如何在CUDA中使用uint4向量正确转换全局内存数组以增加内存吞吐量？
通常有两种技术可以在计算能力1.3 GPU上增加CUDA内核上全局内存的内存吞吐量。 内存访问合并并访问至少4个字节的字。 利用第一种技术，通过相同半翘曲的线程对相同存储器段的访问被合并为更少的事务，而在访问至少4字节的字时，该存储器段有效地从32字节增加到128。
更新：基于talonmies回答的解决方案 。 当存在全局存储器中存在无符号字符时，要访问16字节而不是1字节字，通常通过将存储器arrays转换为uint4来使用uint4向量。 要从uint4向量中获取值，可以将其重新编译为uchar4，如下所示：
#include #include #include __global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; extern __shared__ unsigned char s_array[]; uint4 *uint4_text = reinterpret_cast(d_text); uint4 uint4_var; //memory transaction uint4_var = uint4_text[0]; //recast data to uchar4 uchar4 c0 = *reinterpret_cast(&uint4_var.x); uchar4 c4 = *reinterpret_cast(&uint4_var.y); uchar4 c8 = *reinterpret_cast(&uint4_var.z); uchar4 c12 = *reinterpret_cast(&uint4_var.w); d_out[idx] = c0.y; } int main ( void ) { unsigned char *d_text, *d_out; unsigned char *h_out = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) ); unsigned char *h_text = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) ); int i; for ( i = 0; i < 16; i++ ) h_text[i] = 65 + i; cudaMalloc ( ( void** ) &d_text, 16 * sizeof ( unsigned char ) ); cudaMalloc ( ( void** ) &d_out, 16 * sizeof ( unsigned char ) ); cudaMemcpy ( d_text, h_text, 16 * sizeof ( unsigned char ), cudaMemcpyHostToDevice ); kernel<<>>(d_text, d_out ); cudaMemcpy ( h_out, d_out, 16 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost ); for ( i = 0; i < 16; i++ ) printf("%c\n", h_out[i]); return 0; } 如果我已经理解了你要做的事情，那么逻辑方法是使用C ++ reinterpret_cast机制使编译器生成正确的向量加载指令，然后使用内置字节大小的向量类型uchar4的CUDA来访问每个内部的每个字节。从全局内存加载的四个32位字。 使用这种方法，您真正信任编译器，知道在每个32位寄存器中进行字节访问的最佳方式。
一个完全做作的例子可能如下所示：
#include #include __global__ void kernel(unsigned int *in, unsigned char* out) { int tid = threadIdx.x; uint4* p = reinterpret_cast(in); uint4 i4 = p[tid]; // vector load here uchar4 c0 = *reinterpret_cast(&i4.x); uchar4 c4 = *reinterpret_cast (&i4.y); uchar4 c8 = *reinterpret_cast (&i4.z); uchar4 c12 = *reinterpret_cast (&i4.w); out[tid*4+0] = c0.x; out[tid*4+1] = c4.y; out[tid*4+2] = c8.z; out[tid*4+3] = c12.w; } int main(void) { unsigned int c[8] = { 2021161062, 2021158776, 2020964472, 1920497784, 2021161058, 2021161336, 2020898936, 1702393976 }; unsigned int * _c; cudaMalloc((void **)&_c, sizeof(int)*size_t(8)); cudaMemcpy(_c, c, sizeof(int)*size_t(8), cudaMemcpyHostToDevice); unsigned char * _m; cudaMalloc((void **)&_m, sizeof(unsigned char)*size_t(8)); kernel<<<1,2>>>(_c, _m); unsigned char m[8]; cudaMemcpy(m, _m, sizeof(unsigned char)*size_t(8), cudaMemcpyDeviceToHost); for(int i=0; i<8; i++) fprintf(stdout, "%d %c\n", i, m[i]); return 0; }
这应该产生一个嵌入在提供给内核的无符号整数数组中的可读字符串。
需要注意的是，用于计算1.x目标的open64编译器经常会破坏这种尝试生成向量加载的策略，如果它能够检测到并非向量中的所有单词都被实际使用。 因此，请确保触摸输入向量类型中的所有输入字以确保编译器可以很好地播放。
施放到char *会很好。 你试过了吗？ 如果是这样，发生了什么促使这个问题？
在您的示例中，看起来您可以将s_array转换为int*并从var.x执行单个副本（将j乘以4而不是16）。
如果你需要更灵活地改变单词中的字节，你可以使用__byte_perm()内在函数。 例如，要颠倒整数x字节的顺序，可以执行__byte_perm(x, 0, 0x0123);
通过使用矢量类型甚至单个int来存储字节，您可能无法获得任何东西。 在Fermi上，全局内存事务是128字节宽。 因此，当您的warp命中一个从全局内存加载/存储的指令时，GPU将执行为32个线程提供服务所需的128字节事务。 性能在很大程度上取决于需要多少单独的事务，而不是每个线程如何加载或存储其字节。