将全局复制到共享内存的最佳方法答案

【问题标题】：Best way to copy global into shared memory将全局复制到共享内存的最佳方法
【发布时间】：2015-02-18 03:40:05
【问题描述】：

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

答：

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

或乙：

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

哪个更好？还是有其他更好的方法？

【问题讨论】：

我想说第一个可能会更快，如果实施得当，原因很简单，所有其他线程都会等待第一个线程完成加载。

标签： c++ c cuda

【解决方案1】：

A 更好。

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

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

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

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

【讨论】：

所以我用A 表示的方法更好/更“标准”？
是的，如果CopyFromGlobalToShared() 执行得很好，A 会好很多。

【解决方案2】：

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();   
}

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

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

【讨论】：