說明:html
1.轉載請聯繫本人git
2.代碼在最後矩陣乘法 C = aAB + bC
其中a,b爲常數,A,B,C爲矩陣程序員
CUDA矩陣的優化有多個思路,在本次試驗中我使用了shared memory進行訪問速度的提高,嘗試減小if-else語句的出現,避免串行化,同時作了精度優化以下降錯誤率(結果不怎麼好)。
同時,參考Nvidia給的Samples中0_simple裏的matrixMulCUBLAS相關代碼,思考提高空間。
github
處理速度公式=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
bash
沒有達到理論峯值的緣由是:
程序並不僅是在作浮點數運算或只是在訪問內存;
sgemm中還存在着if-else語句,使得線程存在着divergence;
因爲大小分配的問題存在着Occupancy;
存在着空閒的線程;
以及操做系統的線程調度,和服務器自己的不穩定性等等。服務器
Shared memory的做用在於下降對於全局數據的訪問,充分利用Cuda中線程能夠有獨立的內存空間及寄存器,以及block中線程之間能夠通訊的特色
在shared memory大小定義中,Width要保證不能大於XY對應dim的最小值,另外在測試的時候發現,若是width_size大於32,那麼獲得的結果是全錯(不管XY的dim有多大)暫時不清楚爲何。
函數
在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(),在分配的時候動態設定邊界大小,可是參考調用以後其優化的效果不是很明顯,沒有原做者所說的三倍性能提高,可能和本人的相關知識掌握不足有關。
在sgmm函數的for循環以前,使用 #pragma unroll ,GFLOPS/s提高了10個左右的點,效果比較顯著。
另外參考課件,有考慮過用Parallel Reduction中的連續訪問的方法,可是運行以後程序報錯或者錯誤率很高,暫時沒有找到解決辦法。
優化後的巔峯狀態
CUBLAS 是Nvidia程序員專門優化過的函數,性能表現極好,因爲代碼不開源,暫時不瞭解應該如何調整代碼。
下面的一次測試顯示,在維度爲680的矩陣狀況下,其performance = 1272 GFlops/s , time = 0.154 msec,較我本身的代碼好了三倍有餘。
我的GitHub:Icarusintheworld