CUDA縮減優化示例

我正試圖在CUDA Reduction上實現優化，並且要成功完成，直到第6部分。感謝大家的幫助。爲了獲得CUDA的完整感受，我還需要完成最終優化，如幻燈片＃31中所述，稱爲算法級聯。CUDA縮減優化示例

這個想法本質上是每個線程有512個元素，並在執行縮減之前將所有元素相加。

我嘗試了一種方法，我的每個線程都從內存中訪問連續的512個數字。不幸的是，它的表現最差。我猜測是銀行衝突的一個原因，但還沒有完全弄清楚。你們中的任何一個人能否提出這種行爲的原因？

我還發布了下面的Nvidia提供的示例代碼。

unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x; 
unsigned int gridSize = blockSize*2*gridDim.x; 
sdata[tid] = 0; 
while (i < n) { 
    sdata[tid] += g_idata[i] + g_idata[i+blockSize]; 
    i += gridSize; 
} 
__syncthreads();

有幾個參數沒有定義。我可以推斷blockSize等於每塊的線程數。但我無法推斷變量'gridSize'的重要性。訪問內存的適當方式是什麼，以便我們獲得更好的性能？這是一個跨越訪問的例子嗎？

如果您有任何其他問題，請提前在下面提供幫助和評論。

來源

2016-05-17 Rahul Gunjan

所有這些縮減代碼的完整工作示例在相應的[CUDA示例代碼]（http://docs.nvidia.com/cuda/cuda-samples/index.html#cuda-parallel-reduction）中提供。你不應該猜測任何參數。我懷疑你提供了足夠的信息來解釋你的觀察。如果您的第6部分的實施表現不佳，您可能應該運行CUDA示例代碼並研究差異。 –

這是一個合併訪問的例子。最好的gridDim取決於你的硬件。根據每個線程的寄存器和每個塊的最大線程數，該值應該是硬件上可用多處理器數量的某個乘數。如果你的問題足夠大，8倍的多處理器計數對於開普勒來說是個好選擇，而對於麥克斯韋來說是16倍。

來源

2016-05-17 06:04:53

假設你有blockDim.x = blockSize = 256線程每塊，並且gridDim.x = 32塊在網格中，並且你想減少一個大數組g_idata[8,192,000]。

然後你總共有8192個線程。讓我們用

thread[x][y], x=0..31, y=0..255

來表示這些線程。

每個thread[x][y]加載

g_idata[iter*512*x+y] and g_idata[iter*512*x+256+y], iter = 0 .. 999

到共享存儲器sdata。

對於每個迭代iter，所有8192 threads[x][y]將從GPU內存加載gridSize = 16384元素。

這是合併內存訪問，它是訪問GPU內存的正確方式。

然而，你的方式，其中每個thread[x]讀取data[i*x*512 .. i*(x+1)*512-1], i=0...不是一個好方法。實際上，這是訪問GPU內存的最低效的方式。

來源

2016-05-17 08:50:37 kangshiyin

@RobertCrovella我認爲代碼是Nvidia演示，而不是Rahul的方式。這一個是：'我嘗試了一種方法，我的每個線程都從內存中訪問連續的512個數字。' – kangshiyin

對不起，我誤解了你的答案。刪除了我以前的評論。 –

CUDA縮減優化示例

回答

相關問題