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

[Gurunath Kadam via llvm-dev]

Okay, so as I understand, LLVM inline assembly style is *different than*
the GCC/NVPTX assembly style. So as per LLVM language reference manual
following  constraint codes are supported:


   - b: A 1-bit integer register.
   - c or h: A 16-bit integer register.
   - r: A 32-bit integer register.
   - l or N: A 64-bit integer register.
   - f: A 32-bit float register.
   - d: A 64-bit float register.

Now, I am just wondering if there are equivalent constraints for %smid and
%warpid? As per the NVIDIA documentation
<http://docs.nvidia.com/cuda/cuda-c-programming-guide/#sm-id-and-warp-id> it
is unsafe to rely on this information provided by %smid and %warpid, but
for compiling an existing cuda code I will need to do that.

Any suggestions? Work in progress?

Thank you.

-Guru

On Fri, Oct 14, 2016 at 8:53 AM, Gurunath Kadam <gurunath.kadam at gmail.com>
wrote:

> Hi,
>
> I am sorry for sending this query again here, but maybe I sent it to wrong
> list yesterday.
>
> 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/llvm-dev/attachments/20161014/28d2b524/attachment.html>