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

加速计算 CUDA
1

程序中需要用到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));

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

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

2

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写法发表意见。

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

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

3

[

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

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

4

[
既然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]

5

[

LZ您好：

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

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

请您自重。

6

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

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

7

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

8

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