CUDA 版本矩陣乘

說明:html

1.轉載請聯繫本人git

2.代碼在最後

問題描述

矩陣乘法 C = aAB + bC
其中a,b爲常數,A,B,C爲矩陣程序員

實驗要求

  1. 根據內存大小測不一樣規模矩陣的處理速度(GFLOPS/s),並給出計算公式。
  2. 請計算系統的理論峯值,若是沒有達到理論峯值,嘗試給出緣由。

方法

CUDA矩陣的優化有多個思路,在本次試驗中我使用了shared memory進行訪問速度的提高,嘗試減小if-else語句的出現,避免串行化,同時作了精度優化以下降錯誤率(結果不怎麼好)。
同時,參考Nvidia給的Samples中0_simple裏的matrixMulCUBLAS相關代碼,思考提高空間。
github

實驗

結果及分析

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
bash

沒有達到理論峯值的緣由是:
程序並不僅是在作浮點數運算或只是在訪問內存;
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,較我本身的代碼好了三倍有餘。

結論

  1. shared memory的正確使用可以很是顯著地提高矩陣乘法的性能
  2. if-else語句產生的divergence問題是很是值得關注的,若是不盡可能減小分支語句的使用,並行性能將不會有很好的體現。
  3. CUDA編程的調試難度不亞於OpenMP,以及優化代碼須要細心,按部就班。另外也要謹記Amdahl優化定律,抓重要的代碼進行性能優化。
  4. CUBLAS的代碼表示,可以很是接近峯值是一件很是複雜且辛苦的事情,須要慢慢測量和分析。

參考

《在CUDA中實現任意尺寸的矩陣乘法》

《矩陣乘法—CUDA優化記錄》

《利用cuda的cublas庫實現任意矩陣的乘法》

代碼地址

我的GitHub:Icarusintheworld

相關文章
相關標籤/搜索