algoProg algoProg - 1 month ago 16
C++ Question

LLVM with CUDA inline assembly

I am trying to compile a CUDA code with following inline assembly:

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


The code compiles fine with
nvcc
with a flag
-Xptxas -v
.

When i try to compile it with
clang++
(version 4.0), with corresponding flag
-Xcuda-ptxas -v
(I think this is right, but I maybe mistaken), I get following error:


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



It points to
%smid
.

I think I am suppose to link proper library but I have this too:
L/cuda/install/lib
.

Another possibility is NVPTX asm incompatibility. On this page, it is explained that LLVM has different definitions for all PTX variables (there are some for smid and warpid as well). Now I am lost if the mentioned code has to be separately (not inline) written and compiled as such.

Has anybody dealt with similar issue before? Suggestions are welcomed.

Answer

You need to reference the special register with a double percent sign: %%smid.

The %% escape sequence gets gets converted to a singe percent sign during compilation, so that ptxas sees the correct special register name. The double percent sign version also works under nvcc.

nvcc seems to be more forgiving with escape sequences in inline assembler than clang++ is, and leaves unknown escape sequences untouched rather than emitting an error as clang is in this case.