CUDA中二维结构体怎么分配空间?定义函数和传递参数

程序中需要用到500组数据,500组数据对应一个结构体A,每组数据分别经过5个函数模型后得到了5组计算结果,结果是存放在结构体B和C里面的,在CPU中B和C只需要声明为5个的长度就行了,而在GPU中,需要要把所有的结果500 x 5传递回主机端,然后都打印出来。
这里对B 和 C结构体我采用的是二维结构体。然后分配空间的时候是这样的:

struct vector **de_pos = NULL;
struct vector **de_vel = NULL;

CUDA_CALL(cudaMalloc((void **)&de_pos, sizeof(struct vector *) * 500));
for(q = 0; q < 500; q++)
{
CUDA_CALL(cudaMalloc((void **)&de_pos[q], sizeof(struct vector) * 5));
}

CUDA_CALL(cudaMalloc((void **)&de_vel, sizeof(struct vector *) * 500));
for(q = 0; q < 500; q++)
{
CUDA_CALL(cudaMalloc((void **)&de_vel[q], sizeof(struct vector) * 5));
}

kernel函数是这样定义的以及使用结构里面的成员
global void mykernel(int num, int model, struct tle_ascii *de_sat_data, struct vector **pos, struct vector **vel)
{

double delta = 360.0;
double tsince = 0.0;

struct DD satdata;

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

while(tid < num)
{
tsince = (double)tid * delta;
Convert_function1(de_sat_data[tid], &satdata);
// if(tid == 0)
// printf(“satdata[%d] = %ld \n”, tid, satdata.xno);
for (x = 0; x < 5; x++)
{
tsince = x * delta;

if(model == 0)
{
s0(tsince, &pos[tid], &vel[tid], &satdata);
}
if(model == 1)
{
s1(tsince, &pos[tid], &vel[tid], &satdata);
}
if(model == 2)
{
s2(tsince, &pos[tid], &vel[tid], &satdata);
}
if(model == 3)
{
s3(tsince, &pos[tid], &vel[tid], &satdata);
}
if(model == 4)
{
s4(tsince, &pos[tid], &vel[tid], &satdata);
}
Convert_function2(&pos[tid], &vel[tid]);
pos[tid].v[3] = tsince;
}

tid += blockDim.x * gridDim.x;

}

}

然后调用kernel:
mykernel<<<256, 1024>>>(num, model, de_sat_data, de_pos, de_vel);//我就直接把节结构体指针这样传过去了

然后往主机端传值的时候是这样:

CUDA_CALL(cudaMemcpy(pos, de_pos, sizeof(struct vector *) * num, cudaMemcpyDeviceToHost));

CUDA_CALL(cudaMemcpy(vel, de_vel, sizeof(struct vector *) * num, cudaMemcpyDeviceToHost));

然后是打印部分。
这样哪里有问题?

我在编译的时候,没有报错,但是在执行的时候一直报段错误。段错误,段错误,段错误。。。。。。。
请问斑竹哪里出问题了?

LZ您好:

好久不见,欢迎您重返cudazone论坛。

由于您在1#的文字说明几乎没有说清楚任何所需要的内容和细节,以及凭空提出了“二维结构体”这个概念,造成理解障碍,您的文字说明部分已经被自动跳过。

您的代码这里有一个严重的问题,“CUDA_CALL(cudaMalloc((void **)&de_pos, sizeof(struct vector *) * 500));”这里,de_pos指向的是显存中分配的空间,而下一句您试图在host端将这些空间拿来使用“CUDA_CALL(cudaMalloc((void **)&de_pos[q], sizeof(struct vector) * 5));”,这是不行的,cudaMalloc()的第一个参数要求是保存在host端的。

您的两个循环都有这样的问题,您可以简单修正为struct vector *de_pos[500];如此,de_pos是一个保存在host端的指针数组,后面的循环将不会出错。如果有需要,您可以再将这个指针数组复制到device端使用。

以及因为缺少必要的信息,将先不对您的kernel写法发表意见。

本着先实现,再优化的原则,上述内容不涉及讨论您的实现效率是否够高,以及不涉及是否可以/如何进行优化。

请您先修改和调试您的代码。
祝您调试顺利~

[

本区禁止自顶这种灌水行为,特别是频繁自顶!

特此警告一次,望LZ自重。
如再有发生,将被转移至水区,并作下沉处理。

[
既然LZ玩赖编辑帖子,不妨我也玩玩。(本帖原本含有对LZ最新问题的详细解答,共计近千字,目前已被编辑掉,将择机重新放出。)


LZ您好:

1:根据您4#提供的最新反馈,您已经按照2#的建议进行了修改,并且最前面的cudaMalloc()环节得以成功进行。

根据2#中的建议,您应该是写为“struct vector *de_pos[500];”,这样才能匹配cudaMalloc()的参数类型。但是此时,de_pos是一个存储在host端的指针数组,数组名是指向host端地址的,这样,您的写法:“然后往主机端传值的时候是这样:

CUDA_CALL(cudaMemcpy(pos, de_pos, sizeof(struct vector *) * num, cudaMemcpyDeviceToHost));”
de_pos是指向host端的,而不是指向device端的,所以会报参数错误。

2:以及,即便您修正了1:中的参数错误,(比如在device端再申请一个指针数组,在host端用一个指针指向该数组,并将host端的指针数组的内容复制给device端的该指针数组等),您的这个复制只是将device端的一个指针数组的内容——也就是一堆指向device端地址的指针——复制回host端。

这有什么用呢?a)您复制回host端的这一堆指针又不能在host端直接用,这些地址都是device端的地址;b)如果是之后的kernel需要使用,这个指针数组本身不就在device端么?不用全部复制一遍;c)如果是host端的cuda 的API函数需要使用,根据前面2#的建议,这在host端本身就有一个副本的;d)您只把指针复制回来了,真正的数据都没回来,也无法依靠这一堆指针在host端访问的。

事实上,您应该用这个指针数组的元素(这些元素的指向是真正存放的数据)作为cudaMemcpy()的参数,进行copy,而且按照您的写法需要多次copy才行。(以及,您这样的做法是效率低下的做法,但是现在不讨论优化。)

3:即便这样,根据您4#给出的host端的代码,也一样要出问题的。
您的代码为:

pos = (struct vector *)malloc(sizeof(struct vector) * num);
for(q = 0; q < num; q++)
{
pos[q] = (struct vector *)malloc(sizeof(struct vector) * 5);

}
if(NULL == pos)
{
printf(“struct vector malloc ERROR!\n”);
exit(-1);
}
memset(pos, 0, sizeof(struct vector) * num);

申请了一个指针数组,并通过循环给每个元素申请了空间,使得每个元素都有确切的指向。(以及这里检查pos malloc是否成功的代码位置不对,没检查之前就进行使用了,不过先不讨论这个。)

关键问题在于,您刚刚通过循环初始化好的指针数组,转手就被您memset()成全0了!!
内存就这样泄露了,所有元素就这样浮云了,即便您按照2:中最后的建议循环进行cudaMemcpy(),您host端用于接收数据的都是些无效的指针,一样会挂掉了。

希望您能明白每一步的问题所在,而不是依葫芦画瓢地写着各个语句。

祝您debug顺利~


以及,由于LZ删除了原来的帖子,为了方便网友看帖,特将原4#的截图附上:
[attach]3446[/attach]

[

LZ您好:

尊重是相互的,当斑竹为了您的帖子绞尽脑汁进行人肉DEBUG的时候,请给予基本的耐心和尊重。
有时候一个帖子查看,分析,讨论,总结,写出建议,需要一两个小时的时间,需要动辄写上上千字,查阅多次相关资料,以及上机测试等等。

如果您这点时间都不愿意等,明明就是首页靠前的帖子,刚刚发表一个小时左右,就要自顶,而且自顶以后还自己编辑帖子不认账,那么请您自行解决您的问题,恕我不能回答您。

请您自重。

本帖曾被转移到水区,并做下沉处理。


LZ删除了自己自顶的回帖,并保留了向版主质问的点评,这是极其不厚道的做法。

抱歉了,感谢艾斯版主的警告,以后不会再犯了。问题已经解决,还是很感谢你之前对我的指点。:slight_smile:

既然问题已经解决,将不再就此问题进行讨论。