//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-32267302
// Cuda compilation tools, release 12.0, V12.0.140
// Based on NVVM 7.0.1
//
.version 8.0
.target sm_52
.address_size 64
// .globl mul_by_2
.visible .entry mul_by_2(
.param .u64 mul_by_2_param_0,
.param .u64 mul_by_2_param_1
)
{
.reg .f32 %f<3>;
.reg .b32 %r<8>;
.reg .b64 %rd<8>;
ld.param.u64 %rd1, [mul_by_2_param_0];
ld.param.u64 %rd2, [mul_by_2_param_1];
cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mov.u32 %r1, %tid.x;
mov.u32 %r2, %tid.y;
mov.u32 %r3, %ntid.x;
mov.u32 %r4, %ctaid.x;
mov.u32 %r5, %ntid.y;
mad.lo.s32 %r6, %r5, %r4, %r2;
mad.lo.s32 %r7, %r6, %r3, %r1;
mul.wide.s32 %rd5, %r7, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f1, [%rd6];
add.f32 %f2, %f1, %f1;
add.s64 %rd7, %rd3, %rd5;
st.global.f32 [%rd7], %f2;
ret;
}
// .globl sum_column
.visible .entry sum_column(
.param .u64 sum_column_param_0,
.param .u64 sum_column_param_1,
.param .u32 sum_column_param_2,
.param .u32 sum_column_param_3
)
{
.reg .pred %p<6>;
.reg .f32 %f<30>;
.reg .b32 %r<34>;
.reg .b64 %rd<22>;
ld.param.u64 %rd11, [sum_column_param_0];
ld.param.u64 %rd10, [sum_column_param_1];
ld.param.u32 %r13, [sum_column_param_2];
ld.param.u32 %r14, [sum_column_param_3];
cvta.to.global.u64 %rd1, %rd11;
mov.u32 %r15, %ntid.y;
mov.u32 %r16, %ctaid.x;
mov.u32 %r17, %tid.y;
mad.lo.s32 %r18, %r15, %r16, %r17;
mov.u32 %r19, %ntid.x;
mul.lo.s32 %r1, %r18, %r19;
mov.u32 %r2, %tid.x;
add.s32 %r3, %r1, %r2;
setp.lt.s32 %p1, %r13, 1;
mov.f32 %f29, 0f00000000;
@%p1 bra $L__BB1_7;
add.s32 %r21, %r13, -1;
and.b32 %r33, %r13, 3;
setp.lt.u32 %p2, %r21, 3;
mov.f32 %f29, 0f00000000;
mov.u32 %r32, 0;
@%p2 bra $L__BB1_4;
sub.s32 %r31, %r13, %r33;
mul.wide.s32 %rd12, %r3, 4;
add.s64 %rd20, %rd1, %rd12;
mul.wide.s32 %rd3, %r14, 4;
$L__BB1_3:
ld.global.f32 %f12, [%rd20];
add.f32 %f13, %f29, %f12;
cvt.rzi.s32.f32 %r23, %f13;
cvt.rn.f32.s32 %f14, %r23;
add.s64 %rd13, %rd20, %rd3;
ld.global.f32 %f15, [%rd13];
add.f32 %f16, %f15, %f14;
cvt.rzi.s32.f32 %r24, %f16;
cvt.rn.f32.s32 %f17, %r24;
add.s64 %rd14, %rd13, %rd3;
ld.global.f32 %f18, [%rd14];
add.f32 %f19, %f18, %f17;
cvt.rzi.s32.f32 %r25, %f19;
cvt.rn.f32.s32 %f20, %r25;
add.s64 %rd15, %rd14, %rd3;
add.s64 %rd20, %rd15, %rd3;
ld.global.f32 %f21, [%rd15];
add.f32 %f22, %f21, %f20;
cvt.rzi.s32.f32 %r26, %f22;
add.s32 %r32, %r32, 4;
cvt.rn.f32.s32 %f29, %r26;
add.s32 %r31, %r31, -4;
setp.ne.s32 %p3, %r31, 0;
@%p3 bra $L__BB1_3;
$L__BB1_4:
setp.eq.s32 %p4, %r33, 0;
@%p4 bra $L__BB1_7;
mad.lo.s32 %r27, %r32, %r14, %r2;
add.s32 %r28, %r27, %r1;
mul.wide.s32 %rd16, %r28, 4;
add.s64 %rd21, %rd1, %rd16;
mul.wide.s32 %rd7, %r14, 4;
$L__BB1_6:
.pragma "nounroll";
ld.global.f32 %f23, [%rd21];
add.f32 %f24, %f29, %f23;
cvt.rzi.s32.f32 %r29, %f24;
cvt.rn.f32.s32 %f29, %r29;
add.s64 %rd21, %rd21, %rd7;
add.s32 %r33, %r33, -1;
setp.ne.s32 %p5, %r33, 0;
@%p5 bra $L__BB1_6;
$L__BB1_7:
cvta.to.global.u64 %rd17, %rd10;
mul.wide.s32 %rd18, %r3, 4;
add.s64 %rd19, %rd17, %rd18;
st.global.f32 [%rd19], %f29;
ret;
}