將全局複製到共享內存的最佳方法

比方說，我有一個需要隨機訪問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

將全局複製到共享內存的最佳方法

回答

相關問題