这是一个矩阵相乘的kernel函数
1 __global__ static void MatrixMulKernel(float* Md,float* Nd,float* Pd,int Width)
2 {
3 //共享存储器保存从从全局存储器中加载的数据
4 __shared__ float Mds[TILE_WIDTH][TILE_WIDTH]; //TILE_WIDTH=16
5 __shared__ float Nds[TILE_WIDTH][TILE_WIDTH];
6
7 //计算Pd和Md中元素的行索引
8 int bx=blockIdx.x;
9 int by=blockIdx.y;
10 int tx=threadIdx.x;
11 int ty=threadIdx.y;
12 //Pd的行和列
13 int Row = by*TILE_WIDTH+ty;
14 int Col = bx*TILE_WIDTH+tx;
15
16 float Pvalue = 0.0;
17 //第k小块
18 for (int k=0;k<Width/TILE_WIDTH;k++) //width=128;
19 {
20 //通过协作把Md和Nd的块加载到共享存储器中
21 Mds[ty][tx]=Md[Row*Width+k*TILE_WIDTH+tx];
22 Nds[ty][tx]=Nd[(k*TILE_WIDTH+ty)*Width+Col];
23 __syncthreads(); //等待块中其他线程同步
24
25 for(int m=0;m<TILE_WIDTH;m++)
26 Pvalue +=Mds[ty][m]*Nds[m][tx];
27 __syncthreads(); //等待其他线程计算完,因为Pvalue要用到下一个块的计算
28 }
29 //每个线程负责计算P中的一个元素
30 Pd[Row*Width+Col]=Pvalue;
31 }
其中,Md为n1*128矩阵的元素,Nd为128*n2的矩阵元素,Pd就为计算出的n1*n2矩阵的元素。当n1*n2比较大时,为pd分配显存时就会提示out of memory。
所以,对于比较大的矩阵相乘有没有解决办法?
|