<div dir="ltr"><span style="font-family:arial,sans-serif;font-size:13px">That did the trick! Using </span><span style="font-family:arial,sans-serif;font-size:13px"> llvm.nvvm.read.ptx.sreg.tid.x(</span><span style="font-family:arial,sans-serif;font-size:13px">) compiles down to a mov from %tid.x</span><br>
</div><div class="gmail_extra"><br><br><div class="gmail_quote">On Fri, Mar 1, 2013 at 2:18 PM, Justin Holewinski <span dir="ltr"><<a href="mailto:justin.holewinski@gmail.com" target="_blank">justin.holewinski@gmail.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr">The identifier <span style="font-family:arial,sans-serif;font-size:13px">INT_PTX_SREG_TID_X is the name of an instruction as the back-end sees it, and has very little to do with the name you should use in your IR. Your best bet is to look at the include/llvm/IR/IntrinsicsNVVM.td file and see the definitions for each intrinsic. Then, the name mapping is just:</span><div>
<span style="font-family:arial,sans-serif;font-size:13px"><br></span></div><div><span style="font-family:arial,sans-serif;font-size:13px">int_foo_bar -> llvm.foo.bar()</span></div><div><br></div><div>
<font face="arial, sans-serif">int_ prefix becomes llvm., and all underscores turn into periods.</font></div><div><font face="arial, sans-serif"><br></font></div><div><font face="arial, sans-serif"><br></font></div>
<div><font face="arial, sans-serif">Ex:</font></div><div><font face="arial, sans-serif"><br></font></div><div><span style="font-family:arial,sans-serif;font-size:13px">int_nvvm_read_ptx_sreg_tid_x -> llvm.nvvm.read.ptx.sreg.tid.x()</span><font face="arial, sans-serif"><br>
</font></div><div><span style="font-family:arial,sans-serif;font-size:13px"><br></span></div></div><div class="gmail_extra"><div><div class="h5"><br><br><div class="gmail_quote">On Fri, Mar 1, 2013 at 3:51 PM, Pete Couperus <span dir="ltr"><<a href="mailto:pjcoup@gmail.com" target="_blank">pjcoup@gmail.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Hi Timothy,<br>
<br>
I'm not sure what you mean by this working for other intrinsics, but<br>
in this case, I think you want the intrinsic name<br>
llvm.nvvm.read.ptx.sreg.tid.x.<br>
<br>
For me, this looks like:<br>
%x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()<br>
<span><font color="#888888"><br>
Pete<br>
</font></span><div><div><br>
<br>
On Fri, Mar 1, 2013 at 11:51 AM, Timothy Baldridge <<a href="mailto:tbaldridge@gmail.com" target="_blank">tbaldridge@gmail.com</a>> wrote:<br>
> I'm building this with llvm-c, and accessing these intrinsics via calling<br>
> the intrinsic as if it were a function.<br>
><br>
> class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :<br>
> NVPTXInst<(outs regclassOut:$dst), (ins),<br>
> OpStr,<br>
> [(set regclassOut:$dst, (IntOp))]>;<br>
><br>
> def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,<br>
> int_nvvm_read_ptx_sreg_tid_x>;<br>
><br>
> This method of accessing intrinsics works just fine for other intrinsics<br>
> (for instance sqrt). Should I be declaring these as extern global variables?<br>
><br>
> Thanks,<br>
><br>
> Timothy<br>
><br>
><br>
> On Fri, Mar 1, 2013 at 12:44 PM, Dmitry Mikushin <<a href="mailto:dmitry@kernelgen.org" target="_blank">dmitry@kernelgen.org</a>><br>
> wrote:<br>
>><br>
>> Timothy,<br>
>><br>
>> Those calls to compute grid intrinsics are definitely wrong. In ptx code<br>
>> they should end up into reading special registers, rather than function<br>
>> calls. Try to take some working example and figure out the LLVM IR<br>
>> differences between it and the result of your compiler.<br>
>><br>
>> - D.<br>
>><br>
>> ----- Original message -----<br>
>> > I've written a compiler that outputs PTX code, the result seems fairly<br>
>> > reasonable, but I'm not sure the intrinsics are getting compiled<br>
>> > correctly.<br>
>> ><br>
>> > In addition, when I try load the module using CUDA, I get an<br>
>> > error: CUDA_ERROR_NO_BINARY_FOR_GPU. I'm running this on a 2012 MBP<br>
>> > with a 640M GPU.<br>
>> ><br>
>> > PTX Code (for a mandelbrot calculation):<br>
>> ><br>
>> > //<br>
>> > // Generated by LLVM NVPTX Back-End<br>
>> > //<br>
>> ><br>
>> > .version 3.1<br>
>> > .target sm_10, texmode_independent<br>
>> > .address_size 64<br>
>> ><br>
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X<br>
>> > (<br>
>> ><br>
>> > )<br>
>> > ;<br>
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y<br>
>> > (<br>
>> ><br>
>> > )<br>
>> > ;<br>
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X<br>
>> > (<br>
>> ><br>
>> > )<br>
>> > ;<br>
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_X<br>
>> > (<br>
>> ><br>
>> > )<br>
>> > ;<br>
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_Y<br>
>> > (<br>
>> ><br>
>> > )<br>
>> > ;<br>
>> ><br>
>> > // .globl examples_2E_mandelbrot_2F_square<br>
>> > .func (.reg .b64 func_retval0) examples_2E_mandelbrot_2F_square(<br>
>> > .reg .b64 examples_2E_mandelbrot_2F_square_param_0<br>
>> > )<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>
>> > mov.f64 %fl0, examples_2E_mandelbrot_2F_square_param_0;<br>
>> > mul.f64 %fl0, %fl0, %fl0;<br>
>> > mov.f64 func_retval0, %fl0;<br>
>> > ret;<br>
>> > }<br>
>> ><br>
>> > // .globl<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration<br>
>> > .func (.reg .b64 func_retval0)<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration( .reg .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_0, .reg<br>
>> > .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_1, .reg<br>
>> > .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_2, .reg<br>
>> > .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_3, .reg<br>
>> > .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_4 )<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>
>> > mov.f64 %fl0,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_0;<br>
>> > mov.f64<br>
>> > %fl1, examples_2E_mandelbrot_2F_calc_2D_iteration_param_3;<br>
>> > div.rn.f64 %fl0, %fl0, %fl1; mov.f64 %fl2,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_1;<br>
>> > mul.f64<br>
>> > %fl1, %fl0, 0d400C000000000000; mov.f64 %fl0,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_2;<br>
>> > add.f64<br>
>> > %fl1, %fl1, 0dC004000000000000; mov.f64 %fl3,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_iteration_param_4;<br>
>> > div.rn.f64<br>
>> > %fl2, %fl2, %fl3; add.f64 %fl2,<br>
>> > %fl2, %fl2;<br>
>> > add.f64 %fl3, %fl2, 0dBFF0000000000000;<br>
>> > mov.f64 %fl2, 0d0000000000000000;<br>
>> > mov.f64 %fl5, %fl2;<br>
>> > mov.f64 %fl4, %fl2;<br>
>> > bra.uni BB1_1;<br>
>> > BB1_2:<br>
>> > add.f64 %fl2, %fl2, 0d3FF0000000000000;<br>
>> > sub.f64 %fl6, %fl6, %fl7;<br>
>> > add.f64 %fl6, %fl6, %fl1;<br>
>> > add.f64 %fl5, %fl5, %fl5;<br>
>> > mul.f64 %fl4, %fl5, %fl4;<br>
>> > add.f64 %fl4, %fl4, %fl3;<br>
>> > mov.f64 %fl5, %fl6;<br>
>> > BB1_1:<br>
>> > mul.f64 %fl6, %fl5, %fl5;<br>
>> > mul.f64 %fl7, %fl4, %fl4;<br>
>> > add.f64 %fl8, %fl6, %fl7;<br>
>> > setp.lt.f64 %p0, %fl8, 0d4010000000000000;<br>
>> > setp.lt.f64 %p1, %fl2, %fl0;<br>
>> > and.pred %p0, %p0, %p1;<br>
>> > @!%p0 bra BB1_3;<br>
>> > bra.uni BB1_2;<br>
>> > BB1_3:<br>
>> > mov.f64 func_retval0, %fl2;<br>
>> > ret;<br>
>> > }<br>
>> ><br>
>> > // .globl<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx .func (.reg .b64<br>
>> > func_retval0) examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx(<br>
>> > .reg .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0,<br>
>> > .reg .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1,<br>
>> > .reg .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2,<br>
>> > .reg .b64<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3<br>
>> > )<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>
>> > mov.b64 %rl0,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0;<br>
>> > mov.f64 %fl2,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1;<br>
>> > mov.f64 %fl1,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2;<br>
>> > mov.f64 %fl0,<br>
>> > examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3;<br>
>> > // Callseq Start 0<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_CTAID_X,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r0, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 0<br>
>> > // Callseq Start 1<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_NTID_X,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r1, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 1<br>
>> > // Callseq Start 2<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_TID_X,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r2, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 2<br>
>> > mad.lo.s32 %r0, %r0, %r1, %r2;<br>
>> > cvt.rn.f64.s32 %fl3, %r0;<br>
>> > // Callseq Start 3<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_CTAID_Y,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r0, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 3<br>
>> > // Callseq Start 4<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_NTID_Y,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r1, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 4<br>
>> > // Callseq Start 5<br>
>> > {<br>
>> > .reg .b32 temp_param_reg;<br>
>> > // <end>}<br>
>> > .reg .b32 retval0;<br>
>> > call.uni (retval0),<br>
>> > INT_PTX_SREG_TID_X,<br>
>> > (<br>
>> > );<br>
>> > mov.b32 %r2, retval0;<br>
>> ><br>
>> > //{<br>
>> > }// Callseq End 5<br>
>> > mad.lo.s32 %r0, %r0, %r1, %r2;<br>
>> > cvt.rn.f64.s32 %fl4, %r0;<br>
>> > mul.f64 %fl5, %fl4, %fl2;<br>
>> > add.f64 %fl5, %fl5, %fl3;<br>
>> > cvt.rzi.s64.f64 %rl1, %fl5;<br>
>> > shl.b64 %rl1, %rl1, 3;<br>
>> > add.s64 %rl1, %rl0, %rl1;<br>
>> > div.rn.f64 %fl2, %fl3, %fl2;<br>
>> > mul.f64 %fl2, %fl2, 0d400C000000000000;<br>
>> > add.f64 %fl2, %fl2, 0dC004000000000000;<br>
>> > div.rn.f64 %fl1, %fl4, %fl1;<br>
>> > add.f64 %fl1, %fl1, %fl1;<br>
>> > add.f64 %fl3, %fl1, 0dBFF0000000000000;<br>
>> > mov.f64 %fl1, 0d0000000000000000;<br>
>> > mov.f64 %fl5, %fl1;<br>
>> > mov.f64 %fl4, %fl1;<br>
>> > bra.uni BB2_1;<br>
>> > BB2_2:<br>
>> > add.f64 %fl1, %fl1, 0d3FF0000000000000;<br>
>> > sub.f64 %fl6, %fl6, %fl7;<br>
>> > add.f64 %fl6, %fl6, %fl2;<br>
>> > add.f64 %fl5, %fl5, %fl5;<br>
>> > mul.f64 %fl4, %fl5, %fl4;<br>
>> > add.f64 %fl4, %fl4, %fl3;<br>
>> > mov.f64 %fl5, %fl6;<br>
>> > BB2_1:<br>
>> > mul.f64 %fl6, %fl5, %fl5;<br>
>> > mul.f64 %fl7, %fl4, %fl4;<br>
>> > add.f64 %fl8, %fl6, %fl7;<br>
>> > setp.lt.f64 %p0, %fl8, 0d4010000000000000;<br>
>> > setp.lt.f64 %p1, %fl1, %fl0;<br>
>> > and.pred %p0, %p0, %p1;<br>
>> > @!%p0 bra BB2_3;<br>
>> > bra.uni BB2_2;<br>
>> > BB2_3:<br>
>> > div.rn.f64 %fl0, %fl1, %fl0;<br>
>> > st.global.f64 [%rl1], %fl0;<br>
>> > mov.b64 func_retval0, %rl0;<br>
>> > ret;<br>
>> > }<br>
>><br>
><br>
><br>
><br>
> --<br>
> “One of the main causes of the fall of the Roman Empire was that–lacking<br>
> zero–they had no way to indicate successful termination of their C<br>
> programs.”<br>
> (Robert Firth)<br>
><br>
</div></div><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>
><br>
<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>
</div></div></blockquote></div><br><br clear="all"><div><br></div></div></div><span class="HOEnZb"><font color="#888888">-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div>
</font></span></div>
</blockquote></div><br><br clear="all"><div><br></div>-- <br>“One of the main causes of the fall of the Roman Empire was that–lacking zero–they had no way to indicate successful termination of their C programs.”<br>(Robert Firth)
</div>