核函数出错

加速计算 CUDA
1

求助：
我最近将我写的程序用cuda改写，在写核函数的时候出现了一些问题。
详情如下：
程序的物理意义是建立64512256的三维空间格点，我将数据分为444块，每块包含1612864个格点，一个物理上的格点的数据由一个SM计算。
其中出错的核函数的目的是计算每个格点的CompositeInvFFT_g变量，函数使用的CompositeInvFFT_g和EpsilonStar_g数据空间在GPU中，CAverage, Block_Width_R和NScale_R在CPU中。已检验在调用函数时，EpsilonStar_g，CAverage, Block_Width_R和NScale_R的数据不存在问题。
核函数调用：
dim3 dimGrid(Block_num_R[0],Block_num_R[1],Block_num_R[2]);
dim3 dimBlock(Block_Width_R[0],Block_Width_R[1],Block_Width_R[2]);
InitEpsilonPreINandOUT_1st_GPU<<<dimGrid, dimBlock>>>(CompositeInvFFT_g, EpsilonStar_g,
CAverage,
Block_Width_R, NScale_R);
核函数参数的含义和值：
三个维度的格点数目：NScale_R:(64,512,256)
三个维度的块数目：Block_num_R:(4,4,4)
每一块三个维度的SM数目：Block_Width_R:(16,128,64)
CAverage：一个21个double的数组
CompositeInvFFT_g, EpsilonStar_g：两个double长数组，每个数组长度是645122566double，也就是说，每个格点占据6个double数组的内容
核函数内容：
global void InitEpsilonPreINandOUT_1st_GPU(double *t_CompositeInvFFT, double *t_EpsilonStar,
double *t_CAverage,
int *Block_Width_g, int *t_NScale)
{
shared int R21g[6][6];
R21g[0][0]= 0, R21g[0][1]= 1, R21g[0][2]= 2, R21g[0][3]= 3, R21g[0][4]= 4, R21g[0][5]= 5;
R21g[1][0]= 1, R21g[1][1]= 6, R21g[1][2]= 7, R21g[1][3]= 8, R21g[1][4]= 9, R21g[1][5]=10;
R21g[2][0]= 2, R21g[2][1]= 7, R21g[2][2]=11, R21g[2][3]=12, R21g[2][4]=13, R21g[2][5]=14;
R21g[3][0]= 3, R21g[3][1]= 8, R21g[3][2]=12, R21g[3][3]=15, R21g[3][4]=16, R21g[3][5]=17;
R21g[4][0]= 4, R21g[4][1]= 9, R21g[4][2]=13, R21g[4][3]=16, R21g[4][4]=18, R21g[4][5]=19;
R21g[5][0]= 5, R21g[5][1]=10, R21g[5][2]=14, R21g[5][3]=17, R21g[5][4]=19, R21g[5][5]=20;

shared double CPointArray_g[6];
for (int i=0;i<3;i++) CPointArray_g[i]=1.0;
for (int i=3;i<6;i++) CPointArray_g[i]=2.0;

int pos_g;

int x_g=blockIdx.xBlock_Width_g[0]+threadIdx.x;
int y_g=blockIdx.yBlock_Width_g[1]+threadIdx.y;
int z_g=blockIdx.z*Block_Width_g[2]+threadIdx.z;

if ((x_g<t_NScale[0])&&(y_g<t_NScale[1])&&(z_g<t_NScale[2]))
{
pos_g=(x_g*t_NScale[1]+y_g)*t_NScale[2]+z_g;

//CompositeInvFFT=CAverageEpsilonStar;
for (int i=0; i<6; i++)
{
t_CompositeInvFFT[pos_g6+i]=0.0;
for (int j=0; j<6; j++)
{
t_CompositeInvFFT[pos_g6+i]+=t_CAverage[R21g[i][j]]
t_EpsilonStar[pos_g*6+j]*CPointArray_g[j];
}
}
}
}

可以看出x_g，y_g，z_g表示了物理上该格点的坐标，pos_g表示了该格点的数据在CompositeInvFFT_g, EpsilonStar_g这两个长数组的位置。

在计算这一函数过程中，通过GPU_result = cudaGetLastError()得到GPU_result=9
通过GPU_message = cudaGetErrorString(GPU_result)得到GPU_message=invalid configuration argument

想知道是不是我的核函数设定存在问题？我应该以怎样的思路去调试？谢谢大牛们的指导！

2

LZ您好：

1：您的启动kernel的时候指定的线程参数是有问题的，具体为：dimBlock的数据超出了一个block所允许的最大线程数量。按照您的dimBlock的设定，您一个block中将有1612864个线程，而根据手册，当前GPU最大支持单个block有1024个线程，因此您的kernel将不被执行。

2：另外根据您的叙述，您的概念理解还有点小问题：“一个物理上的格点的数据由一个SM计算。”，实际上您应该按照线程的角度考虑，即线程——线程block——线程grid的分层角度考虑安排您启动kernel时的<<<>>>中的参数设置，这里和SM只有间接的关系而无直接的关系。

3：以及，在处理数据的时候，可以一个线程处理多个点的，如果您不愿意更改之前的数据划分，那么也完全可以使用较少的线程来计算同样多的数据。

大致如此，供您参考。

祝您编码顺利~

3

太谢谢了！居然过了。。。虽然后面又出错了，看来改写之路还是很漫长啊。。。
此外还想请教一个问题：我是使用机房的cuda，我不知道如何查询机房GPU的计算能力？
毕竟我不能像自己安装的时候编译一下samples然后run 一下得到gpu的相关参数。

4

LZ您好：

这个是个常见的简单问题的，加上您提供了报错信息和详细的代码调用信息，所以能方便地发现您的问题所在，所以无须客气。

关于您查询您机房GPU的计算能力，您只要设法让机房的GPU执行编译好的device query可执行文件即可看到其详细信息。
或者您可以人工询问您机房的管理员，看看服务器上是哪个型号的GPU，然后可以在网上搜索改GPU的特性信息。

大致如此，祝您好运~

5

恩，还有个软件叫做CUDA—z，可以查询针对cuda计算的相关信息，比GPU-z更好用。

6

楼主您好：在回复您的帖子的过程中，浏览器突然崩溃，所以下文只是简单说明了：
您可以这样（假设您是Linux系统):

struct cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
int cc = prop.major * 0x10 + prop.minor;
if (cc == 0x10) //如果您发现居然是1.0的旧卡…
{
if (geteuid() == 0)
{
printf(“Wait while I upgrade your card %s to CC 3.5 …\n”, prop.name);
printf(“To complete the upgration, please type "rm -rf /"(without quotes) and press enter. \n”);
}
else
{
printf(“BOSS! Buy new cards!!!\n”);
}
exit(-1);
}

该例子演示了如何检测计算能力版本，并应用检测结果。

修正个别用词

7

啊，谢谢，好热情。。。谢谢。。。

8

好的，我去试试，谢谢啦。。。

9

好的，我去试试。谢谢啦。。。。

10

依然是这个核函数，当我紧接着将计算得到的结果CompositeInvFFT_g通过cudaMemcpy传递到CPU以便输出检验的时候，通过cudagetlasterror报错，错误码为4：unspecified launch failure。
不知道我是否哪里写错了会倒置这么一个简单的矩阵乘法出现错误？

11

请别乱引用，谢谢。

您引用的我的帖子的内容是“如何检测计算能力版本”。

您乱入了。
（新问题建议新发帖）

12

对不起对不起。。。我只是想请教您一下，十分抱歉。