cuda - 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”表示浮点数，“r”表示整数。

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

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

Answer 1

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

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

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