3

假设我有一个包含 32 个线程的块，需要随机访问 1024 个元素的数组。我想通过最初将块从全局转移到共享来减少全局内存调用的数量。我有两个想法：

A：

my_kernel()
{
    CopyFromGlobalToShared(1024 / 32 elements);
    UseSharedMemory();
}

或 B：

my_kernel()
{
    if (first thread in block)
    {
        CopyFromGlobalToShared(all elements);
    }
    UseSharedMemory();
}

哪个更好？或者还有其他更好的方法吗？

4

2 回答 2

5

A 更好。

与 CPU 相比，GPU 具有更高的内存带宽。然而，只有在 GPU 中运行的线程遵循一定的模式时，才能实现峰值带宽。

这种模式需要合并 mem 访问。这意味着您需要使用多个线程来访问全局内存中的顺序地址，并特别注意对齐。

您可以在 CUDA 文档中找到有关合并访问全局内存的更多详细信息。

http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#coalesced-access-global-memory

于 2013-09-26T14:04:27.037 回答
1

A 更好，但取决于元素的大小 - 不一定是最好的！

您想要的是每个线程访问 32 位大小的相邻字（64 位也可能在较新的硬件上工作）。如果元素尺寸更大，您可能更喜欢花哨一点：

//assumes sizeof(T) is multiple of sizeof(int)
//assumes single-dimention kernel
//assumes it is launched for all threads in block
template <typename T>
__device__ void memCopy(T* dest, T* src, size_t size) {
    int* iDest = (int*)dest;
    int* iSrc = (int*)src;
    for(size_t i = threadIdx.x; i<size*sizeof(T)/sizeof(int); i+=blockDim.x)
        iDest[i] = iSrc[i];
    __syncthreads();   
}

注意：这适用于所有内存传输（全局->全局、全局->共享、共享->全局）。即使是没有统一内存寻址的旧设备也是如此，因为函数将被内联。

如果您对更大的元素使用数组结构方法，则不会出现问题。

于 2015-02-17T19:40:05.050 回答