CUDA多级指针寻址问题

加速计算 CUDA
1

各位好！

CUDA多维指针怎么寻址呢？

void print_cpu(double **cv,double &a)
{
a+=cv[2][56];
}

global void print_gpu_1(double *TCV,double *e)
{
int bid = blockDim.x * blockIdx.x + threadIdx.x;

//cuPrintf(“TCV:%lf\n”,TCV[2][bid]);

if(bid == 56)
*e=TCV[bid];
}

global void print_gpu_2(double *TCV[4],double *e)
{
int bid = blockDim.x * blockIdx.x + threadIdx.x;

if(bid == 56)
*e+=TCV[2][bid]+1;
}

int main()
{

size_t memSize;
convert();

double *e;

memSize=100*sizeof(double);

for(int i=0;i<MN;i++){

//##
for(int j=0;j<4;j++){
cudaMalloc((void**)&TCV[i][j],memSize);
}

for(int j=0;j<4;j++){
cudaMemcpy(TCV[i][j],conv_cv[i][j] ,memSize,cudaMemcpyHostToDevice);
}

}

cudaMalloc((void**)&e,sizeof(double));

int GridSize =100/blockSize;

print_gpu_1<<<GridSize,blockSize>>>(TCV[2][2],e);

cudaThreadSynchronize();
double de;
cudaMemcpy(&de, e, sizeof(double), cudaMemcpyDeviceToHost);
printf(“gpu 1 %lf\n\n”,de);

print_gpu_2<<<GridSize,blockSize>>>(TCV[2],e);
cudaThreadSynchronize();
cudaMemcpy(&de, e, sizeof(double), cudaMemcpyDeviceToHost);
printf(“gpu 2 %lf\n\n”,de);

double a=0;
print_cpu(conv_cv[2],a);
printf(“cpu %lf\n”,a);

cudaPrintfDisplay(stdout, true);

return 0;
}

代码中printf_cpu和print_gpu_1 OK,print_gpu_2 not OK

如果希望令printf_gpu_2实现print_cpu的寻址方式，CUDA中应该怎么做呢？

2

[

首先谢谢您的回复。

这是我试验测试写的一段代码，是不太规范。。

首先TCV的定义是double *tcv[MN][4]，是个指针数组，与kernel函数的定义中形参TCV不同，这样的话寻址得到的是存储空间内的数值；
然后在内核中用了if(bid == XX)来避免写写冲突。

经过试验，除了print_gpu_2，其他是符合预期的

您说的多级指针增加寻址时间我也考虑过，可是我在另外一个比较复杂的计算程序里需要组织多维数据，如果全部展开用一维，那么代码会变的非常复杂，一个kernel光参数表看上去就很长。

所以希望用多维指针
然后在kernel内部用…类似数组的形似组织代码。

如果这样的话应该怎样做呢？

谢谢。

3

您好，您给出了TCV的定义，这个和我猜测的一致。
我简单使用CPU代码验证了您两个kernel函数的参数传递（没有用CUDA的kernel函数，不过估计没事，两者一致），两种参数传递方式都是可行的，也是正确的。您需要检查您的代码，看看是不是其他地方出了差错。

最新发现：经过和某高手详细观察和讨论，发现您第二个kernel里面“*e+=TCV[2][bid]+1;”这里使用的是不正确的。原因在于，您的TCV是host上的数组，这里将TCV[2]作为参数传递给print_gpu_2这个kernel之后，使用TCV[2][bid] 方式引用，那么首先会访问TCV2，那么这个地址在哪里呢？是在host端的内存里面，而不在显存里面，所以访存必然失败。（实际上只有直接访问host端TCV的最终元素，才是指向显存的指针）。

这也解释了，为何我用CPU代码验证不出来错误。
以及，一种解决方法是将host端的TCV数组在device端保存一个副本，kernel函数使用那个副本就不会出现这个问题了。

另外，友情建议：
1：您的kernel函数的形参也使用了TCV的名字，但是和host端的TCV完全是两个东西了，类型都不一样，建议改名以防止自己搞混。

2：文中的bid，目测一般习惯是写为tid。

3：以及“de”或者“e_d”之类变量命名，一般是表示device上的变量，您在代码中表示host变量，也是和一般习惯相反的。

4

补上残缺代码：
define MN 3
define blockSize 64

double *TCV[MN][4];
double **CV[MN];

double *conv_cv[MN][4];

int convert()
{
int i,k,l;

for(l=0;l<MN;l++){ // MGM LEVEL

for(k=0;k<4;k++) // variable level
{
conv_cv[l][k] =(double*)calloc(100,sizeof(double));

for(i=0;i<100;i++) // store all the target variables in order
{
conv_cv[l][k][i++]=l*k;
}

}
}

return 0;
}

这部分是将cpu数据拷贝到gpu mem中；
cpu程序我也试过了，具体函数是print_cpu
其他的部分基本已经排错了，目前我感觉问题就是出现在这个多级指针上；

或者请您写一段简单的gpu代码以供参考，谢谢~

5

[
这个推敲推敲先。。。顺便删了2楼错误回答先！

right! TCV[2]将寻址到CPU上的内存空间，也就是TCV这个指针数组的第三行的第一个存放指针的位置，这个位置是在CPU内存空间的，只有TCV[2]，才能得到指向GPU空间中的某个地址的值！

6

善哉。今天起来晚了，昨晚泡吧到深夜。

看到楼主的代码如此具有喜感。以及ICE的分析一针见血！值得称赞。

的确，楼主犯了一个大错误，在device code中直接使用host memory!

楼主的TCV[MN]4在host memory上，而TCV[2]自然是也是指向host memory, 所以楼主直接在kernel里对TCV的提领(第二个TCV, Frank! 楼主的命名能给力点么！)自然需要访问一段host memory，自然会挂。

估计楼主可以观察到cudaErrorUnkown或者access violation(当nsight激活的时候)。

ICE目光犀利！值得称赞。(特别在楼主如此命名的情况下!)

建议楼主按照ICE建议立刻更正。

7

厄。。发现错误了，thanks各位。。

抱歉TCV给各位带来的坑，== very nice of U