ptx 中的原始数据类型

Question

__device__ __inline__ double ld_gbl_cg(const double *addr) {
          double return_value;
          asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr));
          return return_value;
        }

以上代码来自这里： CUDA disable L1 cache only for one variable

根据作者的说法，“d”表示float，“r”表示int。

我想写一小段内联 asm 代码，我想知道其余原始类型变量（如 unsigned short、unsigned long long、float-32 等）的符号是什么，我无法从ptx isa。

我用字母“l”代表unsigned long long，对吗？

Answer 1

您可以找到它们here，但为了完整起见，字母对应于底层 PTX 寄存器类型：

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

所以unsigned long long 映射到“l”（对于 64 位整数 PTX 寄存器）。