Search code examples
c++cudallvm-clang

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.

Solution

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

    The %% escape sequence gets converted to a single 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 does in this case.