primitive data type in 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;
        }

The above code is from here: CUDA disable L1 cache only for one variable

According to the author, "d" means float, "r" means int.

I want to write a small piece of inline asm code, I want to know whats the symbol for rest of the primitive type variables (like unsigned short, unsigned long long, float-32, etc), I cannot find that from ptx isa.

I use letter "l" to represent unsigned long long, is that correct?

Answer 1

You can find them here, but for the sake of completeness, the letters correspond to the underlying PTX register types:

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

So an unsigned long long maps to "l" (for a 64 bit integer PTX register).