为什么CUDA编程指南中char3的对齐大小是1？

Question

我发现

char1

的对齐为1，

char2

的对齐为2，

char3

的对齐为1，

char4

的对齐为4。 为什么 char3 的对齐大小是 1 而不是 3？

我还需要一些帮助来理解“对齐要求”的概念。

请参阅此CUDA编程指南（附录B中B.3.1节的表B-1）

非常感谢，

易

Answer 1

想想如果你有一个

char3

的数组会发生什么 - > 1 的对齐意味着一个不连续的数组。

对于

char4

但是，您没有任何此类填充问题，并且您可能希望确保

char4

可以被视为 32 位 int，因此 4 字节对齐很有帮助。

Answer 2

对于像我这样仍然对这个问题感到困惑的人： 全局内存访问的大小始终为 1、2、4、8、16 字节等。因此，如果数据是 3 字节，则必须填充以对齐。如果是这样，那么

char3

实际上是底层的

char4

，我们正在使用额外的字节来提高效率。那么对齐是4。

还有另外一种方式，就是CUDA实际使用的方式。他们没有牺牲空间来换取时间，而是反其道而行之。 A

char3

并不是底层的

char4

，而是 3

char1

！

使用 CUDA 11.7 和 CC 8.6，我编写了以下代码：

__device__ char2 ch2[100];
__device__ char3 ch3[100];
__device__ char4 ch4[100];
__global__ void test(char3 chval3,char2 chval2,char4 chval4){
    auto x=threadIdx.x;
    ch3[x]=chval3;
    ch2[x]=chval2;
    ch4[x]=chval4;
}

NVCC 生成的 ptx 代码如下：

ld.param.u8     %rs1, [_Z4test5char35char25char4_param_0+2];
ld.param.u8     %rs2, [_Z4test5char35char25char4_param_0+1];
ld.param.u8     %rs3, [_Z4test5char35char25char4_param_0];
mov.u32     %r1, %tid.x;
mul.wide.u32    %rd1, %r1, 3;
mov.u64     %rd2, ch3;
add.s64     %rd3, %rd2, %rd1;
st.global.u8    [%rd3], %rs3;
st.global.u8    [%rd3+1], %rs2;
st.global.u8    [%rd3+2], %rs1;
mul.wide.u32    %rd4, %r1, 2;
mov.u64     %rd5, ch2;
add.s64     %rd6, %rd5, %rd4;
ld.param.u8     %rs4, [_Z4test5char35char25char4_param_1+1];
ld.param.u8     %rs5, [_Z4test5char35char25char4_param_1];
st.global.v2.u8     [%rd6], {%rs5, %rs4};
mul.wide.u32    %rd7, %r1, 4;
mov.u64     %rd8, ch4;
add.s64     %rd9, %rd8, %rd7;
ld.param.u8     %rs6, [_Z4test5char35char25char4_param_2+3];
ld.param.u8     %rs7, [_Z4test5char35char25char4_param_2+2];
ld.param.u8     %rs8, [_Z4test5char35char25char4_param_2+1];
ld.param.u8     %rs9, [_Z4test5char35char25char4_param_2];
st.global.v4.u8     [%rd9], {%rs9, %rs8, %rs7, %rs6};
ret;

我们可以看到，对

char2

和

char4

的操作被编译为单个指令，如

st.global.v2

、

st.global.v4

，而

char3

实际上被视为3个

char1

。

那么回到问题，当

char3

的对齐是 1 时？因为实际上是 3

char1

，不需要对齐。它不是为了填充，相反，如果我们在这里填充的话，对齐方式将是4。