CUDA内联PTX中如何通过b128指令加载int4类型数据？

错误原因 #

原代码中result是int4向量类型，而"=r"约束仅支持标量寄存器（32位/64位），不兼容128位向量类型，因此编译器报错。

解决方案 #

方案一：使用内存操作数约束 #

直接让PTX指令将128位数据加载到result的内存空间，利用"=m"内存约束适配向量类型：

__device__ int4 LoadVolatile(int4* that) {
    int4 result;
    asm("ld.volatile.global.b128 [%0], [%1];" : "=m"(result) : "l"((void*)that));
    return result;
}

方案二：拆分到64位寄存器再组装 #

利用int4在内存中是连续4个32位整数的特性，将其映射为两个64位变量，通过"=l"约束接收加载的数据：

__device__ int4 LoadVolatile(int4* that) {
    int4 result;
    asm("ld.volatile.global.b128 {%0, %1}, [%2];" 
        : "=l"(*(long long*)&result.x), "=l"(*(long long*)&result.z) 
        : "l"((void*)that));
    return result;
}

原理说明 #

PTX的ld.volatile.global.b128指令支持两种目标操作数：

• 内存地址（[%0]形式）：直接将128位数据写入指定内存位置，对应方案一的"=m"约束。
• 一对64位寄存器（{%0, %1}形式）：将128位数据拆分到两个64位寄存器中，对应方案二的两个"=l"约束，再通过指针转换还原为int4。

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