c - CUDA 中的强制对齐？

Question

有没有办法创建一个 2D 数组 a[][]，其中每个 a[i] 本身都被迫与 CUDA 中的某些其他数据类型对齐？

我想做这样的事情:

__shared__ unsigned char a[20][8];// where a[i] is aligned to 8-byte boundary;

double t=*((double *)(a[2]));

或者甚至类似:

__shared__ unsigned char a[20][9];// where a[i] is aligned to 8-byte boundary;

double t=*((double *)(a[2]));

Answer 1

您显示的两个代码片段在 CUDA 中都是非法的。

所有当前硬件都需要将类型与字大小边界对齐。在您的示例中，仅从 a[0]、a[8]、a[16] 等读取才有效，因为硬件需要任何双读都发生在 8 字节对齐边界处。任何其他操作都会导致设备产生运行时错误。即使您尝试使用结构，编译器也会通过填充强制对齐。如果您想要不同的、有效对齐方式，该语言支持 __align__ specifier对于结构。

这是一种冗长的拒绝方式。