CUDA，找到最大使用減少，錯誤

這裏是我的代碼試圖做減少找到一個塊中的50值數組的最大值。我已經將數組填充到64.CUDA，找到最大使用減少，錯誤

對於線程1-31，我有正確的maxVal打印輸出，但對於線程32-49，它是一個完全隨機數。我不知道我做錯了什麼。

btw。我認爲我不需要在展開每一行的同時，但顯然我必須。有關這個的任何建議？

在此先感謝您的幫助。

//block size = 50 


__syncthreads(); 

if (tid<32){ 

    cptmp[tid]=(cptmp[tid]< cptmp[tid+32]) ? cptmp[tid+32] : cptmp[tid];__syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+16]) ? cptmp[tid+16] : cptmp[tid];__syncthreads(); 
    cptmp[tid]=(cptmp[tid]< cptmp[tid+8]) ? cptmp[tid+8] : cptmp[tid]; __syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+4]) ? cptmp[tid+4] : cptmp[tid]; __syncthreads(); 
    cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid]; __syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid]; __syncthreads(); 

} 

__syncthreads(); 

//if (tid==0) { 
    maxVal=cptmp[0]; 
    if(bix==0 && biy==0) cuPrintf(" max:%f x:%d y:%d\n", maxVal, blockIdx.x, blockIdx.y); 
//}

來源

2011-06-28 Kiarash

很難弄清楚發生了什麼，而不描述你的線程/塊/網格結構，以及你如何計算'tid'。 –

我的塊尺寸是50，所以tid = threadIdx.x在0-49之間。網格大小是（40，一個很大的數字）。每個塊都在cptmp數組中找到最大值，它與64的大小共享。maxVal也是共享的。 – Kiarash

這是一個更高效的（至少在費米GPU上）和使用volatile的正確代碼。將T替換爲您的類型（或使用模板）：

if (tid<32) { 
    volatile T *c = cptmp; 
    T t = c[tid]; 
    c[tid] = t = (t < c[tid+32]) ? c[tid+32] : t; 
    c[tid] = t = (t < c[tid+16]) ? c[tid+16] : t; 
    c[tid] = t = (t < c[tid+ 8]) ? c[tid+ 8] : t; 
    c[tid] = t = (t < c[tid+ 4]) ? c[tid+ 4] : t; 
    c[tid] = t = (t < c[tid+ 2]) ? c[tid+ 2] : t; 
    c[tid] = t = (t < c[tid+ 1]) ? c[tid+ 1] : t; 
}

爲什麼更高效？那麼，爲了在__syncthreads()的情況下的正確性，我們必須使用一個易失性指針來共享內存。但是這會迫使編譯器「尊重」所有讀寫共享內存 - 它不能優化和保存寄存器中的任何內容。因此，通過明確始終將c[tid]保存在臨時文件t中，我們會爲每行代碼保存一個共享內存負載。由於Fermi是一種只能使用寄存器作爲指令操作數的加載/存儲架構，這意味着我們每行保存一條指令或總共6條指令（總體上約爲25％，我預計）。

在舊的T10/GT200架構和更早的版本中，您的代碼（帶有易失性且不帶__syncthreads（））將同樣有效，因爲該架構可以直接從共享內存獲取每條指令的一個操作數。

此代碼應該是等價的，如果你喜歡if超過?:：

if (tid<32) { 
    volatile T *c = cptmp; 
    T t = c[tid]; 
    if (t < c[tid+32]) c[tid] = t = c[tid+32]; 
    if (t < c[tid+16]) c[tid] = t = c[tid+16]; 
    if (t < c[tid+ 8]) c[tid] = t = c[tid+ 8]; 
    if (t < c[tid+ 4]) c[tid] = t = c[tid+ 4]; 
    if (t < c[tid+ 2]) c[tid] = t = c[tid+ 2]; 
    if (t < c[tid+ 1]) c[tid] = t = c[tid+ 1]; 
}

來源

2011-06-30 04:06:58 harrism

非常感謝你的哈利。雖然我對volatile的概念不是很滿意。也可以使用類似的方法進行求和嗎？如果可能，請告訴我如何？ – Kiarash

請參閱NVIDIA CUDA SDK中的「縮小」示例。求和的基本思想是'c [tid] = t = c [tid + 32] + t;' – harrism

感謝harrism。這工作得很好。對於記錄，我把這件作品放在這裏：\t if（tid <32）{0} {0} {0} {0} \t \t float t = c [tid]; \t \t \t c [tid] = t = c [tid + 32] + t; \t \t c [tid] = t = c [tid + 16] + t; \t \t c [tid] = t = c [tid + 8] + t; \t \t c [tid] = t = c [tid + 4] + t; \t \t c [tid] = t = c [tid + 2] + t; \t \t c [tid] = t = c [tid + 1] + t; \t} – Kiarash

不要在發散的代碼中使用__syncthreads()！來自給定塊的所有線程或無線程應該在同一位置到達每個__syncthreads()。

來自單個warp（32個線程）的所有線程都是隱式同步的，所以您不需要__syncthreads()將它們放在一起。但是，如果您擔心一個線程的共享內存寫入可能不會被同一個warp的另一個線程看到，請使用__threadfence_block()。

詳細說明__threadfence_block()的重要性。考慮以下兩行：

cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid]; 
cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid];

可能編譯成這樣的事情：

int tmp; //assuming that cptmp is an array of int-s 
tmp=cptmp[tid]; 
tmp=(tmp<cptmp[tid+2])?cptmp[tid+2]:tmp; 
tmp=(tmp<cptmp[tid+1])?cptmp[tid+1]:tmp; 
cptmp[tid]=tmp;

雖然這將是一個單線程代碼是正確的，它顯然失敗了CUDA。

爲了避免這樣的優化，您要麼將cptmp數組聲明爲volatile，要麼在行之間添加此__threadfence_block()。該函數確保在函數存在之前，同一塊的所有線程都可以看到當前線程的共享內存。

存在類似的__threadfence()函數來確保全局內存可見性。

來源

2011-06-28 20:10:41 CygnusX1

非常感謝CygnusX1提供的這些信息。我使用了__threadfence_block（）方法和volatile技巧，顯然它們都可以完美工作。但是你知道哪一個有更好的表現嗎？ – Kiarash

性能差異可能是可以忽略的。我會說'__threadfence_block（）'給你更多的自由：你可以準確地說明共享內存寫入的位置必須可見。針對變量的Statig'volatile'強制共享內存在任何地方寫入，即使在使用臨時寄存器的地方可能是好的並且是有益的。 – CygnusX1

Fences是實際的指令，所以它們比volatile更昂貴 - 假設你的代碼不會強制從共享內存讀取過多的內容。查看我的答案以獲得更高效的代碼（如果我可以適應並格式化，它會將其置於評論中）。 – harrism

對於大家誰將來在這個線程絆倒，像我一樣，這裏是除了harrism回答一個建議 - 它可能是從性能的角度來看值得考慮的隨機操作，所以更新的代碼，以獲得最大程度的發揮使用單經64種元素的應該是這樣的：

auto localMax = max(c[tid], c[tid + 32]);  
for (auto i = 16; i >= 1; i /= 2) 
{ 
    localMax = max(localMax, __shfl_xor(localMax, i)); 
} 
c[tid] = localMax;

只有兩個讀取並從全局內存一個寫需要，所以它非常整齊。

來源

2014-02-13 00:54:33

CUDA，找到最大使用減少，錯誤

回答

相關問題