說明:
1.轉載請聯絡本人
2.程式碼在最後問題描述
矩陣乘法 C = aAB + bC
其中a,b為常數,A,B,C為矩陣
實驗要求
- 根據記憶體大小測不同規模矩陣的處理速度(GFLOPS/s),並給出計算公式。
- 請計算系統的理論峰值,如果沒有達到理論峰值,嘗試給出原因。
方法
CUDA矩陣的優化有多個思路,在本次試驗中我使用了shared memory進行訪問速度的提升,嘗試減少if-else語句的出現,避免序列化,同時做了精度優化以降低錯誤率(結果不怎麼好)。
同時,參考Nvidia給的Samples中0_simple裡的matrixMulCUBLAS相關程式碼,思考提升空間。
實驗
結果及分析
1.假設矩陣維度為n
處理速度公式=2*n/1000000000/time;
頻寬計算公式:
= ( sizeof(int)*dim + sizeof(int)*n + sizeof(float)*n
+ sizeof(float)*dim*2)/1000000000/time;
系統理論峰值(即浮點數理論峰值)
叢集理論浮點峰值
= CPU主頻(GHz)× CPU每時鐘週期執行浮點運算次數 × 節點數 × 8(每節點雙路四核)
=4.2*4*8=134.4GFLOPS/s
峰值頻寬: B=F×D/8=2133MHz*64bit/8=17.064GHz
沒有達到理論峰值的原因是:
程式並不只是在做浮點數運算或只是在訪問記憶體;
sgemm中還存在著if-else語句,使得執行緒存在著divergence;
由於大小分配的問題存在著Occupancy;
存在著空閒的執行緒;
以及作業系統的執行緒排程,和伺服器本身的不穩定性等等。
2.優化過程
2.1嘗試shared memory
Shared memory的作用在於降低對於全域性資料的訪問,充分利用Cuda中執行緒可以有獨立的記憶體空間及暫存器,以及block中執行緒之間可以通訊的特點
在shared memory大小定義中,Width要保證不能大於XY對應dim的最小值,另外在測試的時候發現,如果width_size大於32,那麼得到的結果是全錯(無論XY的dim有多大)暫時不清楚為什麼。
2.2嘗試減少if-else語句
在Sgemm函式中,if-else語句主要用於進行邊界判斷。
這是因為在分配block大小的時候,矩陣的維度可能不能剛好被32整除。例如dim=500時,不進行邊界判斷會引起很多問題。
一個有效的解決方案是,利用ceil的取整函式,在for迴圈中有效限制i的上界。使得對矩陣維度的限制沒有那麼大。
在程式碼中對grid, block 定義如下
dim3 block(DIM_THREAD_BLOCK_Y, DIM_THREAD_BLOCK_Y);
dim3 grid((size_t)ceil( ((float)N) / ((float)block.x) ),
(size_t)ceil( ((float)N) / ((float)block.y)) );
//取整函式ceil
複製程式碼
當然,經過反覆測試表明,矩陣的維度若能被32整除,其效能表現要比不能整除的要好。
另外在搜尋查詢的時候看到有一個方式是利用了cudaMallocPitch(),在分配的時候動態設定邊界大小,但是參考呼叫之後其優化的效果不是很明顯,沒有原作者所說的三倍效能提升,可能和本人的相關知識掌握不足有關。
2.3嘗試採用for迴圈展開
在sgmm函式的for迴圈之前,使用 #pragma unroll ,GFLOPS/s提升了10個左右的點,效果比較顯著。
另外參考課件,有考慮過用Parallel Reduction中的連續訪問的方法,但是執行之後程式報錯或者錯誤率很高,暫時沒有找到解決辦法。
優化後的巔峰狀態
2.4 執行CUBLAS對比
CUBLAS 是Nvidia程式設計師專門優化過的函式,效能表現極好,由於程式碼不開源,暫時不瞭解應該如何調整程式碼。
下面的一次測試顯示,在維度為680的矩陣情況下,其performance = 1272 GFlops/s , time = 0.154 msec,較我自己的程式碼好了三倍有餘。
結論
- shared memory的正確使用能夠非常顯著地提升矩陣乘法的效能
- if-else語句產生的divergence問題是非常值得關注的,如果不盡量減少分支語句的使用,並行效能將不會有很好的體現。
- CUDA程式設計的除錯難度不亞於OpenMP,以及優化程式碼需要細心,循序漸進。另外也要謹記Amdahl優化定律,抓重要的程式碼進行效能優化。
- CUBLAS的程式碼表示,能夠非常接近峰值是一件非常複雜且辛苦的事情,需要慢慢測量和分析。
參考
程式碼地址
個人GitHub:Icarusintheworld