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

Question

Я использую 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]); 
     } 

Answer 1

ГПУ требует, чтобы все данные, естественно, выровнены (т.е. 16-бит данных 16-битный выровнены, 32-битовые данные 32-битный выровнены, 64-битовые данные 64-битный выровнены, и т.д.). Uchar4 представляет собой структуру из четырех символов без знака, которая 32-битная, выровненная с использованием атрибута выравнивания. Поэтому он может быть загружен одним 32-битным доступом. С другой стороны, массив из четырех символов без знака не гарантирует 32-битного выравнивания и, следовательно, не может быть загружен одной 32-разрядной нагрузкой. Союз выравнивается на основе строгого выравнивания, необходимого для любой из составных частей.

определяемые пользователем типы данных могут быть приведены в соответствие с атрибутом __align__, который описан в CUDA Programming Guide