321 lines
11 KiB
Plaintext
321 lines
11 KiB
Plaintext
//
|
|
// Generated by LLVM NVPTX Back-End
|
|
//
|
|
|
|
.version 3.1
|
|
.target sm_35, texmode_independent
|
|
.address_size 64
|
|
|
|
// .globl __aos_to_soa4_float1
|
|
// @__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
|
|
)
|
|
{
|
|
.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.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];
|
|
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];
|
|
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.f32 %f0, [__soa_to_aos4_float1_param_0];
|
|
ld.param.u64 %rl1, [__soa_to_aos4_float1_param_5];
|
|
ld.param.f32 %f1, [__soa_to_aos4_float1_param_1];
|
|
ld.param.u64 %rl2, [__soa_to_aos4_float1_param_6];
|
|
ld.param.f32 %f2, [__soa_to_aos4_float1_param_2];
|
|
ld.param.u64 %rl3, [__soa_to_aos4_float1_param_7];
|
|
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 mandelbrot_scanline
|
|
.entry mandelbrot_scanline(
|
|
.param .f32 mandelbrot_scanline_param_0,
|
|
.param .f32 mandelbrot_scanline_param_1,
|
|
.param .f32 mandelbrot_scanline_param_2,
|
|
.param .f32 mandelbrot_scanline_param_3,
|
|
.param .u32 mandelbrot_scanline_param_4,
|
|
.param .u32 mandelbrot_scanline_param_5,
|
|
.param .u32 mandelbrot_scanline_param_6,
|
|
.param .u32 mandelbrot_scanline_param_7,
|
|
.param .u32 mandelbrot_scanline_param_8,
|
|
.param .u64 .ptr .align 4 mandelbrot_scanline_param_9
|
|
) // @mandelbrot_scanline
|
|
{
|
|
.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.u32 %r6, [mandelbrot_scanline_param_5];
|
|
ld.param.u32 %r5, [mandelbrot_scanline_param_7];
|
|
mov.u32 %r7, %ctaid.y;
|
|
mul.lo.s32 %r0, %r7, %r5;
|
|
mad.lo.s32 %r1, %r7, %r5, %r5;
|
|
setp.lt.s32 %p0, %r1, %r6;
|
|
selp.b32 %r1, %r1, %r6, %p0;
|
|
setp.ge.s32 %p0, %r0, %r1;
|
|
@%p0 bra BB4_13;
|
|
// BB#1: // %for_test28.preheader.lr.ph
|
|
ld.param.f32 %f0, [mandelbrot_scanline_param_0];
|
|
ld.param.f32 %f1, [mandelbrot_scanline_param_1];
|
|
ld.param.f32 %f2, [mandelbrot_scanline_param_2];
|
|
ld.param.f32 %f3, [mandelbrot_scanline_param_3];
|
|
ld.param.u32 %r1, [mandelbrot_scanline_param_4];
|
|
ld.param.u32 %r4, [mandelbrot_scanline_param_6];
|
|
ld.param.u32 %r2, [mandelbrot_scanline_param_8];
|
|
ld.param.u64 %rl0, [mandelbrot_scanline_param_9];
|
|
mov.u32 %r8, %ctaid.x;
|
|
mul.lo.s32 %r3, %r8, %r4;
|
|
mad.lo.s32 %r4, %r8, %r4, %r4;
|
|
setp.lt.s32 %p0, %r4, %r1;
|
|
selp.b32 %r4, %r4, %r1, %p0;
|
|
setp.gt.s32 %p0, %r2, 0;
|
|
not.b32 %r6, %r6;
|
|
add.s32 %r7, %r7, 1;
|
|
mul.lo.s32 %r5, %r7, %r5;
|
|
not.b32 %r5, %r5;
|
|
setp.gt.s32 %p1, %r6, %r5;
|
|
selp.b32 %r5, %r6, %r5, %p1;
|
|
not.b32 %r5, %r5;
|
|
BB4_2: // %for_test28.preheader
|
|
// =>This Loop Header: Depth=1
|
|
// Child Loop BB4_15 Depth 2
|
|
// Child Loop BB4_8 Depth 2
|
|
// Child Loop BB4_11 Depth 3
|
|
setp.ge.s32 %p1, %r3, %r4;
|
|
@%p1 bra BB4_12;
|
|
// BB#3: // %for_loop30.lr.ph
|
|
// in Loop: Header=BB4_2 Depth=1
|
|
mul.lo.s32 %r6, %r0, %r1;
|
|
mov.u32 %r7, %r3;
|
|
@%p0 bra BB4_4;
|
|
bra.uni BB4_15;
|
|
BB4_4: // in Loop: Header=BB4_2 Depth=1
|
|
cvt.rn.f32.s32 %f4, %r0;
|
|
fma.rn.f32 %f4, %f4, %f3, %f2;
|
|
mov.u32 %r7, %r3;
|
|
BB4_8: // %for_loop.i.lr.ph.us
|
|
// Parent Loop BB4_2 Depth=1
|
|
// => This Loop Header: Depth=2
|
|
// Child Loop BB4_11 Depth 3
|
|
mov.u32 %r9, %tid.x;
|
|
mov.u32 %r8, WARP_SZ;
|
|
add.s32 %r10, %r8, -1;
|
|
and.b32 %r10, %r10, %r9;
|
|
add.s32 %r11, %r10, %r7;
|
|
cvt.rn.f32.s32 %f5, %r11;
|
|
fma.rn.f32 %f5, %f5, %f1, %f0;
|
|
mov.u32 %r10, 0;
|
|
mov.pred %p1, 0;
|
|
mov.pred %p3, -1;
|
|
mov.pred %p4, %p0;
|
|
mov.pred %p2, %p1;
|
|
mov.f32 %f7, %f5;
|
|
mov.f32 %f6, %f4;
|
|
BB4_11: // %for_loop.i.us
|
|
// Parent Loop BB4_2 Depth=1
|
|
// Parent Loop BB4_8 Depth=2
|
|
// => This Inner Loop Header: Depth=3
|
|
and.pred %p4, %p3, %p4;
|
|
mul.f32 %f8, %f7, %f7;
|
|
fma.rn.f32 %f9, %f6, %f6, %f8;
|
|
setp.gtu.f32 %p3, %f9, 0f40800000;
|
|
and.pred %p3, %p4, %p3;
|
|
or.pred %p2, %p3, %p2;
|
|
xor.pred %p5, %p2, %p4;
|
|
mov.pred %p3, %p1;
|
|
@!%p5 bra BB4_10;
|
|
bra.uni BB4_9;
|
|
BB4_9: // %not_all_continued_or_breaked.i.us
|
|
// in Loop: Header=BB4_11 Depth=3
|
|
mul.f32 %f9, %f6, %f6;
|
|
not.pred %p3, %p2;
|
|
and.pred %p3, %p4, %p3;
|
|
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;
|
|
BB4_10: // %for_step.i.us
|
|
// in Loop: Header=BB4_11 Depth=3
|
|
add.s32 %r12, %r10, 1;
|
|
selp.b32 %r10, %r12, %r10, %p3;
|
|
setp.lt.s32 %p4, %r10, %r2;
|
|
and.pred %p5, %p3, %p4;
|
|
@%p5 bra BB4_11;
|
|
// BB#5: // %mandel___vyfvyfvyi.exit.us
|
|
// in Loop: Header=BB4_8 Depth=2
|
|
setp.ge.s32 %p1, %r11, %r4;
|
|
@%p1 bra BB4_7;
|
|
// BB#6: // %if_then.us
|
|
// in Loop: Header=BB4_8 Depth=2
|
|
add.s32 %r11, %r8, 1073741823;
|
|
and.b32 %r9, %r11, %r9;
|
|
add.s32 %r11, %r7, %r6;
|
|
add.s32 %r9, %r11, %r9;
|
|
shl.b32 %r9, %r9, 2;
|
|
cvt.s64.s32 %rl1, %r9;
|
|
add.s64 %rl1, %rl1, %rl0;
|
|
st.u32 [%rl1], %r10;
|
|
BB4_7: // %if_exit.us
|
|
// in Loop: Header=BB4_8 Depth=2
|
|
add.s32 %r7, %r8, %r7;
|
|
setp.lt.s32 %p1, %r7, %r4;
|
|
@%p1 bra BB4_8;
|
|
bra.uni BB4_12;
|
|
BB4_15: // %mandel___vyfvyfvyi.exit
|
|
// Parent Loop BB4_2 Depth=1
|
|
// => This Inner Loop Header: Depth=2
|
|
mov.u32 %r9, %tid.x;
|
|
mov.u32 %r8, WARP_SZ;
|
|
add.s32 %r10, %r8, -1;
|
|
and.b32 %r10, %r10, %r9;
|
|
add.s32 %r10, %r10, %r7;
|
|
setp.lt.s32 %p1, %r10, %r4;
|
|
@%p1 bra BB4_16;
|
|
bra.uni BB4_14;
|
|
BB4_16: // %if_then
|
|
// in Loop: Header=BB4_15 Depth=2
|
|
add.s32 %r10, %r8, 1073741823;
|
|
and.b32 %r9, %r10, %r9;
|
|
add.s32 %r10, %r7, %r6;
|
|
add.s32 %r9, %r10, %r9;
|
|
shl.b32 %r9, %r9, 2;
|
|
cvt.s64.s32 %rl1, %r9;
|
|
add.s64 %rl1, %rl1, %rl0;
|
|
mov.u32 %r9, 0;
|
|
st.u32 [%rl1], %r9;
|
|
BB4_14: // %if_exit
|
|
// in Loop: Header=BB4_15 Depth=2
|
|
add.s32 %r7, %r8, %r7;
|
|
setp.lt.s32 %p1, %r7, %r4;
|
|
@%p1 bra BB4_15;
|
|
BB4_12: // %for_exit31
|
|
// in Loop: Header=BB4_2 Depth=1
|
|
add.s32 %r0, %r0, 1;
|
|
setp.eq.s32 %p1, %r0, %r5;
|
|
@%p1 bra BB4_13;
|
|
bra.uni BB4_2;
|
|
BB4_13: // %for_exit
|
|
ret;
|
|
}
|
|
|