关于矩阵相乘

加速计算 CUDA
1

大家好：
我刚刚开始学习cuda，希望从最基础的矩阵相乘开始。
我写了一个在CPU上的矩阵相乘的C语言程序（附件），MN＝P，M是一个48的矩阵，N是一个8*2的矩阵。以二维数组的形式是可以实现的。
但我看关于CUDA的书籍，是将矩阵以一维的形式存储的（我是这么理解的，不知道对不对），为了用threadIdx.x和 threadIdx.y，以及M，N的尺寸进行矩阵元素的寻址。
在解决实际问题时，我还需要将二维的矩阵转化为一维的矩阵吗？
如果不用转化为一维矩阵，我需要如何实现矩阵元素的寻址。（我的思路是用指针实现，但每次还要将矩阵给个指针。有没有其它的方法。）
谢谢。

# include "stdio.h"

int main(void)
{        
   // part 1
   // allocate and initialize the matrices M,N,P
   // I/O to rea dthe input matries M and N
   printf("\n--------- this is part 1------------\n");
   // 4 matrix 
   int M_row=4;
   int M_col=8;
   int N_row=M_col;
   int N_col=2;
   
   // 4 for ite
   int i=0;
   int j=0;
   int k=0;
   float tmp=0;

   float M[M_row][M_col];
   float N[N_row][N_col]; 
   float P[M_row][N_col];

   // init M 
   printf("M=\n");
   for(i=0;i<M_row;i++)
   {
   for (j=0;j<M_col;j++)
   {                       
   M[i][j]=(i+1)+0.01*(j+1);
   printf("%3.3f,   ",M[i][j]);
   }
   printf("\n");
   }
   
   // init N 
   printf("N=\n");
   for(i=0;i<N_row;i++)
   {
   for (j=0;j<N_col;j++)
   {                       
   N[i][j]=(i+1)+0.001*(j+1);
   printf("%3.3f,   ",N[i][j]);
   }
   printf("\n");
   }
   
   // part 2
   // Matrix multiplic on CPU
   printf("\n ------------this is part 2------------\n");
   // cal P
   printf("M * N = P\n");
   printf("\n");
   for (i=0;i<M_row;i++)
   {
   for (j=0;j<N_col;j++)
   {
   tmp=0;
   for (k=0;k<M_col;k++)
   {
   tmp=tmp+M[i][k]*N[k][j];
   }
   P[i][j]=tmp;
   }
   }
   printf("\n--------------this is part 3 -----------\n");
   for (i=0;i<M_row;i++)
   {
   for (j=0;j<N_col;j++)
   {
   printf("%3.3f, ",P[i][j]);
   }
   printf("\n");
   }
   return 0;
}
//mod : Tue Apr  9 14:48:05 HKT 2013
//mod : Tue Apr  9 07:06:19 CST 2013
//mod :Mon Apr  8 06:53:31 CST 2013
2

楼主您好，您的问题问的好。

的确，我们可以使用float a[M][N];的形式，那么为何几乎所有的例子中都是用一维数组，然后进行坐标变换呢？

那是因为，float a[M][N]这种，在CUAD C/C++和它配套的host compiler VC中，均需要要求M, N是编译时刻确定的常数。

而一旦M,N是变量，则您无法使用。（需要C99标准的VLA支持，但cuda c和vc均不支持VLA)。

所以，除了固定的矩阵大小的情况，任何为了能处理通用的矩阵的乘法，您只能写成一维的，然后将宽度和高度信息传入，然后自己计算。（请注意，这并不会带来额外的性能损失）。

所以，您的疑惑就此解开。
以及，不排除以后cuda c支持VLA(以及其指针形式变种), 以后可能您可以直接使用，例如：
global void meatball(float *p, int width, int height)
{
float (array)[width] = (float ()[width])p; //建立指针，width是运行时才可知的
//以下直接使用:
//array[ j ] [i ]形式
}

该操作需要未来版本的cuda toolkit支持。您目前还不能这么用。（只是给您提前说明）。

3

感谢“横扫千军”的回复。
学习了。可能还需要再仔细想想（因为比较菜）。但在大体的意思已经明白了。

4

您客气了，服务您是我的荣幸。