全局内存数据对齐问题~

加速计算 CUDA
1 
typedef unsigned int  uint32;
typedef unsigned char byte;

__device__ byte data[1024];

__device__ inline void Copy128(byte *dest,const byte *src)
{

  ((uint32*)dest)[0]=((uint32*)src)[0];
  ((uint32*)dest)[1]=((uint32*)src)[1];
  ((uint32*)dest)[2]=((uint32*)src)[2];
  ((uint32*)dest)[3]=((uint32*)src)[3];

}

__global__ kernel(){
   byte v[4][4];
   copy128((byte *)v,data);
}

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

2

楼主你好：

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

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

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

感谢您的来访。

3

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

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

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

感谢您的来访。

4

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

static H_S[256];
__device__ S[256];

main(){
   for(int i = 0;i < 256;i++)
   H_S[i] = '\0'+i;

   //然后调用
cudaStatus = cudaMemcpyToSymbol(S,H_S,256,0,cudaMemcpyHostToDevice);
}
5

楼主您好，

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

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

6

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

typedef unsigned char byte;
static byte H_S[256];
__device__ byte S[256];

main(){
   for(int i = 0;i < 256;i++)
   H_S[i] = '\0'+i;

   //然后调用
cudaStatus = cudaMemcpyToSymbol(S,H_S,256,0,cudaMemcpyHostToDevice);
}
7

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

8

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

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

9

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

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

请勿疑神疑鬼。

10

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

11

楼主您好，

nsight可以在nsight菜单的选项里修改最多显示的数组里的元素数目的，默认是64个。如果您需要256个。请设置成256.

感谢夜晚来访。

12

原来是这样~感谢版主指点

13

服务您是我们的荣幸，欢迎再次光临。