[cfe-dev] LLVM/CLANG: CUDA compilation fail for inline assembly code

[Gurunath Kadam via cfe-dev]

I am trying to compile LonestarGPU-rev2.0
<http://iss.ices.utexas.edu/?p=projects/galois/lonestargpu/download>
benchmark suite with LLVM/CLANG.

This suite has a following piece of code (more info here
<https://devtalk.nvidia.com/default/topic/481465/cuda-programming-and-performance/any-way-to-know-on-which-sm-a-thread-is-running-/2/?offset=21#4996171>
):

- static __device__ uint get_smid(void) {
- uint ret;
- asm("mov.u32 %0, %smid;" : "=r"(ret) );
- return ret;
- }
The original make file has nvcc compiler with a flag -Xptxas -v. It
compiles with nvcc.

LLVM has -Xcuda-ptxas <arg>, which I believe is the comparable command for
compiling PTX code. I get following error when I try compiling (clang 4.0).:


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


It points to %smid.

I have been trying to figure out what is this error is but NVIDIA PTX has
this <http://docs.nvidia.com/cuda/cuda-c-programming-guide/>.

Is this a bug or something?

Thanks.

-Guru
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20161013/a6a4a76a/attachment.html>