// // Generated by LLVM NVPTX Back-End // .version 3.1 .target sm_20, texmode_independent .address_size 64 // .globl __land_id // @__land_id .func (.param .b32 func_retval0) __land_id( ) { .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>; // BB#0: mov.u32 %r0, %laneid; st.param.b32 [func_retval0+0], %r0; ret; } // .globl __vselect_i8 .func (.param .align 1 .b8 func_retval0[1]) __vselect_i8( .param .align 1 .b8 __vselect_i8_param_0[1], .param .align 1 .b8 __vselect_i8_param_1[1], .param .align 4 .b8 __vselect_i8_param_2[4] ) // @__vselect_i8 { .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>; // BB#0: ld.param.u32 %r0, [__vselect_i8_param_2]; setp.eq.s32 %p0, %r0, 0; ld.param.u8 %rc0, [__vselect_i8_param_0]; ld.param.u8 %rc1, [__vselect_i8_param_1]; selp.b16 %rc0, %rc0, %rc1, %p0; st.param.b8 [func_retval0+0], %rc0; ret; } // .globl __vselect_i16 .func (.param .align 2 .b8 func_retval0[2]) __vselect_i16( .param .align 2 .b8 __vselect_i16_param_0[2], .param .align 2 .b8 __vselect_i16_param_1[2], .param .align 4 .b8 __vselect_i16_param_2[4] ) // @__vselect_i16 { .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>; // BB#0: ld.param.u32 %r0, [__vselect_i16_param_2]; setp.eq.s32 %p0, %r0, 0; ld.param.u16 %rs0, [__vselect_i16_param_0]; ld.param.u16 %rs1, [__vselect_i16_param_1]; selp.b16 %rs0, %rs0, %rs1, %p0; st.param.b16 [func_retval0+0], %rs0; ret; } // .globl __vselect_i64 .func (.param .align 8 .b8 func_retval0[8]) __vselect_i64( .param .align 8 .b8 __vselect_i64_param_0[8], .param .align 8 .b8 __vselect_i64_param_1[8], .param .align 4 .b8 __vselect_i64_param_2[4] ) // @__vselect_i64 { .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>; // BB#0: ld.param.u32 %r0, [__vselect_i64_param_2]; setp.eq.s32 %p0, %r0, 0; ld.param.u64 %rl0, [__vselect_i64_param_0]; ld.param.u64 %rl1, [__vselect_i64_param_1]; selp.b64 %rl0, %rl0, %rl1, %p0; st.param.b64 [func_retval0+0], %rl0; ret; } // .globl __aos_to_soa4_float1 .func __aos_to_soa4_float1( .param .align 4 .b8 __aos_to_soa4_float1_param_0[4], .param .align 4 .b8 __aos_to_soa4_float1_param_1[4], .param .align 4 .b8 __aos_to_soa4_float1_param_2[4], .param .align 4 .b8 __aos_to_soa4_float1_param_3[4], .param .b64 __aos_to_soa4_float1_param_4, .param .b64 __aos_to_soa4_float1_param_5, .param .b64 __aos_to_soa4_float1_param_6, .param .b64 __aos_to_soa4_float1_param_7 ) // @__aos_to_soa4_float1 { .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>; // BB#0: ld.param.u64 %rl0, [__aos_to_soa4_float1_param_4]; ld.param.u64 %rl1, [__aos_to_soa4_float1_param_5]; ld.param.u64 %rl2, [__aos_to_soa4_float1_param_6]; ld.param.u64 %rl3, [__aos_to_soa4_float1_param_7]; ld.param.f32 %f0, [__aos_to_soa4_float1_param_0]; ld.param.f32 %f1, [__aos_to_soa4_float1_param_1]; ld.param.f32 %f2, [__aos_to_soa4_float1_param_2]; ld.param.f32 %f3, [__aos_to_soa4_float1_param_3]; st.f32 [%rl0], %f0; st.f32 [%rl1], %f1; st.f32 [%rl2], %f2; st.f32 [%rl3], %f3; ret; } // .globl __soa_to_aos4_float1 .func __soa_to_aos4_float1( .param .align 4 .b8 __soa_to_aos4_float1_param_0[4], .param .align 4 .b8 __soa_to_aos4_float1_param_1[4], .param .align 4 .b8 __soa_to_aos4_float1_param_2[4], .param .align 4 .b8 __soa_to_aos4_float1_param_3[4], .param .b64 __soa_to_aos4_float1_param_4, .param .b64 __soa_to_aos4_float1_param_5, .param .b64 __soa_to_aos4_float1_param_6, .param .b64 __soa_to_aos4_float1_param_7 ) // @__soa_to_aos4_float1 { .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>; // BB#0: ld.param.u64 %rl0, [__soa_to_aos4_float1_param_4]; ld.param.u64 %rl1, [__soa_to_aos4_float1_param_5]; ld.param.u64 %rl2, [__soa_to_aos4_float1_param_6]; ld.param.u64 %rl3, [__soa_to_aos4_float1_param_7]; ld.param.f32 %f0, [__soa_to_aos4_float1_param_0]; ld.param.f32 %f1, [__soa_to_aos4_float1_param_1]; ld.param.f32 %f2, [__soa_to_aos4_float1_param_2]; ld.param.f32 %f3, [__soa_to_aos4_float1_param_3]; st.f32 [%rl0], %f0; st.f32 [%rl1], %f1; st.f32 [%rl2], %f2; st.f32 [%rl3], %f3; ret; } // .globl __aos_to_soa3_float1 .func __aos_to_soa3_float1( .param .align 4 .b8 __aos_to_soa3_float1_param_0[4], .param .align 4 .b8 __aos_to_soa3_float1_param_1[4], .param .align 4 .b8 __aos_to_soa3_float1_param_2[4], .param .b64 __aos_to_soa3_float1_param_3, .param .b64 __aos_to_soa3_float1_param_4, .param .b64 __aos_to_soa3_float1_param_5 ) // @__aos_to_soa3_float1 { .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>; // BB#0: ld.param.u64 %rl0, [__aos_to_soa3_float1_param_3]; ld.param.u64 %rl1, [__aos_to_soa3_float1_param_4]; ld.param.u64 %rl2, [__aos_to_soa3_float1_param_5]; ld.param.f32 %f0, [__aos_to_soa3_float1_param_0]; ld.param.f32 %f1, [__aos_to_soa3_float1_param_1]; ld.param.f32 %f2, [__aos_to_soa3_float1_param_2]; st.f32 [%rl0], %f0; st.f32 [%rl1], %f1; st.f32 [%rl2], %f2; ret; } // .globl __soa_to_aos3_float1 .func __soa_to_aos3_float1( .param .align 4 .b8 __soa_to_aos3_float1_param_0[4], .param .align 4 .b8 __soa_to_aos3_float1_param_1[4], .param .align 4 .b8 __soa_to_aos3_float1_param_2[4], .param .b64 __soa_to_aos3_float1_param_3, .param .b64 __soa_to_aos3_float1_param_4, .param .b64 __soa_to_aos3_float1_param_5 ) // @__soa_to_aos3_float1 { .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>; // BB#0: ld.param.u64 %rl0, [__soa_to_aos3_float1_param_3]; ld.param.u64 %rl1, [__soa_to_aos3_float1_param_4]; ld.param.u64 %rl2, [__soa_to_aos3_float1_param_5]; ld.param.f32 %f0, [__soa_to_aos3_float1_param_0]; ld.param.f32 %f1, [__soa_to_aos3_float1_param_1]; ld.param.f32 %f2, [__soa_to_aos3_float1_param_2]; st.f32 [%rl0], %f0; st.f32 [%rl1], %f1; st.f32 [%rl2], %f2; ret; } // .globl __rsqrt_varying_double .func (.param .align 8 .b8 func_retval0[8]) __rsqrt_varying_double( .param .align 8 .b8 __rsqrt_varying_double_param_0[8] ) // @__rsqrt_varying_double { .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>; // BB#0: ld.param.f64 %fl0, [__rsqrt_varying_double_param_0]; rsqrt.approx.f64 %fl0, %fl0; st.param.f64 [func_retval0+0], %fl0; ret; } // .globl mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E_ .func mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E_( .param .b64 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_0, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_1, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_2, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_3, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_4, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_5, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_6, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_7, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_8, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_9, .param .b32 mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_10 ) // @mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E_ { .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>; // BB#0: // %allocas ld.param.u64 %rl0, [mandelbrot_scanline___unfunfunfunfuniuniuniuniuniun_3C_uni_3E__param_0]; ld.f32 %f0, [%rl0]; ld.f32 %f1, [%rl0+4]; ld.f32 %f2, [%rl0+8]; ld.f32 %f3, [%rl0+12]; ld.u32 %r4, [%rl0+16]; ld.u32 %r8, [%rl0+20]; ld.u32 %r9, [%rl0+24]; ld.u32 %r7, [%rl0+28]; ld.u32 %r0, [%rl0+32]; ld.u32 %r5, [%rl0+48]; setp.lt.s32 %p0, %r5, 0; mov.u32 %r11, %ctaid.x; mul.lo.s32 %r1, %r11, %r9; mad.lo.s32 %r2, %r11, %r9, %r9; setp.lt.s32 %p1, %r2, %r4; selp.b32 %r2, %r2, %r4, %p1; mov.u32 %r10, %ctaid.y; mul.lo.s32 %r3, %r10, %r7; mad.lo.s32 %r6, %r10, %r7, %r7; setp.lt.s32 %p1, %r6, %r8; selp.b32 %r6, %r6, %r8, %p1; @%p0 bra BB9_4; // BB#1: // %for_test101.preheader setp.ge.s32 %p0, %r3, %r6; @%p0 bra BB9_31; // BB#2: // %for_test112.preheader.lr.ph setp.gt.s32 %p0, %r0, 0; selp.b32 %r4, -1, 0, %p0; and.b32 %r6, %r4, %r5; not.b32 %r8, %r8; add.s32 %r9, %r10, 1; mul.lo.s32 %r7, %r7, %r9; not.b32 %r7, %r7; setp.gt.s32 %p0, %r8, %r7; selp.b32 %r7, %r8, %r7, %p0; not.b32 %r7, %r7; BB9_3: // %for_test112.preheader // =>This Loop Header: Depth=1 // Child Loop BB9_29 Depth 2 // Child Loop BB9_28 Depth 2 // Child Loop BB9_23 Depth 3 setp.ge.s32 %p0, %r1, %r2; @%p0 bra BB9_30; // BB#21: // %for_loop114.lr.ph // in Loop: Header=BB9_3 Depth=1 setp.lt.s32 %p0, %r6, 0; mov.u32 %r8, %r1; @%p0 bra BB9_22; bra.uni BB9_29; BB9_22: // in Loop: Header=BB9_3 Depth=1 cvt.rn.f32.s32 %f4, %r3; fma.rn.f32 %f4, %f3, %f4, %f2; mov.u32 %r8, %r1; BB9_28: // %for_loop.i.lr.ph.us // Parent Loop BB9_3 Depth=1 // => This Loop Header: Depth=2 // Child Loop BB9_23 Depth 3 mov.u32 %r9, %laneid; add.s32 %r9, %r9, %r8; cvt.rn.f32.s32 %f5, %r9; fma.rn.f32 %f5, %f1, %f5, %f0; mov.u32 %r9, 0; mov.u32 %r12, %r4; mov.u32 %r10, %r9; mov.u32 %r11, %r9; mov.f32 %f7, %f5; mov.f32 %f6, %f4; BB9_23: // %for_loop.i.us // Parent Loop BB9_3 Depth=1 // Parent Loop BB9_28 Depth=2 // => This Inner Loop Header: Depth=3 and.b32 %r13, %r12, %r5; mul.f32 %f8, %f7, %f7; fma.rn.f32 %f9, %f6, %f6, %f8; setp.gtu.f32 %p0, %f9, 0f40800000; selp.b32 %r14, %r12, 0, %p0; or.b32 %r10, %r14, %r10; and.b32 %r14, %r10, %r5; shr.u32 %r14, %r14, 31; shr.u32 %r13, %r13, 31; setp.eq.s32 %p0, %r14, %r13; @%p0 bra BB9_24; bra.uni BB9_25; BB9_24: // in Loop: Header=BB9_23 Depth=3 mov.u32 %r12, %r9; bra.uni BB9_26; BB9_25: // %not_all_continued_or_breaked.i.us // in Loop: Header=BB9_23 Depth=3 mul.f32 %f9, %f6, %f6; not.b32 %r13, %r10; and.b32 %r12, %r12, %r13; sub.f32 %f8, %f8, %f9; add.f32 %f8, %f5, %f8; add.f32 %f7, %f7, %f7; fma.rn.f32 %f6, %f6, %f7, %f4; mov.f32 %f7, %f8; BB9_26: // %for_step.i.us // in Loop: Header=BB9_23 Depth=3 setp.ne.s32 %p0, %r12, 0; selp.u32 %r13, 1, 0, %p0; add.s32 %r11, %r11, %r13; setp.lt.s32 %p0, %r11, %r0; selp.b32 %r12, %r12, 0, %p0; and.b32 %r13, %r12, %r5; setp.lt.s32 %p0, %r13, 0; @%p0 bra BB9_23; // BB#27: // %if_exit156.us // in Loop: Header=BB9_28 Depth=2 add.s32 %r8, %r8, 32; setp.lt.s32 %p0, %r8, %r2; @%p0 bra BB9_28; bra.uni BB9_30; BB9_29: // %if_exit156 // Parent Loop BB9_3 Depth=1 // => This Inner Loop Header: Depth=2 add.s32 %r8, %r8, 32; setp.lt.s32 %p0, %r8, %r2; @%p0 bra BB9_29; BB9_30: // %for_exit115 // in Loop: Header=BB9_3 Depth=1 add.s32 %r3, %r3, 1; setp.eq.s32 %p0, %r3, %r7; @%p0 bra BB9_31; bra.uni BB9_3; BB9_4: // %for_test.preheader setp.ge.s32 %p0, %r3, %r6; @%p0 bra BB9_31; // BB#5: // %for_test40.preheader.lr.ph ld.u64 %rl0, [%rl0+40]; setp.gt.s32 %p0, %r0, 0; selp.b32 %r5, -1, 0, %p0; not.b32 %r6, %r8; add.s32 %r8, %r10, 1; mul.lo.s32 %r8, %r7, %r8; not.b32 %r8, %r8; setp.gt.s32 %p0, %r6, %r8; selp.b32 %r6, %r6, %r8, %p0; not.b32 %r6, %r6; mul.lo.s32 %r7, %r10, %r7; mul.lo.s32 %r7, %r7, %r4; mad.lo.s32 %r7, %r11, %r9, %r7; mov.u32 %r13, 0; BB9_6: // %for_test40.preheader // =>This Loop Header: Depth=1 // Child Loop BB9_19 Depth 2 // Child Loop BB9_12 Depth 2 // Child Loop BB9_13 Depth 3 setp.ge.s32 %p0, %r1, %r2; @%p0 bra BB9_17; // BB#7: // %for_loop42.lr.ph // in Loop: Header=BB9_6 Depth=1 setp.lt.s32 %p0, %r5, 0; mov.u32 %r8, %r7; mov.u32 %r9, %r1; @%p0 bra BB9_8; bra.uni BB9_19; BB9_8: // in Loop: Header=BB9_6 Depth=1 cvt.rn.f32.s32 %f4, %r3; mul.lo.s32 %r8, %r3, %r4; fma.rn.f32 %f4, %f3, %f4, %f2; mov.u32 %r9, %r1; BB9_12: // %for_loop.i206.lr.ph.us // Parent Loop BB9_6 Depth=1 // => This Loop Header: Depth=2 // Child Loop BB9_13 Depth 3 mov.u32 %r10, %laneid; add.s32 %r12, %r10, %r9; cvt.rn.f32.s32 %f5, %r12; fma.rn.f32 %f5, %f1, %f5, %f0; mov.u32 %r15, %r5; mov.u32 %r14, %r13; mov.u32 %r11, %r13; mov.f32 %f7, %f5; mov.f32 %f6, %f4; BB9_13: // %for_loop.i206.us // Parent Loop BB9_6 Depth=1 // Parent Loop BB9_12 Depth=2 // => This Inner Loop Header: Depth=3 mul.f32 %f8, %f7, %f7; fma.rn.f32 %f9, %f6, %f6, %f8; setp.gtu.f32 %p0, %f9, 0f40800000; selp.b32 %r16, %r15, 0, %p0; or.b32 %r14, %r16, %r14; shr.u32 %r16, %r14, 31; shr.u32 %r17, %r15, 31; setp.eq.s32 %p0, %r16, %r17; @%p0 bra BB9_14; bra.uni BB9_15; BB9_14: // in Loop: Header=BB9_13 Depth=3 mov.u32 %r15, %r13; bra.uni BB9_16; BB9_15: // %not_all_continued_or_breaked.i220.us // in Loop: Header=BB9_13 Depth=3 mul.f32 %f9, %f6, %f6; not.b32 %r16, %r14; and.b32 %r15, %r15, %r16; sub.f32 %f8, %f8, %f9; add.f32 %f8, %f5, %f8; add.f32 %f7, %f7, %f7; fma.rn.f32 %f6, %f6, %f7, %f4; mov.f32 %f7, %f8; BB9_16: // %for_step.i189.us // in Loop: Header=BB9_13 Depth=3 setp.ne.s32 %p0, %r15, 0; selp.u32 %r16, 1, 0, %p0; add.s32 %r11, %r11, %r16; setp.lt.s32 %p0, %r11, %r0; selp.b32 %r15, %r15, 0, %p0; setp.lt.s32 %p0, %r15, 0; @%p0 bra BB9_13; // BB#9: // %mandel___vyfvyfvyi.exit221.us // in Loop: Header=BB9_12 Depth=2 setp.ge.s32 %p0, %r12, %r2; @%p0 bra BB9_11; // BB#10: // %if_then.us // in Loop: Header=BB9_12 Depth=2 add.s32 %r12, %r9, %r8; add.s32 %r10, %r12, %r10; shl.b32 %r10, %r10, 2; cvt.s64.s32 %rl1, %r10; add.s64 %rl1, %rl1, %rl0; st.u32 [%rl1], %r11; BB9_11: // %if_exit.us // in Loop: Header=BB9_12 Depth=2 add.s32 %r9, %r9, 32; setp.lt.s32 %p0, %r9, %r2; @%p0 bra BB9_12; bra.uni BB9_17; BB9_19: // %mandel___vyfvyfvyi.exit221 // Parent Loop BB9_6 Depth=1 // => This Inner Loop Header: Depth=2 mov.u32 %r10, %laneid; add.s32 %r11, %r10, %r9; setp.lt.s32 %p0, %r11, %r2; @%p0 bra BB9_20; bra.uni BB9_18; BB9_20: // %if_then // in Loop: Header=BB9_19 Depth=2 add.s32 %r10, %r10, %r8; shl.b32 %r10, %r10, 2; cvt.s64.s32 %rl1, %r10; add.s64 %rl1, %rl1, %rl0; mov.u32 %r10, 0; st.u32 [%rl1], %r10; BB9_18: // %if_exit // in Loop: Header=BB9_19 Depth=2 add.s32 %r9, %r9, 32; add.s32 %r8, %r8, 32; setp.lt.s32 %p0, %r9, %r2; @%p0 bra BB9_19; BB9_17: // %for_exit43 // in Loop: Header=BB9_6 Depth=1 add.s32 %r3, %r3, 1; add.s32 %r7, %r7, %r4; setp.eq.s32 %p0, %r3, %r6; @%p0 bra BB9_31; bra.uni BB9_6; BB9_31: // %for_exit ret; }