0
__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?

Community
  • 1
  • 1
user2188453
  • 1,105
  • 1
  • 12
  • 26

1 Answers1

2

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).

talonmies
  • 70,661
  • 34
  • 192
  • 269
  • Thanks for this, and does int64 use the same symbol/reg type as uint64 here? – user2188453 Aug 21 '13 at 21:09
  • 1
    The constraint letter is the same between an `s64` register and a `u64` register, but the actual register type is different. If you study the link that @talonmies gave, you'll notice immediately after the table excerpted above, there is an example that passes a parameter to an `s64` register using the `l` constraint letter. Likewise you'll see examples where the `r` constraint letter is used to pass values to/from `s32` registers. – Robert Crovella Aug 22 '13 at 01:05