c - CUDA中的强制对齐？

Question

有没有办法创建一个二维数组 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

使用工会怎么样？

union my_union_type {
    char   a[8];
    double t;
};

my_union_type var[20];

char   a = var.a[2][3];
double t = var.t[2];

这不等同于您的第二个示例，因为（正如 talonmies 指出的那样）您的示例是非法的并且会生成无效代码，而我的代码会在每个元素之后插入 7 个字节的正确对齐方式。但我认为它最接近您的要求。

Answer 2

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

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

这是一个冗长的说不的方式。