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

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

A：

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

或B：

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

哪个更好？ 还是有另一种更好的方法吗？

Answer 1

A更好。

与CPU相比，GPU具有更高的内存带宽。 但是，只有当GPU中运行的线程遵循某种模式时，才能实现峰值带宽。

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

您可以在CUDA文档中找到有关Coalesced Access to Global Memory的更多详细信息。

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

Answer 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  __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 

注意：这适用于所有内存传输（global-> global，global-> shared，shared-> global）。 即使是没有统一内存寻址的旧设备也是如此，因为该function将被内联。

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