優化CUDA內核不規則的內存訪問

我有以下CUDA內核，這似乎很「艱難」來優化：優化CUDA內核不規則的內存訪問

__global__ void DataLayoutTransformKernel(cuDoubleComplex* d_origx, cuDoubleComplex* d_origx_remap, int n, int filter_size, int ai) 
{ 
    for(int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < filter_size; idx+=blockDim.x * gridDim.x) 
    { 
     int index = (idx * ai) & (n-1); 
     d_origx_remap[idx] = d_origx[index]; 
    } 
} 

//Parameters were defined before 
int permute[loops] = {29165143,3831769,17603771,9301169,32350975, ...} 
int n = 33554432; 
int filter_size = 1783157; 

for(int i=0; i<loops; i++) 
{ 
    DataLayoutTransformKernel<<<dimGrid, dimBlock, 0, stream[i]>>>((cuDoubleComplex*) d_origx,(cuDoubleComplex*)d_origx_remap+i*filter_size, n, filter_size, permute[i]); 

}

內核的目的是從非正規到正規重新排序d_origx[]數據佈局（d_origx_remap）。內核以不同的訪問級別（ai）啓動幾次。

這裏面臨的挑戰是在引用d_origx[index]數組時不規則的存儲器訪問模式。我的想法是使用共享內存。但是對於這種情況，使用共享內存合併全局內存訪問似乎非常困難。

有沒有人有如何優化這個內核的建議？

來源

2013-12-11 user3089852

也許你可以隱藏一些延遲，通過啓動內核以交錯的方式同時處理重新映射的緩衝區和變換內核。第一次迭代只是重映射內核。第二次迭代同時啓動重新映射內核和處理第一次重新映射結果的內核等。也許還要考慮將此功能直接滾動到下一個內核中（讓內核將它的值在跨步位置）。 –

Trove庫是一個支持AoS支持的CUDA/C++庫，可能爲隨機AoS訪問提供接近最佳性能。從GitHub頁面看來，trove比16字節結構的天真方法要快2倍左右。

https://github.com/BryanCatanzaro/trove

Random access performance using Trove compared to the naive direct access approach

來源

2013-12-11 06:54:48

只要提到Trove適用於計算能力'3.0'及以上。 – JackOLantern

我不知道，你可以做很多工作來優化你的代碼。

根本沒有線程合作，所以我會說共享內存不是要走的路。

您可以嘗試使用const和__restrict__關鍵字改變

__global__ void DataLayoutTransformKernel(cuDoubleComplex* d_origx, cuDoubleComplex* d_origx_remap, int n, int filter_size, int ai)

到

__global__ void DataLayoutTransformKernel(const cuDoubleComplex* __restrict__ d_origx, cuDoubleComplex* __restrict__ d_origx_remap, const int n, const int filter_size, const int ai)

即。特別是__restrict__將使nvcc執行一些優化，請參見CUDA C編程指南的第B.2節。對於開普勒體系結構，const和__restrict關鍵字可由編譯器標記爲通過只讀數據高速緩存加載，請參閱Kepler architecture whitepaper。

來源

2013-12-11 18:29:07 JackOLantern

優化CUDA內核不規則的內存訪問

回答

相關問題