Почему статический объявленный массив из 4 символов без знака создает ld.global.u8 при выборке памяти?

Я использую CUDA 5.5 и нахожу поведение компилятора немного странным: если я попытаюсь обратиться к структуре, в которой только данные состоят из 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]);
        }