(2)那麼下面就是不使用shared
memory的並行化演算法的思路。簡單地來說，就是將上述可並行化的部分傳遞給gpu，使用cuda來計算。**如下:
void matrixmulondevice(float*
m, float* n, float* p, intwidth)
int size = width * width * sizeof(float);
float* md,
nd,pd;
//設定呼叫核心函式時的執行緒數目
dim3 dimblock(width,
width);
dim3 dimgrid(1,
1);//在裝置儲存器上給m和n矩陣分配空間，並將資料複製到裝置儲存器中
cudamalloc(&md,
size);
cudamemcpy(md,
m, size, cudamemcpyhosttodevice);
cudamalloc(&nd,
size);
cudamemcpy(nd,
n, size, cudamemcpyhosttodevice);
//在裝置儲存器上給p矩陣分配空間
cudamalloc(&pd,
size);
//核心函式呼叫，將在後續部分說明
//只使用了乙個執行緒塊(dimgrid)，此執行緒塊中有width*width個執行緒
matrixmulkernel<<
dimblock>>>(md,
nd,pd，width);
// 從裝置中讀取p矩陣的資料
cudamemcpy(p,
pd, size, cudamemcpydevicetohost);
// 釋放裝置儲存器中的空間
cudafree(md);
cudafree(nd);
cudafree (pd);
__global__ void matrixmulkernel(float*
md,float* nd,
float* pd, intwidth)
// 2維的執行緒id號
int tx = threadidx.x;
int ty = threadidx.y;
// pvalue用來儲存被每個執行緒計算完成後的矩陣的元素
float pvalue = 0;
//每個執行緒計算乙個元素
for(intk
= 0; k < width; ++k)
float
melement = md[ty
* width + k];
float
nelement = nd[k
* width + tx];
pvalue
+= melement * nelement;
// 將計算結果寫入裝置儲存器中
pd[ty*
width + tx]
= pvalue;
(3)這裡面有乙個很重要的問題就是所有的資料都放在了global
memory裡面，資料的傳輸速度會受到頻寬的限制，並且矩陣的大小受到計算能力的限制。所以可以將儲存器優化，演算法優化，使它更有效率地執行。
思路如下：
將整個大的width*width的矩陣分成tile_width*tile_width的小矩陣，一共可分為(width/tile_width)*(width/tile_width)個。當然此處要注意tile_width和計算能力的大小。
然後是將這些小矩陣分別進行計算，將計算結果放入shared
memory中，如上面的第二張圖。簡單來說就是m的藍色部分乘以n的藍色部分，結果放入黃色部分中；再用m的橙色部分乘以n的橙色部分，結果加入黃色部分中，依次類推。
定義blockspergrid和threadsperblock：
dim3
dimblock(tile_width,tile_width);
dim3
dimgrid(width/tile_width,width/tile_width);
呼叫核心函式：
matrixmulkernel<<>>(md,nd,pd,width);
核心函式如下：
__global__ void matrixmulkernel(float*
md,float* nd,
float* pd, intwidth)
//獲得執行緒塊號
intbx=
blockidx.x;
intby
= blockidx.y;
//獲得塊內的執行緒號
inttx=
threadidx.x;
intty=
threadidx.y;
//pvalue：執行緒計算完成後的子矩陣元素——自動變數
float pvalue = 0;
//迴圈，遍歷m和n的所有子矩陣
for (intm
= 0; m < width/tile_width; ++m)
//獲取指向當前矩陣m子矩陣的指標msub
float* mdsub =
getsubmatrix(md,m,by,width);
//獲取指向當前矩陣n的子矩陣的指標nsub
float* ndsub = getsubmatrix(nd,bx,m,width);
//共享儲存器空間宣告
__shared__ float mds[tile_width][tile_width];
__shared__ float nds和getmatrixelement(mdsub,tx,ty,width)【獲取子矩陣中某個元素的位址】的示意圖:
接下來是儲存結果：
//獲取指向矩陣p的子矩陣的指標
matrix psub
= getsubmatrix(p,bx,by);
//向全域性儲存器寫入執行緒塊計算後的結果子矩陣
//每個執行緒寫入乙個元素
setmatrixelement(psub,tx,ty,pvalue);
總的來說，核心函式的思路就是：
①獲得執行緒塊號以及塊內的執行緒號
②申明乙個自動變數以記錄執行緒計算完成之後的子矩陣元素
③從第乙個子矩陣開始，遍歷m和n的所有子矩陣，以完成如下操作：
④獲取指向當前矩陣m、n子矩陣的指標msub和nsub
⑤宣告兩個tile_width*tile_width大小的共享儲存器空間，分別對應m的子矩陣和n的子矩陣
⑥分別載入m和n的子矩陣的元素到共享儲存器中
⑦同步，確保所有的元素都已經載入共享儲存器中
⑧開始計算，每個執行緒計算執行緒塊內子矩陣的乙個元素
⑨同步，確保計算過程已經完成
⑩獲取指向矩陣p的子矩陣的指標
⑪向全域性儲存器中寫入計算結果後的子矩陣
但是小白現在還有乙個問題，按理說，pvalue應該是儲存在register中才對，那麼，向全域性儲存器中寫入計算結果的值是cuda自動從暫存器中取值然後再寫入嗎，但又有聽說register是不可定址的？所以表示不清楚，只好留待以後**吧。
CUDA矩陣相乘
include include include include 並行矩陣乘法kernel函式 global void matrixmultiply int m d,int n d,int p d,int width p d row width col p 每個執行緒計算p d矩陣的乙個元素 生成隨機...
CUDA系列三 矩陣相乘
本博文主要講解下基於cuda的矩陣相乘，cuda特別擅長的就是矩陣乘法，而且也比較容易實現。通過矩陣乘法的實現，可以比較容易理解cuda的核心思想。網上也有很多基於cuda實現的矩陣乘法，但是感覺都不完成，要不就是有錯，本文給出的 都是經過驗證可行的，希望能夠幫助到大家。矩陣乘法實現方式一 矩陣乘法...
CUDA 共享儲存器實現矩陣相乘
共享儲存器使用 shared 限定詞分配。cuda程式設計介面 執行初始化與裝置儲存器 cuda程式設計介面 使用nvcc編譯器的相容性 cuda程式設計介面 如何用nvcc編譯cuda程式 cuda程式設計模型 儲存器層次和異構程式設計 cuda程式設計模型 核心與執行緒層次概述 正如在前面的文章...