b128 加载的内联 ptx 语法

Question

我想从内存中加载一些 int4 大小的东西（又名 16 字节结构），但是 b128 似乎没有限制。

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

这给了我错误：

错误：asm 操作数必须具有标量类型

asm("ld.volatile.global.b128 %0, [%1];" : "=r"(result) : "l"((void*)that));

                                          ^

但是没有 128 位限制，Inline PTX assembly
中的表格 只显示：

“h”= .u16 reg
“r”= .u32 reg
“l” = .u64 reg
“f”=.f32 reg
“d”=.f64 reg

b128 没有限制。

如何更改上述代码以将

int4

加载到

result

中？

Answer 1

自从大约 20 年前首次发布 CUDA 以来，PTX 就支持 C++ 向量类型的向量化加载和存储。在这种情况下，您想要使用“V4”版本的负载，在内联 PTX 中将返回

int4

viz 的四个成员：

__device__ int4 LoadVolatile(const int4* that) 
{ 
    int4 result; 
    asm("ld.volatile.global.v4.u32 {%0, %1, %2, %3}, [%4];" : 
        "=r"(result.x), 
        "=r"(result.y), 
        "=r"(result.z), 
        "=r"(result.w) : "l"((void*)that)); 
    return result; 
}

b128没有限制

128 位整数类型（相对于 128 位向量类型）是 CUDA 中的一个新事物，只有最新的架构支持它们。虽然 128 位整数加载在这里是错误的解决方案，但 CUDA 中的内联 PTX 工具似乎尚未更新以包含 128 位整数类型约束。如果这对您来说真的很重要，我建议向 NVIDIA 提交功能请求。