CUDA内存拷贝力对准

我写了一个测试来说明我的问题，该代码尝试16个字节复制到无 - 4字节对齐的存储空间，但DEST自动修改CUDA内存拷贝力对准

#include <cuda.h> 
#include <cuda_runtime.h> 
#include <stdio.h> 

__global__ 
void Copy128(char *dest,const char *src) 
{ 
    ((int*)dest)[0]=((int*)src)[0]; 
    ((int*)dest)[1]=((int*)src)[1]; 
    ((int*)dest)[2]=((int*)src)[2]; 
    ((int*)dest)[3]=((int*)src)[3]; 
} 
__global__ 
void fill_src(char *src) 
{ 
    for(int i=0; i<16; i++) 
     src[i] = i+1; // starts from 1 
} 

int main() 
{ 
    char* dest; 
    cudaMalloc(&dest, 17); 

    char* src; 
    cudaMalloc(&src, 16); 

    fill_src<<<1, 1>>>((char*)src); // fill some value for debugging 

    // copy to dest+1 which is not aligned to 4 
    Copy128<<<1, 1>>>(dest + 1, src); 

    getchar(); 
}

在VS2013中调试代码，如图所示，目标内存为0x40A8000 ，但实际上它复制到0x40A8000 。

问题是dest如果它没有对齐到4字节会自动修改。它被无声地修改，我花了几个小时找到这个bug。我知道最好是使用良好对齐的内存，但我正在写一些rar解压缩程序，解压缩一些字节，然后连接一些字节，它不能总是对齐。

我想我会使用uint64像Copy256一样的功能。这是正常行为，内存是力量对齐？任何可以关闭此功能的编译标志？或者我应该一个接一个地复制字节？

环境：CUDA 6.5，Win7-32bit，VS2013

来源

2016-04-19 aj3423

当我运行您的示例代码时，由于未对齐的内存访问，在Copy128内核中出现非法写入错误，这正是应该发生的情况。我不明白你在这里试图做什么 – talonmies

除了x86 CPU之外，GPU上的所有内存访问必须自然对齐，即与访问的大小对齐，例如， 4字节访问必须与4字节边界对齐。所以在GPU上，内存访问的这种对齐对于*功能正确性*是必要的，而不仅仅是在x86上的性能。这在CUDA文档中提到。对于未对齐的副本，您无需逐个字节地复制较大的对象，只需对最终案例使用较窄的访问权限，并将大量副本用于大部分传输。 – njuffa

- 这是正常的行为，该内存的力对齐？ 是：引自here，“驻留在全局内存中或由驱动程序或运行时API的内存分配例程之一返回的变量的任何地址总是对齐至少256个字节”。

任何可以关闭此功能的编译标志？ 我想不会，这可能是与硬件相关的

还是应该由一个复制字节一个？ 如果你处理（非常）未对齐的内存，这是你唯一的选择，以避免错位的商店（如上面评论）。但是，您应该尝试在编译时或运行时检测内存操作是否对齐，然后使用手头最宽的加载/存储（int4会导致ldg指令，这会提供更好的方法带宽）

来源

2016-04-19 11:33:54

CUDA内存拷贝力对准

回答

相关问题