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

比方说，我有一个需要随机访问1024个元素数组的32个线程块。我想通过将块从全局转移到共享来减少全局内存调用的次数。我有两个想法去做：将全局复制到共享内存的最佳方法

答：

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

或B：

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

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

来源

2013-09-26 Boyko Perfanov

我想说的第一个可能会更快，如果正确实施，原因很简单，所有其他线程将等待第一个线程完成加载。 – JackOLantern

A更好。

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

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

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

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

来源

2013-09-26 14:04:27 kangshiyin

那么我表示为'A'的方法是更好还是更“标准”呢？ –

是的，如果'CopyFromGlobalToShared（）'得到很好的实现，A会好得多。 – kangshiyin

A是更好的，但是根据元素的大小 - 不necessairly最好的！

你想要的是每个线程访问一个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(); 
}

注：这将为所有存储器的传输工作（全球 - >全球，全球 - >共享，shared->全球）。即使是没有统一内存寻址的旧设备也是如此，因为该功能将被内联。

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

来源

2015-02-17 19:40:05 CygnusX1

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

回答

相关问题