<div dir="ltr">I've written a compiler that outputs PTX code, the result seems fairly reasonable, but I'm not sure the intrinsics are getting compiled correctly. <br><br>In addition, when I try load the module using CUDA, I get an error: CUDA_ERROR_NO_BINARY_FOR_GPU. I'm running this on a 2012 MBP with a 640M GPU. <br>
<br>PTX Code (for a mandelbrot calculation):<br><br><div>//</div><div>// Generated by LLVM NVPTX Back-End</div><div>//</div><div><br></div><div>.version 3.1</div><div>.target sm_10, texmode_independent</div><div>.address_size 64</div>
<div><br></div><div>.func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X</div><div>(</div><div><br></div><div>)</div><div>;</div><div>.func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y</div><div>(</div><div><br></div><div>
)</div><div>;</div><div>.func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X</div><div>(</div><div><br></div><div>)</div><div>;</div><div>.func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_X</div><div>(</div><div><br></div><div>
)</div><div>;</div><div>.func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_Y</div><div>(</div><div><br></div><div>)</div><div>;</div><div><br></div><div> // .globl examples_2E_mandelbrot_2F_square</div><div>.func (.reg .b64 func_retval0) examples_2E_mandelbrot_2F_square(</div>
<div> .reg .b64 examples_2E_mandelbrot_2F_square_param_0</div><div>)</div><div>{</div><div> .reg .pred %p<396>;</div><div> .reg .s16 %rc<396>;</div><div> .reg .s16 %rs<396>;</div>
<div> .reg .s32 %r<396>;</div><div> .reg .s64 %rl<396>;</div><div> .reg .f32 %f<396>;</div><div> .reg .f64 %fl<396>;</div><div><br></div><div> mov.f64 %fl0, examples_2E_mandelbrot_2F_square_param_0;</div>
<div> mul.f64 %fl0, %fl0, %fl0;</div><div> mov.f64 func_retval0, %fl0;</div><div> ret;</div><div>}</div><div><br></div><div> // .globl examples_2E_mandelbrot_2F_calc_2D_iteration</div>
<div>.func (.reg .b64 func_retval0) examples_2E_mandelbrot_2F_calc_2D_iteration(</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_iteration_param_0,</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_iteration_param_1,</div>
<div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_iteration_param_2,</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_iteration_param_3,</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_iteration_param_4</div>
<div>)</div><div>{</div><div> .reg .pred %p<396>;</div><div> .reg .s16 %rc<396>;</div><div> .reg .s16 %rs<396>;</div><div> .reg .s32 %r<396>;</div><div> .reg .s64 %rl<396>;</div>
<div> .reg .f32 %f<396>;</div><div> .reg .f64 %fl<396>;</div><div><br></div><div> mov.f64 %fl0, examples_2E_mandelbrot_2F_calc_2D_iteration_param_0;</div><div> mov.f64 %fl1, examples_2E_mandelbrot_2F_calc_2D_iteration_param_3;</div>
<div> div.rn.f64 %fl0, %fl0, %fl1;</div><div> mov.f64 %fl2, examples_2E_mandelbrot_2F_calc_2D_iteration_param_1;</div><div> mul.f64 %fl1, %fl0, 0d400C000000000000;</div><div> mov.f64 %fl0, examples_2E_mandelbrot_2F_calc_2D_iteration_param_2;</div>
<div> add.f64 %fl1, %fl1, 0dC004000000000000;</div><div> mov.f64 %fl3, examples_2E_mandelbrot_2F_calc_2D_iteration_param_4;</div><div> div.rn.f64 %fl2, %fl2, %fl3;</div><div> add.f64 %fl2, %fl2, %fl2;</div>
<div> add.f64 %fl3, %fl2, 0dBFF0000000000000;</div><div> mov.f64 %fl2, 0d0000000000000000;</div><div> mov.f64 %fl5, %fl2;</div><div> mov.f64 %fl4, %fl2;</div><div>
bra.uni BB1_1;</div><div>BB1_2:</div><div> add.f64 %fl2, %fl2, 0d3FF0000000000000;</div><div> sub.f64 %fl6, %fl6, %fl7;</div><div> add.f64 %fl6, %fl6, %fl1;</div>
<div> add.f64 %fl5, %fl5, %fl5;</div><div> mul.f64 %fl4, %fl5, %fl4;</div><div> add.f64 %fl4, %fl4, %fl3;</div><div> mov.f64 %fl5, %fl6;</div><div>BB1_1:</div><div>
mul.f64 %fl6, %fl5, %fl5;</div><div> mul.f64 %fl7, %fl4, %fl4;</div><div> add.f64 %fl8, %fl6, %fl7;</div><div> setp.lt.f64 %p0, %fl8, 0d4010000000000000;</div><div>
setp.lt.f64 %p1, %fl2, %fl0;</div><div> and.pred %p0, %p0, %p1;</div><div> @!%p0 bra BB1_3;</div><div> bra.uni BB1_2;</div><div>BB1_3:</div><div> mov.f64 func_retval0, %fl2;</div>
<div> ret;</div><div>}</div><div><br></div><div> // .globl examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx</div><div>.func (.reg .b64 func_retval0) examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx(</div>
<div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0,</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1,</div><div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2,</div>
<div> .reg .b64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3</div><div>)</div><div>{</div><div> .reg .pred %p<396>;</div><div> .reg .s16 %rc<396>;</div><div> .reg .s16 %rs<396>;</div>
<div> .reg .s32 %r<396>;</div><div> .reg .s64 %rl<396>;</div><div> .reg .f32 %f<396>;</div><div> .reg .f64 %fl<396>;</div><div><br></div><div> mov.b64 %rl0, examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0;</div>
<div> mov.f64 %fl2, examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1;</div><div> mov.f64 %fl1, examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2;</div><div> mov.f64 %fl0, examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3;</div>
<div> // Callseq Start 0</div><div> {</div><div> .reg .b32 temp_param_reg;</div><div> // <end>}</div><div> .reg .b32 retval0;</div><div> call.uni (retval0), </div><div> INT_PTX_SREG_CTAID_X, </div>
<div> (</div><div> );</div><div> mov.b32 %r0, retval0;</div><div> </div><div> //{</div><div> }// Callseq End 0</div><div> // Callseq Start 1</div><div> {</div><div> .reg .b32 temp_param_reg;</div>
<div> // <end>}</div><div> .reg .b32 retval0;</div><div> call.uni (retval0), </div><div> INT_PTX_SREG_NTID_X, </div><div> (</div><div> );</div><div> mov.b32 %r1, retval0;</div>
<div> </div><div> //{</div><div> }// Callseq End 1</div><div> // Callseq Start 2</div><div> {</div><div> .reg .b32 temp_param_reg;</div><div> // <end>}</div><div> .reg .b32 retval0;</div>
<div> call.uni (retval0), </div><div> INT_PTX_SREG_TID_X, </div><div> (</div><div> );</div><div> mov.b32 %r2, retval0;</div><div> </div><div> //{</div><div> }// Callseq End 2</div>
<div> mad.lo.s32 %r0, %r0, %r1, %r2;</div><div> cvt.rn.f64.s32 %fl3, %r0;</div><div> // Callseq Start 3</div><div> {</div><div> .reg .b32 temp_param_reg;</div><div> // <end>}</div>
<div> .reg .b32 retval0;</div><div> call.uni (retval0), </div><div> INT_PTX_SREG_CTAID_Y, </div><div> (</div><div> );</div><div> mov.b32 %r0, retval0;</div><div> </div><div>
//{</div><div> }// Callseq End 3</div><div> // Callseq Start 4</div><div> {</div><div> .reg .b32 temp_param_reg;</div><div> // <end>}</div><div> .reg .b32 retval0;</div>
<div> call.uni (retval0), </div><div> INT_PTX_SREG_NTID_Y, </div><div> (</div><div> );</div><div> mov.b32 %r1, retval0;</div><div> </div><div> //{</div><div> }// Callseq End 4</div>
<div> // Callseq Start 5</div><div> {</div><div> .reg .b32 temp_param_reg;</div><div> // <end>}</div><div> .reg .b32 retval0;</div><div> call.uni (retval0), </div><div> INT_PTX_SREG_TID_X, </div>
<div> (</div><div> );</div><div> mov.b32 %r2, retval0;</div><div> </div><div> //{</div><div> }// Callseq End 5</div><div> mad.lo.s32 %r0, %r0, %r1, %r2;</div><div> cvt.rn.f64.s32 %fl4, %r0;</div>
<div> mul.f64 %fl5, %fl4, %fl2;</div><div> add.f64 %fl5, %fl5, %fl3;</div><div> cvt.rzi.s64.f64 %rl1, %fl5;</div><div> shl.b64 %rl1, %rl1, 3;</div><div> add.s64 %rl1, %rl0, %rl1;</div>
<div> div.rn.f64 %fl2, %fl3, %fl2;</div><div> mul.f64 %fl2, %fl2, 0d400C000000000000;</div><div> add.f64 %fl2, %fl2, 0dC004000000000000;</div><div> div.rn.f64 %fl1, %fl4, %fl1;</div>
<div> add.f64 %fl1, %fl1, %fl1;</div><div> add.f64 %fl3, %fl1, 0dBFF0000000000000;</div><div> mov.f64 %fl1, 0d0000000000000000;</div><div> mov.f64 %fl5, %fl1;</div>
<div> mov.f64 %fl4, %fl1;</div><div> bra.uni BB2_1;</div><div>BB2_2:</div><div> add.f64 %fl1, %fl1, 0d3FF0000000000000;</div><div> sub.f64 %fl6, %fl6, %fl7;</div>
<div> add.f64 %fl6, %fl6, %fl2;</div><div> add.f64 %fl5, %fl5, %fl5;</div><div> mul.f64 %fl4, %fl5, %fl4;</div><div> add.f64 %fl4, %fl4, %fl3;</div><div> mov.f64 %fl5, %fl6;</div>
<div>BB2_1:</div><div> mul.f64 %fl6, %fl5, %fl5;</div><div> mul.f64 %fl7, %fl4, %fl4;</div><div> add.f64 %fl8, %fl6, %fl7;</div><div> setp.lt.f64 %p0, %fl8, 0d4010000000000000;</div>
<div> setp.lt.f64 %p1, %fl1, %fl0;</div><div> and.pred %p0, %p0, %p1;</div><div> @!%p0 bra BB2_3;</div><div> bra.uni BB2_2;</div><div>BB2_3:</div><div> div.rn.f64 %fl0, %fl1, %fl0;</div>
<div> st.global.f64 [%rl1], %fl0;</div><div> mov.b64 func_retval0, %rl0;</div><div> ret;</div><div>}</div><div><br></div><div><br></div></div>