爲什麼全球+共享的速度比全球獨自

我需要一些幫助理解的羅恩·法伯的代碼的行爲：http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/208801731?pgno=2 爲什麼全球+共享的速度比全球獨自

我不理解如何使用共享MEM的是給在非共享內存更快的性能版。即，如果我再添加一些索引計算步驟並使用另一個Rd/Wr循環來訪問共享內存，那麼如何比單獨使用全局內存更快？在這兩種情況下，相同的數字或Rd/Wr循環訪問全局內存。每個內核實例只能訪問一次數據。數據仍然使用全局內存進/出。內核實例的數量是相同的。寄存器數量看起來是一樣的。如何添加更多的處理步驟使其更快。（我們沒有減去任何流程步驟。）基本上我們正在做更多的工作，並且它正在更快地完成。

共享內存訪問速度比全球快得多，但它不是零（或負值）。我錯過了什麼？

的「慢」的代碼：

__global__ void reverseArrayBlock(int *d_out, int *d_in) { 
int inOffset = blockDim.x * blockIdx.x; 
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 
int in = inOffset + threadIdx.x; 
int out = outOffset + (blockDim.x - 1 - threadIdx.x); 
d_out[out] = d_in[in]; 
}

的「快」的代碼：

__global__ void reverseArrayBlock(int *d_out, int *d_in) { 
extern __shared__ int s_data[]; 

int inOffset = blockDim.x * blockIdx.x; 
int in = inOffset + threadIdx.x; 

// Load one element per thread from device memory and store it 
// *in reversed order* into temporary shared memory 
s_data[blockDim.x - 1 - threadIdx.x] = d_in[in]; 

// Block until all threads in the block have written their data to shared mem 
__syncthreads(); 

// write the data from shared memory in forward order, 
// but to the reversed block offset as before 
int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 
int out = outOffset + threadIdx.x; 
d_out[out] = s_data[threadIdx.x]; 
}

來源

2012-08-13 Doug

你在使用什麼卡？它可以產生顯着的差異。 – 2012-08-13 18:13:08

由於顯着不同，較舊的卡僅支持在「正向順序」模式下有效讀取全局內存。較新的卡片不應該受此影響。（我相信這在2.x發佈時有所變化） – 2012-08-13 19:32:16

該文章寫於2008年。硬件是費米前。我認爲這是'老'卡。 – Doug 2012-08-13 19:45:40

早期支持CUDA的設備（計算能力< 1.2）不會將您的「慢」版本中的d_out [out]寫入視爲合併寫入。這些設備只會在「最好」的情況下合併內存訪問，在這種情況下，第二個半字節中的第i個線程訪問第i個字。結果，將發出16個內存事務來服務每半個warp的d_out [out]寫入，而不是僅僅一次內存事務。

從計算能力1.2開始，CUDA中的內存合併規則變得更放鬆了。因此，「慢」版本中的d_out [out]寫入也會合並，並且不再需要使用共享內存作爲便箋。

代碼示例的源代碼是2008年6月編寫的文章「CUDA，大衆超級計算：第5部分」。具有計算能力1.2的CUDA支持的設備僅在2009年上市，所以作者該文章清楚地討論了具有計算能力的設備< 1.2。

有關更多詳細信息，請參閱NVIDIA CUDA C Programming Guide中的F.3.2.1小節。

來源

2012-08-13 19:37:46

這是因爲共享的存儲器更接近計算單元，因此等待時間和峯值帶寬將不是這個計算的瓶頸（至少在矩陣乘法的情況下）

但最重要的是，最重要的原因是瓷磚中的很多數字被重複使用很多線程。所以如果你從全球進行訪問，你會多次檢索這些數字。將它們寫入共享內存將消除浪費的帶寬使用

來源

2012-08-13 18:21:20 lurscher

這裏沒有數據重用。它是從一個獨特的全球讀取的，它被寫入另一個獨特的全球一次。在這兩個代碼塊中都會發生相同數量的全局訪問。即相同的全局帶寬。 – Doug 2012-08-13 18:29:35

在查看全局內存訪問時，慢代碼會向前讀取並向後寫入。快速代碼都可以正向讀取和寫入。我認爲快速代碼的速度更快，因爲緩存層次結構在某種程度上以降序訪問全局內存（朝着更高的內存地址）進行了優化。

CPU執行一些推測性提取，在數據被程序觸及之前，它們將從更高的存儲器地址填充緩存行。也許在GPU上發生類似的事情。

來源

2012-08-13 19:23:03

CC 2.0和3.0數據高速緩存層次結構不預取。預取通過軟件說明得到支持（參見PTX手冊）。 SM到L1的訪問都是使用128B的事務。 L1到L2訪問使用32B事務。未緩存的L2（未緩存在L1中）可以使用32B事務完成。 – 2012-08-14 02:48:36

爲什麼全球+共享的速度比全球獨自

回答

相關問題