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

[Gurunath Kadam via llvm-dev]

Hi,

I am just following up on my question.

Meanwhile I looked up the 'include/llvm/IR/IntrinsicsNVVM.td' and found
there is a definition for smid (declare i32 @llvm.nvvm.read.ptx.sreg.smid).

At this page <http://llvm.org/docs/NVPTXUsage.html>, there is code for
kernel.ll. I am lost here because this looks like LLVM code completely and
not inline asm.

I also tried a cpp-only program with simple asm (so now this is gcc
compatible code) and found that LLVM can compile it (may be this was a too
simplified test program).

Can someone please point me in the right direction?

Thank you.

-Guru

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

> 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/20161015/eff1777d/attachment-0001.html>