1

我正在使用 CUDA 5.5，我发现编译器的行为有点奇怪，如果我尝试解决只有数据是 4 个无符号字符的结构，它会触发 4 次加载 u8。相反，如果我使用联合并加载 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]);
        }
4

1 回答 1

6

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

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

于 2013-06-12T15:39:36.947 回答