CUDA 矩陣轉置

在CUDA SDK中有可以找到很多常用的代碼, 比如矩陣轉置, CUDA的記憶體顯示, 都比較具有代表性

#define BLOCK_DIM 16</p><p>__global__ void transpose(float *odata, float *idata, int width, int height)<br />{<br />__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];</p><p>unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;<br />unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;<br />if((xIndex < width) && (yIndex < height))<br />{<br />unsigned int index_in = yIndex * width + xIndex;<br />block[threadIdx.y][threadIdx.x] = idata[index_in];<br />}</p><p>__syncthreads();</p><p>xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;<br />yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;<br />if((xIndex < height) && (yIndex < width))<br />{<br />unsigned int index_out = yIndex * height + xIndex;<br />odata[index_out] = block[threadIdx.x][threadIdx.y];<br />}<br />}</p><p>

 

      這個是個典型的利用 shared memory 來使記憶體讀寫加速的例子, CUDA SDK 中有個 transpose_native, 是採用比較原始的辦法, 將每列的第一個元素讀出來湊成一行, 這個辦法顯而易見是比較的慢的,比上面這個程式要慢10倍以上.

      使用16x16的塊大小來做運算，個人認為有兩個原因：

1）在X86平台上由於 cache line 可以產生同樣的效果, 但是GPU中為了放更多的運算單元, 並沒有製造很多的 Cache Line, 但是查看nvidia 的GPU,發現每次資料轉送的基本數量是 256bit, 或者 512bit, 假設是 512bit, 則每次是傳輸 16 個單精確度浮點數, 所以對應的你也至少要每次出來一行中的 16 個單精確度浮點數據, 才能充分利用顯存頻寬

2）如果使用8X8: 64 threads/block. 每個SM至多接受768 threads, 即12 blocks。但是, SM至多接受8 blocks, 所以實際上僅有512threads!如果使用16X16: 256 threads/block.每個SM至多接受768 threads, 即3 blocks 只要其它計算資源許可，可以滿負荷工作，SM滿負荷工作效率最高！