any way to know on which SM a thread is running?


I am trying to compile a program with LLVM/CLANG which has this piece of code:

static __device__ uint get_smid(void) {
     uint ret;
     asm("mov.u32 %0, %smid;" : "=r"(ret) );
     return ret;

I get following error:

../../include/cutil_subset.h:23:25: error: invalid % escape in inline assembly string
     asm("mov.u32 %0, %smid;" : "=r"(ret) );

Any idea what’s going on?

You can go through this document, which I came across and I loved it. It tells how the scheduling of blocks occur in Fermi architecture.

I’m not sure it still matters after almost a year, but for the record (and anyone getting here via a search for the error message):
The correct syntax for the asm statement is

asm("mov.u32 %0, %%smid;" : "=r"(ret) );

The double “%%” turns into a single % during parameter substitution, which then forms the correct %smid special register name. nvcc is more forgiving, leaving unknown escape sequences untouched, but clang complains about them.

1 Like