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

Accelerated Computing CUDA CUDA Programming and Performance
21

Hi,

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?

22

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.
https://users.ices.utexas.edu/~sreepai/fermi-tbs/

23

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.