Dear LLVM,<br><br>Could someone please commit the attached patch?<br><br>Description: adding curly braces in PTX code, when they are specified in LLVM inline asm for definitions scoping.<br><br><a href="http://llvm.org/bugs/show_bug.cgi?id=13322" target="_blank">http://llvm.org/bugs/show_bug.cgi?id=13322</a><br>
<br>Staged for commit, as per discussed below.<br><br>Thanks<br>- Dima Mikushin.<br><br><div class="gmail_quote">2012/7/11 Dmitry N. Mikushin <span dir="ltr"><<a href="mailto:maemarcus@gmail.com" target="_blank">maemarcus@gmail.com</a>></span><br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Let me propose a fix:<br><br>--- a/llvm/tools/clang/lib/Basic/Targets.cpp (revision 157736)<br>+++ b/llvm/tools/clang/lib/Basic/Targets.cpp (working copy)<br>
@@ -966,6 +966,10 @@<br> AddrSpaceMap = &NVPTXAddrSpaceMap;<br>
// Define available target features<br> // These must be defined in sorted order!<br>+<br>+ // {} in inline assembly are local scope specifiers, not assembly variant<br>+ // specifiers.<br>+ NoAsmVariants = true;<br>
}<br> virtual void getTargetDefines(const LangOptions &Opts,<br> MacroBuilder &Builder) const {<br><br>Is it correct? Could you please review,<br><br>Thanks,<br>- D.<br>
<br>
<div class="gmail_quote">2012/7/11 Dmitry N. Mikushin <span dir="ltr"><<a href="mailto:maemarcus@gmail.com" target="_blank">maemarcus@gmail.com</a>></span><div><div class="h5"><br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
Yes, sure, good idea, because might be also Clang-related.<br><br><a href="http://llvm.org/bugs/show_bug.cgi?id=13322" target="_blank">http://llvm.org/bugs/show_bug.cgi?id=13322</a><div><div><br>
<br><div class="gmail_quote">2012/7/11 Chad Rosier <span dir="ltr"><<a href="mailto:mcrosier@apple.com" target="_blank">mcrosier@apple.com</a>></span><br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div style="word-wrap:break-word">Dmitry,<div>You might be better served by filing this as a bug (<a href="http://llvm.org/bugs/" target="_blank">http://llvm.org/bugs/</a>). Please include a test case and the steps to reproduce (i.e., what you've provided below).</div>
<div><br></div><div> Chad</div><div><br><div><div><div><div>On Jul 10, 2012, at 3:15 PM, Dmitry N. Mikushin wrote:</div><br></div></div><blockquote type="cite"><div><div>Hi,<br><br>Looks like "{" and "}" are lost when trying to use the combination of Clang and NVPTX, which may result into clash of definitions of the function-scope and asm-scope. Here is an example:<br>
<br>
> cat <a href="http://test.cu/" target="_blank">test.cu</a><br>__attribute__((device)) __attribute__((nv_linkonce_odr)) __inline__ int __any(int a) {<br> int result;<br> asm __volatile__ ("{ \n\t"<br> ".reg .pred \t%%p1; \n\t"<br>
".reg .pred \t%%p2; \n\t"<br> "setp.ne.u32 \t%%p1, %1, 0; \n\t"<br> "vote.any.pred \t%%p2, %%p1; \n\t"<br> "selp.s32 \t%0, 1, 0, %%p2; \n\t"<br> "}" : "=r"(result) : "r"(a));<br>
return result;<br>}<br><br>> clang -cc1 -emit-llvm -fcuda-is-device -triple ptx64-unknown-unknown <a href="http://test.cu/" target="_blank">test.cu</a> -o test.ll<br>> cat test.ll<br>; ModuleID = '<a href="http://test.cu/" target="_blank">test.cu</a>'<br>
target datalayout = "e-p:64:64-i64:64:64-f64:64:64-n1:8:16:32:64"<br>target triple = "ptx64-unknown-unknown"<br><br>define ptx_device i32 @_Z5__anyi(i32 %a) nounwind inlinehint {<br>entry:<br> %a.addr = alloca i32, align 4<br>
%result = alloca i32, align 4<br> store i32 %a, i32* %a.addr, align 4<br> %0 = load i32* %a.addr, align 4<br> %1 = call i32 asm sideeffect "$( \0A\09.reg .pred \09%p1; \0A\09.reg .pred \09%p2; \0A\09setp.ne.u32 \09%p1, $1, 0; \0A\09vote.any.pred \09%p2, %p1; \0A\09selp.s32 \09$0, 1, 0, %p2; \0A\09$)", "=r,r"(i32 %0) nounwind, !srcloc !0<br>
store i32 %1, i32* %result, align 4<br> %2 = load i32* %result, align 4<br> ret i32 %2<br>}<br><br>!0 = metadata !{i32 127, i32 132, i32 166, i32 200, i32 242, i32 285, i32 327}<br><br>> llc -march=nvptx64 test.ll -o test.ptx<br>
> cat test.ptx<br>//<br>// Generated by LLVM NVPTX Back-End<br>//<br><br>.version 3.0<br>.target sm_10, texmode_independent<br>.address_size 64<br><br><br> // .globl _Z5__anyi<br>.visible .global .align 4 .b8 __local_depot0[8];<br>
<br>.func (.reg .b32 func_retval0) _Z5__anyi(<br> .reg .b32 _Z5__anyi_param_0<br>) // @_Z5__anyi<br>{<br> .reg .b64 %SP;<br> .reg .b64 %SPL;<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>
mov.u64 %SP, __local_depot0;<br> mov.b32 %r0, _Z5__anyi_param_0;<br> st.global.u32 [%SP+0], %r0;<br> // inline asm<br> <br> .reg .pred %p1; <br> .reg .pred %p2; <br> setp.ne.u32 %p1, %r0, 0; <br>
vote.any.pred %p2, %p1; <br> selp.s32 %r0, 1, 0, %p2; <br> <br> // inline asm<br> st.global.u32 [%SP+4], %r0;<br> mov.b32 func_retval0, %r0;<br> ret;<br>}<br><br>> ptxas test.ptx -o test.cubin<br>
ptxas test.ptx, line 33; error : Duplicate definition of variable '%p1'<br>ptxas test.ptx, line 34; error : Duplicate definition of variable '%p2'<br>ptxas test.ptx, line 36; error : Instruction 'vote' requires .target sm_12 or higher<br>
ptxas fatal : Ptx assembly aborted due to errors<br><br>- D.<br><br></div></div>
_______________________________________________<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></blockquote></div><br></div></div>
</blockquote></div><br>
</div></div></blockquote></div></div></div><br>
</blockquote></div><br>