为什么静态声明的4个无符号字符数组在获取内存时会产生ld.global.u8？

Question

我正在使用CUDA 5.5并且我发现编译器行为有点奇怪,如果我尝试寻址只有数据是4个无符号字符的结构,它会触发4次加载u8.相反,如果我使用union并加载uchar4,它会产生所需的nc.v4.u8加载

此代码生成ld.global.u8%rs5,[%r32];

        const int wu = 4;
        struct data {
            uchar_t v[wu];           
            CUDA_CALLABLE_MEMBER uchar_t &operator[] (int i) {
                return v[i];
            }
        } fetch[rows];

        for (int i = 0; i < rows; i++) {
            fetch[i] = *((data*)&src[offsetSrc + i*strideSrc]);
        }

所以我必须解决这个问题,建立一个产生所需要的联盟:ld.global.nc.v4.u8 {%rs49,%rs50,%rs51,%rs52},[%r37];

       const int wu = 4;
       struct data {
            union {
                uchar_t v[wu];
                uchar4 v4;
            };
            CUDA_CALLABLE_MEMBER uchar_t &operator[] (int i) {
                return v[i];
            }
        } fetch[rows];

        for (int i = 0; i < rows; i++) {
            fetch[i].v4 = *((uchar4*)&src[offsetSrc + i*strideSrc]);
        }

Answer 1

GPU要求所有数据自然对齐(即16位数据是16位对齐,32位数据是32位对齐,64位数据是64位对齐等).uchar4是四个无符号字符的结构,通过使用alignment属性进行32位对齐.因此,它可以加载一个32位访问.另一方面,四个无符号字符的数组不能保证具有32位对齐,因此不能加载单个32位负载.根据任何组成部分所需的最严格的对齐来对齐联合.

用户定义的数据类型可以与__align__属性对齐,这在CUDA编程指南中有描述