CUDA中能否将全局内存128bit数据直接加载到寄存器？

可行性确认 #

128bit全局内存加载完全可以实现，且没有额外硬件开销。GPU单寄存器位宽确实为32bit，但int4等CUDA向量类型在硬件层面会映射为4个连续的32bit寄存器，128bit加载指令会一次性把连续的128bit数据写入这4个寄存器，无需手动拆分。

你当前写法失效的核心原因 #

编译器生成4次32bit加载的问题，几乎都是以下几个原因导致的：

• 结构体未显式指定16字节对齐：128bit全局加载要求目标地址必须16字节对齐，如果你没有给mystruct_t加对齐修饰，编译器无法保证数组中每个结构体的首地址满足对齐要求，就会自动降级为4次32bit加载。
• 编译器优化未开启：必须开启-O2及更高等级优化，编译器才会自动生成向量加载指令。
• 存在指针别名问题：如果内核的数组指针没有加__restrict__修饰，编译器无法确认指针不存在地址重叠，会保守选择不做向量化优化。
• 结构体定义存在语法错误：你给出的结构体定义中成员间用逗号分隔，正确写法应该用分号，语法错误也可能导致编译器对齐判断失效。

修正后的实现代码 #

第一步：修正结构体定义，显式指定对齐 #

typedef struct __align__(16) {
    uint32_t a;
    uint32_t b;
    uint32_t c;
    uint32_t d;
} mystruct_t;

第二步：内核写法优化 #

__global__ void kernel(const mystruct_t* __restrict__ input_array, mystruct_t* __restrict__ output_array, int array_size)
{
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (global_idx >= array_size) return;
    
    // 方法1：普通向量加载，支持后续写回
    int4 vec = *reinterpret_cast<const int4*>(input_array + global_idx);
    // 方法2：只读数据推荐用__ldg加载，走只读数据缓存性能更高
    // int4 vec = __ldg(reinterpret_cast<const int4*>(input_array + global_idx));
    
    uint32_t a = vec.x;
    uint32_t b = vec.y;
    uint32_t c = vec.z;
    uint32_t d = vec.w;
    
    // 此处添加你的计算逻辑
    // ...
    
    // 计算完成后同样可以用128bit向量存回
    *reinterpret_cast<int4*>(output_array + global_idx) = vec;
}

验证方式 #

编译时添加-ptx参数生成PTX汇编代码，搜索是否存在ld.global.v4.u32或者ld.global.nc.v4.u32（使用__ldg时）指令，存在即代表编译器已经生成了128bit向量加载。如果出现4条独立的ld.global.u32指令，则需要检查上述优化项是否都已配置正确。

额外说明 #

该方案保留了AoS的存储布局，同时实现了完美合并的内存访问，warp单次连续读取512字节，带宽利用率接近峰值，完全避免了你担心的SoA布局带来的L2缓存局部性下降问题。如果你的硬件算力在sm_50及以上，即便出现偶发的非对齐访问，硬件也支持非对齐128bit加载，仅会有极小的性能损耗，远优于4次非合并32bit加载的性能。

内容的提问来源于stack exchange，提问作者Quim