ptx • Lines: 1272//
// 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 fill_value
// _ZZ10matrix_mulE4ds_A has been demoted
// _ZZ10matrix_mulE4ds_B has been demoted
.global .align 4 .b8 __cudart_i2opi_f[24] = {65, 144, 67, 60, 153, 149, 98, 219, 192, 221, 52, 245, 209, 87, 39, 252, 41, 21, 68, 78, 110, 131, 249, 162};
.visible .entry fill_value(
.param .u64 fill_value_param_0,
.param .u32 fill_value_param_1,
.param .f32 fill_value_param_2
)
{
.reg .pred %p<2>;
.reg .f32 %f<2>;
.reg .b32 %r<6>;
.reg .b64 %rd<5>;
ld.param.u64 %rd1, [fill_value_param_0];
ld.param.u32 %r2, [fill_value_param_1];
ld.param.f32 %f1, [fill_value_param_2];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r3, %r4, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;
cvta.to.global.u64 %rd2, %rd1;
mul.wide.s32 %rd3, %r1, 4;
add.s64 %rd4, %rd2, %rd3;
st.global.f32 [%rd4], %f1;
$L__BB0_2:
ret;
}
// .globl vector_arithmatic
.visible .entry vector_arithmatic(
.param .u64 vector_arithmatic_param_0,
.param .u64 vector_arithmatic_param_1,
.param .u64 vector_arithmatic_param_2,
.param .u32 vector_arithmatic_param_3,
.param .u32 vector_arithmatic_param_4
)
{
.reg .pred %p<7>;
.reg .f32 %f<16>;
.reg .b32 %r<7>;
.reg .b64 %rd<11>;
ld.param.u64 %rd4, [vector_arithmatic_param_0];
ld.param.u64 %rd5, [vector_arithmatic_param_1];
ld.param.u64 %rd6, [vector_arithmatic_param_2];
ld.param.u32 %r3, [vector_arithmatic_param_3];
ld.param.u32 %r2, [vector_arithmatic_param_4];
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %ctaid.x;
mov.u32 %r6, %tid.x;
mad.lo.s32 %r1, %r5, %r4, %r6;
setp.ge.s32 %p1, %r1, %r3;
@%p1 bra $L__BB1_11;
cvta.to.global.u64 %rd7, %rd4;
mul.wide.s32 %rd8, %r1, 4;
add.s64 %rd1, %rd7, %rd8;
cvta.to.global.u64 %rd9, %rd5;
add.s64 %rd2, %rd9, %rd8;
cvta.to.global.u64 %rd10, %rd6;
add.s64 %rd3, %rd10, %rd8;
setp.gt.s32 %p2, %r2, 2;
@%p2 bra $L__BB1_4;
setp.eq.s32 %p5, %r2, 1;
@%p5 bra $L__BB1_9;
setp.eq.s32 %p6, %r2, 2;
@%p6 bra $L__BB1_8;
bra.uni $L__BB1_10;
$L__BB1_8:
ld.global.f32 %f7, [%rd1];
ld.global.f32 %f8, [%rd2];
sub.f32 %f9, %f7, %f8;
st.global.f32 [%rd3], %f9;
bra.uni $L__BB1_11;
$L__BB1_4:
setp.eq.s32 %p3, %r2, 3;
@%p3 bra $L__BB1_7;
setp.ne.s32 %p4, %r2, 4;
@%p4 bra $L__BB1_10;
ld.global.f32 %f1, [%rd1];
ld.global.f32 %f2, [%rd2];
div.rn.f32 %f3, %f1, %f2;
st.global.f32 [%rd3], %f3;
bra.uni $L__BB1_11;
$L__BB1_10:
ld.global.f32 %f13, [%rd1];
ld.global.f32 %f14, [%rd2];
add.f32 %f15, %f13, %f14;
st.global.f32 [%rd3], %f15;
bra.uni $L__BB1_11;
$L__BB1_9:
ld.global.f32 %f10, [%rd1];
ld.global.f32 %f11, [%rd2];
add.f32 %f12, %f10, %f11;
st.global.f32 [%rd3], %f12;
bra.uni $L__BB1_11;
$L__BB1_7:
ld.global.f32 %f4, [%rd1];
ld.global.f32 %f5, [%rd2];
mul.f32 %f6, %f4, %f5;
st.global.f32 [%rd3], %f6;
$L__BB1_11:
ret;
}
// .globl clip
.visible .entry clip(
.param .u64 clip_param_0,
.param .u64 clip_param_1,
.param .u32 clip_param_2,
.param .f32 clip_param_3,
.param .f32 clip_param_4
)
{
.reg .pred %p<4>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;
ld.param.u64 %rd2, [clip_param_0];
ld.param.u64 %rd3, [clip_param_1];
ld.param.u32 %r2, [clip_param_2];
ld.param.f32 %f2, [clip_param_3];
ld.param.f32 %f3, [clip_param_4];
mov.u32 %r3, %ntid.x;
mov.u32 %r4, %ctaid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r4, %r3, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB2_6;
cvta.to.global.u64 %rd4, %rd2;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f1, [%rd6];
setp.lt.f32 %p2, %f1, %f2;
cvta.to.global.u64 %rd7, %rd3;
add.s64 %rd1, %rd7, %rd5;
@%p2 bra $L__BB2_5;
bra.uni $L__BB2_2;
$L__BB2_5:
st.global.f32 [%rd1], %f2;
bra.uni $L__BB2_6;
$L__BB2_2:
setp.gt.f32 %p3, %f1, %f3;
@%p3 bra $L__BB2_4;
bra.uni $L__BB2_3;
$L__BB2_4:
st.global.f32 [%rd1], %f3;
bra.uni $L__BB2_6;
$L__BB2_3:
st.global.f32 [%rd1], %f1;
$L__BB2_6:
ret;
}
// .globl element_op
.visible .entry element_op(
.param .u64 element_op_param_0,
.param .u64 element_op_param_1,
.param .u32 element_op_param_2,
.param .u32 element_op_param_3,
.param .f32 element_op_param_4
)
{
.local .align 4 .b8 __local_depot3[28];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<56>;
.reg .f32 %f<266>;
.reg .b32 %r<196>;
.reg .f64 %fd<7>;
.reg .b64 %rd<89>;
mov.u64 %SPL, __local_depot3;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd23, [element_op_param_0];
ld.param.u64 %rd24, [element_op_param_1];
ld.param.u32 %r57, [element_op_param_2];
ld.param.u32 %r56, [element_op_param_3];
ld.param.f32 %f48, [element_op_param_4];
add.u64 %rd25, %SP, 0;
add.u64 %rd1, %SPL, 0;
mov.u32 %r58, %ntid.x;
mov.u32 %r59, %ctaid.x;
mov.u32 %r60, %tid.x;
mad.lo.s32 %r1, %r59, %r58, %r60;
setp.ge.s32 %p1, %r1, %r57;
@%p1 bra $L__BB3_70;
cvta.to.global.u64 %rd26, %rd23;
mul.wide.s32 %rd27, %r1, 4;
add.s64 %rd2, %rd26, %rd27;
cvta.to.global.u64 %rd28, %rd24;
add.s64 %rd3, %rd28, %rd27;
add.s64 %rd4, %rd1, 24;
setp.gt.s32 %p2, %r56, 4;
@%p2 bra $L__BB3_8;
bra.uni $L__BB3_2;
$L__BB3_8:
setp.gt.s32 %p3, %r56, 7;
@%p3 bra $L__BB3_12;
setp.eq.s32 %p7, %r56, 5;
@%p7 bra $L__BB3_30;
setp.eq.s32 %p8, %r56, 6;
@%p8 bra $L__BB3_27;
setp.eq.s32 %p9, %r56, 7;
@%p9 bra $L__BB3_24;
bra.uni $L__BB3_70;
$L__BB3_24:
ld.global.f32 %f80, [%rd2];
setp.lt.f32 %p21, %f80, 0f00800000;
mul.f32 %f81, %f80, 0f4B000000;
selp.f32 %f6, %f81, %f80, %p21;
selp.f32 %f82, 0fC1B80000, 0f00000000, %p21;
mov.b32 %r68, %f6;
add.s32 %r69, %r68, -1059760811;
and.b32 %r70, %r69, -8388608;
sub.s32 %r71, %r68, %r70;
mov.b32 %f83, %r71;
cvt.rn.f32.s32 %f84, %r70;
mov.f32 %f85, 0f34000000;
fma.rn.f32 %f86, %f84, %f85, %f82;
add.f32 %f87, %f83, 0fBF800000;
mov.f32 %f88, 0f3E1039F6;
mov.f32 %f89, 0fBE055027;
fma.rn.f32 %f90, %f89, %f87, %f88;
mov.f32 %f91, 0fBDF8CDCC;
fma.rn.f32 %f92, %f90, %f87, %f91;
mov.f32 %f93, 0f3E0F2955;
fma.rn.f32 %f94, %f92, %f87, %f93;
mov.f32 %f95, 0fBE2AD8B9;
fma.rn.f32 %f96, %f94, %f87, %f95;
mov.f32 %f97, 0f3E4CED0B;
fma.rn.f32 %f98, %f96, %f87, %f97;
mov.f32 %f99, 0fBE7FFF22;
fma.rn.f32 %f100, %f98, %f87, %f99;
mov.f32 %f101, 0f3EAAAA78;
fma.rn.f32 %f102, %f100, %f87, %f101;
mov.f32 %f103, 0fBF000000;
fma.rn.f32 %f104, %f102, %f87, %f103;
mul.f32 %f105, %f87, %f104;
fma.rn.f32 %f106, %f105, %f87, %f87;
mov.f32 %f107, 0f3F317218;
fma.rn.f32 %f258, %f86, %f107, %f106;
setp.lt.u32 %p22, %r68, 2139095040;
@%p22 bra $L__BB3_26;
mov.f32 %f108, 0f7F800000;
fma.rn.f32 %f258, %f6, %f108, %f108;
$L__BB3_26:
mul.f32 %f109, %f258, 0f3EDE5BD9;
setp.eq.f32 %p23, %f6, 0f00000000;
selp.f32 %f110, 0fFF800000, %f109, %p23;
st.global.f32 [%rd3], %f110;
bra.uni $L__BB3_70;
$L__BB3_2:
setp.gt.s32 %p10, %r56, 1;
@%p10 bra $L__BB3_5;
bra.uni $L__BB3_3;
$L__BB3_5:
setp.eq.s32 %p11, %r56, 2;
@%p11 bra $L__BB3_55;
setp.eq.s32 %p12, %r56, 3;
@%p12 bra $L__BB3_42;
setp.eq.s32 %p13, %r56, 4;
@%p13 bra $L__BB3_33;
bra.uni $L__BB3_70;
$L__BB3_33:
ld.global.f32 %f16, [%rd2];
mul.f32 %f165, %f16, 0f3F22F983;
cvt.rni.s32.f32 %r187, %f165;
cvt.rn.f32.s32 %f166, %r187;
mov.f32 %f167, 0fBFC90FDA;
fma.rn.f32 %f168, %f166, %f167, %f16;
mov.f32 %f169, 0fB3A22168;
fma.rn.f32 %f170, %f166, %f169, %f168;
mov.f32 %f171, 0fA7C234C5;
fma.rn.f32 %f259, %f166, %f171, %f170;
abs.f32 %f18, %f16;
setp.ltu.f32 %p27, %f18, 0f47CE4780;
@%p27 bra $L__BB3_41;
setp.eq.f32 %p28, %f18, 0f7F800000;
@%p28 bra $L__BB3_40;
bra.uni $L__BB3_35;
$L__BB3_40:
mov.f32 %f174, 0f00000000;
mul.rn.f32 %f259, %f16, %f174;
mov.u32 %r187, 0;
bra.uni $L__BB3_41;
$L__BB3_12:
setp.eq.s32 %p4, %r56, 8;
@%p4 bra $L__BB3_21;
setp.eq.s32 %p5, %r56, 9;
@%p5 bra $L__BB3_18;
setp.ne.s32 %p6, %r56, 10;
@%p6 bra $L__BB3_70;
ld.global.f32 %f1, [%rd2];
setp.ltu.f32 %p16, %f1, 0f00000000;
@%p16 bra $L__BB3_17;
bra.uni $L__BB3_16;
$L__BB3_17:
mov.u32 %r61, 0;
st.global.u32 [%rd3], %r61;
bra.uni $L__BB3_70;
$L__BB3_3:
setp.eq.s32 %p14, %r56, 0;
@%p14 bra $L__BB3_69;
setp.eq.s32 %p15, %r56, 1;
@%p15 bra $L__BB3_68;
bra.uni $L__BB3_70;
$L__BB3_68:
ld.global.f32 %f236, [%rd2];
mul.f32 %f237, %f236, %f48;
st.global.f32 [%rd3], %f237;
bra.uni $L__BB3_70;
$L__BB3_30:
ld.global.f32 %f11, [%rd2];
abs.f32 %f12, %f11;
setp.ltu.f32 %p25, %f12, 0f3F19999A;
@%p25 bra $L__BB3_32;
bra.uni $L__BB3_31;
$L__BB3_32:
mul.f32 %f155, %f11, %f11;
mov.f32 %f156, 0fBD563CAE;
mov.f32 %f157, 0f3C80F082;
fma.rn.f32 %f158, %f157, %f155, %f156;
mov.f32 %f159, 0f3E085941;
fma.rn.f32 %f160, %f158, %f155, %f159;
mov.f32 %f161, 0fBEAAA9ED;
fma.rn.f32 %f162, %f160, %f155, %f161;
mov.f32 %f163, 0f00000000;
fma.rn.f32 %f164, %f162, %f155, %f163;
fma.rn.f32 %f14, %f164, %f11, %f11;
st.global.f32 [%rd3], %f14;
bra.uni $L__BB3_70;
$L__BB3_27:
ld.global.f32 %f10, [%rd2];
setp.ltu.f32 %p24, %f10, 0f00000000;
@%p24 bra $L__BB3_29;
bra.uni $L__BB3_28;
$L__BB3_29:
mov.f32 %f129, 0f3F000000;
mov.f32 %f130, 0f3BBB989D;
fma.rn.f32 %f131, %f10, %f130, %f129;
mov.f32 %f132, 0f3FB8AA3B;
mov.f32 %f133, 0f437C0000;
cvt.sat.f32.f32 %f134, %f131;
mov.f32 %f135, 0f4B400001;
fma.rm.f32 %f136, %f134, %f133, %f135;
add.f32 %f137, %f136, 0fCB40007F;
neg.f32 %f138, %f137;
fma.rn.f32 %f139, %f10, %f132, %f138;
mov.f32 %f140, 0f32A57060;
fma.rn.f32 %f141, %f10, %f140, %f139;
mov.b32 %r74, %f136;
shl.b32 %r75, %r74, 23;
mov.b32 %f142, %r75;
ex2.approx.ftz.f32 %f143, %f141;
mul.f32 %f144, %f143, %f142;
add.f32 %f145, %f144, 0f3F800000;
div.rn.f32 %f146, %f144, %f145;
st.global.f32 [%rd3], %f146;
bra.uni $L__BB3_70;
$L__BB3_55:
ld.global.f32 %f35, [%rd2];
mul.f32 %f215, %f35, 0f3F22F983;
cvt.rni.s32.f32 %r195, %f215;
cvt.rn.f32.s32 %f216, %r195;
mov.f32 %f217, 0fBFC90FDA;
fma.rn.f32 %f218, %f216, %f217, %f35;
mov.f32 %f219, 0fB3A22168;
fma.rn.f32 %f220, %f216, %f219, %f218;
mov.f32 %f221, 0fA7C234C5;
fma.rn.f32 %f263, %f216, %f221, %f220;
abs.f32 %f37, %f35;
setp.ltu.f32 %p46, %f37, 0f47CE4780;
@%p46 bra $L__BB3_63;
setp.eq.f32 %p47, %f37, 0f7F800000;
@%p47 bra $L__BB3_62;
bra.uni $L__BB3_57;
$L__BB3_62:
mov.f32 %f224, 0f00000000;
mul.rn.f32 %f263, %f35, %f224;
mov.u32 %r195, 0;
bra.uni $L__BB3_63;
$L__BB3_42:
ld.global.f32 %f22, [%rd2];
mul.f32 %f194, %f22, 0f3F22F983;
cvt.rni.s32.f32 %r191, %f194;
cvt.rn.f32.s32 %f195, %r191;
mov.f32 %f196, 0fBFC90FDA;
fma.rn.f32 %f197, %f195, %f196, %f22;
mov.f32 %f198, 0fB3A22168;
fma.rn.f32 %f199, %f195, %f198, %f197;
mov.f32 %f200, 0fA7C234C5;
fma.rn.f32 %f260, %f195, %f200, %f199;
abs.f32 %f24, %f22;
setp.ltu.f32 %p36, %f24, 0f47CE4780;
@%p36 bra $L__BB3_50;
setp.eq.f32 %p37, %f24, 0f7F800000;
@%p37 bra $L__BB3_49;
bra.uni $L__BB3_44;
$L__BB3_49:
mov.f32 %f203, 0f00000000;
mul.rn.f32 %f260, %f22, %f203;
mov.u32 %r191, 0;
bra.uni $L__BB3_50;
$L__BB3_21:
ld.global.f32 %f50, [%rd2];
setp.lt.f32 %p18, %f50, 0f00800000;
mul.f32 %f51, %f50, 0f4B000000;
selp.f32 %f2, %f51, %f50, %p18;
selp.f32 %f52, 0fC1B80000, 0f00000000, %p18;
mov.b32 %r64, %f2;
add.s32 %r65, %r64, -1059760811;
and.b32 %r66, %r65, -8388608;
sub.s32 %r67, %r64, %r66;
mov.b32 %f53, %r67;
cvt.rn.f32.s32 %f54, %r66;
mov.f32 %f55, 0f34000000;
fma.rn.f32 %f56, %f54, %f55, %f52;
add.f32 %f57, %f53, 0fBF800000;
mov.f32 %f58, 0f3E1039F6;
mov.f32 %f59, 0fBE055027;
fma.rn.f32 %f60, %f59, %f57, %f58;
mov.f32 %f61, 0fBDF8CDCC;
fma.rn.f32 %f62, %f60, %f57, %f61;
mov.f32 %f63, 0f3E0F2955;
fma.rn.f32 %f64, %f62, %f57, %f63;
mov.f32 %f65, 0fBE2AD8B9;
fma.rn.f32 %f66, %f64, %f57, %f65;
mov.f32 %f67, 0f3E4CED0B;
fma.rn.f32 %f68, %f66, %f57, %f67;
mov.f32 %f69, 0fBE7FFF22;
fma.rn.f32 %f70, %f68, %f57, %f69;
mov.f32 %f71, 0f3EAAAA78;
fma.rn.f32 %f72, %f70, %f57, %f71;
mov.f32 %f73, 0fBF000000;
fma.rn.f32 %f74, %f72, %f57, %f73;
mul.f32 %f75, %f57, %f74;
fma.rn.f32 %f76, %f75, %f57, %f57;
mov.f32 %f77, 0f3F317218;
fma.rn.f32 %f257, %f56, %f77, %f76;
setp.lt.u32 %p19, %r64, 2139095040;
@%p19 bra $L__BB3_23;
mov.f32 %f78, 0f7F800000;
fma.rn.f32 %f257, %f2, %f78, %f78;
$L__BB3_23:
setp.eq.f32 %p20, %f2, 0f00000000;
selp.f32 %f79, 0fFF800000, %f257, %p20;
st.global.f32 [%rd3], %f79;
bra.uni $L__BB3_70;
$L__BB3_18:
ld.global.f32 %f49, [%rd2];
setp.ltu.f32 %p17, %f49, 0f00000000;
@%p17 bra $L__BB3_20;
bra.uni $L__BB3_19;
$L__BB3_20:
mov.u32 %r63, 0;
st.global.u32 [%rd3], %r63;
bra.uni $L__BB3_70;
$L__BB3_69:
ld.global.f32 %f238, [%rd2];
mov.f32 %f239, 0f3F000000;
mov.f32 %f240, 0f3BBB989D;
fma.rn.f32 %f241, %f238, %f240, %f239;
mov.f32 %f242, 0f3FB8AA3B;
mov.f32 %f243, 0f437C0000;
cvt.sat.f32.f32 %f244, %f241;
mov.f32 %f245, 0f4B400001;
fma.rm.f32 %f246, %f244, %f243, %f245;
add.f32 %f247, %f246, 0fCB40007F;
neg.f32 %f248, %f247;
fma.rn.f32 %f249, %f238, %f242, %f248;
mov.f32 %f250, 0f32A57060;
fma.rn.f32 %f251, %f238, %f250, %f249;
mov.b32 %r182, %f246;
shl.b32 %r183, %r182, 23;
mov.b32 %f252, %r183;
ex2.approx.ftz.f32 %f253, %f251;
mul.f32 %f254, %f253, %f252;
st.global.f32 [%rd3], %f254;
bra.uni $L__BB3_70;
$L__BB3_16:
st.global.f32 [%rd3], %f1;
bra.uni $L__BB3_70;
$L__BB3_31:
mul.f32 %f147, %f12, 0f4038AA3B;
ex2.approx.ftz.f32 %f148, %f147;
add.f32 %f149, %f148, 0f3F800000;
mov.f32 %f150, 0f3F800000;
rcp.approx.ftz.f32 %f151, %f149;
mov.f32 %f152, 0fC0000000;
fma.rn.f32 %f153, %f151, %f152, %f150;
setp.ge.f32 %p26, %f12, 0f41102CB4;
selp.f32 %f154, 0f3F800000, %f153, %p26;
mov.b32 %r76, %f154;
mov.b32 %r77, %f11;
and.b32 %r78, %r77, -2147483648;
or.b32 %r79, %r78, %r76;
mov.b32 %f13, %r79;
st.global.f32 [%rd3], %f13;
bra.uni $L__BB3_70;
$L__BB3_28:
neg.f32 %f111, %f10;
mov.f32 %f112, 0f3F000000;
mov.f32 %f113, 0f3BBB989D;
fma.rn.f32 %f114, %f111, %f113, %f112;
mov.f32 %f115, 0f3FB8AA3B;
mov.f32 %f116, 0f437C0000;
cvt.sat.f32.f32 %f117, %f114;
mov.f32 %f118, 0f4B400001;
fma.rm.f32 %f119, %f117, %f116, %f118;
add.f32 %f120, %f119, 0fCB40007F;
neg.f32 %f121, %f120;
fma.rn.f32 %f122, %f111, %f115, %f121;
mov.f32 %f123, 0f32A57060;
fma.rn.f32 %f124, %f111, %f123, %f122;
mov.b32 %r72, %f119;
shl.b32 %r73, %r72, 23;
mov.b32 %f125, %r73;
ex2.approx.ftz.f32 %f126, %f124;
fma.rn.f32 %f127, %f126, %f125, 0f3F800000;
rcp.rn.f32 %f128, %f127;
st.global.f32 [%rd3], %f128;
bra.uni $L__BB3_70;
$L__BB3_19:
mov.u32 %r62, 1065353216;
st.global.u32 [%rd3], %r62;
bra.uni $L__BB3_70;
$L__BB3_35:
mov.b32 %r3, %f16;
bfe.u32 %r81, %r3, 23, 8;
add.s32 %r4, %r81, -128;
shl.b32 %r82, %r3, 8;
or.b32 %r5, %r82, -2147483648;
shr.u32 %r6, %r4, 5;
mov.u64 %rd82, 0;
mov.u32 %r184, 0;
mov.u64 %rd81, __cudart_i2opi_f;
$L__BB3_36:
.pragma "nounroll";
ld.global.nc.u32 %r83, [%rd81];
mad.wide.u32 %rd31, %r83, %r5, %rd82;
shr.u64 %rd82, %rd31, 32;
st.local.u32 [%rd1], %rd31;
add.s64 %rd81, %rd81, 4;
add.s64 %rd1, %rd1, 4;
add.s32 %r184, %r184, 1;
setp.ne.s32 %p29, %r184, 6;
@%p29 bra $L__BB3_36;
st.local.u32 [%rd4], %rd82;
mov.u32 %r84, 4;
sub.s32 %r9, %r84, %r6;
mov.u32 %r85, 6;
sub.s32 %r86, %r85, %r6;
cvta.to.local.u64 %rd33, %rd25;
mul.wide.s32 %rd34, %r86, 4;
add.s64 %rd35, %rd33, %rd34;
ld.local.u32 %r185, [%rd35];
ld.local.u32 %r186, [%rd35+-4];
and.b32 %r12, %r4, 31;
setp.eq.s32 %p30, %r12, 0;
@%p30 bra $L__BB3_39;
mov.u32 %r87, 32;
sub.s32 %r88, %r87, %r12;
shr.u32 %r89, %r186, %r88;
shl.b32 %r90, %r185, %r12;
add.s32 %r185, %r89, %r90;
mul.wide.s32 %rd38, %r9, 4;
add.s64 %rd39, %rd33, %rd38;
ld.local.u32 %r91, [%rd39];
shr.u32 %r92, %r91, %r88;
shl.b32 %r93, %r186, %r12;
add.s32 %r186, %r92, %r93;
$L__BB3_39:
and.b32 %r94, %r3, -2147483648;
shr.u32 %r95, %r186, 30;
shl.b32 %r96, %r185, 2;
or.b32 %r97, %r95, %r96;
shr.u32 %r98, %r97, 31;
shr.u32 %r99, %r185, 30;
add.s32 %r100, %r98, %r99;
neg.s32 %r101, %r100;
setp.eq.s32 %p31, %r94, 0;
selp.b32 %r187, %r100, %r101, %p31;
setp.ne.s32 %p32, %r98, 0;
xor.b32 %r102, %r94, -2147483648;
selp.b32 %r103, %r102, %r94, %p32;
selp.b32 %r104, -1, 0, %p32;
xor.b32 %r105, %r97, %r104;
shl.b32 %r106, %r186, 2;
xor.b32 %r107, %r106, %r104;
cvt.u64.u32 %rd40, %r105;
cvt.u64.u32 %rd41, %r107;
bfi.b64 %rd42, %rd40, %rd41, 32, 32;
cvt.rn.f64.s64 %fd1, %rd42;
mul.f64 %fd2, %fd1, 0d3BF921FB54442D19;
cvt.rn.f32.f64 %f172, %fd2;
setp.eq.s32 %p33, %r103, 0;
neg.f32 %f173, %f172;
selp.f32 %f259, %f172, %f173, %p33;
$L__BB3_41:
mul.f32 %f175, %f259, %f259;
mov.f32 %f176, 0f3B560000;
mov.f32 %f177, 0f3C190000;
fma.rn.f32 %f178, %f177, %f175, %f176;
mov.f32 %f179, 0f3CC70000;
fma.rn.f32 %f180, %f178, %f175, %f179;
mov.f32 %f181, 0f3D5B0000;
fma.rn.f32 %f182, %f180, %f175, %f181;
mov.f32 %f183, 0f3E089438;
fma.rn.f32 %f184, %f182, %f175, %f183;
mov.f32 %f185, 0f3EAAAA88;
fma.rn.f32 %f186, %f184, %f175, %f185;
mul.rn.f32 %f187, %f175, %f259;
fma.rn.f32 %f188, %f186, %f187, %f259;
abs.f32 %f189, %f259;
setp.eq.f32 %p34, %f189, 0f3A00B43C;
selp.f32 %f190, %f259, %f188, %p34;
and.b32 %r109, %r187, 1;
setp.eq.b32 %p35, %r109, 1;
neg.f32 %f191, %f190;
rcp.approx.ftz.f32 %f192, %f191;
selp.f32 %f193, %f192, %f190, %p35;
st.global.f32 [%rd3], %f193;
bra.uni $L__BB3_70;
$L__BB3_57:
mov.b32 %r39, %f35;
bfe.u32 %r149, %r39, 23, 8;
add.s32 %r40, %r149, -128;
shl.b32 %r150, %r39, 8;
or.b32 %r41, %r150, -2147483648;
shr.u32 %r42, %r40, 5;
mov.u64 %rd88, 0;
mov.u32 %r192, 0;
mov.u64 %rd87, __cudart_i2opi_f;
$L__BB3_58:
.pragma "nounroll";
ld.global.nc.u32 %r151, [%rd87];
mad.wide.u32 %rd65, %r151, %r41, %rd88;
shr.u64 %rd88, %rd65, 32;
st.local.u32 [%rd1], %rd65;
add.s64 %rd87, %rd87, 4;
add.s64 %rd1, %rd1, 4;
add.s32 %r192, %r192, 1;
setp.ne.s32 %p48, %r192, 6;
@%p48 bra $L__BB3_58;
st.local.u32 [%rd4], %rd88;
mov.u32 %r152, 4;
sub.s32 %r45, %r152, %r42;
mov.u32 %r153, 6;
sub.s32 %r154, %r153, %r42;
cvta.to.local.u64 %rd67, %rd25;
mul.wide.s32 %rd68, %r154, 4;
add.s64 %rd69, %rd67, %rd68;
ld.local.u32 %r193, [%rd69];
ld.local.u32 %r194, [%rd69+-4];
and.b32 %r48, %r40, 31;
setp.eq.s32 %p49, %r48, 0;
@%p49 bra $L__BB3_61;
mov.u32 %r155, 32;
sub.s32 %r156, %r155, %r48;
shr.u32 %r157, %r194, %r156;
shl.b32 %r158, %r193, %r48;
add.s32 %r193, %r157, %r158;
mul.wide.s32 %rd72, %r45, 4;
add.s64 %rd73, %rd67, %rd72;
ld.local.u32 %r159, [%rd73];
shr.u32 %r160, %r159, %r156;
shl.b32 %r161, %r194, %r48;
add.s32 %r194, %r160, %r161;
$L__BB3_61:
and.b32 %r162, %r39, -2147483648;
shr.u32 %r163, %r194, 30;
shl.b32 %r164, %r193, 2;
or.b32 %r165, %r163, %r164;
shr.u32 %r166, %r165, 31;
shr.u32 %r167, %r193, 30;
add.s32 %r168, %r166, %r167;
neg.s32 %r169, %r168;
setp.eq.s32 %p50, %r162, 0;
selp.b32 %r195, %r168, %r169, %p50;
setp.ne.s32 %p51, %r166, 0;
xor.b32 %r170, %r162, -2147483648;
selp.b32 %r171, %r170, %r162, %p51;
selp.b32 %r172, -1, 0, %p51;
xor.b32 %r173, %r165, %r172;
shl.b32 %r174, %r194, 2;
xor.b32 %r175, %r174, %r172;
cvt.u64.u32 %rd74, %r173;
cvt.u64.u32 %rd75, %r175;
bfi.b64 %rd76, %rd74, %rd75, 32, 32;
cvt.rn.f64.s64 %fd5, %rd76;
mul.f64 %fd6, %fd5, 0d3BF921FB54442D19;
cvt.rn.f32.f64 %f222, %fd6;
setp.eq.s32 %p52, %r171, 0;
neg.f32 %f223, %f222;
selp.f32 %f263, %f222, %f223, %p52;
$L__BB3_63:
and.b32 %r55, %r195, 1;
setp.eq.s32 %p53, %r55, 0;
selp.f32 %f41, %f263, 0f3F800000, %p53;
mul.rn.f32 %f42, %f263, %f263;
mov.f32 %f264, 0fB94D4153;
@%p53 bra $L__BB3_65;
mov.f32 %f226, 0fBAB607ED;
mov.f32 %f227, 0f37CBAC00;
fma.rn.f32 %f264, %f227, %f42, %f226;
$L__BB3_65:
selp.f32 %f228, 0f3C0885E4, 0f3D2AAABB, %p53;
fma.rn.f32 %f229, %f264, %f42, %f228;
selp.f32 %f230, 0fBE2AAAA8, 0fBEFFFFFF, %p53;
fma.rn.f32 %f231, %f229, %f42, %f230;
mov.f32 %f232, 0f00000000;
fma.rn.f32 %f233, %f42, %f41, %f232;
fma.rn.f32 %f265, %f231, %f233, %f41;
and.b32 %r177, %r195, 2;
setp.eq.s32 %p55, %r177, 0;
@%p55 bra $L__BB3_67;
mov.f32 %f235, 0fBF800000;
fma.rn.f32 %f265, %f265, %f235, %f232;
$L__BB3_67:
st.global.f32 [%rd3], %f265;
bra.uni $L__BB3_70;
$L__BB3_44:
mov.b32 %r20, %f22;
bfe.u32 %r115, %r20, 23, 8;
add.s32 %r21, %r115, -128;
shl.b32 %r116, %r20, 8;
or.b32 %r22, %r116, -2147483648;
shr.u32 %r23, %r21, 5;
mov.u64 %rd85, 0;
mov.u32 %r188, 0;
mov.u64 %rd84, __cudart_i2opi_f;
$L__BB3_45:
.pragma "nounroll";
ld.global.nc.u32 %r117, [%rd84];
mad.wide.u32 %rd48, %r117, %r22, %rd85;
shr.u64 %rd85, %rd48, 32;
st.local.u32 [%rd1], %rd48;
add.s64 %rd84, %rd84, 4;
add.s64 %rd1, %rd1, 4;
add.s32 %r188, %r188, 1;
setp.ne.s32 %p38, %r188, 6;
@%p38 bra $L__BB3_45;
st.local.u32 [%rd4], %rd85;
mov.u32 %r118, 4;
sub.s32 %r26, %r118, %r23;
mov.u32 %r119, 6;
sub.s32 %r120, %r119, %r23;
cvta.to.local.u64 %rd50, %rd25;
mul.wide.s32 %rd51, %r120, 4;
add.s64 %rd52, %rd50, %rd51;
ld.local.u32 %r189, [%rd52];
ld.local.u32 %r190, [%rd52+-4];
and.b32 %r29, %r21, 31;
setp.eq.s32 %p39, %r29, 0;
@%p39 bra $L__BB3_48;
mov.u32 %r121, 32;
sub.s32 %r122, %r121, %r29;
shr.u32 %r123, %r190, %r122;
shl.b32 %r124, %r189, %r29;
add.s32 %r189, %r123, %r124;
mul.wide.s32 %rd55, %r26, 4;
add.s64 %rd56, %rd50, %rd55;
ld.local.u32 %r125, [%rd56];
shr.u32 %r126, %r125, %r122;
shl.b32 %r127, %r190, %r29;
add.s32 %r190, %r126, %r127;
$L__BB3_48:
and.b32 %r128, %r20, -2147483648;
shr.u32 %r129, %r190, 30;
shl.b32 %r130, %r189, 2;
or.b32 %r131, %r129, %r130;
shr.u32 %r132, %r131, 31;
shr.u32 %r133, %r189, 30;
add.s32 %r134, %r132, %r133;
neg.s32 %r135, %r134;
setp.eq.s32 %p40, %r128, 0;
selp.b32 %r191, %r134, %r135, %p40;
setp.ne.s32 %p41, %r132, 0;
xor.b32 %r136, %r128, -2147483648;
selp.b32 %r137, %r136, %r128, %p41;
selp.b32 %r138, -1, 0, %p41;
xor.b32 %r139, %r131, %r138;
shl.b32 %r140, %r190, 2;
xor.b32 %r141, %r140, %r138;
cvt.u64.u32 %rd57, %r139;
cvt.u64.u32 %rd58, %r141;
bfi.b64 %rd59, %rd57, %rd58, 32, 32;
cvt.rn.f64.s64 %fd3, %rd59;
mul.f64 %fd4, %fd3, 0d3BF921FB54442D19;
cvt.rn.f32.f64 %f201, %fd4;
setp.eq.s32 %p42, %r137, 0;
neg.f32 %f202, %f201;
selp.f32 %f260, %f201, %f202, %p42;
$L__BB3_50:
add.s32 %r36, %r191, 1;
and.b32 %r37, %r36, 1;
setp.eq.s32 %p43, %r37, 0;
selp.f32 %f28, %f260, 0f3F800000, %p43;
mul.rn.f32 %f29, %f260, %f260;
mov.f32 %f261, 0fB94D4153;
@%p43 bra $L__BB3_52;
mov.f32 %f205, 0fBAB607ED;
mov.f32 %f206, 0f37CBAC00;
fma.rn.f32 %f261, %f206, %f29, %f205;
$L__BB3_52:
selp.f32 %f207, 0f3C0885E4, 0f3D2AAABB, %p43;
fma.rn.f32 %f208, %f261, %f29, %f207;
selp.f32 %f209, 0fBE2AAAA8, 0fBEFFFFFF, %p43;
fma.rn.f32 %f210, %f208, %f29, %f209;
mov.f32 %f211, 0f00000000;
fma.rn.f32 %f212, %f29, %f28, %f211;
fma.rn.f32 %f262, %f210, %f212, %f28;
and.b32 %r143, %r36, 2;
setp.eq.s32 %p45, %r143, 0;
@%p45 bra $L__BB3_54;
mov.f32 %f214, 0fBF800000;
fma.rn.f32 %f262, %f262, %f214, %f211;
$L__BB3_54:
st.global.f32 [%rd3], %f262;
$L__BB3_70:
ret;
}
// .globl compare_memory
.visible .entry compare_memory(
.param .u64 compare_memory_param_0,
.param .u64 compare_memory_param_1,
.param .u64 compare_memory_param_2,
.param .u64 compare_memory_param_3
)
{
.reg .pred %p<3>;
.reg .f32 %f<5>;
.reg .b32 %r<6>;
.reg .b64 %rd<12>;
ld.param.u64 %rd2, [compare_memory_param_0];
ld.param.u64 %rd3, [compare_memory_param_1];
ld.param.u64 %rd5, [compare_memory_param_2];
ld.param.u64 %rd4, [compare_memory_param_3];
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r2, %r1, %r3;
cvt.u64.u32 %rd1, %r4;
setp.ge.u64 %p1, %rd1, %rd5;
@%p1 bra $L__BB4_3;
cvta.to.global.u64 %rd6, %rd2;
shl.b64 %rd7, %rd1, 2;
add.s64 %rd8, %rd6, %rd7;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd7;
ld.global.f32 %f1, [%rd10];
ld.global.f32 %f2, [%rd8];
sub.f32 %f3, %f2, %f1;
abs.f32 %f4, %f3;
setp.leu.f32 %p2, %f4, 0f3089705F;
@%p2 bra $L__BB4_3;
cvta.to.global.u64 %rd11, %rd4;
atom.global.exch.b32 %r5, [%rd11], 5;
$L__BB4_3:
ret;
}
// .globl transpose_naive
.visible .entry transpose_naive(
.param .u64 transpose_naive_param_0,
.param .u64 transpose_naive_param_1,
.param .u32 transpose_naive_param_2,
.param .u32 transpose_naive_param_3
)
{
.reg .pred %p<4>;
.reg .f32 %f<2>;
.reg .b32 %r<13>;
.reg .b64 %rd<9>;
ld.param.u64 %rd1, [transpose_naive_param_0];
ld.param.u64 %rd2, [transpose_naive_param_1];
ld.param.u32 %r3, [transpose_naive_param_2];
ld.param.u32 %r4, [transpose_naive_param_3];
mov.u32 %r5, %ctaid.x;
mov.u32 %r6, %ntid.x;
mov.u32 %r7, %tid.x;
mad.lo.s32 %r1, %r5, %r6, %r7;
mov.u32 %r8, %ntid.y;
mov.u32 %r9, %ctaid.y;
mov.u32 %r10, %tid.y;
mad.lo.s32 %r2, %r9, %r8, %r10;
setp.ge.s32 %p1, %r2, %r3;
setp.ge.s32 %p2, %r1, %r4;
or.pred %p3, %p1, %p2;
@%p3 bra $L__BB5_2;
cvta.to.global.u64 %rd3, %rd1;
mad.lo.s32 %r11, %r2, %r4, %r1;
mul.wide.s32 %rd4, %r11, 4;
add.s64 %rd5, %rd3, %rd4;
ld.global.f32 %f1, [%rd5];
mad.lo.s32 %r12, %r1, %r3, %r2;
cvta.to.global.u64 %rd6, %rd2;
mul.wide.s32 %rd7, %r12, 4;
add.s64 %rd8, %rd6, %rd7;
st.global.f32 [%rd8], %f1;
$L__BB5_2:
ret;
}
// .globl matrix_mul
.visible .entry matrix_mul(
.param .u64 matrix_mul_param_0,
.param .u64 matrix_mul_param_1,
.param .u64 matrix_mul_param_2,
.param .u32 matrix_mul_param_3,
.param .u32 matrix_mul_param_4,
.param .u32 matrix_mul_param_5
)
{
.reg .pred %p<12>;
.reg .f32 %f<63>;
.reg .b32 %r<44>;
.reg .b64 %rd<16>;
// demoted variable
.shared .align 4 .b8 _ZZ10matrix_mulE4ds_A[1024];
// demoted variable
.shared .align 4 .b8 _ZZ10matrix_mulE4ds_B[1024];
ld.param.u64 %rd4, [matrix_mul_param_0];
ld.param.u64 %rd5, [matrix_mul_param_1];
ld.param.u64 %rd6, [matrix_mul_param_2];
ld.param.u32 %r20, [matrix_mul_param_3];
ld.param.u32 %r21, [matrix_mul_param_4];
ld.param.u32 %r22, [matrix_mul_param_5];
mov.u32 %r23, %ctaid.y;
shl.b32 %r24, %r23, 4;
mov.u32 %r41, %tid.y;
add.s32 %r2, %r24, %r41;
mov.u32 %r25, %ctaid.x;
shl.b32 %r3, %r25, 4;
mov.u32 %r40, %tid.x;
add.s32 %r5, %r3, %r40;
setp.lt.s32 %p1, %r22, 1;
mov.f32 %f62, 0f00000000;
@%p1 bra $L__BB6_7;
shl.b32 %r27, %r41, 6;
mov.u32 %r28, _ZZ10matrix_mulE4ds_A;
add.s32 %r8, %r28, %r27;
shl.b32 %r29, %r40, 2;
add.s32 %r6, %r8, %r29;
mov.u32 %r30, _ZZ10matrix_mulE4ds_B;
add.s32 %r31, %r30, %r27;
add.s32 %r7, %r31, %r29;
add.s32 %r9, %r30, %r29;
mad.lo.s32 %r32, %r41, %r21, %r40;
add.s32 %r42, %r32, %r3;
mad.lo.s32 %r33, %r22, %r2, %r40;
cvta.to.global.u64 %rd7, %rd4;
mul.wide.s32 %rd8, %r33, 4;
add.s64 %rd15, %rd7, %rd8;
add.s32 %r34, %r22, 15;
shr.s32 %r35, %r34, 31;
shr.u32 %r36, %r35, 28;
add.s32 %r37, %r34, %r36;
shr.s32 %r11, %r37, 4;
mov.u32 %r43, 0;
cvta.to.global.u64 %rd9, %rd5;
$L__BB6_2:
setp.ge.s32 %p2, %r40, %r22;
setp.ge.s32 %p3, %r2, %r20;
mov.f32 %f61, 0f00000000;
or.pred %p4, %p3, %p2;
mov.f32 %f60, %f61;
@%p4 bra $L__BB6_4;
ld.global.f32 %f60, [%rd15];
$L__BB6_4:
st.shared.f32 [%r6], %f60;
setp.ge.s32 %p5, %r41, %r22;
setp.ge.s32 %p6, %r5, %r21;
or.pred %p7, %p6, %p5;
@%p7 bra $L__BB6_6;
mul.wide.s32 %rd10, %r42, 4;
add.s64 %rd11, %rd9, %rd10;
ld.global.f32 %f61, [%rd11];
$L__BB6_6:
st.shared.f32 [%r7], %f61;
bar.sync 0;
ld.shared.f32 %f12, [%r9];
ld.shared.f32 %f13, [%r8];
fma.rn.f32 %f14, %f13, %f12, %f62;
ld.shared.f32 %f15, [%r9+64];
ld.shared.f32 %f16, [%r8+4];
fma.rn.f32 %f17, %f16, %f15, %f14;
ld.shared.f32 %f18, [%r9+128];
ld.shared.f32 %f19, [%r8+8];
fma.rn.f32 %f20, %f19, %f18, %f17;
ld.shared.f32 %f21, [%r9+192];
ld.shared.f32 %f22, [%r8+12];
fma.rn.f32 %f23, %f22, %f21, %f20;
ld.shared.f32 %f24, [%r9+256];
ld.shared.f32 %f25, [%r8+16];
fma.rn.f32 %f26, %f25, %f24, %f23;
ld.shared.f32 %f27, [%r9+320];
ld.shared.f32 %f28, [%r8+20];
fma.rn.f32 %f29, %f28, %f27, %f26;
ld.shared.f32 %f30, [%r9+384];
ld.shared.f32 %f31, [%r8+24];
fma.rn.f32 %f32, %f31, %f30, %f29;
ld.shared.f32 %f33, [%r9+448];
ld.shared.f32 %f34, [%r8+28];
fma.rn.f32 %f35, %f34, %f33, %f32;
ld.shared.f32 %f36, [%r9+512];
ld.shared.f32 %f37, [%r8+32];
fma.rn.f32 %f38, %f37, %f36, %f35;
ld.shared.f32 %f39, [%r9+576];
ld.shared.f32 %f40, [%r8+36];
fma.rn.f32 %f41, %f40, %f39, %f38;
ld.shared.f32 %f42, [%r9+640];
ld.shared.f32 %f43, [%r8+40];
fma.rn.f32 %f44, %f43, %f42, %f41;
ld.shared.f32 %f45, [%r9+704];
ld.shared.f32 %f46, [%r8+44];
fma.rn.f32 %f47, %f46, %f45, %f44;
ld.shared.f32 %f48, [%r9+768];
ld.shared.f32 %f49, [%r8+48];
fma.rn.f32 %f50, %f49, %f48, %f47;
ld.shared.f32 %f51, [%r9+832];
ld.shared.f32 %f52, [%r8+52];
fma.rn.f32 %f53, %f52, %f51, %f50;
ld.shared.f32 %f54, [%r9+896];
ld.shared.f32 %f55, [%r8+56];
fma.rn.f32 %f56, %f55, %f54, %f53;
ld.shared.f32 %f57, [%r9+960];
ld.shared.f32 %f58, [%r8+60];
fma.rn.f32 %f62, %f58, %f57, %f56;
bar.sync 0;
shl.b32 %r38, %r21, 4;
add.s32 %r42, %r42, %r38;
add.s32 %r41, %r41, 16;
add.s64 %rd15, %rd15, 64;
add.s32 %r40, %r40, 16;
add.s32 %r43, %r43, 1;
setp.lt.s32 %p8, %r43, %r11;
@%p8 bra $L__BB6_2;
$L__BB6_7:
setp.ge.s32 %p9, %r5, %r21;
setp.ge.s32 %p10, %r2, %r20;
or.pred %p11, %p10, %p9;
@%p11 bra $L__BB6_9;
mad.lo.s32 %r39, %r2, %r21, %r5;
cvta.to.global.u64 %rd12, %rd6;
mul.wide.s32 %rd13, %r39, 4;
add.s64 %rd14, %rd12, %rd13;
st.global.f32 [%rd14], %f62;
$L__BB6_9:
ret;
}
// .globl column_reduce
.visible .entry column_reduce(
.param .u64 column_reduce_param_0,
.param .u64 column_reduce_param_1,
.param .u32 column_reduce_param_2,
.param .u32 column_reduce_param_3
)
{
.reg .pred %p<7>;
.reg .f32 %f<25>;
.reg .b32 %r<24>;
.reg .b64 %rd<22>;
ld.param.u64 %rd11, [column_reduce_param_0];
ld.param.u64 %rd10, [column_reduce_param_1];
ld.param.u32 %r11, [column_reduce_param_2];
ld.param.u32 %r12, [column_reduce_param_3];
cvta.to.global.u64 %rd1, %rd11;
mov.u32 %r13, %ntid.x;
mov.u32 %r14, %ctaid.x;
mov.u32 %r15, %tid.x;
mad.lo.s32 %r1, %r14, %r13, %r15;
setp.ge.s32 %p1, %r1, %r12;
@%p1 bra $L__BB7_9;
setp.lt.s32 %p2, %r11, 1;
mov.f32 %f24, 0f00000000;
@%p2 bra $L__BB7_8;
add.s32 %r17, %r11, -1;
and.b32 %r23, %r11, 3;
setp.lt.u32 %p3, %r17, 3;
mov.f32 %f24, 0f00000000;
mov.u32 %r22, 0;
@%p3 bra $L__BB7_5;
sub.s32 %r21, %r11, %r23;
mul.wide.s32 %rd12, %r1, 4;
add.s64 %rd20, %rd1, %rd12;
mul.wide.s32 %rd3, %r12, 4;
$L__BB7_4:
ld.global.f32 %f12, [%rd20];
add.f32 %f13, %f24, %f12;
add.s64 %rd13, %rd20, %rd3;
ld.global.f32 %f14, [%rd13];
add.f32 %f15, %f13, %f14;
add.s64 %rd14, %rd13, %rd3;
ld.global.f32 %f16, [%rd14];
add.f32 %f17, %f15, %f16;
add.s64 %rd15, %rd14, %rd3;
add.s64 %rd20, %rd15, %rd3;
ld.global.f32 %f18, [%rd15];
add.f32 %f24, %f17, %f18;
add.s32 %r22, %r22, 4;
add.s32 %r21, %r21, -4;
setp.ne.s32 %p4, %r21, 0;
@%p4 bra $L__BB7_4;
$L__BB7_5:
setp.eq.s32 %p5, %r23, 0;
@%p5 bra $L__BB7_8;
mad.lo.s32 %r19, %r22, %r12, %r1;
mul.wide.s32 %rd16, %r19, 4;
add.s64 %rd21, %rd1, %rd16;
mul.wide.s32 %rd7, %r12, 4;
$L__BB7_7:
.pragma "nounroll";
ld.global.f32 %f19, [%rd21];
add.f32 %f24, %f24, %f19;
add.s64 %rd21, %rd21, %rd7;
add.s32 %r23, %r23, -1;
setp.ne.s32 %p6, %r23, 0;
@%p6 bra $L__BB7_7;
$L__BB7_8:
cvta.to.global.u64 %rd17, %rd10;
mul.wide.s32 %rd18, %r1, 4;
add.s64 %rd19, %rd17, %rd18;
st.global.f32 [%rd19], %f24;
$L__BB7_9:
ret;
}