1. 程式人生 > >CUDA效能調優(一)--合併訪問&迴圈展開

CUDA效能調優(一)--合併訪問&迴圈展開

1  合併訪問 

       當同一個warp中的所有執行緒都執行同一條指令訪問全域性儲存器中連續的單元時,就獲得最有利的訪問模式。硬體檢測到同一個warp中的這些執行緒訪問全域性儲存器中連續的儲存單元,並將這些單元結合成一個合併的訪問

     合併訪問可以提高DRAM的頻寬利用率,使DRAM在傳輸資料時的速度接近全域性儲存器頻寬的峰值

這裡需要補充一點知識:二維三維陣列的線性對映--多維執行緒,維度對映到線性順序

二維執行緒按照線性順序排列如下圖所示:


也就是說矩陣在記憶體空間中已經被等價線性化為等價的一維陣列,以行優先的方式逐行順序存放。

示例如下:

(1)  在未使用shared memory

的GPU矩陣相乘中的兩種訪問模式


接合訪問模式

在第一次迭代中,訪問0號元素,這些0號元素在global memory中相鄰,硬體接合這些訪問:


非接合訪問模式


在這個例子中得出如下結論:

對於全域性儲存器的訪問,kernel迴圈遍歷一行時,要比遍歷一列的效率低得多。

(2)  在使用shared memory的矩陣相乘中,使用合併方式載入資料

//以合併的模式載入
ds_A[ty][tx] = A[Row*n+t*TILE_WIDTH+tx];
ds_B[ty][tx] = B[(t*TILE_WIDTH + ty)*k+Col];

每個執行緒負責載入一個Md及Nd元素。

m識別左側瓦片的位置。

Md瓦片的一行由TILE_WIDTH個執行緒載入,threadIdx.x變化。

2  指令混合

      在當前裝置中,每個SM的指令處理頻寬有限。每個指令佔用指令處理頻寬,包括浮點計算指令、載入指令和分支指令。消除重複指令可減少指令處理頻寬的壓力,提升kernel函式的整體執行效能。

以下面兩行程式碼為例進行介紹:

  for (int i = 0; i < TILE_WIDTH; ++i)  
            Cvalue += ds_A[ty][i] * ds_B[i][tx]
這兩行程式碼包括一下幾種指令:

(1)迴圈引入額外指令更新計數器k    1次

(2)每次迭代結尾位置執行條件跳轉    1次

(3)使用k計算Mds,Nds索引引入了地址運算指令    2次

(4)浮點乘加計算指令    2次

浮點乘加計算指令僅佔了1/3的指令。由於指令處理頻寬有限,因此這種指令混合將能取得的效能限制在頻寬峰值的1/3以內。

為了改善這用指令混合,採取迴圈展開(Unroll loop)的方法。

將上面的程式碼修改如下:

Cvalue += ds_A[ty][0] * ds_B[0][tx]+ds_A[ty][1] * ds_B[1][tx]
		         +ds_A[ty][2] * ds_B[2][tx]+ds_A[ty][3] * ds_B[3][tx]
		         +ds_A[ty][4] * ds_B[4][tx]+ds_A[ty][5] * ds_B[5][tx]
				 +ds_A[ty][6] * ds_B[6][tx]+ds_A[ty][7] * ds_B[7][tx]
				 +ds_A[ty][8] * ds_B[8][tx]+ds_A[ty][9] * ds_B[9][tx]
				 +ds_A[ty][10] * ds_B[10][tx]+ds_A[ty][11] * ds_B[11][tx]
				 +ds_A[ty][12] * ds_B[12][tx]+ds_A[ty][13] * ds_B[13][tx]
				 +ds_A[ty][14] * ds_B[14][tx]+ds_A[ty][15] * ds_B[15][tx];  

程式碼分析:

Long-multiply-add操作

消除分支指令以及loop計數器更新

索引為常量-編譯器可以使用載入指令的定址模式對應的偏移量,這樣可以消除地址運算指令。

因此,這個很長的表示式的執行速度幾乎可以接近效能的峰值。