[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
Dmitry Mikushin
dmitry at kernelgen.org
Fri Mar 1 11:44:12 PST 2013
Timothy,
Those calls to compute grid intrinsics are definitely wrong. In ptx code they should end up into reading special registers, rather than function calls. Try to take some working example and figure out the LLVM IR differences between it and the result of your compiler.
- D.
----- Original message -----
> 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.
>
> 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.
>
> PTX Code (for a mandelbrot calculation):
>
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.1
> .target sm_10, texmode_independent
> .address_size 64
>
> .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
> (
>
> )
> ;
> .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
> (
>
> )
> ;
> .func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
> (
>
> )
> ;
> .func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_X
> (
>
> )
> ;
> .func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_Y
> (
>
> )
> ;
>
> // .globl examples_2E_mandelbrot_2F_square
> .func (.reg .b64 func_retval0) examples_2E_mandelbrot_2F_square(
> .reg .b64 examples_2E_mandelbrot_2F_square_param_0
> )
> {
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r<396>;
> .reg .s64 %rl<396>;
> .reg .f32 %f<396>;
> .reg .f64 %fl<396>;
>
> mov.f64 %fl0, examples_2E_mandelbrot_2F_square_param_0;
> mul.f64 %fl0, %fl0, %fl0;
> mov.f64 func_retval0, %fl0;
> ret;
> }
>
> // .globl examples_2E_mandelbrot_2F_calc_2D_iteration
> .func (.reg .b64 func_retval0)
> examples_2E_mandelbrot_2F_calc_2D_iteration( .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_0, .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_1, .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_2, .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_3, .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_4 )
> {
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r<396>;
> .reg .s64 %rl<396>;
> .reg .f32 %f<396>;
> .reg .f64 %fl<396>;
>
> mov.f64 %fl0,
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_0; mov.f64
> %fl1, examples_2E_mandelbrot_2F_calc_2D_iteration_param_3;
> div.rn.f64 %fl0, %fl0, %fl1; mov.f64 %fl2,
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_1; mul.f64
> %fl1, %fl0, 0d400C000000000000; mov.f64 %fl0,
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_2; add.f64
> %fl1, %fl1, 0dC004000000000000; mov.f64 %fl3,
> examples_2E_mandelbrot_2F_calc_2D_iteration_param_4; div.rn.f64
> %fl2, %fl2, %fl3; add.f64 %fl2, %fl2, %fl2;
> add.f64 %fl3, %fl2, 0dBFF0000000000000;
> mov.f64 %fl2, 0d0000000000000000;
> mov.f64 %fl5, %fl2;
> mov.f64 %fl4, %fl2;
> bra.uni BB1_1;
> BB1_2:
> add.f64 %fl2, %fl2, 0d3FF0000000000000;
> sub.f64 %fl6, %fl6, %fl7;
> add.f64 %fl6, %fl6, %fl1;
> add.f64 %fl5, %fl5, %fl5;
> mul.f64 %fl4, %fl5, %fl4;
> add.f64 %fl4, %fl4, %fl3;
> mov.f64 %fl5, %fl6;
> BB1_1:
> mul.f64 %fl6, %fl5, %fl5;
> mul.f64 %fl7, %fl4, %fl4;
> add.f64 %fl8, %fl6, %fl7;
> setp.lt.f64 %p0, %fl8, 0d4010000000000000;
> setp.lt.f64 %p1, %fl2, %fl0;
> and.pred %p0, %p0, %p1;
> @!%p0 bra BB1_3;
> bra.uni BB1_2;
> BB1_3:
> mov.f64 func_retval0, %fl2;
> ret;
> }
>
> // .globl
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx .func (.reg .b64
> func_retval0) examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx(
> .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0,
> .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1,
> .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2,
> .reg .b64
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3
> )
> {
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r<396>;
> .reg .s64 %rl<396>;
> .reg .f32 %f<396>;
> .reg .f64 %fl<396>;
>
> mov.b64 %rl0,
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0;
> mov.f64 %fl2,
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1;
> mov.f64 %fl1,
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2;
> mov.f64 %fl0,
> examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3;
> // Callseq Start 0
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_CTAID_X,
> (
> );
> mov.b32 %r0, retval0;
>
> //{
> }// Callseq End 0
> // Callseq Start 1
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_NTID_X,
> (
> );
> mov.b32 %r1, retval0;
>
> //{
> }// Callseq End 1
> // Callseq Start 2
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_TID_X,
> (
> );
> mov.b32 %r2, retval0;
>
> //{
> }// Callseq End 2
> mad.lo.s32 %r0, %r0, %r1, %r2;
> cvt.rn.f64.s32 %fl3, %r0;
> // Callseq Start 3
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_CTAID_Y,
> (
> );
> mov.b32 %r0, retval0;
>
> //{
> }// Callseq End 3
> // Callseq Start 4
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_NTID_Y,
> (
> );
> mov.b32 %r1, retval0;
>
> //{
> }// Callseq End 4
> // Callseq Start 5
> {
> .reg .b32 temp_param_reg;
> // <end>}
> .reg .b32 retval0;
> call.uni (retval0),
> INT_PTX_SREG_TID_X,
> (
> );
> mov.b32 %r2, retval0;
>
> //{
> }// Callseq End 5
> mad.lo.s32 %r0, %r0, %r1, %r2;
> cvt.rn.f64.s32 %fl4, %r0;
> mul.f64 %fl5, %fl4, %fl2;
> add.f64 %fl5, %fl5, %fl3;
> cvt.rzi.s64.f64 %rl1, %fl5;
> shl.b64 %rl1, %rl1, 3;
> add.s64 %rl1, %rl0, %rl1;
> div.rn.f64 %fl2, %fl3, %fl2;
> mul.f64 %fl2, %fl2, 0d400C000000000000;
> add.f64 %fl2, %fl2, 0dC004000000000000;
> div.rn.f64 %fl1, %fl4, %fl1;
> add.f64 %fl1, %fl1, %fl1;
> add.f64 %fl3, %fl1, 0dBFF0000000000000;
> mov.f64 %fl1, 0d0000000000000000;
> mov.f64 %fl5, %fl1;
> mov.f64 %fl4, %fl1;
> bra.uni BB2_1;
> BB2_2:
> add.f64 %fl1, %fl1, 0d3FF0000000000000;
> sub.f64 %fl6, %fl6, %fl7;
> add.f64 %fl6, %fl6, %fl2;
> add.f64 %fl5, %fl5, %fl5;
> mul.f64 %fl4, %fl5, %fl4;
> add.f64 %fl4, %fl4, %fl3;
> mov.f64 %fl5, %fl6;
> BB2_1:
> mul.f64 %fl6, %fl5, %fl5;
> mul.f64 %fl7, %fl4, %fl4;
> add.f64 %fl8, %fl6, %fl7;
> setp.lt.f64 %p0, %fl8, 0d4010000000000000;
> setp.lt.f64 %p1, %fl1, %fl0;
> and.pred %p0, %p0, %p1;
> @!%p0 bra BB2_3;
> bra.uni BB2_2;
> BB2_3:
> div.rn.f64 %fl0, %fl1, %fl0;
> st.global.f64 [%rl1], %fl0;
> mov.b64 func_retval0, %rl0;
> ret;
> }
More information about the llvm-dev
mailing list