cuda實現任意尺寸的矩陣乘法

nvidia gpu常見的塊內線程數最大為1024，當矩陣的行數和列數均小於1024時，我們可以簡單的採用行和列點到點依次相乘構建核函式，即塊內的每個執行緒負責一對元素的乘積計算，然後將所有塊內線程相乘的結果累加求和，得到結果矩陣對應行和列的元素值。

>>code：參照cuda11指導手冊，給出核函式**如下：

// matrix multiplication kernel called by matmul()
__global__ void
matmulkernel
(matrix a, matrix b, matrix c)
}

這裡matrix是乙個結構體，如下：

// matrices are stored in row-major order:
// m(row, col) = *(m.elements + row * m.width + col)
typedef
struct
matrix;

上述核函式簡單、易理解，且未使用塊的共享記憶體，下面將介紹使用子矩陣劃分和共享記憶體的方法實現任意尺寸的矩陣乘法。

思想如下：

設定塊為寬高相等的二維形狀，尺寸大小block_size為32（32*32=1024），從而塊內包含的執行緒個數為1024個，源矩陣a和b可以被na和nb個塊劃分，如下圖所示：

上圖中，每個藍色的塊代表矩陣的劃分得到的子矩陣，每個子矩陣對應乙個執行緒塊，它們的大小均為block_size * block_size, 為了計算c(1, 0)處的值，必須將每個塊內第一行和b裡面對應塊（按矩陣乘法規則相對應）的第0列相乘得到結果，然後求和，塊內則是執行緒求對應元素相乘的結果求和。

值得注意的是，最後一行和最後一列的子矩陣行或列可能小於block_size（圖二中綠框內的子矩陣），因此在核函式計算時，需要作判斷，約束僅計算在行列範圍內的值。

比如說：矩陣a的shape為（33， 32），矩陣b的shape為（32， 35），則矩陣a的第二行第乙個塊的行數為1，矩陣b的第二列第乙個塊的列數為3，核函式內在該執行緒塊下僅計算1行和3列的值，詳見**及注釋。

>>code：**如下：

template
<
int block_size> __global__ void
matrixmulcuda
(float
* c,
float
* a,
float
* b,
int wa,
int wb,
int ha,
int hb)
else
subbh = subaw;if(
(by +1)
* block_size > ha)
//a矩陣最後一行的塊的行數少於block_size
elseif(
(bx +1)
* block_size > wb)
//b矩陣最後一列的塊的列數少於block_size
else
/* 開闢塊內共享記憶體 */
__shared__ float as[block_size]
[block_size]
; __shared__ float bs[block_size]
[block_size]
;/* 為行和列範圍內的子矩陣對應元素賦值 */
if(ty < subah && tx < subaw)
if(ty < subbh && tx < subbw)
__syncthreads()
;//展開迴圈來 編譯以加速 
#pragma unroll
//內迴圈計算每個子矩陣內對應行和列的向量乘積，累加到之前得到的值上
for(
int k =
0; k < subaw; k++)}
__syncthreads()
;}//滿足行和列約束內的元素計算乘積並求和
if(ty < subah && tx < subbw)
}

此時函式的呼叫舉例為：

/* 引數設定 */
dim3 dimsa
(1055
,2137,1
);//矩陣的寬、高和未使用引數1
dim3 dimsb
(108
,1055,1
);//矩陣的寬、高和未使用引數1
const
int block_size =32;
/* 矩陣初始化、記憶體傳遞等常規步驟
....
*//* 呼叫核函式計算 */
dim3 threads
(block_size, block_size)
;dim3 grid
((dimsb.x -1)
/ threads.x +1,
(dimsa.y -1)
/ threads.y +1)
;matrixmulcuda
<<
>
>
(d_c, d_a, d_b,
dimsa.x, dimsb.x, dimsa.y, dimsb.y)
;

當然，也可以對原始矩陣padding使其成為block_size的整數倍來實現，理論上能有更快的計算速度。

cuda實現任意尺寸的矩陣乘法

cuda矩陣相乘 CUDA的矩陣乘法

矩陣乘法矩陣乘法的基本實現

CUDA學習矩陣乘法的並行運算

cuda實現任意尺寸的矩陣乘法

cuda矩陣相乘 CUDA的矩陣乘法

矩陣乘法 矩陣乘法的基本實現

CUDA學習 矩陣乘法的並行運算

相關推薦

矩陣乘法矩陣乘法的基本實現

CUDA學習矩陣乘法的並行運算