全局内存数据对齐问题~

1. typedef unsigned int  uint32;
   
2. typedef unsigned char byte;
   
3. 
   
4. __device__ byte data[1024];
   
5. 
   
6. __device__ inline void Copy128(byte *dest,const byte *src)
   
7. {
   
8. 
   
9.   ((uint32*)dest)[0]=((uint32*)src)[0];
   
10.   ((uint32*)dest)[1]=((uint32*)src)[1];
    
11.   ((uint32*)dest)[2]=((uint32*)src)[2];
    
12.   ((uint32*)dest)[3]=((uint32*)src)[3];
    
13. 
    
14. }
    
15. 
    
16. __global__ kernel(){
    
17.     byte v[4][4];
    
18.    copy128((byte *)v,data);
    
19. }

复制代码

代码大概是这样，调试的时候发现 v中的内容整体后移了两个字节，是什么原因呢？
data 前16字节：                                                             v
0x45                                                                             0x89
0x67                                                                             0x45
0x89                                                                             0x7a
0x45                                                                             0xf1
0x7a                                                                                .
0xf1                                                                                 .
   .
   .                                                                                    .
   .
   .                                                                                   0x56
                                                                                       0x77
0x56                                                                               0xb2
0x77                                                                               0x3c

楼主你好：

请不要定义local memory中的byte[]数组, 然后将首地址cast成int类型的指针用。

当你是使用byte或者char之类的来定义[]数组的时候，你已经明确的告知了编译器你不需要任何对齐方面的保证。编译器有充分的权利将你的byte[]数组的首地址安排到任意位置，而无须考虑对其到4B边界。此时将可能导致不可预测的后果（例如你这种结果，作为4B写入的时候，指令将忽略地址的最低2位，并可能在其它方面导致错误）。

任何未定义的行为都是完全无保证的！

感谢您的来访。

横扫千军 发表于 2013-10-10 22:05
楼主你好：

请不要定义local memory中的byte[]数组, 然后将首地址cast成int类型的指针用。

如果你真的需要使用char[]这种，然后作为int写入，您可以自定义多3B字节。

例如这样：
char v[16 + 3];
char *p = v;
p = (p + 3) & (~3); //如果编译器抱怨此处，请改用&0xffffffffc (32位下), 或&0xfffffffffffffffc(64位xia)
这将会执行一次手工对齐，会导致1字节，2字节，或者最多3字节的空间浪费（当然也可能无浪费，如果原地址已经对其到4B边界了）

请注意范例代码没有经过编译器检验，仅供示意用途。

感谢您的来访。

横扫千军 发表于 2013-10-10 22:11
如果你真的需要使用char[]这种，然后作为int写入，您可以自定义多3B字节。

例如这样：

感谢版主的耐心解答。我还想请问个问题，代码大概如下：调试的时候发现，只拷贝了H_S中的前64个字节到S。我想实现在CPU中计算相关数据，然后传送到Gpu中作为只读的全局变量，应该怎么做比较好呢？

1. static H_S[256];
   
2. __device__ S[256];
   
3. 
   
4. main(){
   
5.      for(int i = 0;i < 256;i++)
   
6.            H_S[i] = '\0'+i;
   
7. 
   
8.      //然后调用
   
9. cudaStatus = cudaMemcpyToSymbol(S,H_S,256,0,cudaMemcpyHostToDevice);
   
10. }  

复制代码

hujun21 发表于 2013-10-11 19:21
感谢版主的耐心解答。我还想请问个问题，代码大概如下：调试的时候发现，只拷贝了H_S中的前64个字节到S。 ...

楼主您好，

您的代码试图使用“默认int"类型，这在C99中和C++中均被废弃，请使用正确的类型声明。

在您修复您的错误之前，无法回答您的任何基于错误代码的任何推论。

玫瑰幻想 发表于 2013-10-11 20:31
楼主您好，

您的代码试图使用“默认int"类型，这在C99中和C++中均被废弃，请使用正确的类型声明。

抱歉，代码写错了，漏了类型声明，应该是下面这个：

1. typedef unsigned char byte;
   
2. static byte H_S[256];
   
3. __device__ byte S[256];
   
4. 
   
5. main(){
   
6.      for(int i = 0;i < 256;i++)
   
7.            H_S[i] = '\0'+i;
   
8. 
   
9.      //然后调用
   
10. cudaStatus = cudaMemcpyToSymbol(S,H_S,256,0,cudaMemcpyHostToDevice);
    
11. }  

复制代码

玫瑰幻想 发表于 2013-10-11 20:31
楼主您好，

您的代码试图使用“默认int"类型，这在C99中和C++中均被废弃，请使用正确的类型声明。

用Nsight调试发现，__device__ byte S[256]  好像只分配64个字节的空间。。。。

hujun21 发表于 2013-10-12 10:02
抱歉，代码写错了，漏了类型声明，应该是下面这个：

楼主您好，经过人脑和上机验证，您的代码毫无问题。

请确保您不是自己的错觉。

hujun21 发表于 2013-10-12 11:55
用Nsight调试发现，__device__ byte S[256]  好像只分配64个字节的空间。。。。

楼主您好，经过人脑和上机验证，

__device__ byte s[256];将分配至少256字节的空间。

请勿疑神疑鬼。

横扫千军 发表于 2013-10-12 12:17
楼主您好，经过人脑和上机验证，

__device__ byte s[256];将分配至少256字节的空间。

感谢版主细心指导啊~确实是我搞错了，问题不出在数据的传输上。但奇怪的是我在vs2012用nsight调试的时候，对S监视，确实只显示了64个值，不过不是大问题，再次感谢~