2399 lines
57 KiB
Text
2399 lines
57 KiB
Text
//
|
|
// Generated by NVIDIA NVVM Compiler
|
|
//
|
|
// Compiler Build ID: CL-36037853
|
|
// Cuda compilation tools, release 12.9, V12.9.86
|
|
// Based on NVVM 7.0.1
|
|
//
|
|
|
|
.version 8.8
|
|
.target sm_75
|
|
.address_size 64
|
|
|
|
// .globl lif_noise
|
|
// _ZZ9lif_sweepE8per_warp has been demoted
|
|
// _ZZ11lif_compactE8per_warp has been demoted
|
|
// _ZZ15lif_bucket_scanE7partial has been demoted
|
|
.extern .shared .align 16 .b8 partial[];
|
|
|
|
.visible .entry lif_noise(
|
|
.param .u64 lif_noise_param_0,
|
|
.param .u64 lif_noise_param_1,
|
|
.param .u64 lif_noise_param_2,
|
|
.param .u32 lif_noise_param_3,
|
|
.param .f64 lif_noise_param_4
|
|
)
|
|
{
|
|
.reg .pred %p<7>;
|
|
.reg .f32 %f<18>;
|
|
.reg .b32 %r<17>;
|
|
.reg .f64 %fd<12>;
|
|
.reg .b64 %rd<12>;
|
|
|
|
|
|
ld.param.u64 %rd2, [lif_noise_param_0];
|
|
ld.param.u64 %rd3, [lif_noise_param_1];
|
|
ld.param.u64 %rd4, [lif_noise_param_2];
|
|
ld.param.u32 %r9, [lif_noise_param_3];
|
|
ld.param.f64 %fd1, [lif_noise_param_4];
|
|
mov.u32 %r10, %ctaid.x;
|
|
mov.u32 %r11, %ntid.x;
|
|
mov.u32 %r12, %tid.x;
|
|
mad.lo.s32 %r1, %r10, %r11, %r12;
|
|
setp.ge.u32 %p1, %r1, %r9;
|
|
@%p1 bra $L__BB0_8;
|
|
|
|
cvta.to.global.u64 %rd5, %rd2;
|
|
cvta.to.global.u64 %rd6, %rd3;
|
|
mul.wide.u32 %rd7, %r1, 4;
|
|
add.s64 %rd8, %rd6, %rd7;
|
|
cvta.to.global.u64 %rd9, %rd4;
|
|
add.s64 %rd10, %rd9, %rd7;
|
|
ld.global.u32 %r13, [%rd8];
|
|
mul.wide.u32 %rd11, %r13, 4;
|
|
add.s64 %rd1, %rd5, %rd11;
|
|
ld.global.f32 %f17, [%rd1];
|
|
ld.global.u32 %r2, [%rd10];
|
|
setp.eq.s32 %p2, %r2, 0;
|
|
@%p2 bra $L__BB0_7;
|
|
|
|
add.s32 %r14, %r2, -1;
|
|
and.b32 %r16, %r2, 3;
|
|
setp.lt.u32 %p3, %r14, 3;
|
|
@%p3 bra $L__BB0_5;
|
|
|
|
sub.s32 %r15, %r2, %r16;
|
|
|
|
$L__BB0_4:
|
|
cvt.f64.f32 %fd2, %f17;
|
|
add.rn.f64 %fd3, %fd2, %fd1;
|
|
cvt.rn.f32.f64 %f10, %fd3;
|
|
cvt.f64.f32 %fd4, %f10;
|
|
add.rn.f64 %fd5, %fd4, %fd1;
|
|
cvt.rn.f32.f64 %f11, %fd5;
|
|
cvt.f64.f32 %fd6, %f11;
|
|
add.rn.f64 %fd7, %fd6, %fd1;
|
|
cvt.rn.f32.f64 %f12, %fd7;
|
|
cvt.f64.f32 %fd8, %f12;
|
|
add.rn.f64 %fd9, %fd8, %fd1;
|
|
cvt.rn.f32.f64 %f17, %fd9;
|
|
add.s32 %r15, %r15, -4;
|
|
setp.ne.s32 %p4, %r15, 0;
|
|
@%p4 bra $L__BB0_4;
|
|
|
|
$L__BB0_5:
|
|
setp.eq.s32 %p5, %r16, 0;
|
|
@%p5 bra $L__BB0_7;
|
|
|
|
$L__BB0_6:
|
|
.pragma "nounroll";
|
|
cvt.f64.f32 %fd10, %f17;
|
|
add.rn.f64 %fd11, %fd10, %fd1;
|
|
cvt.rn.f32.f64 %f17, %fd11;
|
|
add.s32 %r16, %r16, -1;
|
|
setp.ne.s32 %p6, %r16, 0;
|
|
@%p6 bra $L__BB0_6;
|
|
|
|
$L__BB0_7:
|
|
st.global.f32 [%rd1], %f17;
|
|
|
|
$L__BB0_8:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_visual
|
|
.visible .entry lif_visual(
|
|
.param .u64 lif_visual_param_0,
|
|
.param .u64 lif_visual_param_1,
|
|
.param .u64 lif_visual_param_2,
|
|
.param .u64 lif_visual_param_3,
|
|
.param .u64 lif_visual_param_4,
|
|
.param .u32 lif_visual_param_5
|
|
)
|
|
{
|
|
.reg .pred %p<7>;
|
|
.reg .f32 %f<23>;
|
|
.reg .b32 %r<35>;
|
|
.reg .b64 %rd<38>;
|
|
|
|
|
|
ld.param.u64 %rd4, [lif_visual_param_0];
|
|
ld.param.u64 %rd5, [lif_visual_param_1];
|
|
ld.param.u64 %rd6, [lif_visual_param_2];
|
|
ld.param.u64 %rd7, [lif_visual_param_3];
|
|
ld.param.u64 %rd8, [lif_visual_param_4];
|
|
ld.param.u32 %r13, [lif_visual_param_5];
|
|
cvta.to.global.u64 %rd1, %rd8;
|
|
cvta.to.global.u64 %rd2, %rd7;
|
|
mov.u32 %r14, %ntid.x;
|
|
mov.u32 %r15, %ctaid.x;
|
|
mov.u32 %r16, %tid.x;
|
|
mad.lo.s32 %r1, %r15, %r14, %r16;
|
|
setp.ge.u32 %p1, %r1, %r13;
|
|
@%p1 bra $L__BB1_8;
|
|
|
|
cvta.to.global.u64 %rd9, %rd6;
|
|
cvta.to.global.u64 %rd10, %rd4;
|
|
cvta.to.global.u64 %rd11, %rd5;
|
|
mul.wide.u32 %rd12, %r1, 4;
|
|
add.s64 %rd13, %rd11, %rd12;
|
|
ld.global.u32 %r17, [%rd13];
|
|
mul.wide.u32 %rd14, %r17, 4;
|
|
add.s64 %rd3, %rd10, %rd14;
|
|
ld.global.f32 %f22, [%rd3];
|
|
add.s64 %rd15, %rd9, %rd12;
|
|
add.s32 %r18, %r1, 1;
|
|
mul.wide.u32 %rd16, %r18, 4;
|
|
add.s64 %rd17, %rd9, %rd16;
|
|
ld.global.u32 %r2, [%rd17];
|
|
ld.global.u32 %r3, [%rd15];
|
|
setp.ge.u32 %p2, %r3, %r2;
|
|
@%p2 bra $L__BB1_7;
|
|
|
|
add.s32 %r19, %r3, 1;
|
|
max.u32 %r4, %r2, %r19;
|
|
sub.s32 %r20, %r4, %r3;
|
|
and.b32 %r32, %r20, 3;
|
|
setp.eq.s32 %p3, %r32, 0;
|
|
mov.u32 %r33, %r3;
|
|
@%p3 bra $L__BB1_5;
|
|
|
|
mov.u32 %r33, %r3;
|
|
|
|
$L__BB1_4:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd18, %r33, 4;
|
|
add.s64 %rd19, %rd2, %rd18;
|
|
ld.global.u32 %r21, [%rd19];
|
|
mul.wide.u32 %rd20, %r21, 4;
|
|
add.s64 %rd21, %rd1, %rd20;
|
|
ld.global.f32 %f10, [%rd21];
|
|
add.rn.f32 %f22, %f22, %f10;
|
|
add.s32 %r33, %r33, 1;
|
|
add.s32 %r32, %r32, -1;
|
|
setp.ne.s32 %p4, %r32, 0;
|
|
@%p4 bra $L__BB1_4;
|
|
|
|
$L__BB1_5:
|
|
not.b32 %r22, %r3;
|
|
add.s32 %r23, %r4, %r22;
|
|
setp.lt.u32 %p5, %r23, 3;
|
|
@%p5 bra $L__BB1_7;
|
|
|
|
$L__BB1_6:
|
|
mul.wide.u32 %rd22, %r33, 4;
|
|
add.s64 %rd23, %rd2, %rd22;
|
|
ld.global.u32 %r24, [%rd23];
|
|
mul.wide.u32 %rd24, %r24, 4;
|
|
add.s64 %rd25, %rd1, %rd24;
|
|
ld.global.f32 %f11, [%rd25];
|
|
add.rn.f32 %f12, %f22, %f11;
|
|
add.s32 %r25, %r33, 1;
|
|
mul.wide.u32 %rd26, %r25, 4;
|
|
add.s64 %rd27, %rd2, %rd26;
|
|
ld.global.u32 %r26, [%rd27];
|
|
mul.wide.u32 %rd28, %r26, 4;
|
|
add.s64 %rd29, %rd1, %rd28;
|
|
ld.global.f32 %f13, [%rd29];
|
|
add.rn.f32 %f14, %f12, %f13;
|
|
add.s32 %r27, %r33, 2;
|
|
mul.wide.u32 %rd30, %r27, 4;
|
|
add.s64 %rd31, %rd2, %rd30;
|
|
ld.global.u32 %r28, [%rd31];
|
|
mul.wide.u32 %rd32, %r28, 4;
|
|
add.s64 %rd33, %rd1, %rd32;
|
|
ld.global.f32 %f15, [%rd33];
|
|
add.rn.f32 %f16, %f14, %f15;
|
|
add.s32 %r29, %r33, 3;
|
|
mul.wide.u32 %rd34, %r29, 4;
|
|
add.s64 %rd35, %rd2, %rd34;
|
|
ld.global.u32 %r30, [%rd35];
|
|
mul.wide.u32 %rd36, %r30, 4;
|
|
add.s64 %rd37, %rd1, %rd36;
|
|
ld.global.f32 %f17, [%rd37];
|
|
add.rn.f32 %f22, %f16, %f17;
|
|
add.s32 %r33, %r33, 4;
|
|
setp.lt.u32 %p6, %r33, %r2;
|
|
@%p6 bra $L__BB1_6;
|
|
|
|
$L__BB1_7:
|
|
st.global.f32 [%rd3], %f22;
|
|
|
|
$L__BB1_8:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_stimulate
|
|
.visible .entry lif_stimulate(
|
|
.param .u64 lif_stimulate_param_0,
|
|
.param .u64 lif_stimulate_param_1,
|
|
.param .u64 lif_stimulate_param_2,
|
|
.param .u32 lif_stimulate_param_3,
|
|
.param .f64 lif_stimulate_param_4
|
|
)
|
|
{
|
|
.reg .pred %p<7>;
|
|
.reg .f32 %f<18>;
|
|
.reg .b32 %r<17>;
|
|
.reg .f64 %fd<12>;
|
|
.reg .b64 %rd<12>;
|
|
|
|
|
|
ld.param.u64 %rd2, [lif_stimulate_param_0];
|
|
ld.param.u64 %rd3, [lif_stimulate_param_1];
|
|
ld.param.u64 %rd4, [lif_stimulate_param_2];
|
|
ld.param.u32 %r9, [lif_stimulate_param_3];
|
|
ld.param.f64 %fd1, [lif_stimulate_param_4];
|
|
mov.u32 %r10, %ctaid.x;
|
|
mov.u32 %r11, %ntid.x;
|
|
mov.u32 %r12, %tid.x;
|
|
mad.lo.s32 %r1, %r10, %r11, %r12;
|
|
setp.ge.u32 %p1, %r1, %r9;
|
|
@%p1 bra $L__BB2_8;
|
|
|
|
cvta.to.global.u64 %rd5, %rd2;
|
|
cvta.to.global.u64 %rd6, %rd3;
|
|
mul.wide.u32 %rd7, %r1, 4;
|
|
add.s64 %rd8, %rd6, %rd7;
|
|
cvta.to.global.u64 %rd9, %rd4;
|
|
add.s64 %rd10, %rd9, %rd7;
|
|
ld.global.u32 %r13, [%rd8];
|
|
mul.wide.u32 %rd11, %r13, 4;
|
|
add.s64 %rd1, %rd5, %rd11;
|
|
ld.global.f32 %f17, [%rd1];
|
|
ld.global.u32 %r2, [%rd10];
|
|
setp.eq.s32 %p2, %r2, 0;
|
|
@%p2 bra $L__BB2_7;
|
|
|
|
add.s32 %r14, %r2, -1;
|
|
and.b32 %r16, %r2, 3;
|
|
setp.lt.u32 %p3, %r14, 3;
|
|
@%p3 bra $L__BB2_5;
|
|
|
|
sub.s32 %r15, %r2, %r16;
|
|
|
|
$L__BB2_4:
|
|
cvt.f64.f32 %fd2, %f17;
|
|
add.rn.f64 %fd3, %fd2, %fd1;
|
|
cvt.rn.f32.f64 %f10, %fd3;
|
|
cvt.f64.f32 %fd4, %f10;
|
|
add.rn.f64 %fd5, %fd4, %fd1;
|
|
cvt.rn.f32.f64 %f11, %fd5;
|
|
cvt.f64.f32 %fd6, %f11;
|
|
add.rn.f64 %fd7, %fd6, %fd1;
|
|
cvt.rn.f32.f64 %f12, %fd7;
|
|
cvt.f64.f32 %fd8, %f12;
|
|
add.rn.f64 %fd9, %fd8, %fd1;
|
|
cvt.rn.f32.f64 %f17, %fd9;
|
|
add.s32 %r15, %r15, -4;
|
|
setp.ne.s32 %p4, %r15, 0;
|
|
@%p4 bra $L__BB2_4;
|
|
|
|
$L__BB2_5:
|
|
setp.eq.s32 %p5, %r16, 0;
|
|
@%p5 bra $L__BB2_7;
|
|
|
|
$L__BB2_6:
|
|
.pragma "nounroll";
|
|
cvt.f64.f32 %fd10, %f17;
|
|
add.rn.f64 %fd11, %fd10, %fd1;
|
|
cvt.rn.f32.f64 %f17, %fd11;
|
|
add.s32 %r16, %r16, -1;
|
|
setp.ne.s32 %p6, %r16, 0;
|
|
@%p6 bra $L__BB2_6;
|
|
|
|
$L__BB2_7:
|
|
st.global.f32 [%rd1], %f17;
|
|
|
|
$L__BB2_8:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_sweep
|
|
.visible .entry lif_sweep(
|
|
.param .u64 lif_sweep_param_0,
|
|
.param .u64 lif_sweep_param_1,
|
|
.param .u64 lif_sweep_param_2,
|
|
.param .u32 lif_sweep_param_3,
|
|
.param .f64 lif_sweep_param_4,
|
|
.param .f64 lif_sweep_param_5,
|
|
.param .u32 lif_sweep_param_6,
|
|
.param .u64 lif_sweep_param_7,
|
|
.param .u64 lif_sweep_param_8
|
|
)
|
|
{
|
|
.reg .pred %p<19>;
|
|
.reg .b16 %rs<3>;
|
|
.reg .f32 %f<6>;
|
|
.reg .b32 %r<61>;
|
|
.reg .f64 %fd<9>;
|
|
.reg .b64 %rd<33>;
|
|
// demoted variable
|
|
.shared .align 4 .b8 _ZZ9lif_sweepE8per_warp[32];
|
|
|
|
ld.param.u64 %rd7, [lif_sweep_param_0];
|
|
ld.param.u64 %rd8, [lif_sweep_param_1];
|
|
ld.param.u64 %rd9, [lif_sweep_param_2];
|
|
ld.param.u32 %r26, [lif_sweep_param_3];
|
|
ld.param.f64 %fd2, [lif_sweep_param_4];
|
|
ld.param.f64 %fd3, [lif_sweep_param_5];
|
|
ld.param.u32 %r27, [lif_sweep_param_6];
|
|
ld.param.u64 %rd10, [lif_sweep_param_7];
|
|
ld.param.u64 %rd11, [lif_sweep_param_8];
|
|
mov.u32 %r1, %ntid.x;
|
|
mov.u32 %r2, %ctaid.x;
|
|
mul.lo.s32 %r3, %r2, %r1;
|
|
mov.u32 %r4, %tid.x;
|
|
add.s32 %r5, %r3, %r4;
|
|
setp.ge.u32 %p4, %r5, %r26;
|
|
mov.pred %p18, 0;
|
|
@%p4 bra $L__BB3_6;
|
|
|
|
cvta.to.global.u64 %rd12, %rd8;
|
|
cvt.u64.u32 %rd1, %r5;
|
|
add.s64 %rd2, %rd12, %rd1;
|
|
ld.global.u8 %rs1, [%rd2];
|
|
setp.eq.s16 %p5, %rs1, 0;
|
|
cvta.to.global.u64 %rd13, %rd7;
|
|
mul.wide.u32 %rd14, %r5, 4;
|
|
add.s64 %rd3, %rd13, %rd14;
|
|
@%p5 bra $L__BB3_3;
|
|
|
|
add.s16 %rs2, %rs1, -1;
|
|
st.global.u8 [%rd2], %rs2;
|
|
ld.global.f32 %f1, [%rd3];
|
|
cvt.f64.f32 %fd4, %f1;
|
|
mul.rn.f64 %fd5, %fd4, %fd2;
|
|
cvt.rn.f32.f64 %f2, %fd5;
|
|
st.global.f32 [%rd3], %f2;
|
|
bra.uni $L__BB3_6;
|
|
|
|
$L__BB3_3:
|
|
ld.global.f32 %f3, [%rd3];
|
|
cvt.f64.f32 %fd6, %f3;
|
|
mul.rn.f64 %fd7, %fd6, %fd2;
|
|
cvta.to.global.u64 %rd15, %rd9;
|
|
shl.b64 %rd16, %rd1, 2;
|
|
add.s64 %rd17, %rd15, %rd16;
|
|
ld.global.f32 %f4, [%rd17];
|
|
cvt.f64.f32 %fd8, %f4;
|
|
add.rn.f64 %fd1, %fd7, %fd8;
|
|
setp.ltu.f64 %p7, %fd1, %fd3;
|
|
@%p7 bra $L__BB3_5;
|
|
bra.uni $L__BB3_4;
|
|
|
|
$L__BB3_5:
|
|
cvt.rn.f32.f64 %f5, %fd1;
|
|
st.global.f32 [%rd3], %f5;
|
|
bra.uni $L__BB3_6;
|
|
|
|
$L__BB3_4:
|
|
mov.u32 %r28, 0;
|
|
st.global.u32 [%rd3], %r28;
|
|
st.global.u8 [%rd2], %r27;
|
|
mov.pred %p18, -1;
|
|
|
|
$L__BB3_6:
|
|
mov.u32 %r29, -1;
|
|
vote.sync.ballot.b32 %r6, %p18, %r29;
|
|
and.b32 %r30, %r4, 31;
|
|
setp.ne.s32 %p10, %r30, 0;
|
|
@%p10 bra $L__BB3_10;
|
|
|
|
and.b32 %r31, %r4, -32;
|
|
add.s32 %r32, %r3, %r31;
|
|
shr.u32 %r8, %r32, 5;
|
|
add.s32 %r33, %r26, 31;
|
|
shr.u32 %r34, %r33, 5;
|
|
setp.ge.u32 %p11, %r8, %r34;
|
|
@%p11 bra $L__BB3_9;
|
|
|
|
cvta.to.global.u64 %rd18, %rd10;
|
|
mul.wide.u32 %rd19, %r8, 4;
|
|
add.s64 %rd20, %rd18, %rd19;
|
|
st.global.u32 [%rd20], %r6;
|
|
|
|
$L__BB3_9:
|
|
popc.b32 %r35, %r6;
|
|
shr.u32 %r36, %r4, 5;
|
|
mul.wide.u32 %rd21, %r36, 4;
|
|
mov.u64 %rd22, _ZZ9lif_sweepE8per_warp;
|
|
add.s64 %rd23, %rd22, %rd21;
|
|
st.shared.u32 [%rd23], %r35;
|
|
|
|
$L__BB3_10:
|
|
bar.sync 0;
|
|
setp.ne.s32 %p12, %r4, 0;
|
|
@%p12 bra $L__BB3_19;
|
|
|
|
shr.u32 %r9, %r1, 5;
|
|
setp.eq.s32 %p13, %r9, 0;
|
|
mov.u32 %r60, 0;
|
|
@%p13 bra $L__BB3_18;
|
|
|
|
add.s32 %r41, %r9, -1;
|
|
and.b32 %r59, %r9, 3;
|
|
setp.lt.u32 %p14, %r41, 3;
|
|
mov.u32 %r56, 0;
|
|
mov.u32 %r60, %r56;
|
|
@%p14 bra $L__BB3_15;
|
|
|
|
sub.s32 %r54, %r9, %r59;
|
|
mov.u32 %r56, 0;
|
|
mov.u64 %rd25, _ZZ9lif_sweepE8per_warp;
|
|
|
|
$L__BB3_14:
|
|
mul.wide.u32 %rd24, %r56, 4;
|
|
add.s64 %rd26, %rd25, %rd24;
|
|
ld.shared.u32 %r44, [%rd26];
|
|
add.s32 %r45, %r44, %r60;
|
|
ld.shared.u32 %r46, [%rd26+4];
|
|
add.s32 %r47, %r46, %r45;
|
|
ld.shared.u32 %r48, [%rd26+8];
|
|
add.s32 %r49, %r48, %r47;
|
|
ld.shared.u32 %r50, [%rd26+12];
|
|
add.s32 %r60, %r50, %r49;
|
|
add.s32 %r56, %r56, 4;
|
|
add.s32 %r54, %r54, -4;
|
|
setp.ne.s32 %p15, %r54, 0;
|
|
@%p15 bra $L__BB3_14;
|
|
|
|
$L__BB3_15:
|
|
setp.eq.s32 %p16, %r59, 0;
|
|
@%p16 bra $L__BB3_18;
|
|
|
|
mul.wide.u32 %rd27, %r56, 4;
|
|
mov.u64 %rd28, _ZZ9lif_sweepE8per_warp;
|
|
add.s64 %rd32, %rd28, %rd27;
|
|
|
|
$L__BB3_17:
|
|
.pragma "nounroll";
|
|
ld.shared.u32 %r51, [%rd32];
|
|
add.s32 %r60, %r51, %r60;
|
|
add.s64 %rd32, %rd32, 4;
|
|
add.s32 %r59, %r59, -1;
|
|
setp.ne.s32 %p17, %r59, 0;
|
|
@%p17 bra $L__BB3_17;
|
|
|
|
$L__BB3_18:
|
|
cvta.to.global.u64 %rd29, %rd11;
|
|
mul.wide.u32 %rd30, %r2, 4;
|
|
add.s64 %rd31, %rd29, %rd30;
|
|
st.global.u32 [%rd31], %r60;
|
|
|
|
$L__BB3_19:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_scan
|
|
.visible .entry lif_scan(
|
|
.param .u64 lif_scan_param_0,
|
|
.param .u64 lif_scan_param_1,
|
|
.param .u32 lif_scan_param_2,
|
|
.param .u64 lif_scan_param_3,
|
|
.param .u64 lif_scan_param_4,
|
|
.param .u32 lif_scan_param_5,
|
|
.param .u64 lif_scan_param_6
|
|
)
|
|
{
|
|
.reg .pred %p<15>;
|
|
.reg .b32 %r<124>;
|
|
.reg .b64 %rd<53>;
|
|
|
|
|
|
ld.param.u64 %rd8, [lif_scan_param_0];
|
|
ld.param.u64 %rd9, [lif_scan_param_1];
|
|
ld.param.u32 %r43, [lif_scan_param_2];
|
|
ld.param.u64 %rd10, [lif_scan_param_3];
|
|
ld.param.u64 %rd11, [lif_scan_param_4];
|
|
ld.param.u32 %r41, [lif_scan_param_5];
|
|
ld.param.u64 %rd12, [lif_scan_param_6];
|
|
mov.u32 %r44, %tid.x;
|
|
mov.u32 %r45, %ntid.x;
|
|
add.s32 %r46, %r45, %r43;
|
|
add.s32 %r47, %r46, -1;
|
|
div.u32 %r48, %r47, %r45;
|
|
mul.lo.s32 %r49, %r48, %r44;
|
|
add.s32 %r50, %r49, %r48;
|
|
min.u32 %r1, %r49, %r43;
|
|
min.u32 %r2, %r50, %r43;
|
|
setp.ge.u32 %p1, %r1, %r2;
|
|
mov.u32 %r114, 0;
|
|
@%p1 bra $L__BB4_7;
|
|
|
|
not.b32 %r53, %r2;
|
|
mov.u32 %r54, -2;
|
|
sub.s32 %r55, %r54, %r1;
|
|
sub.s32 %r56, %r55, %r53;
|
|
sub.s32 %r3, %r2, %r1;
|
|
and.b32 %r113, %r3, 3;
|
|
setp.lt.u32 %p2, %r56, 3;
|
|
mov.u32 %r114, 0;
|
|
mov.u32 %r109, %r1;
|
|
@%p2 bra $L__BB4_4;
|
|
|
|
sub.s32 %r107, %r3, %r113;
|
|
cvta.to.global.u64 %rd1, %rd8;
|
|
mov.u32 %r114, 0;
|
|
mov.u32 %r109, %r1;
|
|
|
|
$L__BB4_3:
|
|
mul.wide.u32 %rd13, %r109, 4;
|
|
add.s64 %rd14, %rd1, %rd13;
|
|
ld.global.u32 %r58, [%rd14];
|
|
add.s32 %r59, %r58, %r114;
|
|
add.s32 %r60, %r109, 1;
|
|
mul.wide.u32 %rd15, %r60, 4;
|
|
add.s64 %rd16, %rd1, %rd15;
|
|
ld.global.u32 %r61, [%rd16];
|
|
add.s32 %r62, %r61, %r59;
|
|
add.s32 %r63, %r109, 2;
|
|
mul.wide.u32 %rd17, %r63, 4;
|
|
add.s64 %rd18, %rd1, %rd17;
|
|
ld.global.u32 %r64, [%rd18];
|
|
add.s32 %r65, %r64, %r62;
|
|
add.s32 %r66, %r109, 3;
|
|
mul.wide.u32 %rd19, %r66, 4;
|
|
add.s64 %rd20, %rd1, %rd19;
|
|
ld.global.u32 %r67, [%rd20];
|
|
add.s32 %r114, %r67, %r65;
|
|
add.s32 %r109, %r109, 4;
|
|
add.s32 %r107, %r107, -4;
|
|
setp.ne.s32 %p3, %r107, 0;
|
|
@%p3 bra $L__BB4_3;
|
|
|
|
$L__BB4_4:
|
|
setp.eq.s32 %p4, %r113, 0;
|
|
@%p4 bra $L__BB4_7;
|
|
|
|
cvta.to.global.u64 %rd2, %rd8;
|
|
|
|
$L__BB4_6:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd21, %r109, 4;
|
|
add.s64 %rd22, %rd2, %rd21;
|
|
ld.global.u32 %r68, [%rd22];
|
|
add.s32 %r114, %r68, %r114;
|
|
add.s32 %r109, %r109, 1;
|
|
add.s32 %r113, %r113, -1;
|
|
setp.ne.s32 %p5, %r113, 0;
|
|
@%p5 bra $L__BB4_6;
|
|
|
|
$L__BB4_7:
|
|
mul.wide.u32 %rd23, %r44, 4;
|
|
mov.u64 %rd24, partial;
|
|
add.s64 %rd3, %rd24, %rd23;
|
|
st.shared.u32 [%rd3], %r114;
|
|
bar.sync 0;
|
|
setp.lt.u32 %p6, %r45, 2;
|
|
@%p6 bra $L__BB4_12;
|
|
|
|
mov.u32 %r115, 1;
|
|
|
|
$L__BB4_9:
|
|
setp.lt.u32 %p7, %r44, %r115;
|
|
mov.u32 %r116, 0;
|
|
@%p7 bra $L__BB4_11;
|
|
|
|
sub.s32 %r75, %r44, %r115;
|
|
mul.wide.u32 %rd25, %r75, 4;
|
|
add.s64 %rd27, %rd24, %rd25;
|
|
ld.shared.u32 %r116, [%rd27];
|
|
|
|
$L__BB4_11:
|
|
bar.sync 0;
|
|
ld.shared.u32 %r76, [%rd3];
|
|
add.s32 %r77, %r76, %r116;
|
|
st.shared.u32 [%rd3], %r77;
|
|
bar.sync 0;
|
|
shl.b32 %r115, %r115, 1;
|
|
setp.lt.u32 %p8, %r115, %r45;
|
|
@%p8 bra $L__BB4_9;
|
|
|
|
$L__BB4_12:
|
|
cvta.to.global.u64 %rd28, %rd12;
|
|
ld.global.u32 %r26, [%rd28];
|
|
sub.s32 %r79, %r26, %r114;
|
|
ld.shared.u32 %r80, [%rd3];
|
|
add.s32 %r121, %r79, %r80;
|
|
@%p1 bra $L__BB4_19;
|
|
|
|
sub.s32 %r81, %r2, %r1;
|
|
and.b32 %r119, %r81, 3;
|
|
setp.eq.s32 %p10, %r119, 0;
|
|
mov.u32 %r120, %r1;
|
|
@%p10 bra $L__BB4_16;
|
|
|
|
cvta.to.global.u64 %rd4, %rd8;
|
|
cvta.to.global.u64 %rd5, %rd9;
|
|
mov.u32 %r120, %r1;
|
|
|
|
$L__BB4_15:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd29, %r120, 4;
|
|
add.s64 %rd30, %rd5, %rd29;
|
|
st.global.u32 [%rd30], %r121;
|
|
add.s64 %rd31, %rd4, %rd29;
|
|
ld.global.u32 %r84, [%rd31];
|
|
add.s32 %r121, %r84, %r121;
|
|
add.s32 %r120, %r120, 1;
|
|
add.s32 %r119, %r119, -1;
|
|
setp.ne.s32 %p11, %r119, 0;
|
|
@%p11 bra $L__BB4_15;
|
|
|
|
$L__BB4_16:
|
|
mov.u32 %r85, -2;
|
|
sub.s32 %r86, %r85, %r1;
|
|
not.b32 %r87, %r2;
|
|
sub.s32 %r88, %r86, %r87;
|
|
setp.lt.u32 %p12, %r88, 3;
|
|
@%p12 bra $L__BB4_19;
|
|
|
|
cvta.to.global.u64 %rd6, %rd8;
|
|
cvta.to.global.u64 %rd7, %rd9;
|
|
|
|
$L__BB4_18:
|
|
mul.wide.u32 %rd32, %r120, 4;
|
|
add.s64 %rd33, %rd7, %rd32;
|
|
st.global.u32 [%rd33], %r121;
|
|
add.s64 %rd34, %rd6, %rd32;
|
|
ld.global.u32 %r89, [%rd34];
|
|
add.s32 %r90, %r89, %r121;
|
|
add.s32 %r91, %r120, 1;
|
|
mul.wide.u32 %rd35, %r91, 4;
|
|
add.s64 %rd36, %rd7, %rd35;
|
|
st.global.u32 [%rd36], %r90;
|
|
add.s64 %rd37, %rd6, %rd35;
|
|
ld.global.u32 %r92, [%rd37];
|
|
add.s32 %r93, %r92, %r90;
|
|
add.s32 %r94, %r120, 2;
|
|
mul.wide.u32 %rd38, %r94, 4;
|
|
add.s64 %rd39, %rd7, %rd38;
|
|
st.global.u32 [%rd39], %r93;
|
|
add.s64 %rd40, %rd6, %rd38;
|
|
ld.global.u32 %r95, [%rd40];
|
|
add.s32 %r96, %r95, %r93;
|
|
add.s32 %r97, %r120, 3;
|
|
mul.wide.u32 %rd41, %r97, 4;
|
|
add.s64 %rd42, %rd7, %rd41;
|
|
st.global.u32 [%rd42], %r96;
|
|
add.s64 %rd43, %rd6, %rd41;
|
|
ld.global.u32 %r98, [%rd43];
|
|
add.s32 %r121, %r98, %r96;
|
|
add.s32 %r120, %r120, 4;
|
|
setp.lt.u32 %p13, %r120, %r2;
|
|
@%p13 bra $L__BB4_18;
|
|
|
|
$L__BB4_19:
|
|
bar.sync 0;
|
|
add.s32 %r100, %r45, -1;
|
|
setp.ne.s32 %p14, %r44, %r100;
|
|
@%p14 bra $L__BB4_21;
|
|
|
|
ld.shared.u32 %r103, [%rd3];
|
|
cvta.to.global.u64 %rd47, %rd10;
|
|
mul.wide.u32 %rd48, %r41, 4;
|
|
add.s64 %rd49, %rd47, %rd48;
|
|
st.global.u32 [%rd49], %r103;
|
|
cvta.to.global.u64 %rd50, %rd11;
|
|
add.s64 %rd51, %rd50, %rd48;
|
|
st.global.u32 [%rd51], %r26;
|
|
add.s32 %r104, %r103, %r26;
|
|
st.global.u32 [%rd28], %r104;
|
|
|
|
$L__BB4_21:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_compact
|
|
.visible .entry lif_compact(
|
|
.param .u64 lif_compact_param_0,
|
|
.param .u64 lif_compact_param_1,
|
|
.param .u32 lif_compact_param_2,
|
|
.param .u64 lif_compact_param_3
|
|
)
|
|
{
|
|
.reg .pred %p<11>;
|
|
.reg .b32 %r<66>;
|
|
.reg .b64 %rd<25>;
|
|
// demoted variable
|
|
.shared .align 4 .b8 _ZZ11lif_compactE8per_warp[32];
|
|
|
|
ld.param.u64 %rd4, [lif_compact_param_0];
|
|
ld.param.u64 %rd5, [lif_compact_param_1];
|
|
ld.param.u32 %r25, [lif_compact_param_2];
|
|
ld.param.u64 %rd6, [lif_compact_param_3];
|
|
mov.u32 %r27, %ntid.x;
|
|
mov.u32 %r1, %ctaid.x;
|
|
mul.lo.s32 %r28, %r1, %r27;
|
|
mov.u32 %r29, %tid.x;
|
|
add.s32 %r2, %r28, %r29;
|
|
and.b32 %r3, %r29, 31;
|
|
shr.u32 %r4, %r29, 5;
|
|
and.b32 %r30, %r29, -32;
|
|
add.s32 %r31, %r28, %r30;
|
|
shr.u32 %r5, %r31, 5;
|
|
add.s32 %r32, %r25, 31;
|
|
shr.u32 %r33, %r32, 5;
|
|
setp.ge.u32 %p1, %r5, %r33;
|
|
mov.u32 %r56, 0;
|
|
@%p1 bra $L__BB5_2;
|
|
|
|
cvta.to.global.u64 %rd7, %rd4;
|
|
mul.wide.u32 %rd8, %r5, 4;
|
|
add.s64 %rd9, %rd7, %rd8;
|
|
ld.global.u32 %r56, [%rd9];
|
|
|
|
$L__BB5_2:
|
|
setp.ne.s32 %p2, %r3, 0;
|
|
@%p2 bra $L__BB5_4;
|
|
|
|
popc.b32 %r34, %r56;
|
|
mul.wide.u32 %rd10, %r4, 4;
|
|
mov.u64 %rd11, _ZZ11lif_compactE8per_warp;
|
|
add.s64 %rd12, %rd11, %rd10;
|
|
st.shared.u32 [%rd12], %r34;
|
|
|
|
$L__BB5_4:
|
|
bar.sync 0;
|
|
mov.u32 %r35, 1;
|
|
shl.b32 %r36, %r35, %r3;
|
|
and.b32 %r37, %r56, %r36;
|
|
setp.eq.s32 %p3, %r37, 0;
|
|
setp.ge.u32 %p4, %r2, %r25;
|
|
or.pred %p5, %p4, %p3;
|
|
@%p5 bra $L__BB5_13;
|
|
|
|
mov.u32 %r38, -1;
|
|
shl.b32 %r39, %r38, %r3;
|
|
not.b32 %r40, %r39;
|
|
and.b32 %r41, %r56, %r40;
|
|
popc.b32 %r65, %r41;
|
|
setp.eq.s32 %p6, %r4, 0;
|
|
@%p6 bra $L__BB5_12;
|
|
|
|
add.s32 %r44, %r4, -1;
|
|
and.b32 %r64, %r4, 3;
|
|
setp.lt.u32 %p7, %r44, 3;
|
|
mov.u32 %r61, 0;
|
|
@%p7 bra $L__BB5_9;
|
|
|
|
sub.s32 %r59, %r4, %r64;
|
|
mov.u32 %r61, 0;
|
|
mov.u64 %rd14, _ZZ11lif_compactE8per_warp;
|
|
|
|
$L__BB5_8:
|
|
mul.wide.u32 %rd13, %r61, 4;
|
|
add.s64 %rd15, %rd14, %rd13;
|
|
ld.shared.u32 %r46, [%rd15];
|
|
add.s32 %r47, %r46, %r65;
|
|
ld.shared.u32 %r48, [%rd15+4];
|
|
add.s32 %r49, %r48, %r47;
|
|
ld.shared.u32 %r50, [%rd15+8];
|
|
add.s32 %r51, %r50, %r49;
|
|
ld.shared.u32 %r52, [%rd15+12];
|
|
add.s32 %r65, %r52, %r51;
|
|
add.s32 %r61, %r61, 4;
|
|
add.s32 %r59, %r59, -4;
|
|
setp.ne.s32 %p8, %r59, 0;
|
|
@%p8 bra $L__BB5_8;
|
|
|
|
$L__BB5_9:
|
|
setp.eq.s32 %p9, %r64, 0;
|
|
@%p9 bra $L__BB5_12;
|
|
|
|
mul.wide.u32 %rd16, %r61, 4;
|
|
mov.u64 %rd17, _ZZ11lif_compactE8per_warp;
|
|
add.s64 %rd24, %rd17, %rd16;
|
|
|
|
$L__BB5_11:
|
|
.pragma "nounroll";
|
|
ld.shared.u32 %r53, [%rd24];
|
|
add.s32 %r65, %r53, %r65;
|
|
add.s64 %rd24, %rd24, 4;
|
|
add.s32 %r64, %r64, -1;
|
|
setp.ne.s32 %p10, %r64, 0;
|
|
@%p10 bra $L__BB5_11;
|
|
|
|
$L__BB5_12:
|
|
cvta.to.global.u64 %rd18, %rd5;
|
|
mul.wide.u32 %rd19, %r1, 4;
|
|
add.s64 %rd20, %rd18, %rd19;
|
|
ld.global.u32 %r54, [%rd20];
|
|
add.s32 %r55, %r54, %r65;
|
|
cvta.to.global.u64 %rd21, %rd6;
|
|
mul.wide.u32 %rd22, %r55, 4;
|
|
add.s64 %rd23, %rd21, %rd22;
|
|
st.global.u32 [%rd23], %r2;
|
|
|
|
$L__BB5_13:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_bucket_count
|
|
.visible .entry lif_bucket_count(
|
|
.param .u64 lif_bucket_count_param_0,
|
|
.param .u64 lif_bucket_count_param_1,
|
|
.param .u64 lif_bucket_count_param_2,
|
|
.param .u32 lif_bucket_count_param_3,
|
|
.param .u64 lif_bucket_count_param_4,
|
|
.param .u64 lif_bucket_count_param_5,
|
|
.param .u64 lif_bucket_count_param_6
|
|
)
|
|
{
|
|
.reg .pred %p<5>;
|
|
.reg .b32 %r<27>;
|
|
.reg .b64 %rd<28>;
|
|
|
|
|
|
ld.param.u64 %rd6, [lif_bucket_count_param_0];
|
|
ld.param.u64 %rd11, [lif_bucket_count_param_1];
|
|
ld.param.u64 %rd7, [lif_bucket_count_param_2];
|
|
ld.param.u32 %r14, [lif_bucket_count_param_3];
|
|
ld.param.u64 %rd8, [lif_bucket_count_param_4];
|
|
ld.param.u64 %rd9, [lif_bucket_count_param_5];
|
|
ld.param.u64 %rd10, [lif_bucket_count_param_6];
|
|
cvt.u64.u32 %rd1, %r14;
|
|
cvta.to.global.u64 %rd12, %rd11;
|
|
mul.wide.u32 %rd13, %r14, 4;
|
|
add.s64 %rd14, %rd12, %rd13;
|
|
mov.u32 %r1, %ntid.x;
|
|
mov.u32 %r15, %ctaid.x;
|
|
mov.u32 %r2, %tid.x;
|
|
mad.lo.s32 %r16, %r15, %r1, %r2;
|
|
shr.u32 %r25, %r16, 5;
|
|
ld.global.nc.u32 %r4, [%rd14];
|
|
setp.ge.u32 %p1, %r25, %r4;
|
|
@%p1 bra $L__BB6_5;
|
|
|
|
cvta.to.global.u64 %rd2, %rd8;
|
|
cvta.to.global.u64 %rd15, %rd7;
|
|
shl.b64 %rd16, %rd1, 2;
|
|
add.s64 %rd17, %rd15, %rd16;
|
|
ld.global.nc.u32 %r5, [%rd17];
|
|
and.b32 %r6, %r2, 31;
|
|
mov.u32 %r17, %nctaid.x;
|
|
mul.lo.s32 %r18, %r17, %r1;
|
|
shr.u32 %r7, %r18, 5;
|
|
cvta.to.global.u64 %rd3, %rd6;
|
|
cvta.to.global.u64 %rd4, %rd9;
|
|
cvta.to.global.u64 %rd5, %rd10;
|
|
|
|
$L__BB6_2:
|
|
add.s32 %r19, %r25, %r5;
|
|
mul.wide.u32 %rd18, %r19, 4;
|
|
add.s64 %rd19, %rd3, %rd18;
|
|
ld.global.nc.u32 %r20, [%rd19];
|
|
mul.wide.u32 %rd20, %r20, 4;
|
|
add.s64 %rd21, %rd2, %rd20;
|
|
add.s32 %r21, %r20, 1;
|
|
mul.wide.u32 %rd22, %r21, 4;
|
|
add.s64 %rd23, %rd2, %rd22;
|
|
ld.global.nc.u32 %r22, [%rd21];
|
|
add.s32 %r26, %r22, %r6;
|
|
ld.global.nc.u32 %r10, [%rd23];
|
|
setp.ge.u32 %p2, %r26, %r10;
|
|
@%p2 bra $L__BB6_4;
|
|
|
|
$L__BB6_3:
|
|
mul.wide.u32 %rd24, %r26, 4;
|
|
add.s64 %rd25, %rd4, %rd24;
|
|
ld.global.nc.u32 %r23, [%rd25];
|
|
mul.wide.u32 %rd26, %r23, 4;
|
|
add.s64 %rd27, %rd5, %rd26;
|
|
atom.global.add.u32 %r24, [%rd27], 1;
|
|
add.s32 %r26, %r26, 32;
|
|
setp.lt.u32 %p3, %r26, %r10;
|
|
@%p3 bra $L__BB6_3;
|
|
|
|
$L__BB6_4:
|
|
add.s32 %r25, %r25, %r7;
|
|
setp.lt.u32 %p4, %r25, %r4;
|
|
@%p4 bra $L__BB6_2;
|
|
|
|
$L__BB6_5:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_bucket_place
|
|
.visible .entry lif_bucket_place(
|
|
.param .u64 lif_bucket_place_param_0,
|
|
.param .u64 lif_bucket_place_param_1,
|
|
.param .u64 lif_bucket_place_param_2,
|
|
.param .u32 lif_bucket_place_param_3,
|
|
.param .u64 lif_bucket_place_param_4,
|
|
.param .u64 lif_bucket_place_param_5,
|
|
.param .u64 lif_bucket_place_param_6,
|
|
.param .u64 lif_bucket_place_param_7,
|
|
.param .u64 lif_bucket_place_param_8,
|
|
.param .u64 lif_bucket_place_param_9,
|
|
.param .u64 lif_bucket_place_param_10
|
|
)
|
|
{
|
|
.reg .pred %p<6>;
|
|
.reg .b16 %rs<2>;
|
|
.reg .b32 %r<38>;
|
|
.reg .b64 %rd<59>;
|
|
|
|
|
|
ld.param.u64 %rd11, [lif_bucket_place_param_0];
|
|
ld.param.u64 %rd20, [lif_bucket_place_param_1];
|
|
ld.param.u64 %rd12, [lif_bucket_place_param_2];
|
|
ld.param.u32 %r14, [lif_bucket_place_param_3];
|
|
ld.param.u64 %rd13, [lif_bucket_place_param_4];
|
|
ld.param.u64 %rd14, [lif_bucket_place_param_5];
|
|
ld.param.u64 %rd15, [lif_bucket_place_param_6];
|
|
ld.param.u64 %rd16, [lif_bucket_place_param_7];
|
|
ld.param.u64 %rd17, [lif_bucket_place_param_8];
|
|
ld.param.u64 %rd18, [lif_bucket_place_param_9];
|
|
ld.param.u64 %rd19, [lif_bucket_place_param_10];
|
|
cvt.u64.u32 %rd1, %r14;
|
|
cvta.to.global.u64 %rd21, %rd20;
|
|
mul.wide.u32 %rd22, %r14, 4;
|
|
add.s64 %rd23, %rd21, %rd22;
|
|
mov.u32 %r1, %ntid.x;
|
|
mov.u32 %r15, %ctaid.x;
|
|
mov.u32 %r16, %tid.x;
|
|
mad.lo.s32 %r17, %r15, %r1, %r16;
|
|
shr.u32 %r35, %r17, 5;
|
|
ld.global.nc.u32 %r3, [%rd23];
|
|
setp.ge.u32 %p1, %r35, %r3;
|
|
@%p1 bra $L__BB7_7;
|
|
|
|
cvta.to.global.u64 %rd24, %rd12;
|
|
shl.b64 %rd25, %rd1, 2;
|
|
add.s64 %rd26, %rd24, %rd25;
|
|
ld.global.nc.u32 %r4, [%rd26];
|
|
mov.u32 %r18, %nctaid.x;
|
|
mul.lo.s32 %r19, %r18, %r1;
|
|
shr.u32 %r5, %r19, 5;
|
|
cvta.to.global.u64 %rd2, %rd11;
|
|
cvta.to.global.u64 %rd3, %rd16;
|
|
cvta.to.global.u64 %rd4, %rd17;
|
|
cvta.to.global.u64 %rd5, %rd15;
|
|
cvta.to.global.u64 %rd6, %rd18;
|
|
cvta.to.global.u64 %rd7, %rd19;
|
|
cvta.to.global.u64 %rd29, %rd13;
|
|
cvta.to.global.u64 %rd52, %rd14;
|
|
|
|
$L__BB7_2:
|
|
add.s32 %r20, %r35, %r4;
|
|
mul.wide.u32 %rd27, %r20, 4;
|
|
add.s64 %rd28, %rd2, %rd27;
|
|
ld.global.nc.u32 %r21, [%rd28];
|
|
mul.wide.u32 %rd30, %r21, 4;
|
|
add.s64 %rd31, %rd29, %rd30;
|
|
add.s32 %r22, %r21, 1;
|
|
mul.wide.u32 %rd32, %r22, 4;
|
|
add.s64 %rd33, %rd29, %rd32;
|
|
and.b32 %r24, %r16, 31;
|
|
ld.global.nc.u32 %r25, [%rd31];
|
|
add.s32 %r36, %r25, %r24;
|
|
ld.global.nc.u32 %r8, [%rd33];
|
|
setp.ge.u32 %p2, %r36, %r8;
|
|
@%p2 bra $L__BB7_6;
|
|
|
|
$L__BB7_3:
|
|
shr.u32 %r27, %r36, 6;
|
|
cvt.u64.u32 %rd8, %r27;
|
|
mul.wide.u32 %rd34, %r27, 8;
|
|
add.s64 %rd35, %rd3, %rd34;
|
|
and.b32 %r28, %r36, 63;
|
|
cvt.u64.u32 %rd9, %r28;
|
|
mov.u64 %rd36, 1;
|
|
shl.b64 %rd37, %rd36, %r28;
|
|
ld.global.nc.u64 %rd10, [%rd35];
|
|
and.b64 %rd38, %rd10, %rd37;
|
|
setp.eq.s64 %p3, %rd38, 0;
|
|
mov.u32 %r37, 0;
|
|
@%p3 bra $L__BB7_5;
|
|
|
|
cvt.u32.u64 %r29, %rd9;
|
|
mov.u64 %rd39, -1;
|
|
shl.b64 %rd40, %rd39, %r29;
|
|
not.b64 %rd41, %rd40;
|
|
and.b64 %rd42, %rd10, %rd41;
|
|
popc.b64 %r30, %rd42;
|
|
shl.b64 %rd43, %rd8, 2;
|
|
add.s64 %rd44, %rd4, %rd43;
|
|
ld.global.nc.u32 %r31, [%rd44];
|
|
add.s32 %r32, %r30, %r31;
|
|
add.s32 %r37, %r32, 1;
|
|
|
|
$L__BB7_5:
|
|
mul.wide.u32 %rd45, %r36, 2;
|
|
add.s64 %rd46, %rd5, %rd45;
|
|
ld.global.nc.u16 %rs1, [%rd46];
|
|
cvt.u64.u16 %rd47, %rs1;
|
|
cvt.u64.u32 %rd48, %r36;
|
|
bfi.b64 %rd49, %rd48, %rd47, 32, 32;
|
|
mul.wide.u32 %rd50, %r37, 65536;
|
|
or.b64 %rd51, %rd49, %rd50;
|
|
mul.wide.u32 %rd53, %r36, 4;
|
|
add.s64 %rd54, %rd52, %rd53;
|
|
ld.global.nc.u32 %r33, [%rd54];
|
|
mul.wide.u32 %rd55, %r33, 4;
|
|
add.s64 %rd56, %rd6, %rd55;
|
|
atom.global.add.u32 %r34, [%rd56], 1;
|
|
mul.wide.u32 %rd57, %r34, 8;
|
|
add.s64 %rd58, %rd7, %rd57;
|
|
st.global.u64 [%rd58], %rd51;
|
|
add.s32 %r36, %r36, 32;
|
|
setp.lt.u32 %p4, %r36, %r8;
|
|
@%p4 bra $L__BB7_3;
|
|
|
|
$L__BB7_6:
|
|
add.s32 %r35, %r35, %r5;
|
|
setp.lt.u32 %p5, %r35, %r3;
|
|
@%p5 bra $L__BB7_2;
|
|
|
|
$L__BB7_7:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_bucket_scan
|
|
.visible .entry lif_bucket_scan(
|
|
.param .u64 lif_bucket_scan_param_0,
|
|
.param .u64 lif_bucket_scan_param_1,
|
|
.param .u64 lif_bucket_scan_param_2,
|
|
.param .u32 lif_bucket_scan_param_3
|
|
)
|
|
{
|
|
.reg .pred %p<7>;
|
|
.reg .b32 %r<25>;
|
|
.reg .b64 %rd<19>;
|
|
// demoted variable
|
|
.shared .align 4 .b8 _ZZ15lif_bucket_scanE7partial[1024];
|
|
|
|
ld.param.u64 %rd2, [lif_bucket_scan_param_0];
|
|
ld.param.u64 %rd3, [lif_bucket_scan_param_1];
|
|
ld.param.u64 %rd4, [lif_bucket_scan_param_2];
|
|
ld.param.u32 %r11, [lif_bucket_scan_param_3];
|
|
mov.u32 %r1, %ntid.x;
|
|
mov.u32 %r2, %ctaid.x;
|
|
mov.u32 %r3, %tid.x;
|
|
mad.lo.s32 %r4, %r2, %r1, %r3;
|
|
setp.ge.u32 %p1, %r4, %r11;
|
|
mov.u32 %r22, 0;
|
|
@%p1 bra $L__BB8_2;
|
|
|
|
cvta.to.global.u64 %rd5, %rd2;
|
|
mul.wide.u32 %rd6, %r4, 4;
|
|
add.s64 %rd7, %rd5, %rd6;
|
|
ld.global.nc.u32 %r22, [%rd7];
|
|
|
|
$L__BB8_2:
|
|
mul.wide.u32 %rd8, %r3, 4;
|
|
mov.u64 %rd9, _ZZ15lif_bucket_scanE7partial;
|
|
add.s64 %rd1, %rd9, %rd8;
|
|
st.shared.u32 [%rd1], %r22;
|
|
bar.sync 0;
|
|
setp.lt.u32 %p2, %r1, 2;
|
|
@%p2 bra $L__BB8_7;
|
|
|
|
mov.u32 %r23, 1;
|
|
|
|
$L__BB8_4:
|
|
setp.lt.u32 %p3, %r3, %r23;
|
|
mov.u32 %r24, 0;
|
|
@%p3 bra $L__BB8_6;
|
|
|
|
sub.s32 %r15, %r3, %r23;
|
|
mul.wide.u32 %rd10, %r15, 4;
|
|
add.s64 %rd12, %rd9, %rd10;
|
|
ld.shared.u32 %r24, [%rd12];
|
|
|
|
$L__BB8_6:
|
|
bar.sync 0;
|
|
ld.shared.u32 %r16, [%rd1];
|
|
add.s32 %r17, %r16, %r24;
|
|
st.shared.u32 [%rd1], %r17;
|
|
bar.sync 0;
|
|
shl.b32 %r23, %r23, 1;
|
|
setp.lt.u32 %p4, %r23, %r1;
|
|
@%p4 bra $L__BB8_4;
|
|
|
|
$L__BB8_7:
|
|
@%p1 bra $L__BB8_9;
|
|
|
|
ld.shared.u32 %r18, [%rd1];
|
|
sub.s32 %r19, %r18, %r22;
|
|
cvta.to.global.u64 %rd13, %rd3;
|
|
mul.wide.u32 %rd14, %r4, 4;
|
|
add.s64 %rd15, %rd13, %rd14;
|
|
st.global.u32 [%rd15], %r19;
|
|
|
|
$L__BB8_9:
|
|
add.s32 %r20, %r1, -1;
|
|
setp.ne.s32 %p6, %r3, %r20;
|
|
@%p6 bra $L__BB8_11;
|
|
|
|
ld.shared.u32 %r21, [%rd1];
|
|
cvta.to.global.u64 %rd16, %rd4;
|
|
mul.wide.u32 %rd17, %r2, 4;
|
|
add.s64 %rd18, %rd16, %rd17;
|
|
st.global.u32 [%rd18], %r21;
|
|
|
|
$L__BB8_11:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_exscan
|
|
.visible .entry lif_exscan(
|
|
.param .u64 lif_exscan_param_0,
|
|
.param .u64 lif_exscan_param_1,
|
|
.param .u32 lif_exscan_param_2
|
|
)
|
|
{
|
|
.reg .pred %p<14>;
|
|
.reg .b32 %r<108>;
|
|
.reg .b64 %rd<36>;
|
|
|
|
|
|
ld.param.u64 %rd4, [lif_exscan_param_0];
|
|
ld.param.u64 %rd5, [lif_exscan_param_1];
|
|
ld.param.u32 %r43, [lif_exscan_param_2];
|
|
cvta.to.global.u64 %rd1, %rd4;
|
|
cvta.to.global.u64 %rd2, %rd5;
|
|
mov.u32 %r1, %ntid.x;
|
|
add.s32 %r44, %r1, %r43;
|
|
add.s32 %r45, %r44, -1;
|
|
div.u32 %r46, %r45, %r1;
|
|
mov.u32 %r2, %tid.x;
|
|
mul.lo.s32 %r47, %r46, %r2;
|
|
add.s32 %r48, %r47, %r46;
|
|
min.u32 %r3, %r47, %r43;
|
|
min.u32 %r4, %r48, %r43;
|
|
setp.ge.u32 %p1, %r3, %r4;
|
|
mov.u32 %r98, 0;
|
|
@%p1 bra $L__BB9_6;
|
|
|
|
not.b32 %r51, %r4;
|
|
mov.u32 %r52, -2;
|
|
sub.s32 %r53, %r52, %r3;
|
|
sub.s32 %r54, %r53, %r51;
|
|
sub.s32 %r5, %r4, %r3;
|
|
and.b32 %r97, %r5, 3;
|
|
setp.lt.u32 %p2, %r54, 3;
|
|
mov.u32 %r98, 0;
|
|
mov.u32 %r93, %r3;
|
|
@%p2 bra $L__BB9_4;
|
|
|
|
sub.s32 %r91, %r5, %r97;
|
|
mov.u32 %r98, 0;
|
|
mov.u32 %r93, %r3;
|
|
|
|
$L__BB9_3:
|
|
mul.wide.u32 %rd6, %r93, 4;
|
|
add.s64 %rd7, %rd1, %rd6;
|
|
ld.global.nc.u32 %r56, [%rd7];
|
|
add.s32 %r57, %r56, %r98;
|
|
add.s32 %r58, %r93, 1;
|
|
mul.wide.u32 %rd8, %r58, 4;
|
|
add.s64 %rd9, %rd1, %rd8;
|
|
ld.global.nc.u32 %r59, [%rd9];
|
|
add.s32 %r60, %r59, %r57;
|
|
add.s32 %r61, %r93, 2;
|
|
mul.wide.u32 %rd10, %r61, 4;
|
|
add.s64 %rd11, %rd1, %rd10;
|
|
ld.global.nc.u32 %r62, [%rd11];
|
|
add.s32 %r63, %r62, %r60;
|
|
add.s32 %r64, %r93, 3;
|
|
mul.wide.u32 %rd12, %r64, 4;
|
|
add.s64 %rd13, %rd1, %rd12;
|
|
ld.global.nc.u32 %r65, [%rd13];
|
|
add.s32 %r98, %r65, %r63;
|
|
add.s32 %r93, %r93, 4;
|
|
add.s32 %r91, %r91, -4;
|
|
setp.ne.s32 %p3, %r91, 0;
|
|
@%p3 bra $L__BB9_3;
|
|
|
|
$L__BB9_4:
|
|
setp.eq.s32 %p4, %r97, 0;
|
|
@%p4 bra $L__BB9_6;
|
|
|
|
$L__BB9_5:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd14, %r93, 4;
|
|
add.s64 %rd15, %rd1, %rd14;
|
|
ld.global.nc.u32 %r66, [%rd15];
|
|
add.s32 %r98, %r66, %r98;
|
|
add.s32 %r93, %r93, 1;
|
|
add.s32 %r97, %r97, -1;
|
|
setp.ne.s32 %p5, %r97, 0;
|
|
@%p5 bra $L__BB9_5;
|
|
|
|
$L__BB9_6:
|
|
mul.wide.u32 %rd16, %r2, 4;
|
|
mov.u64 %rd17, partial;
|
|
add.s64 %rd3, %rd17, %rd16;
|
|
st.shared.u32 [%rd3], %r98;
|
|
bar.sync 0;
|
|
setp.lt.u32 %p6, %r1, 2;
|
|
@%p6 bra $L__BB9_11;
|
|
|
|
mov.u32 %r99, 1;
|
|
|
|
$L__BB9_8:
|
|
setp.lt.u32 %p7, %r2, %r99;
|
|
mov.u32 %r100, 0;
|
|
@%p7 bra $L__BB9_10;
|
|
|
|
sub.s32 %r69, %r2, %r99;
|
|
mul.wide.u32 %rd18, %r69, 4;
|
|
add.s64 %rd20, %rd17, %rd18;
|
|
ld.shared.u32 %r100, [%rd20];
|
|
|
|
$L__BB9_10:
|
|
bar.sync 0;
|
|
ld.shared.u32 %r70, [%rd3];
|
|
add.s32 %r71, %r70, %r100;
|
|
st.shared.u32 [%rd3], %r71;
|
|
bar.sync 0;
|
|
shl.b32 %r99, %r99, 1;
|
|
setp.lt.u32 %p8, %r99, %r1;
|
|
@%p8 bra $L__BB9_8;
|
|
|
|
$L__BB9_11:
|
|
ld.shared.u32 %r72, [%rd3];
|
|
sub.s32 %r105, %r72, %r98;
|
|
@%p1 bra $L__BB9_17;
|
|
|
|
sub.s32 %r73, %r4, %r3;
|
|
and.b32 %r103, %r73, 3;
|
|
setp.eq.s32 %p10, %r103, 0;
|
|
mov.u32 %r104, %r3;
|
|
@%p10 bra $L__BB9_15;
|
|
|
|
mov.u32 %r104, %r3;
|
|
|
|
$L__BB9_14:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd21, %r104, 4;
|
|
add.s64 %rd22, %rd2, %rd21;
|
|
st.global.u32 [%rd22], %r105;
|
|
add.s64 %rd23, %rd1, %rd21;
|
|
ld.global.nc.u32 %r74, [%rd23];
|
|
add.s32 %r105, %r74, %r105;
|
|
add.s32 %r104, %r104, 1;
|
|
add.s32 %r103, %r103, -1;
|
|
setp.ne.s32 %p11, %r103, 0;
|
|
@%p11 bra $L__BB9_14;
|
|
|
|
$L__BB9_15:
|
|
mov.u32 %r75, -2;
|
|
sub.s32 %r76, %r75, %r3;
|
|
not.b32 %r77, %r4;
|
|
sub.s32 %r78, %r76, %r77;
|
|
setp.lt.u32 %p12, %r78, 3;
|
|
@%p12 bra $L__BB9_17;
|
|
|
|
$L__BB9_16:
|
|
mul.wide.u32 %rd24, %r104, 4;
|
|
add.s64 %rd25, %rd2, %rd24;
|
|
st.global.u32 [%rd25], %r105;
|
|
add.s64 %rd26, %rd1, %rd24;
|
|
ld.global.nc.u32 %r79, [%rd26];
|
|
add.s32 %r80, %r79, %r105;
|
|
add.s32 %r81, %r104, 1;
|
|
mul.wide.u32 %rd27, %r81, 4;
|
|
add.s64 %rd28, %rd2, %rd27;
|
|
st.global.u32 [%rd28], %r80;
|
|
add.s64 %rd29, %rd1, %rd27;
|
|
ld.global.nc.u32 %r82, [%rd29];
|
|
add.s32 %r83, %r82, %r80;
|
|
add.s32 %r84, %r104, 2;
|
|
mul.wide.u32 %rd30, %r84, 4;
|
|
add.s64 %rd31, %rd2, %rd30;
|
|
st.global.u32 [%rd31], %r83;
|
|
add.s64 %rd32, %rd1, %rd30;
|
|
ld.global.nc.u32 %r85, [%rd32];
|
|
add.s32 %r86, %r85, %r83;
|
|
add.s32 %r87, %r104, 3;
|
|
mul.wide.u32 %rd33, %r87, 4;
|
|
add.s64 %rd34, %rd2, %rd33;
|
|
st.global.u32 [%rd34], %r86;
|
|
add.s64 %rd35, %rd1, %rd33;
|
|
ld.global.nc.u32 %r88, [%rd35];
|
|
add.s32 %r105, %r88, %r86;
|
|
add.s32 %r104, %r104, 4;
|
|
setp.lt.u32 %p13, %r104, %r4;
|
|
@%p13 bra $L__BB9_16;
|
|
|
|
$L__BB9_17:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_bucket_rebase
|
|
.visible .entry lif_bucket_rebase(
|
|
.param .u64 lif_bucket_rebase_param_0,
|
|
.param .u64 lif_bucket_rebase_param_1,
|
|
.param .u64 lif_bucket_rebase_param_2,
|
|
.param .u32 lif_bucket_rebase_param_3
|
|
)
|
|
{
|
|
.reg .pred %p<2>;
|
|
.reg .b32 %r<9>;
|
|
.reg .b64 %rd<12>;
|
|
|
|
|
|
ld.param.u64 %rd1, [lif_bucket_rebase_param_0];
|
|
ld.param.u64 %rd2, [lif_bucket_rebase_param_1];
|
|
ld.param.u64 %rd3, [lif_bucket_rebase_param_2];
|
|
ld.param.u32 %r3, [lif_bucket_rebase_param_3];
|
|
mov.u32 %r1, %ctaid.x;
|
|
mov.u32 %r4, %ntid.x;
|
|
mov.u32 %r5, %tid.x;
|
|
mad.lo.s32 %r2, %r1, %r4, %r5;
|
|
setp.ge.u32 %p1, %r2, %r3;
|
|
@%p1 bra $L__BB10_2;
|
|
|
|
cvta.to.global.u64 %rd4, %rd1;
|
|
mul.wide.u32 %rd5, %r2, 4;
|
|
add.s64 %rd6, %rd4, %rd5;
|
|
cvta.to.global.u64 %rd7, %rd3;
|
|
mul.wide.u32 %rd8, %r1, 4;
|
|
add.s64 %rd9, %rd7, %rd8;
|
|
ld.global.nc.u32 %r6, [%rd9];
|
|
ld.global.u32 %r7, [%rd6];
|
|
add.s32 %r8, %r6, %r7;
|
|
st.global.u32 [%rd6], %r8;
|
|
cvta.to.global.u64 %rd10, %rd2;
|
|
add.s64 %rd11, %rd10, %rd5;
|
|
st.global.u32 [%rd11], %r8;
|
|
|
|
$L__BB10_2:
|
|
ret;
|
|
|
|
}
|
|
// .globl lif_bucket_apply
|
|
.visible .entry lif_bucket_apply(
|
|
.param .u64 lif_bucket_apply_param_0,
|
|
.param .u64 lif_bucket_apply_param_1,
|
|
.param .u64 lif_bucket_apply_param_2,
|
|
.param .u64 lif_bucket_apply_param_3,
|
|
.param .u64 lif_bucket_apply_param_4,
|
|
.param .u32 lif_bucket_apply_param_5,
|
|
.param .f64 lif_bucket_apply_param_6,
|
|
.param .f64 lif_bucket_apply_param_7
|
|
)
|
|
{
|
|
.reg .pred %p<133>;
|
|
.reg .b16 %rs<16>;
|
|
.reg .f32 %f<61>;
|
|
.reg .b32 %r<223>;
|
|
.reg .f64 %fd<188>;
|
|
.reg .b64 %rd<166>;
|
|
|
|
|
|
ld.param.u64 %rd44, [lif_bucket_apply_param_0];
|
|
ld.param.u64 %rd45, [lif_bucket_apply_param_1];
|
|
ld.param.u64 %rd46, [lif_bucket_apply_param_2];
|
|
ld.param.u64 %rd47, [lif_bucket_apply_param_3];
|
|
ld.param.u64 %rd48, [lif_bucket_apply_param_4];
|
|
ld.param.u32 %r101, [lif_bucket_apply_param_5];
|
|
ld.param.f64 %fd66, [lif_bucket_apply_param_6];
|
|
ld.param.f64 %fd67, [lif_bucket_apply_param_7];
|
|
cvta.to.global.u64 %rd1, %rd48;
|
|
cvta.to.global.u64 %rd2, %rd47;
|
|
mov.u32 %r102, %ntid.x;
|
|
mov.u32 %r103, %ctaid.x;
|
|
mov.u32 %r104, %tid.x;
|
|
mad.lo.s32 %r1, %r103, %r102, %r104;
|
|
setp.ge.u32 %p17, %r1, %r101;
|
|
@%p17 bra $L__BB11_154;
|
|
|
|
cvta.to.global.u64 %rd49, %rd45;
|
|
cvt.u64.u32 %rd3, %r1;
|
|
mul.wide.u32 %rd50, %r1, 4;
|
|
add.s64 %rd51, %rd49, %rd50;
|
|
ld.global.nc.u32 %r2, [%rd51];
|
|
setp.eq.s32 %p18, %r2, 0;
|
|
@%p18 bra $L__BB11_154;
|
|
|
|
cvta.to.global.u64 %rd52, %rd46;
|
|
shl.b64 %rd53, %rd3, 2;
|
|
add.s64 %rd54, %rd52, %rd53;
|
|
ld.global.nc.u32 %r3, [%rd54];
|
|
setp.lt.u32 %p19, %r2, 702;
|
|
@%p19 bra $L__BB11_9;
|
|
|
|
mov.u32 %r191, 701;
|
|
|
|
$L__BB11_4:
|
|
add.s32 %r106, %r191, %r3;
|
|
mul.wide.u32 %rd55, %r106, 8;
|
|
add.s64 %rd56, %rd2, %rd55;
|
|
ld.global.u64 %rd4, [%rd56];
|
|
mov.u32 %r192, %r191;
|
|
bra.uni $L__BB11_5;
|
|
|
|
$L__BB11_162:
|
|
ld.global.u64 %rd165, [%rd5];
|
|
st.global.u64 [%rd6], %rd165;
|
|
add.s32 %r192, %r192, -701;
|
|
|
|
$L__BB11_5:
|
|
add.s32 %r6, %r192, %r3;
|
|
add.s32 %r107, %r6, -701;
|
|
mul.wide.u32 %rd57, %r107, 8;
|
|
add.s64 %rd5, %rd2, %rd57;
|
|
setp.lt.u32 %p21, %r192, 701;
|
|
mov.pred %p125, 0;
|
|
@%p21 bra $L__BB11_7;
|
|
|
|
ld.global.u64 %rd58, [%rd5];
|
|
setp.gt.u64 %p125, %rd58, %rd4;
|
|
|
|
$L__BB11_7:
|
|
mul.wide.u32 %rd59, %r6, 8;
|
|
add.s64 %rd6, %rd2, %rd59;
|
|
@%p125 bra $L__BB11_162;
|
|
|
|
st.global.u64 [%rd6], %rd4;
|
|
add.s32 %r191, %r191, 1;
|
|
setp.lt.u32 %p22, %r191, %r2;
|
|
@%p22 bra $L__BB11_4;
|
|
|
|
$L__BB11_9:
|
|
setp.lt.u32 %p23, %r2, 302;
|
|
@%p23 bra $L__BB11_16;
|
|
|
|
mov.u32 %r193, 301;
|
|
|
|
$L__BB11_11:
|
|
add.s32 %r109, %r193, %r3;
|
|
mul.wide.u32 %rd60, %r109, 8;
|
|
add.s64 %rd61, %rd2, %rd60;
|
|
ld.global.u64 %rd7, [%rd61];
|
|
mov.u32 %r194, %r193;
|
|
bra.uni $L__BB11_12;
|
|
|
|
$L__BB11_161:
|
|
ld.global.u64 %rd164, [%rd8];
|
|
st.global.u64 [%rd9], %rd164;
|
|
add.s32 %r194, %r194, -301;
|
|
|
|
$L__BB11_12:
|
|
add.s32 %r10, %r194, %r3;
|
|
add.s32 %r110, %r10, -301;
|
|
mul.wide.u32 %rd62, %r110, 8;
|
|
add.s64 %rd8, %rd2, %rd62;
|
|
setp.lt.u32 %p25, %r194, 301;
|
|
mov.pred %p126, 0;
|
|
@%p25 bra $L__BB11_14;
|
|
|
|
ld.global.u64 %rd63, [%rd8];
|
|
setp.gt.u64 %p126, %rd63, %rd7;
|
|
|
|
$L__BB11_14:
|
|
mul.wide.u32 %rd64, %r10, 8;
|
|
add.s64 %rd9, %rd2, %rd64;
|
|
@%p126 bra $L__BB11_161;
|
|
|
|
st.global.u64 [%rd9], %rd7;
|
|
add.s32 %r193, %r193, 1;
|
|
setp.lt.u32 %p26, %r193, %r2;
|
|
@%p26 bra $L__BB11_11;
|
|
|
|
$L__BB11_16:
|
|
setp.lt.u32 %p27, %r2, 133;
|
|
@%p27 bra $L__BB11_23;
|
|
|
|
mov.u32 %r195, 132;
|
|
|
|
$L__BB11_18:
|
|
add.s32 %r112, %r195, %r3;
|
|
mul.wide.u32 %rd65, %r112, 8;
|
|
add.s64 %rd66, %rd2, %rd65;
|
|
ld.global.u64 %rd10, [%rd66];
|
|
mov.u32 %r196, %r195;
|
|
bra.uni $L__BB11_19;
|
|
|
|
$L__BB11_160:
|
|
ld.global.u64 %rd163, [%rd11];
|
|
st.global.u64 [%rd12], %rd163;
|
|
add.s32 %r196, %r196, -132;
|
|
|
|
$L__BB11_19:
|
|
add.s32 %r14, %r196, %r3;
|
|
add.s32 %r113, %r14, -132;
|
|
mul.wide.u32 %rd67, %r113, 8;
|
|
add.s64 %rd11, %rd2, %rd67;
|
|
setp.lt.u32 %p29, %r196, 132;
|
|
mov.pred %p127, 0;
|
|
@%p29 bra $L__BB11_21;
|
|
|
|
ld.global.u64 %rd68, [%rd11];
|
|
setp.gt.u64 %p127, %rd68, %rd10;
|
|
|
|
$L__BB11_21:
|
|
mul.wide.u32 %rd69, %r14, 8;
|
|
add.s64 %rd12, %rd2, %rd69;
|
|
@%p127 bra $L__BB11_160;
|
|
|
|
st.global.u64 [%rd12], %rd10;
|
|
add.s32 %r195, %r195, 1;
|
|
setp.lt.u32 %p30, %r195, %r2;
|
|
@%p30 bra $L__BB11_18;
|
|
|
|
$L__BB11_23:
|
|
setp.lt.u32 %p31, %r2, 58;
|
|
@%p31 bra $L__BB11_30;
|
|
|
|
mov.u32 %r197, 57;
|
|
|
|
$L__BB11_25:
|
|
add.s32 %r115, %r197, %r3;
|
|
mul.wide.u32 %rd70, %r115, 8;
|
|
add.s64 %rd71, %rd2, %rd70;
|
|
ld.global.u64 %rd13, [%rd71];
|
|
mov.u32 %r198, %r197;
|
|
bra.uni $L__BB11_26;
|
|
|
|
$L__BB11_159:
|
|
ld.global.u64 %rd162, [%rd14];
|
|
st.global.u64 [%rd15], %rd162;
|
|
add.s32 %r198, %r198, -57;
|
|
|
|
$L__BB11_26:
|
|
add.s32 %r18, %r198, %r3;
|
|
add.s32 %r116, %r18, -57;
|
|
mul.wide.u32 %rd72, %r116, 8;
|
|
add.s64 %rd14, %rd2, %rd72;
|
|
setp.lt.u32 %p33, %r198, 57;
|
|
mov.pred %p128, 0;
|
|
@%p33 bra $L__BB11_28;
|
|
|
|
ld.global.u64 %rd73, [%rd14];
|
|
setp.gt.u64 %p128, %rd73, %rd13;
|
|
|
|
$L__BB11_28:
|
|
mul.wide.u32 %rd74, %r18, 8;
|
|
add.s64 %rd15, %rd2, %rd74;
|
|
@%p128 bra $L__BB11_159;
|
|
|
|
st.global.u64 [%rd15], %rd13;
|
|
add.s32 %r197, %r197, 1;
|
|
setp.lt.u32 %p34, %r197, %r2;
|
|
@%p34 bra $L__BB11_25;
|
|
|
|
$L__BB11_30:
|
|
setp.lt.u32 %p35, %r2, 24;
|
|
@%p35 bra $L__BB11_37;
|
|
|
|
mov.u32 %r199, 23;
|
|
|
|
$L__BB11_32:
|
|
add.s32 %r118, %r199, %r3;
|
|
mul.wide.u32 %rd75, %r118, 8;
|
|
add.s64 %rd76, %rd2, %rd75;
|
|
ld.global.u64 %rd16, [%rd76];
|
|
mov.u32 %r200, %r199;
|
|
bra.uni $L__BB11_33;
|
|
|
|
$L__BB11_158:
|
|
ld.global.u64 %rd161, [%rd17];
|
|
st.global.u64 [%rd18], %rd161;
|
|
add.s32 %r200, %r200, -23;
|
|
|
|
$L__BB11_33:
|
|
add.s32 %r22, %r200, %r3;
|
|
add.s32 %r119, %r22, -23;
|
|
mul.wide.u32 %rd77, %r119, 8;
|
|
add.s64 %rd17, %rd2, %rd77;
|
|
setp.lt.u32 %p37, %r200, 23;
|
|
mov.pred %p129, 0;
|
|
@%p37 bra $L__BB11_35;
|
|
|
|
ld.global.u64 %rd78, [%rd17];
|
|
setp.gt.u64 %p129, %rd78, %rd16;
|
|
|
|
$L__BB11_35:
|
|
mul.wide.u32 %rd79, %r22, 8;
|
|
add.s64 %rd18, %rd2, %rd79;
|
|
@%p129 bra $L__BB11_158;
|
|
|
|
st.global.u64 [%rd18], %rd16;
|
|
add.s32 %r199, %r199, 1;
|
|
setp.lt.u32 %p38, %r199, %r2;
|
|
@%p38 bra $L__BB11_32;
|
|
|
|
$L__BB11_37:
|
|
setp.lt.u32 %p39, %r2, 11;
|
|
@%p39 bra $L__BB11_44;
|
|
|
|
mov.u32 %r201, 10;
|
|
|
|
$L__BB11_39:
|
|
add.s32 %r121, %r201, %r3;
|
|
mul.wide.u32 %rd80, %r121, 8;
|
|
add.s64 %rd81, %rd2, %rd80;
|
|
ld.global.u64 %rd19, [%rd81];
|
|
mov.u32 %r202, %r201;
|
|
bra.uni $L__BB11_40;
|
|
|
|
$L__BB11_157:
|
|
ld.global.u64 %rd160, [%rd20];
|
|
st.global.u64 [%rd21], %rd160;
|
|
add.s32 %r202, %r202, -10;
|
|
|
|
$L__BB11_40:
|
|
add.s32 %r26, %r202, %r3;
|
|
add.s32 %r122, %r26, -10;
|
|
mul.wide.u32 %rd82, %r122, 8;
|
|
add.s64 %rd20, %rd2, %rd82;
|
|
setp.lt.u32 %p41, %r202, 10;
|
|
mov.pred %p130, 0;
|
|
@%p41 bra $L__BB11_42;
|
|
|
|
ld.global.u64 %rd83, [%rd20];
|
|
setp.gt.u64 %p130, %rd83, %rd19;
|
|
|
|
$L__BB11_42:
|
|
mul.wide.u32 %rd84, %r26, 8;
|
|
add.s64 %rd21, %rd2, %rd84;
|
|
@%p130 bra $L__BB11_157;
|
|
|
|
st.global.u64 [%rd21], %rd19;
|
|
add.s32 %r201, %r201, 1;
|
|
setp.lt.u32 %p42, %r201, %r2;
|
|
@%p42 bra $L__BB11_39;
|
|
|
|
$L__BB11_44:
|
|
setp.lt.u32 %p43, %r2, 5;
|
|
@%p43 bra $L__BB11_51;
|
|
|
|
mov.u32 %r203, 4;
|
|
|
|
$L__BB11_46:
|
|
add.s32 %r124, %r203, %r3;
|
|
mul.wide.u32 %rd85, %r124, 8;
|
|
add.s64 %rd86, %rd2, %rd85;
|
|
ld.global.u64 %rd22, [%rd86];
|
|
mov.u32 %r204, %r203;
|
|
bra.uni $L__BB11_47;
|
|
|
|
$L__BB11_156:
|
|
ld.global.u64 %rd159, [%rd23];
|
|
st.global.u64 [%rd24], %rd159;
|
|
add.s32 %r204, %r204, -4;
|
|
|
|
$L__BB11_47:
|
|
add.s32 %r30, %r204, %r3;
|
|
add.s32 %r125, %r30, -4;
|
|
mul.wide.u32 %rd87, %r125, 8;
|
|
add.s64 %rd23, %rd2, %rd87;
|
|
setp.lt.u32 %p45, %r204, 4;
|
|
mov.pred %p131, 0;
|
|
@%p45 bra $L__BB11_49;
|
|
|
|
ld.global.u64 %rd88, [%rd23];
|
|
setp.gt.u64 %p131, %rd88, %rd22;
|
|
|
|
$L__BB11_49:
|
|
mul.wide.u32 %rd89, %r30, 8;
|
|
add.s64 %rd24, %rd2, %rd89;
|
|
@%p131 bra $L__BB11_156;
|
|
|
|
st.global.u64 [%rd24], %rd22;
|
|
add.s32 %r203, %r203, 1;
|
|
setp.lt.u32 %p46, %r203, %r2;
|
|
@%p46 bra $L__BB11_46;
|
|
|
|
$L__BB11_51:
|
|
setp.lt.u32 %p47, %r2, 2;
|
|
@%p47 bra $L__BB11_58;
|
|
|
|
mov.u32 %r205, 1;
|
|
|
|
$L__BB11_53:
|
|
add.s32 %r206, %r205, %r3;
|
|
mul.wide.u32 %rd90, %r206, 8;
|
|
add.s64 %rd91, %rd2, %rd90;
|
|
ld.global.u64 %rd25, [%rd91];
|
|
add.s32 %r207, %r205, -1;
|
|
bra.uni $L__BB11_54;
|
|
|
|
$L__BB11_155:
|
|
ld.global.u64 %rd158, [%rd26];
|
|
st.global.u64 [%rd27], %rd158;
|
|
add.s32 %r207, %r207, -1;
|
|
mov.u32 %r206, %r37;
|
|
|
|
$L__BB11_54:
|
|
add.s32 %r37, %r206, -1;
|
|
mul.wide.u32 %rd92, %r37, 8;
|
|
add.s64 %rd26, %rd2, %rd92;
|
|
setp.eq.s32 %p49, %r207, -1;
|
|
mov.pred %p132, 0;
|
|
@%p49 bra $L__BB11_56;
|
|
|
|
ld.global.u64 %rd93, [%rd26];
|
|
setp.gt.u64 %p132, %rd93, %rd25;
|
|
|
|
$L__BB11_56:
|
|
mul.wide.u32 %rd94, %r206, 8;
|
|
add.s64 %rd27, %rd2, %rd94;
|
|
@%p132 bra $L__BB11_155;
|
|
|
|
st.global.u64 [%rd27], %rd25;
|
|
add.s32 %r205, %r205, 1;
|
|
setp.lt.u32 %p50, %r205, %r2;
|
|
@%p50 bra $L__BB11_53;
|
|
|
|
$L__BB11_58:
|
|
cvta.to.global.u64 %rd95, %rd44;
|
|
add.s64 %rd28, %rd95, %rd53;
|
|
ld.global.f32 %f60, [%rd28];
|
|
setp.eq.f64 %p51, %fd67, 0d0000000000000000;
|
|
@%p51 bra $L__BB11_90;
|
|
bra.uni $L__BB11_59;
|
|
|
|
$L__BB11_90:
|
|
mov.b64 %rd117, %fd67;
|
|
setp.gt.s64 %p71, %rd117, -1;
|
|
add.s32 %r57, %r2, -1;
|
|
and.b32 %r217, %r2, 3;
|
|
@%p71 bra $L__BB11_122;
|
|
bra.uni $L__BB11_91;
|
|
|
|
$L__BB11_122:
|
|
setp.lt.u32 %p106, %r57, 3;
|
|
mov.u32 %r220, 0;
|
|
@%p106 bra $L__BB11_145;
|
|
|
|
add.s32 %r218, %r3, 3;
|
|
sub.s32 %r77, %r217, %r2;
|
|
mov.u32 %r220, 0;
|
|
|
|
$L__BB11_124:
|
|
add.s32 %r172, %r218, -3;
|
|
mul.wide.u32 %rd138, %r172, 8;
|
|
add.s64 %rd139, %rd2, %rd138;
|
|
ld.global.u64 %rd39, [%rd139];
|
|
cvt.u32.u64 %r173, %rd39;
|
|
shr.u32 %r174, %r173, 16;
|
|
and.b32 %r80, %r174, 32767;
|
|
setp.eq.s32 %p107, %r80, 0;
|
|
mov.f64 %fd178, 0d3FF0000000000000;
|
|
@%p107 bra $L__BB11_126;
|
|
|
|
add.s32 %r175, %r80, -1;
|
|
mul.wide.u32 %rd140, %r175, 4;
|
|
add.s64 %rd141, %rd1, %rd140;
|
|
ld.global.nc.f32 %f40, [%rd141];
|
|
cvt.f64.f32 %fd178, %f40;
|
|
|
|
$L__BB11_126:
|
|
cvt.u16.u64 %rs11, %rd39;
|
|
cvt.rn.f64.s16 %fd129, %rs11;
|
|
mul.rn.f64 %fd130, %fd178, %fd129;
|
|
mul.rn.f64 %fd131, %fd130, %fd66;
|
|
cvt.f64.f32 %fd132, %f60;
|
|
add.rn.f64 %fd179, %fd131, %fd132;
|
|
setp.gt.f64 %p108, %fd179, %fd67;
|
|
@%p108 bra $L__BB11_129;
|
|
|
|
setp.nan.f64 %p109, %fd179, %fd67;
|
|
mov.f64 %fd179, 0d7FF8000000000000;
|
|
@%p109 bra $L__BB11_129;
|
|
|
|
mov.f64 %fd179, %fd67;
|
|
|
|
$L__BB11_129:
|
|
add.s32 %r176, %r218, -2;
|
|
mul.wide.u32 %rd142, %r176, 8;
|
|
add.s64 %rd143, %rd2, %rd142;
|
|
ld.global.u64 %rd40, [%rd143];
|
|
cvt.u32.u64 %r177, %rd40;
|
|
shr.u32 %r178, %r177, 16;
|
|
and.b32 %r81, %r178, 32767;
|
|
setp.eq.s32 %p110, %r81, 0;
|
|
mov.f64 %fd180, 0d3FF0000000000000;
|
|
@%p110 bra $L__BB11_131;
|
|
|
|
add.s32 %r179, %r81, -1;
|
|
mul.wide.u32 %rd144, %r179, 4;
|
|
add.s64 %rd145, %rd1, %rd144;
|
|
ld.global.nc.f32 %f41, [%rd145];
|
|
cvt.f64.f32 %fd180, %f41;
|
|
|
|
$L__BB11_131:
|
|
cvt.u16.u64 %rs12, %rd40;
|
|
cvt.rn.f64.s16 %fd135, %rs12;
|
|
mul.rn.f64 %fd136, %fd180, %fd135;
|
|
mul.rn.f64 %fd137, %fd136, %fd66;
|
|
cvt.rn.f32.f64 %f42, %fd179;
|
|
cvt.f64.f32 %fd138, %f42;
|
|
add.rn.f64 %fd181, %fd137, %fd138;
|
|
setp.gt.f64 %p111, %fd181, %fd67;
|
|
@%p111 bra $L__BB11_134;
|
|
|
|
setp.nan.f64 %p112, %fd181, %fd67;
|
|
mov.f64 %fd181, 0d7FF8000000000000;
|
|
@%p112 bra $L__BB11_134;
|
|
|
|
mov.f64 %fd181, %fd67;
|
|
|
|
$L__BB11_134:
|
|
add.s32 %r180, %r218, -1;
|
|
mul.wide.u32 %rd146, %r180, 8;
|
|
add.s64 %rd147, %rd2, %rd146;
|
|
ld.global.u64 %rd41, [%rd147];
|
|
cvt.u32.u64 %r181, %rd41;
|
|
shr.u32 %r182, %r181, 16;
|
|
and.b32 %r82, %r182, 32767;
|
|
setp.eq.s32 %p113, %r82, 0;
|
|
mov.f64 %fd182, 0d3FF0000000000000;
|
|
@%p113 bra $L__BB11_136;
|
|
|
|
add.s32 %r183, %r82, -1;
|
|
mul.wide.u32 %rd148, %r183, 4;
|
|
add.s64 %rd149, %rd1, %rd148;
|
|
ld.global.nc.f32 %f43, [%rd149];
|
|
cvt.f64.f32 %fd182, %f43;
|
|
|
|
$L__BB11_136:
|
|
cvt.u16.u64 %rs13, %rd41;
|
|
cvt.rn.f64.s16 %fd141, %rs13;
|
|
mul.rn.f64 %fd142, %fd182, %fd141;
|
|
mul.rn.f64 %fd143, %fd142, %fd66;
|
|
cvt.rn.f32.f64 %f44, %fd181;
|
|
cvt.f64.f32 %fd144, %f44;
|
|
add.rn.f64 %fd183, %fd143, %fd144;
|
|
setp.gt.f64 %p114, %fd183, %fd67;
|
|
@%p114 bra $L__BB11_139;
|
|
|
|
setp.nan.f64 %p115, %fd183, %fd67;
|
|
mov.f64 %fd183, 0d7FF8000000000000;
|
|
@%p115 bra $L__BB11_139;
|
|
|
|
mov.f64 %fd183, %fd67;
|
|
|
|
$L__BB11_139:
|
|
mul.wide.u32 %rd150, %r218, 8;
|
|
add.s64 %rd151, %rd2, %rd150;
|
|
ld.global.u64 %rd42, [%rd151];
|
|
cvt.u32.u64 %r184, %rd42;
|
|
shr.u32 %r185, %r184, 16;
|
|
and.b32 %r83, %r185, 32767;
|
|
setp.eq.s32 %p116, %r83, 0;
|
|
mov.f64 %fd184, 0d3FF0000000000000;
|
|
@%p116 bra $L__BB11_141;
|
|
|
|
add.s32 %r186, %r83, -1;
|
|
mul.wide.u32 %rd152, %r186, 4;
|
|
add.s64 %rd153, %rd1, %rd152;
|
|
ld.global.nc.f32 %f45, [%rd153];
|
|
cvt.f64.f32 %fd184, %f45;
|
|
|
|
$L__BB11_141:
|
|
cvt.u16.u64 %rs14, %rd42;
|
|
cvt.rn.f64.s16 %fd147, %rs14;
|
|
mul.rn.f64 %fd148, %fd184, %fd147;
|
|
mul.rn.f64 %fd149, %fd148, %fd66;
|
|
cvt.rn.f32.f64 %f46, %fd183;
|
|
cvt.f64.f32 %fd150, %f46;
|
|
add.rn.f64 %fd185, %fd149, %fd150;
|
|
setp.gt.f64 %p117, %fd185, %fd67;
|
|
@%p117 bra $L__BB11_144;
|
|
|
|
setp.nan.f64 %p118, %fd185, %fd67;
|
|
mov.f64 %fd185, 0d7FF8000000000000;
|
|
@%p118 bra $L__BB11_144;
|
|
|
|
mov.f64 %fd185, %fd67;
|
|
|
|
$L__BB11_144:
|
|
cvt.rn.f32.f64 %f60, %fd185;
|
|
add.s32 %r218, %r218, 4;
|
|
add.s32 %r220, %r220, 4;
|
|
add.s32 %r187, %r77, %r220;
|
|
setp.ne.s32 %p119, %r187, 0;
|
|
@%p119 bra $L__BB11_124;
|
|
|
|
$L__BB11_145:
|
|
setp.eq.s32 %p120, %r217, 0;
|
|
@%p120 bra $L__BB11_153;
|
|
|
|
add.s32 %r221, %r220, %r3;
|
|
|
|
$L__BB11_147:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd154, %r221, 8;
|
|
add.s64 %rd155, %rd2, %rd154;
|
|
ld.global.u64 %rd43, [%rd155];
|
|
cvt.u32.u64 %r188, %rd43;
|
|
shr.u32 %r189, %r188, 16;
|
|
and.b32 %r90, %r189, 32767;
|
|
setp.eq.s32 %p121, %r90, 0;
|
|
mov.f64 %fd186, 0d3FF0000000000000;
|
|
@%p121 bra $L__BB11_149;
|
|
|
|
add.s32 %r190, %r90, -1;
|
|
mul.wide.u32 %rd156, %r190, 4;
|
|
add.s64 %rd157, %rd1, %rd156;
|
|
ld.global.nc.f32 %f47, [%rd157];
|
|
cvt.f64.f32 %fd186, %f47;
|
|
|
|
$L__BB11_149:
|
|
cvt.u16.u64 %rs15, %rd43;
|
|
cvt.rn.f64.s16 %fd153, %rs15;
|
|
mul.rn.f64 %fd154, %fd186, %fd153;
|
|
mul.rn.f64 %fd155, %fd154, %fd66;
|
|
cvt.f64.f32 %fd156, %f60;
|
|
add.rn.f64 %fd187, %fd155, %fd156;
|
|
setp.gt.f64 %p122, %fd187, %fd67;
|
|
@%p122 bra $L__BB11_152;
|
|
|
|
setp.nan.f64 %p123, %fd187, %fd67;
|
|
mov.f64 %fd187, 0d7FF8000000000000;
|
|
@%p123 bra $L__BB11_152;
|
|
|
|
mov.f64 %fd187, %fd67;
|
|
|
|
$L__BB11_152:
|
|
cvt.rn.f32.f64 %f60, %fd187;
|
|
add.s32 %r221, %r221, 1;
|
|
add.s32 %r217, %r217, -1;
|
|
setp.ne.s32 %p124, %r217, 0;
|
|
@%p124 bra $L__BB11_147;
|
|
bra.uni $L__BB11_153;
|
|
|
|
$L__BB11_59:
|
|
add.s32 %r128, %r2, -1;
|
|
and.b32 %r212, %r2, 3;
|
|
setp.lt.u32 %p52, %r128, 3;
|
|
mov.u32 %r210, 0;
|
|
@%p52 bra $L__BB11_82;
|
|
|
|
add.s32 %r208, %r3, 3;
|
|
sub.s32 %r41, %r212, %r2;
|
|
mov.u32 %r210, 0;
|
|
|
|
$L__BB11_61:
|
|
add.s32 %r130, %r208, -3;
|
|
mul.wide.u32 %rd97, %r130, 8;
|
|
add.s64 %rd98, %rd2, %rd97;
|
|
ld.global.u64 %rd29, [%rd98];
|
|
cvt.u32.u64 %r131, %rd29;
|
|
shr.u32 %r132, %r131, 16;
|
|
and.b32 %r44, %r132, 32767;
|
|
setp.eq.s32 %p53, %r44, 0;
|
|
mov.f64 %fd158, 0d3FF0000000000000;
|
|
@%p53 bra $L__BB11_63;
|
|
|
|
add.s32 %r133, %r44, -1;
|
|
mul.wide.u32 %rd99, %r133, 4;
|
|
add.s64 %rd100, %rd1, %rd99;
|
|
ld.global.nc.f32 %f22, [%rd100];
|
|
cvt.f64.f32 %fd158, %f22;
|
|
|
|
$L__BB11_63:
|
|
cvt.u16.u64 %rs1, %rd29;
|
|
cvt.rn.f64.s16 %fd69, %rs1;
|
|
mul.rn.f64 %fd70, %fd158, %fd69;
|
|
mul.rn.f64 %fd71, %fd70, %fd66;
|
|
cvt.f64.f32 %fd72, %f60;
|
|
add.rn.f64 %fd159, %fd71, %fd72;
|
|
setp.gt.f64 %p54, %fd159, %fd67;
|
|
@%p54 bra $L__BB11_66;
|
|
|
|
setp.nan.f64 %p55, %fd159, %fd67;
|
|
mov.f64 %fd159, 0d7FF8000000000000;
|
|
@%p55 bra $L__BB11_66;
|
|
|
|
mov.f64 %fd159, %fd67;
|
|
|
|
$L__BB11_66:
|
|
add.s32 %r134, %r208, -2;
|
|
mul.wide.u32 %rd101, %r134, 8;
|
|
add.s64 %rd102, %rd2, %rd101;
|
|
ld.global.u64 %rd30, [%rd102];
|
|
cvt.u32.u64 %r135, %rd30;
|
|
shr.u32 %r136, %r135, 16;
|
|
and.b32 %r45, %r136, 32767;
|
|
setp.eq.s32 %p56, %r45, 0;
|
|
mov.f64 %fd160, 0d3FF0000000000000;
|
|
@%p56 bra $L__BB11_68;
|
|
|
|
add.s32 %r137, %r45, -1;
|
|
mul.wide.u32 %rd103, %r137, 4;
|
|
add.s64 %rd104, %rd1, %rd103;
|
|
ld.global.nc.f32 %f23, [%rd104];
|
|
cvt.f64.f32 %fd160, %f23;
|
|
|
|
$L__BB11_68:
|
|
cvt.u16.u64 %rs2, %rd30;
|
|
cvt.rn.f64.s16 %fd75, %rs2;
|
|
mul.rn.f64 %fd76, %fd160, %fd75;
|
|
mul.rn.f64 %fd77, %fd76, %fd66;
|
|
cvt.rn.f32.f64 %f24, %fd159;
|
|
cvt.f64.f32 %fd78, %f24;
|
|
add.rn.f64 %fd161, %fd77, %fd78;
|
|
setp.gt.f64 %p57, %fd161, %fd67;
|
|
@%p57 bra $L__BB11_71;
|
|
|
|
setp.nan.f64 %p58, %fd161, %fd67;
|
|
mov.f64 %fd161, 0d7FF8000000000000;
|
|
@%p58 bra $L__BB11_71;
|
|
|
|
mov.f64 %fd161, %fd67;
|
|
|
|
$L__BB11_71:
|
|
add.s32 %r138, %r208, -1;
|
|
mul.wide.u32 %rd105, %r138, 8;
|
|
add.s64 %rd106, %rd2, %rd105;
|
|
ld.global.u64 %rd31, [%rd106];
|
|
cvt.u32.u64 %r139, %rd31;
|
|
shr.u32 %r140, %r139, 16;
|
|
and.b32 %r46, %r140, 32767;
|
|
setp.eq.s32 %p59, %r46, 0;
|
|
mov.f64 %fd162, 0d3FF0000000000000;
|
|
@%p59 bra $L__BB11_73;
|
|
|
|
add.s32 %r141, %r46, -1;
|
|
mul.wide.u32 %rd107, %r141, 4;
|
|
add.s64 %rd108, %rd1, %rd107;
|
|
ld.global.nc.f32 %f25, [%rd108];
|
|
cvt.f64.f32 %fd162, %f25;
|
|
|
|
$L__BB11_73:
|
|
cvt.u16.u64 %rs3, %rd31;
|
|
cvt.rn.f64.s16 %fd81, %rs3;
|
|
mul.rn.f64 %fd82, %fd162, %fd81;
|
|
mul.rn.f64 %fd83, %fd82, %fd66;
|
|
cvt.rn.f32.f64 %f26, %fd161;
|
|
cvt.f64.f32 %fd84, %f26;
|
|
add.rn.f64 %fd163, %fd83, %fd84;
|
|
setp.gt.f64 %p60, %fd163, %fd67;
|
|
@%p60 bra $L__BB11_76;
|
|
|
|
setp.nan.f64 %p61, %fd163, %fd67;
|
|
mov.f64 %fd163, 0d7FF8000000000000;
|
|
@%p61 bra $L__BB11_76;
|
|
|
|
mov.f64 %fd163, %fd67;
|
|
|
|
$L__BB11_76:
|
|
mul.wide.u32 %rd109, %r208, 8;
|
|
add.s64 %rd110, %rd2, %rd109;
|
|
ld.global.u64 %rd32, [%rd110];
|
|
cvt.u32.u64 %r142, %rd32;
|
|
shr.u32 %r143, %r142, 16;
|
|
and.b32 %r47, %r143, 32767;
|
|
setp.eq.s32 %p62, %r47, 0;
|
|
mov.f64 %fd164, 0d3FF0000000000000;
|
|
@%p62 bra $L__BB11_78;
|
|
|
|
add.s32 %r144, %r47, -1;
|
|
mul.wide.u32 %rd111, %r144, 4;
|
|
add.s64 %rd112, %rd1, %rd111;
|
|
ld.global.nc.f32 %f27, [%rd112];
|
|
cvt.f64.f32 %fd164, %f27;
|
|
|
|
$L__BB11_78:
|
|
cvt.u16.u64 %rs4, %rd32;
|
|
cvt.rn.f64.s16 %fd87, %rs4;
|
|
mul.rn.f64 %fd88, %fd164, %fd87;
|
|
mul.rn.f64 %fd89, %fd88, %fd66;
|
|
cvt.rn.f32.f64 %f28, %fd163;
|
|
cvt.f64.f32 %fd90, %f28;
|
|
add.rn.f64 %fd165, %fd89, %fd90;
|
|
setp.gt.f64 %p63, %fd165, %fd67;
|
|
@%p63 bra $L__BB11_81;
|
|
|
|
setp.nan.f64 %p64, %fd165, %fd67;
|
|
mov.f64 %fd165, 0d7FF8000000000000;
|
|
@%p64 bra $L__BB11_81;
|
|
|
|
mov.f64 %fd165, %fd67;
|
|
|
|
$L__BB11_81:
|
|
cvt.rn.f32.f64 %f60, %fd165;
|
|
add.s32 %r208, %r208, 4;
|
|
add.s32 %r210, %r210, 4;
|
|
add.s32 %r145, %r41, %r210;
|
|
setp.ne.s32 %p65, %r145, 0;
|
|
@%p65 bra $L__BB11_61;
|
|
|
|
$L__BB11_82:
|
|
setp.eq.s32 %p66, %r212, 0;
|
|
@%p66 bra $L__BB11_153;
|
|
|
|
add.s32 %r211, %r210, %r3;
|
|
|
|
$L__BB11_84:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd113, %r211, 8;
|
|
add.s64 %rd114, %rd2, %rd113;
|
|
ld.global.u64 %rd33, [%rd114];
|
|
cvt.u32.u64 %r146, %rd33;
|
|
shr.u32 %r147, %r146, 16;
|
|
and.b32 %r54, %r147, 32767;
|
|
setp.eq.s32 %p67, %r54, 0;
|
|
mov.f64 %fd166, 0d3FF0000000000000;
|
|
@%p67 bra $L__BB11_86;
|
|
|
|
add.s32 %r148, %r54, -1;
|
|
mul.wide.u32 %rd115, %r148, 4;
|
|
add.s64 %rd116, %rd1, %rd115;
|
|
ld.global.nc.f32 %f29, [%rd116];
|
|
cvt.f64.f32 %fd166, %f29;
|
|
|
|
$L__BB11_86:
|
|
cvt.u16.u64 %rs5, %rd33;
|
|
cvt.rn.f64.s16 %fd93, %rs5;
|
|
mul.rn.f64 %fd94, %fd166, %fd93;
|
|
mul.rn.f64 %fd95, %fd94, %fd66;
|
|
cvt.f64.f32 %fd96, %f60;
|
|
add.rn.f64 %fd167, %fd95, %fd96;
|
|
setp.gt.f64 %p68, %fd167, %fd67;
|
|
@%p68 bra $L__BB11_89;
|
|
|
|
setp.nan.f64 %p69, %fd167, %fd67;
|
|
mov.f64 %fd167, 0d7FF8000000000000;
|
|
@%p69 bra $L__BB11_89;
|
|
|
|
mov.f64 %fd167, %fd67;
|
|
|
|
$L__BB11_89:
|
|
cvt.rn.f32.f64 %f60, %fd167;
|
|
add.s32 %r211, %r211, 1;
|
|
add.s32 %r212, %r212, -1;
|
|
setp.eq.s32 %p70, %r212, 0;
|
|
@%p70 bra $L__BB11_153;
|
|
bra.uni $L__BB11_84;
|
|
|
|
$L__BB11_91:
|
|
setp.lt.u32 %p72, %r57, 3;
|
|
mov.u32 %r215, 0;
|
|
@%p72 bra $L__BB11_114;
|
|
|
|
add.s32 %r213, %r3, 3;
|
|
sub.s32 %r60, %r217, %r2;
|
|
mov.u32 %r215, 0;
|
|
|
|
$L__BB11_93:
|
|
add.s32 %r151, %r213, -3;
|
|
mul.wide.u32 %rd118, %r151, 8;
|
|
add.s64 %rd119, %rd2, %rd118;
|
|
ld.global.u64 %rd34, [%rd119];
|
|
cvt.u32.u64 %r152, %rd34;
|
|
shr.u32 %r153, %r152, 16;
|
|
and.b32 %r63, %r153, 32767;
|
|
setp.eq.s32 %p73, %r63, 0;
|
|
mov.f64 %fd168, 0d3FF0000000000000;
|
|
@%p73 bra $L__BB11_95;
|
|
|
|
add.s32 %r154, %r63, -1;
|
|
mul.wide.u32 %rd120, %r154, 4;
|
|
add.s64 %rd121, %rd1, %rd120;
|
|
ld.global.nc.f32 %f31, [%rd121];
|
|
cvt.f64.f32 %fd168, %f31;
|
|
|
|
$L__BB11_95:
|
|
cvt.u16.u64 %rs6, %rd34;
|
|
cvt.rn.f64.s16 %fd99, %rs6;
|
|
mul.rn.f64 %fd100, %fd168, %fd99;
|
|
mul.rn.f64 %fd101, %fd100, %fd66;
|
|
cvt.f64.f32 %fd102, %f60;
|
|
add.rn.f64 %fd23, %fd101, %fd102;
|
|
setp.gt.f64 %p74, %fd23, %fd67;
|
|
mov.f64 %fd169, %fd23;
|
|
@%p74 bra $L__BB11_98;
|
|
|
|
setp.nan.f64 %p75, %fd23, %fd67;
|
|
mov.f64 %fd169, 0d7FF8000000000000;
|
|
@%p75 bra $L__BB11_98;
|
|
|
|
setp.geu.f64 %p76, %fd23, %fd67;
|
|
setp.eq.f64 %p77, %fd23, 0d0000000000000000;
|
|
and.pred %p78, %p76, %p77;
|
|
selp.f64 %fd169, %fd23, %fd67, %p78;
|
|
|
|
$L__BB11_98:
|
|
add.s32 %r155, %r213, -2;
|
|
mul.wide.u32 %rd122, %r155, 8;
|
|
add.s64 %rd123, %rd2, %rd122;
|
|
ld.global.u64 %rd35, [%rd123];
|
|
cvt.u32.u64 %r156, %rd35;
|
|
shr.u32 %r157, %r156, 16;
|
|
and.b32 %r64, %r157, 32767;
|
|
setp.eq.s32 %p79, %r64, 0;
|
|
mov.f64 %fd170, 0d3FF0000000000000;
|
|
@%p79 bra $L__BB11_100;
|
|
|
|
add.s32 %r158, %r64, -1;
|
|
mul.wide.u32 %rd124, %r158, 4;
|
|
add.s64 %rd125, %rd1, %rd124;
|
|
ld.global.nc.f32 %f32, [%rd125];
|
|
cvt.f64.f32 %fd170, %f32;
|
|
|
|
$L__BB11_100:
|
|
cvt.u16.u64 %rs7, %rd35;
|
|
cvt.rn.f64.s16 %fd105, %rs7;
|
|
mul.rn.f64 %fd106, %fd170, %fd105;
|
|
mul.rn.f64 %fd107, %fd106, %fd66;
|
|
cvt.rn.f32.f64 %f33, %fd169;
|
|
cvt.f64.f32 %fd108, %f33;
|
|
add.rn.f64 %fd28, %fd107, %fd108;
|
|
setp.gt.f64 %p80, %fd28, %fd67;
|
|
mov.f64 %fd171, %fd28;
|
|
@%p80 bra $L__BB11_103;
|
|
|
|
setp.nan.f64 %p81, %fd28, %fd67;
|
|
mov.f64 %fd171, 0d7FF8000000000000;
|
|
@%p81 bra $L__BB11_103;
|
|
|
|
setp.geu.f64 %p82, %fd28, %fd67;
|
|
setp.eq.f64 %p83, %fd28, 0d0000000000000000;
|
|
and.pred %p84, %p82, %p83;
|
|
selp.f64 %fd171, %fd28, %fd67, %p84;
|
|
|
|
$L__BB11_103:
|
|
add.s32 %r159, %r213, -1;
|
|
mul.wide.u32 %rd126, %r159, 8;
|
|
add.s64 %rd127, %rd2, %rd126;
|
|
ld.global.u64 %rd36, [%rd127];
|
|
cvt.u32.u64 %r160, %rd36;
|
|
shr.u32 %r161, %r160, 16;
|
|
and.b32 %r65, %r161, 32767;
|
|
setp.eq.s32 %p85, %r65, 0;
|
|
mov.f64 %fd172, 0d3FF0000000000000;
|
|
@%p85 bra $L__BB11_105;
|
|
|
|
add.s32 %r162, %r65, -1;
|
|
mul.wide.u32 %rd128, %r162, 4;
|
|
add.s64 %rd129, %rd1, %rd128;
|
|
ld.global.nc.f32 %f34, [%rd129];
|
|
cvt.f64.f32 %fd172, %f34;
|
|
|
|
$L__BB11_105:
|
|
cvt.u16.u64 %rs8, %rd36;
|
|
cvt.rn.f64.s16 %fd111, %rs8;
|
|
mul.rn.f64 %fd112, %fd172, %fd111;
|
|
mul.rn.f64 %fd113, %fd112, %fd66;
|
|
cvt.rn.f32.f64 %f35, %fd171;
|
|
cvt.f64.f32 %fd114, %f35;
|
|
add.rn.f64 %fd33, %fd113, %fd114;
|
|
setp.gt.f64 %p86, %fd33, %fd67;
|
|
mov.f64 %fd173, %fd33;
|
|
@%p86 bra $L__BB11_108;
|
|
|
|
setp.nan.f64 %p87, %fd33, %fd67;
|
|
mov.f64 %fd173, 0d7FF8000000000000;
|
|
@%p87 bra $L__BB11_108;
|
|
|
|
setp.geu.f64 %p88, %fd33, %fd67;
|
|
setp.eq.f64 %p89, %fd33, 0d0000000000000000;
|
|
and.pred %p90, %p88, %p89;
|
|
selp.f64 %fd173, %fd33, %fd67, %p90;
|
|
|
|
$L__BB11_108:
|
|
mul.wide.u32 %rd130, %r213, 8;
|
|
add.s64 %rd131, %rd2, %rd130;
|
|
ld.global.u64 %rd37, [%rd131];
|
|
cvt.u32.u64 %r163, %rd37;
|
|
shr.u32 %r164, %r163, 16;
|
|
and.b32 %r66, %r164, 32767;
|
|
setp.eq.s32 %p91, %r66, 0;
|
|
mov.f64 %fd174, 0d3FF0000000000000;
|
|
@%p91 bra $L__BB11_110;
|
|
|
|
add.s32 %r165, %r66, -1;
|
|
mul.wide.u32 %rd132, %r165, 4;
|
|
add.s64 %rd133, %rd1, %rd132;
|
|
ld.global.nc.f32 %f36, [%rd133];
|
|
cvt.f64.f32 %fd174, %f36;
|
|
|
|
$L__BB11_110:
|
|
cvt.u16.u64 %rs9, %rd37;
|
|
cvt.rn.f64.s16 %fd117, %rs9;
|
|
mul.rn.f64 %fd118, %fd174, %fd117;
|
|
mul.rn.f64 %fd119, %fd118, %fd66;
|
|
cvt.rn.f32.f64 %f37, %fd173;
|
|
cvt.f64.f32 %fd120, %f37;
|
|
add.rn.f64 %fd38, %fd119, %fd120;
|
|
setp.gt.f64 %p92, %fd38, %fd67;
|
|
mov.f64 %fd175, %fd38;
|
|
@%p92 bra $L__BB11_113;
|
|
|
|
setp.nan.f64 %p93, %fd38, %fd67;
|
|
mov.f64 %fd175, 0d7FF8000000000000;
|
|
@%p93 bra $L__BB11_113;
|
|
|
|
setp.geu.f64 %p94, %fd38, %fd67;
|
|
setp.eq.f64 %p95, %fd38, 0d0000000000000000;
|
|
and.pred %p96, %p94, %p95;
|
|
selp.f64 %fd175, %fd38, %fd67, %p96;
|
|
|
|
$L__BB11_113:
|
|
cvt.rn.f32.f64 %f60, %fd175;
|
|
add.s32 %r213, %r213, 4;
|
|
add.s32 %r215, %r215, 4;
|
|
add.s32 %r166, %r60, %r215;
|
|
setp.ne.s32 %p97, %r166, 0;
|
|
@%p97 bra $L__BB11_93;
|
|
|
|
$L__BB11_114:
|
|
setp.eq.s32 %p98, %r217, 0;
|
|
@%p98 bra $L__BB11_153;
|
|
|
|
add.s32 %r216, %r215, %r3;
|
|
|
|
$L__BB11_116:
|
|
.pragma "nounroll";
|
|
mul.wide.u32 %rd134, %r216, 8;
|
|
add.s64 %rd135, %rd2, %rd134;
|
|
ld.global.u64 %rd38, [%rd135];
|
|
cvt.u32.u64 %r167, %rd38;
|
|
shr.u32 %r168, %r167, 16;
|
|
and.b32 %r73, %r168, 32767;
|
|
setp.eq.s32 %p99, %r73, 0;
|
|
mov.f64 %fd176, 0d3FF0000000000000;
|
|
@%p99 bra $L__BB11_118;
|
|
|
|
add.s32 %r169, %r73, -1;
|
|
mul.wide.u32 %rd136, %r169, 4;
|
|
add.s64 %rd137, %rd1, %rd136;
|
|
ld.global.nc.f32 %f38, [%rd137];
|
|
cvt.f64.f32 %fd176, %f38;
|
|
|
|
$L__BB11_118:
|
|
cvt.u16.u64 %rs10, %rd38;
|
|
cvt.rn.f64.s16 %fd123, %rs10;
|
|
mul.rn.f64 %fd124, %fd176, %fd123;
|
|
mul.rn.f64 %fd125, %fd124, %fd66;
|
|
cvt.f64.f32 %fd126, %f60;
|
|
add.rn.f64 %fd43, %fd125, %fd126;
|
|
setp.gt.f64 %p100, %fd43, %fd67;
|
|
mov.f64 %fd177, %fd43;
|
|
@%p100 bra $L__BB11_121;
|
|
|
|
setp.nan.f64 %p101, %fd43, %fd67;
|
|
mov.f64 %fd177, 0d7FF8000000000000;
|
|
@%p101 bra $L__BB11_121;
|
|
|
|
setp.geu.f64 %p102, %fd43, %fd67;
|
|
setp.eq.f64 %p103, %fd43, 0d0000000000000000;
|
|
and.pred %p104, %p102, %p103;
|
|
selp.f64 %fd177, %fd43, %fd67, %p104;
|
|
|
|
$L__BB11_121:
|
|
cvt.rn.f32.f64 %f60, %fd177;
|
|
add.s32 %r216, %r216, 1;
|
|
add.s32 %r217, %r217, -1;
|
|
setp.eq.s32 %p105, %r217, 0;
|
|
@%p105 bra $L__BB11_153;
|
|
bra.uni $L__BB11_116;
|
|
|
|
$L__BB11_153:
|
|
st.global.f32 [%rd28], %f60;
|
|
|
|
$L__BB11_154:
|
|
ret;
|
|
|
|
}
|
|
|