Hi Justin,<br><br>Thanks for explanation!<br><br>Not really a blocker, but a bit of extra work, as I'm going to convert /opt/cuda/nvvm/ci_include.h to LLVM IR, in order to fuse it with non-C frontend language at IR-level. This header contains __inline__-s.<br>
<br>- D.<br><br><div class="gmail_quote">2012/6/13 Justin Holewinski <span dir="ltr"><<a href="mailto:justin.holewinski@gmail.com" target="_blank">justin.holewinski@gmail.com</a>></span><br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div class="gmail_quote"><div><div class="h5">On Tue, Jun 12, 2012 at 6:11 PM, Dmitry N. Mikushin <span dir="ltr"><<a href="mailto:maemarcus@gmail.com" target="_blank">maemarcus@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
Dear LLVM NVPTX maintainers,<br><br>Just to have the issue recorded, I don't know how important it is:<br><br>clang generates linkonce_odr out of
__inline__, and NVPTX generates .weak out of linkonce_odr (how it
happens - a big question, btw, because I can't find anything related in
NVPTX asm printer - does it chain to some other printer?), and finally
ptxas (both 4.2 and 5) fails to compile it to cubin. Below is the test case:<br>
<br>> cat <a href="http://test3.cu/" target="_blank">test3.cu</a><br><br>__inline__ __attribute__((device)) __attribute__((used)) void test()<br>{<br> return;<br>}<br><br>> clang -cc1 -emit-llvm -triple ptx64-unknown-unknown -fcuda-is-device <a href="http://test3.cu/" target="_blank">test3.cu</a> -o test3.ll<br>
> cat test3.ll<br>; ModuleID = '<a href="http://test3.cu/" target="_blank">test3.cu</a>'<br>target datalayout = "e-p:64:64-i64:64:64-f64:64:<div>64-n1:8:16:32:64"<br>target triple = "ptx64-unknown-unknown"<br>
<br>
@llvm.used = appending global [1 x i8*] [i8* bitcast (void ()* @_Z4testv to i8*)], section "llvm.metadata"<br><br>define linkonce_odr ptx_device void @_Z4testv() nounwind inlinehint {<br>entry:<br> ret void<br>
}<br><br>> llc -march=nvptx64 -mcpu=sm_20 test3.ll -o test3.ptx<br>> cat test3.ptx<br><br>//<br>// Generated by LLVM NVPTX Back-End<br>//<br><br>.version 3.0<br>.target sm_20, texmode_independent<br>.address_size 64<br>
<br><br> .weak _Z4testv<br>.func _Z4testv(<br><br>) // @_Z4testv<br>{<br> .reg .pred %p<396>;<br> .reg .s16 %rc<396>;<br> .reg .s16 %rs<396>;<br> .reg .s32 %r<396>;<br>
.reg .s64 %rl<396>;<br> .reg .f32 %f<396>;<br> .reg .f64 %fl<396>;<br><br>// BB#0: // %entry<br> ret;<br>}<br><br>1) ptxas @ CUDA 4.2:<br><br>> ptxas -arch=sm_20 -m64 test3.ptx -o -<br>
<br>ptxas test3.ptx, line 10; fatal : Parsing error near '.weak': syntax error<br>ptxas fatal : Ptx assembly aborted due to errors<br><br>2) ptxas @ CUDA 5:<br><br>> ~/cuda/bin/ptxas -arch=sm_20 -m64 test3.ptx -o -<br>
<br>ptxas test3.ptx, line 10; error : Feature '.weak directive' requires PTX ISA .version 3.1 or later<br>ptxas test3.ptx, line 10; fatal : Parsing error near '_Z4testv': syntax error<br>ptxas fatal : Ptx assembly aborted due to errors<br>
<br>3) ptxas @ CUDA 5, changed .version to 3.1: still error, because according to 3.1 PTX spec, .weak must be followed by .func:<br><br>> ~/cuda/bin/ptxas -arch=sm_20 -m64 test3.ptx -o -<br><br>ptxas test3.ptx, line 10; fatal : Parsing error near '_Z4testv': syntax error<br>
ptxas fatal : Ptx assembly aborted due to errors<br></div></blockquote><div><br></div></div></div><div>Thanks for the report. Unfortunately, this does not appear to have a trivial fix. As you mentioned, it is not the NVPTX back-end itself that is emitting the ".weak", but the default MCAsmStreamer implementation. Setting WeakDefDirective in the NVPTXMCAsmInfo class seems to trigger an emission of ".weak_directive", which doesn't help things. Setting LinkOnceDirective helps the ".weak" case, but there is code in LLVM that causes a special label to be produced when LinkOnceDirective is set, which again messes with the PTX assembler.</div>
<div><br></div><div>This is a case where we really need a custom MCAsmStreamer, but this will take a bit of time. Is this a blocker for you?</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div>
<br>Best,<br>- Dima.</div>
<br>_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:LLVMdev@cs.uiuc.edu" target="_blank">LLVMdev@cs.uiuc.edu</a> <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
<br></blockquote></div><span class="HOEnZb"><font color="#888888"><br><br clear="all"><div><br></div>-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div><br>
</font></span></blockquote></div><br>