矩阵自乘的问题

加速计算 CUDA
1

///////////////////////////// Aderivative的定义////////////////////////////////////

ADerivative = newfloat*[m_receiverm_tnum]; //30300

  for( i=0; i<m_receiver*m_tnum; ++i)

  {

         ADerivative=new float[m_GridnumX*m_GridnumZ];//30*50

  }

  for( i=0; i<m_receiver*m_tnum; ++i)

         for( j=0; j<m_GridnumX*m_GridnumZ; ++j)

         {

                ADerivative[j]=0.0;

         }

/////////////////////////////////////////////////////////////////////////////

#defineM 30*300

#define N 30*50

float *g_ADerivative[N];

float *g_C1[N];

cudaMalloc((void**)&g_ADerivative,NMsizeof(float));

/////////// Aderivative是 new 的 不连续的地址，可不可以直接cudamemcpy////////////

cudaMemcpy(g_ADerivative,ADerivative,m_receiverm_tnumm_GridnumXm_GridnumZsizeof(float),cudaMemcpyHostToDevice);

////////////////////Aderivative的转置 * Aderivative////////////////////////////////

   MatrixMutil0(g_ADerivative,g_C1,m_receiver*m_tnum,m_GridnumX*m_GridnumZ,1,BLOCK);

void MatrixMutil0(float **Mx,float **result,intMX,int MZ,floatalpha,int thread)

{

   dim3 THread(thread,thread);

   dim3 BLock(MZ/thread,MZ/thread);

   g_MatrixMutil0<<<BLock,THread>>>(Mx,result,MX,MZ,alpha);



   cudaError_t cudaStatus;

   cudaStatus = cudaDeviceSynchronize();

   if(cudaStatus != cudaSuccess) {

   fprintf(stderr, "cudaDeviceSynchronizereturned error code %d after Kernel!\n", cudaStatus);}

}

global void g_MatrixMutil0(float**Mx,float **result,intMX,int MZ,floatalpha) {

   int x =blockIdx.x * blockDim.x + threadIdx.x;

   int y =blockIdx.y * blockDim.y + threadIdx.y;

   int tx =threadIdx.x;

   int ty =threadIdx.y;

   int bx =blockIdx.x;

   int by =blockIdx.y;

   floatsub;

   inti,j,k,a,b;



   for(i=0;i<MX/30;i++)

   {



          float*Asub = GetSubMatrix(Mx,i,bx);

          //获取指向当前矩阵N的子矩阵的指针Nsub

          float*A_sub = GetSubMatrix(Mx,i,by);



          __shared__ float AT[BLOCK][BLOCK];

          __shared__ float A[BLOCK][BLOCK];



          AT[ty][tx]=*(A_sub+tx*BLOCK+ty);

          A[ty][tx] =*(Asub+tx*BLOCK+ty);



          __syncthreads();



          for(k=0;k<BLOCK;k++)

          {

                 sub+=AT[ty][k]*A[k][tx]*alpha;

          }

          __syncthreads();

   }



   result[y][x]=sub;

}

device float* GetSubMatrix( float**A,int row ,int col)

{

   return&A[row*BLOCK][col*BLOCK];

}

程序运行报错：cudaDeviceSynchronize returnederror code 30 after Kernel!

但是有结果，结果不正确。

程序是数据传输问题，还是矩阵自乘部分的问题。烦请大神ice和横扫千军指导，谢谢！！

2

30号错误是cuda error unknown,

所以您的问题：“程序是数据传输问题，还是矩阵自乘部分的问题。”
的回答是：
“您的矩阵自乘的kernel有问题，而且是非法访存或者越界”。

您可以进一步通过start cuda debugging(启用memory check的情况下），一步定位到具体出错的行。

感谢来访。

3

LZ您好：

我稍微补充一下2#横扫斑竹。

您在host端申请的ADerivative是指针数组，并多次分段申请了存储空间（您的代码可能丢失了部分字符，我猜测是如此）。

而您向device端的g_ADerivative复制的时候，是一次性复制全部逻辑所需大小的，这个是不行的。非连续的地址内容，无法直接copy。以及您的float *g_ADerivative[N];float *g_C1[N];这两个数据类型似乎也和您的用法对不上。

请您考虑这个问题。

祝您编码顺利~

4

向device端的g_ADerivative复制的时候可不可以这样写：

for( i=0; i<m_receiverm_tnum; ++i)
{ cudaMemcpy(g_ADerivative[i],ADerivative[i],m_GridnumXm_GridnumZ*sizeof(float),cudaMemcpyHostToDevice);
}

至于float *g_ADerivative[N]，是传入GPU的矩阵;float *g_C1[N]是矩阵相乘的结果。对应核函数的参数float **Mx,float **result。

5

LZ您好：

1：可以分段copy。

2：您在copy的时候，使用的存放指针的那个数组需要保存在host端，否则cudaMemcpy()这个runtime函数参数会报错。

3：您还需要定义一个存在host端的指向指针数组的指针，并cudaMalloc这个指针数组的大小，并将host端保存的指针数组复制过去，然后使用这个指向指针数组的指针作为kernel的参数。

大致如此，供您参考。

您还可以参考近期的一个讨论帖：

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=7056&extra=page%3D1

祝您编码顺利~

6

我的程序就是按照那个帖子写的。
float g_A[N];
cudaMalloc((void*)&g_A,MNsizeof(float));
for( i=0; i<M; ++i)
cudaMemcpy(g_A,A,N*sizeof(float),cudaMemcpyHostToDevice);
//A是在主机上的二维指针数组
传递参数的时候直接使用g_A。好像会报错。有错误的话ice可不可以写个例程指出来

7

LZ您好：

上述引用的讨论帖的6#就是现成的正确的例程，经 无与伦比 网友编译和运行检验通过的，请您参照。

祝您好运~