CUDA中的地址错位

Question

谁能告诉我CUDA内核中的以下代码有什么问题:

__constant__ unsigned char MT[256] = {
    0xde, 0x6f, 0x6f, 0xb1, 0xde, 0x6f, 0x6f, 0xb1, 0x91, 0xc5, 0xc5, 0x54, 0x91, 0xc5, 0xc5, 0x54,....};

typedef unsinged int U32;

__global__ void Kernel (unsigned int  *PT, unsigned int  *CT, unsigned int  *rk)
{

    long int i;
    __shared__ unsigned char sh_MT[256];    

    for (i = 0; i < 64; i += 4)
        ((U32*)sh_MT)[threadIdx.x + i] = ((U32*)MT)[threadIdx.x + i];

    __shared__ unsigned int sh_rkey[4];
    __shared__ unsigned int sh_state_pl[4];
    __shared__ unsigned int sh_state_ct[4];

    sh_state_pl[threadIdx.x] = PT[threadIdx.x];
    sh_rkey[threadIdx.x] = rk[threadIdx.x];
    __syncthreads();


    sh_state_ct[threadIdx.x] = ((U32*)sh_MT)[sh_state_pl[threadIdx.x]]^\
    ((U32*)(sh_MT+3))[((sh_state_pl[(1 + threadIdx.x) % 4] >> 8) & 0xff)] ^ \
    ((U32*)(sh_MT+2))[((sh_state_pl[(2 + threadIdx.x) % 4] >> 16) & 0xff)] ^\
    ((U32*)(sh_MT+1))[((sh_state_pl[(3 + threadIdx.x) % 4] >> 24) & 0xff )];


    CT[threadIdx.x] = sh_state_ct[threadIdx.x];
}

在这行代码中,

((U32*)(sh_MT+3))......

CUDA调试器给出了错误消息: 未对齐的地址

我该如何解决这个错误？

我在MVSC中使用CUDA 7,我使用1个Block和4个线程来执行内核函数,如下所示:

__device__ unsigned int *state;
__device__ unsigned int *key;
__device__ unsigned int *ct;
.
.
main()
{
cudaMalloc((void**)&state, 16);
cudaMalloc((void**)&ct, 16);
cudaMalloc((void**)&key, 16);
//cudamemcpy(copy some values to => state , ct, key);   
Kernel << <1, 4 >> >(state, ct, key); 
}

请记住,我不能改变我的"MT Table"类型.提前感谢任何建议或答案.

Answer 1

正如错误消息所示,指针未与处理器所需的边界对齐.

从CUDA编程指南,第5.3.2节:

全局存储器指令支持读取或写入大小等于1,2,4,8或16字节的字.当且仅当数据类型的大小为1,2,4,8或16字节并且数据自然地存在时,对驻留在全局存储器中的数据的任何访问(通过变量或指针)编译成单个全局存储器指令.对齐(即,其地址是该大小的倍数).

这是调试器试图告诉您的内容:基本上,您不应该从未在32位边界对齐的地址取消引用指向32位值的指针.

你可以做到(U32*)(sh_MT)并且(U32*)(sh_MT+4)很好,但不是(U32*)(sh_MT+3)或者这样.

您可能必须单独读取字节并将它们连接在一起.