【问题标题】：CUDA allocation alignment is 256 bytes - seriously?CUDA 分配对齐是 256 字节 - 认真吗？
【发布时间】：2012-12-14 12:18:44
【问题描述】：

在“CUDA C 编程指南 5.0”中，p73 (also here) 说“驻留在全局内存中或由驱动程序或运行时 API 的内存分配例程之一返回的变量的任何地址始终至少对齐256 字节”。我不知道这句话的确切含义。谁能给我举个例子？非常感谢。

一个衍生问题：那么，如何分配基本元素（如 int）或自定义元素的一维数组呢？数组的起始地址会是256B的倍数，而数组中每个元素的地址不一定是256B的倍数？

【问题讨论】：

任何 cudaMalloc 的起始地址都是 256 的倍数。256 是十六进制的两个符号，它们在 addrexx 中为零；所以你可以得到像 0x0456ad00 但不是 0x0456ad80 这样的地址。

【解决方案1】：

使用任何 CUDA 运行时的设备内存分配函数分配的指针，例如 cudaMalloc 或 cudaMallocPitch，保证是 256 字节对齐的，即地址是 256 的倍数。

考虑以下示例：

char *ptr1, *ptr2;

int bytes = 1;

cudaMalloc((void**)&ptr1,bytes);
cudaMalloc((void**)&ptr2,bytes);

假设ptr1中返回的地址是256的某个倍数，那么ptr2中返回的地址至少是(ptr1 + 256)。

这是分配内存的设备所施加的限制。大多数情况下，指针是出于性能目的而对齐的。（一些 NVIDIA 的家伙应该能够判断是否还有其他原因）。

重要：

指针对齐方式并不总是 256。在我的设备（GTX460M）上是512，可以通过cudaDeviceProp::textureAlignment字段获取设备指针对齐方式。

指针对齐也是将指针绑定到纹理的要求。

【讨论】：