在gpu上使用结构体的问题

[chenbo888666999]

大家好，我在host上定义了一个结构体，然后将一个结构体变量拷贝到device上，但是输出结构显示拷贝没有成功，具体的代码片段如下：
＃define MAX_SUB_NUM 32
typedef struct
{
short int haplo[MAX_SUB_NUM];
double prob;
int cur;
}STRUCT;

STRUCT obj1,obj2;//定义host上的结构体变量

__device__ STRUCT dobj;//定义device上的结构体变量

void print_struct(STRUCT *b)
{
int i;
for(i=0;i<MAX_SUB_NUM;i++)
printf("%d ",b->haplo[i]);
printf("\n");
printf("%f\n",b->prob);
printf("%d\n",b->cur);
}

int main()
{
int i;
for(i=0;i<MAX_SUB_NUM;i++)//初始化host上的结构体变量obj1
obj1.haplo[i]=i;
obj1.prob=0.8;
obj1.cur=123;

size_t memsize=sizeof(obj1);
cudaMemcpy(&dobj,&obj1,memsize,cudaMemcpyHostToDevice);//将obj1拷贝到设备上

cudaMemcpy(&obj2,&dobj,memsize,cudaMemcpyDeviceToHost);//将dobj拷贝到主机上

print_struct(&obj);//打印，检验拷贝是否成功
}
但是打印的结果全是0，也就是说从host到device的拷贝没成功。我想到了结构体的内存对齐问题，尝试了显示的进行对齐，如：typedef struct __align__(16) ,但还是不行，请大家帮忙看看，怎么才能拷贝成功。

[chenbo888666999]

刚才测试了一下，二者的对齐规则确实不一样，但是怎么才能使他们的对齐规则一样呢？

[无心人_过过小日子]

1)cuda上不推荐用结构。
2)一种笨办法是自己加dummy成员，使所有成员的边界在8字节上。

[chenbo888666999]

关于编译选项-maxrregcount，我有些不解。我从程序的.cubin文件中，看到的各级设备存储器使用情况如下：
kernel1:
lmem = 32
smem = 56
reg = 25
kernel2：
lmem = 1536
smem = 68
reg = 34
我的blocksize=256，blocknum=32，即每个块有256个线程，共32个线程块，我的GPU是Geforce9800GT，每个线程块有16K个寄存器，按照计算，每个线程最多可用64个寄存器，但是在编译时，如果我不加编译选项-maxrregcount 32 限制，kernel计算时间只有0.1ms，应该是没有启动；但是加上选项-maxrregcount 32，kernel就能运行2s，请问大家这是为什呀？
（2）还有在我的程序中使用了结构体（用过__align__限制），kernel计算完后，需要将结构拷贝到主机，但是打印的结构全是0，我仔细的检查了程序，没有发现什么错误，但也不敢肯定程序没错，请问出现这种情况，会是结构体对齐问题吗？

[无心人_过过小日子]

如果你使用了double,则一个reg占用2个32位reg的资源。

[chenbo888666999]

[Quote=引用 5 楼 l7331014 的回复:]

如果你使用了double,则一个reg占用2个32位reg的资源。
[/Quote]
哦，你说得对，我的确用了double型变量。
另外，上次你说结构体里让每个成员的边界8对齐，我这样做了，但打印的结果（结果是从gpu上的一个动态分配的结构体数组拷贝到host上的memory，然后打印）还是全为0，这几天我一直都在检查程序，还是觉得程序没错，但又不知道问题出在哪儿？
下面是我结构体的定义，麻烦各位看看，有没有什么问题,谢谢。
typedef struct __align__(16)
{
short int haplo[32];
double prob;
long long cur;//本来是int cur，但为了8对齐就改为long long
}HaploList;

typedef struct __align__(16)
{
usigned int hapcur[2];
double prob;
long long cur;
}HconfList;

typedef struct __align__(16)
{
long long hapcur;
long long cur;
}TempList;

[无心人_过过小日子]

__device__ STRUCT dobj;//定义device上的结构体变量
定义成指针。用cudamalloc分配dobj.

[tengwl]

使用cudaMemcpyFromSymbol和cudaMemcpyToSymbol试试看

[chenbo888666999]

[Quote=引用 9 楼 tengwl 的回复:]

使用cudaMemcpyFromSymbol和cudaMemcpyToSymbol试试看
[/Quote]
呵呵，你说的对，复制device上的全局数组得用cudaMemcpyFromSymbol和cudaMemcpyToSymbol。
那如果定义全局二维数组：__device__ STRUCT dobj[32][64]，它也是表示32*64的数组吗（也和cpu上一样以行为主序吗）？

[无心人_过过小日子]

[Quote=引用 10 楼 chenbo888666999 的回复:]
引用 9 楼 tengwl 的回复:

使用cudaMemcpyFromSymbol和cudaMemcpyToSymbol试试看

呵呵，你说的对，复制device上的全局数组得用cudaMemcpyFromSymbol和cudaMemcpyToSymbol。
那如果定义全局二维数组：__device__ STRUCT dobj[32][64]，它也是表示32*64的数组吗（也和cpu上……
[/Quote]

gpu上全局数组好像是作为"常量"处理的。空间开辟在常量存储器中，有大小限制(64K).

[无心人_过过小日子]

[Quote=引用 8 楼 chenbo888666999 的回复:]
引用 7 楼 l7331014 的回复:

__device__ STRUCT dobj;//定义device上的结构体变量
定义成指针。用cudamalloc分配dobj.

我是这样定义的：
STRUCT *dobj;
cudaMalloc((void**)&dobj,memsize);
之后将参数传给kernel函数。
说到着，我有一个疑惑，本来我是这样定义的： _……
[/Quote]

1)gpu上的gmem地址是一个从0开始的32位数。
2)&dobj[0]是主机中地址。dobj中存放cudamalloc中返回值的话，是gpu上的地址。
记住gpu地址空间和主机地址空间是独立的。

[tengwl]

[Quote=引用 11 楼 l7331014 的回复:]
引用 10 楼 chenbo888666999 的回复:
引用 9 楼 tengwl 的回复:

使用cudaMemcpyFromSymbol和cudaMemcpyToSymbol试试看

呵呵，你说的对，复制device上的全局数组得用cudaMemcpyFromSymbol和cudaMemcpyToSymbol。
那如果定义全局二维数组：__device__ STRUCT dobj……
[/Quote]

GPU上的全局内存就是正常的__global__内存，不是常量内存，常量内存使用_constant_修饰的。

[tengwl]

[Quote=引用 10 楼 chenbo888666999 的回复:]
引用 9 楼 tengwl 的回复:

使用cudaMemcpyFromSymbol和cudaMemcpyToSymbol试试看

呵呵，你说的对，复制device上的全局数组得用cudaMemcpyFromSymbol和cudaMemcpyToSymbol。
那如果定义全局二维数组：__device__ STRUCT dobj[32][64]，它也是表示32*64的数组吗（也和cpu上……
[/Quote]

是的，表示在device上的二维数组，另外补充一点，使用__device__修饰的全局静态内存global内存和用cudamalloc申请的内存存放的位置一样，不同的是cudamalloc申请的内存可以使用cudafree释放，而前者不可以

[chenbo888666999]

[Quote=引用 7 楼 l7331014 的回复:]

__device__ STRUCT dobj;//定义device上的结构体变量
定义成指针。用cudamalloc分配dobj.
[/Quote]
我是这样定义的：
STRUCT *dobj;
cudaMalloc((void**)&dobj,memsize);
之后将参数传给kernel函数。
说到着，我有一个疑惑，本来我是这样定义的： __device__ STRUCT *dobj（全局的）;然后用cudaMalloc分配空间，按说这样就可以在kernel和device函数中直接使用dobj，而不用传递参数，但是我在用VS2005的EmuDebug模式调试时，发现如果加了__device__限制符，调试时看的分配的空间地址为OX0，但是不加__device__的话，看到的dobj的地址就不为0。我想问的是，EmuDebug模式下，所有device上的设备空间不是都映射到host空间吗？那为什么还会有这种差别，这二者还有那些差别？

（2）我发现如果定义__device__ STRUCT dobj[SIZE]（全局的）;这样在kernel和device函数中就可以直接使用该数组了，但是在拷贝会host内存时会出现问题：
cudaMemcpy(hostmem,&dobj[0],cpy_memsize,cudaMemcpyDeviceToHost);发现这样拷贝不行
但是如果把dobj定义为指针就没问题，请问&dobj[0]和dobj在device上难道表示的是不同的地址吗，还是有其他原因？

[无心人_过过小日子]

[Quote=引用 13 楼 tengwl 的回复:]
GPU上的全局内存就是正常的__global__内存，不是常量内存，常量内存使用_constant_修饰的。
[/Quote]

忘了这个前缀了....一直当常量来认识的.
看书不认真,该检讨.呵呵.

[chenbo888666999]

唉，我实在找不出程序的错误呀，大家都是怎么调试cuda程序的，我看有人说用cuda-gdb，但是我的cuda-gdb似乎有问题，用的时候会报错：
cuda-gdb: error while loading shared libraries: libtermcap.so.2: cannot open shared object file: No such file or directory
不知道大家有人用过cuda-gdb吗？或是推荐其他调试工具，谢谢。

[无心人_过过小日子]

[Quote=引用 18 楼 chenbo888666999 的回复:]
唉，我实在找不出程序的错误呀，大家都是怎么调试cuda程序的，我看有人说用cuda-gdb，但是我的cuda-gdb似乎有问题，用的时候会报错：
cuda-gdb: error while loading shared libraries: libtermcap.so.2: cannot open shared object file: No such file or directory
不知……
[/Quote]

检查一下是否有可能的死循环(浮点数的边界判断等)。

[chenbo888666999]

现在又遇到一个奇怪的问题：我的程序里使用了double型数据，如果程序在Geforce9800 GT（计算能力为1.1）上运行，大概运行了8s结束，但是结果不对。如果在C1060（计算能力为1.3）上运行程序出不来（就是程序结束不了），操作系统为linux，不知道这是问什么？我有一点疑惑：
（1）看编程手册上说，计算能力为1.1的设备不支持double型，那我的程序在Geforce9800 GT上运行为什么没有报错？还是说编译器自动将double型数据转换成了float型？
（2）为什么我的程序在计算能力低的Geforce 9800GT上能运行结束，但在计算能力高的C1060上反而结束不了？希望大家给点建议和提示，谢谢。

[tengwl]

[Quote=引用 16 楼 chenbo888666999 的回复:]
现在又遇到一个奇怪的问题：我的程序里使用了double型数据，如果程序在Geforce9800 GT（计算能力为1.1）上运行，大概运行了8s结束，但是结果不对。如果在C1060（计算能力为1.3）上运行程序出不来（就是程序结束不了），操作系统为linux，不知道这是问什么？我有一点疑惑：
（1）看编程手册上说，计算能力为1.1的设备不支持double型，那我的程序在Geforce9800 GT……
[/Quote]

1）在Geforce9800 GT上编译器自动将double型数据转换成了float型。
2）估计是你的程序本身问题，不是平台的问题

[tengwl]

支持LS的

[feehung]

都用指针吧，这样分配和释放空间意义都更清楚：

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#define MAX_SUB_NUM 32

typedef struct //__align__(4)
{
int haplo[MAX_SUB_NUM];
double prob;
int cur;
}S;

S *obj1,*obj2;
__device__ S *dobj;

void print_struct(S *b)
{

int i;
for(i=0;i<MAX_SUB_NUM;i++)
printf("%d ",b->haplo[i]);
printf("\n");
printf("%f\n",b->prob);
printf("%d\n",b->cur);
}

int main()
{
int i;
size_t size;

obj2 = (S*)malloc(sizeof(*obj1));
obj1 = (S*)malloc(sizeof(*obj1));

for(i=0;i<MAX_SUB_NUM;i++)//...host.......obj1
{
obj1->haplo[i] = i;

}


obj1->prob = 0.8;
obj1->cur = 123;

size = sizeof(*obj1);


cudaMalloc((void**)&dobj,size);

cudaMemcpy(dobj,obj1,size,cudaMemcpyHostToDevice);//.obj1......
cudaMemcpy(obj2,dobj,size,cudaMemcpyDeviceToHost);//.dobj......

print_struct(obj2);

cudaFree(dobj);
free(obj1);
free(obj2);

}