CUDA 结构体中包含指针的数据传输问题

最近在做一个CUDA并行计算的设计,遇到的一些问题,求版主大人能给看一下。。。

代码中存在这样一个结构体:

struct Layer 
{
   int Cell;
   double **Weight; 
   double *Threshold;
};
struct Net
{
   Layer *NetLayer;
}ANNet;

这里的ANNet是一个全局变量,在用到的时候已经分配了内存进行了初始化,使用正常,

我在调用Kernel函数时要用到Net结构体,所以将数据传输到Device端。

数据传输的代码如下:

Net *dev_net = NULL;
cudaError_t result2 = cudaMalloc((void**)&dev_net,Size*sizeof(Net));
if (result2 != cudaSuccess)
   return false;
Net nt[Size] = {ANNet};

cudaError_t result5 = cudaMemcpy(dev_net,nt,Size*sizeof(Net),cudaMemcpyHostToDevice);
if (result5 != cudaSuccess)
   return false;

这里涉及到的大小都是Size。

到这里貌似都是正常的,但是在调试Kernel函数时,遇到了使用dev_net时,出现如图这样的错误,说是读取内存失败,

[attach]3155[/attach]

看过这个帖子了http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=1086

但是没有看懂,初步认为是我的结构体里面又包含有指针导致的,但是不知道该怎么去更改,

这难道是那种常见的分配了结构体或者对象数组的空间,而没有分配数组元素的空间的错误???

这种问题使用CUDA该怎么改呢?着急呀。。。。。。

难道是要在分配dev_net内存的同时分配它的下一级dev_net[i].NetLayer的内存呢?????
是不是这样呢?

 cudaError_t result2 = cudaMalloc((void**)&dev_net,Size*sizeof(Net));
 if (result2 != cudaSuccess)
   return false;

	for (int i=0;i<ParticleSize;i++)
	{
   cudaError_t result8 = cudaMalloc((void**)&dev_net[i].NetLayer,Layer_Num*sizeof(Layer));
   if (result8 != cudaSuccess)
   return false;
	}
cudaError_t result5 = cudaMemcpy(dev_net,nt,ParticleSize*sizeof(Net),cudaMemcpyHostToDevice);
	if (result5 != cudaSuccess)
   return false;
	//是不是要再接着把数组的分量元素数据传输过去

	for (int i=0;i<ParticleSize;i++)
	{
   cudaError_t result9 =     cudaMemcpy(dev_net[i].NetLayer,nt[i].NetLayer,Layer_Num*sizeof(Layer),cudaMemcpyHostToDevice);
		if (result9 != cudaSuccess)
   return false;
	}

可是在调试的时候,运行到这段代码的时候就会出现异常。

[attach]3156[/attach]

LZ您好,我来大致解读一下您的代码:

1:1#第一段代码大致含义为:您定义了一个名为Layer的结构体类型,该结构体有3个元素,分别为 int类型变量,double类型的二重指针,double类型的指针;您又定义了一个名为Net的结构体类型,含有一个元素,为Layer类型的指针,并声明了该类型的一个实例 ANNet。

2:根据您1#的说法,ANNet是全局变量,已经分配内存做了初始化。我认为,您申请了一个Layer类型的结构体的空间,并使用ANNet的元素(Layer指针类型)NetLayer指向该结构体。

3:您提到,您的kernel要使用Net类型的结构体,您采用了cudaMemcpy的形式。您的第二段代码定义了一个Net类型的结构体指针 dev_net,并使用cudaMalloc在global memory中申请了size个Net结构体的空间,并使用dev_net指向该空间的首地址。dev_net指针自身存储于host端,但是其指向地址空间是global memory的空间。

4:您的第三段代码表示,您在host端定义了一个Net类型的结构体数组nt,有Size个元素,并使用Net类型结构体的实例ANNet对该数组的第一个变量进行了初始化赋值。之后您使用了cudaMemcpy将nt的内容复制给dev_net。(共Size个Net类型结构体的空间)

如果上述分析没问题的话,您实际上是把host端的ANNet的内容复制给了device端的dev_net所指向的结构体数组的第一个元素。注意到这里只是ANNet的值的原样复制,ANNet的元素——指针NetLayer指向的应该是host端的内存地址,这个地址device端是无法使用的。

您可以在host端 先定义一个Layer类型的结构体Layer1,然后用cudaMalloc等函数给Layer1的两个指针指向合适的global memory的空间。
然后在host端定义一个Layer类型的指针Layer_ptr,使用cudaMalloc申请一个global memory上的Layer类型的结构体,并使用Layer_ptr指向该结构体。
使用cudaMemcpy将Layer1的内容复制给Layer_ptr指向的那个device端的结构体。此时,Layer_ptr指向一个device端的Layer类型的结构体,该结构体元素中的指针亦指向device端的global memory。
此时,在host端定义一个Net类型的结构体变量ANNet1,将Layer_ptr的值赋值给ANNet的元素(Layer类型的指针)。此时ANNet的元素是指向device端显存空间的,并且其指向的Layer类型的结构体的元素继续指向device端显存空间的合适位置。

此时您可以将ANNet1直接作为kernel的参数,这养ANNet1的内容将在kernel启动的时候,自动复制给kernel(作为参数压入),您就可以使用了。

或者您也可以,声明一个Net类型的结构体指针(指针存放于host端),cudaMalloc一个Net类型结构体的空间(device端)并使用这个指针指向该cudaMalloc的空间,将此指针作为参数传给kernel。

如果您需要多个Net类型的结构体,可以使用结构体数组,但是需要保证每个结构体中元素指向的地址空间及各级指向的地址空间都是合适的device端的地址空间,此时使用cudaMalloc申请一个device端的结构体数组,并使用cudaMemcpy将host端的该结构体数组复制到device端,只使用一个指向该数组的指针作为kernel的参数。

大致用法如此,供您参考。

祝您编码顺利~

LZ您好,cudaMalloc里面用的指针需要是保存在host端的指针,dev_net这个指针自身是在host端保存的,但是其指向的空间是device端,所以dev_net[ i ].Netlayer是一个保存在device端的指针。

您可以用host端的指针保存申请到的空间的地址,然后cudaMemcpy将地址值复制过去。

大致如上,祝您好运~


调整了一下行文格式,避免被论坛转为转义字符。

使用cudaMemcpy将Layer1的内容复制给Layer_ptr指向的那个device端的结构体。能不能具体写一下这句话的代码啊?是不是
cudaMemcpy(Layer1.weight,Layer_ptr.weight,sizeof(double),cudaMemcpyDeviceToDevice);
cudaMemcpy(Layer1.Threshold,Layer_ptr.Threshold,sizeof(double),cudaMemcpyDeviceToDevice);
谢谢!

ICE在3#的帖子都回复你1500多字了!请认真重新再次反复阅读吧。

一千五百字啊!他对你够意思了!!

[attach]3160[/attach]

非常感谢ice版主的详细全面的讲解,谢谢,您的说明帮了我的大忙,最近一直忙着毕

业论文了,表示感谢有点晚了,还请千军版主见谅。

对于这个问题,ice版主的讲解非常详细清晰,希望也能帮助到其他人,另外我要说的是,

这样的设计没有问题,但是牵扯到太多的一维指针和二维指针,还是建议不要用的好,

一是容易出错,二是太多的数据链接会导致用在读取数据上的时间太多而降低性能。

改成一维数组是最好的了,只要注意一下下标索引的偏移就好了。

再次感谢ice版主。

LZ您好,恭喜您解决了这个问题,我和千军斑竹都为您感到高兴。

以及您最后总结的关于此用法的经验和感想也是很有价值的,感谢您回馈论坛。

最后祝您顺利搞定各项毕业事宜,早日为国家贡献一份力量~

意思是这样,但是你不能这样写,一般是先定义之后开辟空间,再将它赋值给它的上一层下面的指针。
我这里由于设计问题,需要的是开辟ParticleSize 个Net型的GPU空间变量,即指针是我要传递给Kernel函数的参数。

        Net *dev_net = NULL;        Net *my_net = NULL;
   result = cudaMalloc((void**)&dev_net,ParticleSize*sizeof(Net));        if (result != cudaSuccess)                return false;        my_net = (Net *)malloc(ParticleSize*sizeof(Net));

但是我的设计中Net结构是这样的:

struct Layer {        int Cell;        float **Weight;         float *Threshold;  };struct Net {        Layer *NetLayer; }ANNet;

下面的指针使用太多了。所以在使用my_net来初始化dev_net的时候首先要保证my_net下面的指针都是指向GPU的,不然的话你的dev_net下面的指针仍然是指向CPU的,所以在Kernel函数 中就会出现dev_net内存无法读取的错误。

我最终的设计是换成了一位数组,使用下标偏移来控制。。

        for (int i=0;i<ParticleSize;i++)                my_net[i].NetLayer = GenerateCUDALayer();
   result = cudaMemcpy(dev_net,my_net,ParticleSize*sizeof(Net),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;

其中设计为:(供参考)

//用来为device端指针分配内存,返回一个存放在host端的指针,它指向device端的global memoryLayer* GenerateCUDALayer()
{        //这里稍有问题        //思路:直接只分配内存,不涉及内容数据复制了,直接分配出四个float类型device变量        //在分配出一个带有三个层的层结构体指针,将值直接复制给他,        //这样就能得到所有的指针均为d端的一个变量。
   cudaError_t result;
   Layer *l = NULL;        Layer *dev_l = NULL;
   //处理下一级的指针
/*************************                 这里是分界线      
*****************************************/        //输入层        float **dev_wt0 = NULL;//[HN][IN]        float *dev_wt0_0 = NULL;        float *wt0_0 = NULL;     //用于初始化*dev_wt0_0的赋值        float **wt0=NULL;
   float dev_th0[HN];
   result = cudaMalloc((void**)&dev_wt0,HN*sizeof(float*));        if (result != cudaSuccess)                return false;        result = cudaMalloc((void**)&dev_wt0_0,IN*sizeof(float));        if (result != cudaSuccess)                return false;                wt0_0 = (float *)malloc(IN*sizeof(float));        for (int i=0;i<IN;i++)                wt0_0[i] = (float)i;   //用于后面测试的数据                result = cudaMemcpy(dev_wt0_0,wt0_0,IN*sizeof(float),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   wt0 = (float **)malloc(HN*sizeof(float*));        for (int i=0;i<HN;i++)                wt0[i] = dev_wt0_0 + i*IN;
   result = cudaMemcpy((float *)dev_wt0,(float *)wt0,HN*sizeof(float 
*),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   result = cudaMalloc((void**)&dev_th0,HN*sizeof(float));        if (result != cudaSuccess)                return false;
/*************************                    这里是分界线      
*****************************************/        //隐含层        float **dev_wt1 = NULL;//[HN][IN]        float *dev_wt1_0 = NULL;        float *wt1_0 = NULL;        float **wt1=NULL;
   float dev_th1[HN];
   result = cudaMalloc((void**)&dev_wt1,HN*sizeof(float*));        if (result != cudaSuccess)                return false;        result = cudaMalloc((void**)&dev_wt1_0,HN*sizeof(float));        if (result != cudaSuccess)                return false;
   wt1_0 = (float *)malloc(HN*sizeof(float));        for (int i=0;i<HN;i++)                wt1_0[i] = (float)i;   //用于后面测试的数据
   result = cudaMemcpy(dev_wt1_0,wt1_0,HN*sizeof(float),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   wt1 = (float **)malloc(HN*sizeof(float*));        for (int i=0;i<HN;i++)                wt1[i] = dev_wt1_0 + i*IN;
   result = cudaMemcpy((float *)dev_wt1,(float *)wt1,HN*sizeof(float 
*),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   result = cudaMalloc((void**)&dev_th1,HN*sizeof(float));        if (result != cudaSuccess)                return false;
/*************************         这里是分界线      
*****************************************/        //输出层        float **dev_wt2 = NULL;//[ON][HN]        float *dev_wt2_0 = NULL;        float *wt2_0 = NULL;        float **wt2=NULL;
   float dev_th2[ON];
   result = cudaMalloc((void**)&dev_wt2,ON*sizeof(float*));        if (result != cudaSuccess)                return false;        result = cudaMalloc((void**)&dev_wt2_0,ON*sizeof(float));        if (result != cudaSuccess)                return false;

   wt2_0 = (float *)malloc(ON*sizeof(float));        for (int i=0;i<ON;i++)                wt2_0[i] = (float)i;   //用于后面测试的数据
   result = cudaMemcpy(dev_wt2_0,wt2_0,ON*sizeof(float),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;

   wt2 = (float **)malloc(ON*sizeof(float*));        for (int i=0;i<ON;i++)                wt2[i] = dev_wt2_0 + i*HN;
   result = cudaMemcpy((float *)dev_wt2,(float *)wt2,ON*sizeof(float 
*),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   result = cudaMalloc((void**)&dev_th2,ON*sizeof(float));        if (result != cudaSuccess)                return false;
/*************************           这里是分界线      
*****************************************/        l=(Layer *)malloc(Layer_Num*sizeof(Layer));
   l[0].Cell = IN;                l[1].Cell = HN;        l[1].Threshold = dev_th1;        l[1].Weight = dev_wt1;
   //赋值输出层的指针        l[2].Cell = ON;        l[2].Threshold = dev_th2;        l[2].Weight = dev_wt2;
   result = cudaMalloc((void**)&dev_l,Layer_Num*sizeof(Layer));        if (result != cudaSuccess)                return false;
   result = cudaMemcpy(dev_l,l,Layer_Num*sizeof(Layer),cudaMemcpyHostToDevice);        if (result != cudaSuccess)                return false;
   return dev_l;}

木有办法呀,代码有点乱,我也调不好呀,凑合理解一下就好了

善哉!!

LZ费心了!

感谢您详细地提供了您的代码,并详细举例讲述了此问题,同时分享了您的经验。
论坛的发展离不开您这样的热心网友!

在此奉上我和横扫千军斑竹的一份感激。

欢迎您常来论坛~祝您一切顺利!

ice版主您好,另外我还有一个问题,就是,不知到为什么我的Kernel函数中的参数传递过去的数据

其中包括数组的,其中数组怎么都只有64个数据?我的其中一个Weight数组要有1240个数据呢?

我截个图。。

嗯嗯,我吃饭回来看看,您先发即可。

版主费心了,万分感谢呀。。。

首先是指针*dev_net,我是用Net my_net[ParticleSize]来进行数据赋值的,即把my_net的数据复制给dev_net,

其中my_net的数据如图:

(其中的Weight数组为1240大小的,肯定大于64,而在调试CUDA时,这里传递过来的数据却只有64个)

[attach]3164[/attach][attach]3165[/attach]

另外Particles *dev_p也是如此:

[attach]3167[/attach][attach]3166[/attach]

难道是在global memory上面申请空间时,对数组的大小有限制?在学习的时候也没有遇到过呢??

并不是只传输过来64个数据,只是这样查看只会显示出64个数据而已,实际上完整赋值了,

我一步一步地进行了一遍调试,发现正好是在循环结束时控制变量变成了1240,

只是用16进制表示而已,是0x000004d8。

另外不知道为什么VS在调试时会这样显示呢?极其容易让人误解,以为数据传输不完整呢。

那又怎么样才能完整查看呢?

谢谢LZ、横扫千军和ICE版主,非常感激各位对小菜鸟的帮忙。
祝各位一切顺利。

洋哥爱猫猫 发表于 2013-5-26 12:19
难道是要在分配dev_net内存的同时分配它的下一级dev_net.NetLayer的内存呢?????
是不是这 …

LZ您好,cudaMalloc里面用的指针需要是保存在host端的指针,dev_net这个指针自身是在host端保存的,但是其指向的空间是device端,所以dev_net[ i ].Netlayer是一个保存在device端的指针。

您可以用host端的指针保存申请到的空间的地址,然后cudaMemcpy将地址值复制过去。

想问LZ,在这里您是按照ICE斑竹提供的方法做的吗?

LZ您好,global memory申请空间确实有限制的,但是应该远大于您当前申请的数量。

以及,您说的这个问题,可以在nsight的设置里面修改:
[attach]3168[/attach]

修改Max array expansion elements一项即可。

以及如您发现的那样,这个不影响实际执行的。

祝您好运~

大致是按照ice版主的方法实现的,如果Layer下面没再有指针了,就可以直接按照ice版主的意思来设置了,
就是先申请一个global memory空间上的*dev_Layer,若还要求有数据则是再有一个CPU上定义开辟空间的my_Layer,

使用cudaMemcpy,将值进行复制,

另外在从Net这个级别进行操作,即

CPU定义开辟空间的一个变量my_Net,将上面的dev_Layer直接赋值给my_Net的NetLayer指针,

在开辟global memory空间上的*dev_Net,再使用cudaMemcpy进行值的复制。

这样的dev_Net就可以直接作为参数传递了,其下一层的指针NetLayer也是指向GPU空间的指针了。

问题在于我的设计Layer下面还有

double **Weight;

double *Threshold;

这就要求依次进行如上操作了,我上面的代码就是如此操作的,看起来是非常繁琐复杂的,

所以我果断放弃了,我从整体设计上进行了修改,不用指针了改用数组,使用偏移来控制。

嗯,这样呀,好好学习了一把,谢ice版主咯。。

原来这样,非常感谢LZ您的回复,谢谢啦。。。