mov.b64%rd33, _ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0;
mov.u64 %rd1, %rd33;
ld.param.u64 %rd34, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0+48];
cvta.to.global.u64 %rd2, %rd34;
mov.u32 %r301, %ctaid.x;
shl.b32 %r1, %r301, 4;
ld.param.u64 %rd36, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0+136];
cvta.to.global.u64 %rd4, %rd36;
mov.u32 %r302, %ctaid.y;
shl.b32 %r2, %r302, 4;
mov.u32 %r3, %tid.x;
shr.u32 %r4, %r3, 1;
and.b32 %r303, %r3, 1;
shl.b32 %r5, %r303, 3;
add.s32 %r6, %r5, %r1;
ld.param.u32 %r7, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0+20];
mad.lo.s32 %r304, %r7, %r4, %r6;
cvt.s64.s32%rd5, %r304;
mul.wide.s32 %rd37, %r304, 2;
add.s64 %rd6, %rd2, %rd37;
ld.param.u32 %r8, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0];
sub.s32 %r305, %r8, %r6;
ld.param.u32 %r9, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0+8];
setp.gt.s32%p2, %r9, 0;
setp.gt.s32%p3, %r305, 0;
and.pred %p4, %p3, %p2;
selp.u32%r306, 1, 0, %p4;
setp.gt.s32%p5, %r9, 16;
and.pred %p6, %p3, %p5;
selp.u16%rs1, 1, 0, %p6;
mul.wide.u16 %r10, %rs1, 2;
selp.u32%r307, -1, 0, %p6;
bfi.b32 %r481, %r307, %r306, 1, 1;
bfe.u32 %r308, %r3, 1, 4;
mad.lo.s32 %r12, %r308, 24, %r5;
setp.lt.s32%p7, %r12, 768;
@%p7 bra BB0_2;
ld.param.u32 %r310, [%rd1+4];
ld.param.u32 %r13, [%rd1+108];
add.s32 %r16, %r5, %r2;
mad.lo.s32 %r311, %r13, %r4, %r16;
cvt.s64.s32%rd7, %r311;
mul.wide.s32 %rd45, %r311, 2;
add.s64 %rd8, %rd4, %rd45;
sub.s32 %r312, %r310, %r16;
setp.gt.s32%p9, %r312, 0;
and.pred %p11, %p9, %p2;
selp.u32%r313, 1, 0, %p11;
and.pred %p13, %p9, %p5;
selp.u16%rs2, 1, 0, %p13;
mul.wide.u16 %r17, %rs2, 2;
selp.u32%r314, -1, 0, %p13;
bfi.b32 %r480, %r314, %r313, 1, 1;
@%p7 bra BB0_4;
setp.lt.s32%p14, %r9, 32;
@%p14 bra BB0_7;
setp.lt.s32%p18, %r4, %r9;
selp.b32%r321, %r481, %r10, %p18;
add.s32 %r322, %r4, 16;
setp.lt.s32%p19, %r322, %r9;
and.b32 %r323, %r321, 1;
selp.b32%r481, %r321, %r323, %p19;
selp.b32%r324, %r480, %r17, %p18;
and.b32 %r325, %r324, 1;
selp.b32%r480, %r324, %r325, %p19;
bra.uni BB0_8;
and.b32 %r330, %r481, 1;
setp.eq.b32%p20, %r330, 1;
mov.u32 %r571, 0;
mov.u32 %r559, %r571;
mov.u32 %r560, %r571;
mov.u32 %r561, %r571;
mov.u32 %r562, %r571;
@!%p20 bra BB0_10;
bra.uni BB0_9;
ld.global.v4.u32 {%r561, %r562, %r559, %r560}, [%rd6];
ld.param.u32 %r33, [%rd1+32];
cvt.s64.s32%rd53, %r33;
add.s64 %rd9, %rd5, %rd53;
shl.b64 %rd54, %rd9, 1;
add.s64 %rd10, %rd2, %rd54;
and.b32 %r339, %r481, 2;
setp.eq.s32%p21, %r339, 0;
mov.u32 %r563, %r571;
mov.u32 %r564, %r571;
mov.u32 %r565, %r571;
mov.u32 %r566, %r571;
@%p21 bra BB0_12;
and.b32 %r348, %r480, 1;
setp.eq.b32%p22, %r348, 1;
mov.u32 %r567, %r571;
mov.u32 %r568, %r571;
mov.u32 %r569, %r571;
mov.u32 %r570, %r571;
@!%p22 bra BB0_14;
bra.uni BB0_13;
ld.global.v4.u32 {%r569, %r570, %r567, %r568}, [%rd8];
ld.param.u32 %r50, [%rd1+120];
cvt.s64.s32%rd55, %r50;
add.s64 %rd56, %rd7, %rd55;
shl.b64 %rd57, %rd56, 1;
add.s64 %rd11, %rd4, %rd57;
and.b32 %r357, %r480, 2;
setp.eq.s32%p23, %r357, 0;
mov.u32 %r572, %r571;
mov.u32 %r573, %r571;
mov.u32 %r574, %r571;
@%p23 bra BB0_16;
shl.b32 %r362, %r12, 1;
mov.u32 %r363, _ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE$__cuda_local_var_29658_57_non_const_shared_storage;
add.s32 %r59, %r363, %r362;
st.shared.v4.u32 [%r59], {%r561, %r562, %r559, %r560};
st.shared.v4.u32 [%r59+768], {%r565, %r566, %r563, %r564};
st.shared.v4.u32 [%r59+1536], {%r569, %r570, %r567, %r568};
st.shared.v4.u32 [%r59+2304], {%r573, %r574, %r571, %r572};
bar.sync 0;
shr.u32 %r60, %r3, 5;
shl.b32 %r61, %r60, 4;
setp.lt.u32%p24, %r61, 768;
@%p24 bra BB0_18;
mul.wide.s32 %rd65, %r61, 2;
cvt.u64.u32 %temp, %r363;
cvta.shared.u64 %rd66, %temp;
add.s64 %rd67, %rd66, %rd65;
add.s64 %rd12, %rd67, 1536;
mov.u32 %r374, 24;
wmma.load.a.sync.col.m16n16k16.f16 {%r590, %r589, %r588, %r587, %r586, %r585, %r584, %r583}, [%rd66], %r374;
wmma.load.b.sync.row.m16n16k16.f16 {%r582, %r581, %r580, %r579, %r578, %r577, %r576, %r575}, [%rd12], %r374;
add.s32 %r599, %r9, -32;
mov.u32 %r609, 0;
setp.lt.s32%p25, %r599, 1;
@%p25 bra BB0_19;
mov.u32 %r610, %r609;
mov.u32 %r611, %r609;
mov.u32 %r612, %r609;
mov.u32 %r613, %r609;
mov.u32 %r614, %r609;
mov.u32 %r615, %r609;
mov.u32 %r616, %r609;
setp.lt.s32%p34, %r599, -31;
@%p34 bra BB0_35;
mov.u32 %r477, %tid.x;
shr.u32 %r476, %r477, 5;
shl.b32 %r475, %r476, 4;
add.s32 %r434, %r475, 384;
mul.wide.s32 %rd96, %r434, 2;
add.s64 %rd97, %rd66, %rd96;
add.s64 %rd26, %rd97, 1536;
add.s64 %rd99, %rd66, 768;
wmma.load.a.sync.col.m16n16k16.f16 {%r257, %r256, %r255, %r254, %r250, %r251, %r253, %r252}, [%rd99], %r374;
wmma.load.b.sync.row.m16n16k16.f16 {%r265, %r264, %r263, %r262, %r258, %r259, %r261, %r260}, [%rd26], %r374;
mov.b32 %f63, %r615;
mov.b32 %f64, %r616;
mov.b32 %f65, %r613;
mov.b32 %f66, %r614;
mov.b32 %f67, %r611;
mov.b32 %f68, %r612;
mov.b32 %f69, %r609;
mov.b32 %f70, %r610;
wmma.mma.sync.col.row.m16n16k16.f32.f32 {%f10, %f11, %f12, %f13, %f14, %f15, %f16, %f17}, {%r590, %r589, %r588, %r587, %r586, %r585, %r584, %r583}, {%r582, %r581, %r580, %r579, %r578, %r577, %r576, %r575}, {%f70, %f69, %f68, %f67, %f66, %f65, %f64, %f63};
bar.sync 0;
bar.sync 0;
wmma.mma.sync.col.row.m16n16k16.f32.f32 {%f71, %f72, %f73, %f74, %f75, %f76, %f77, %f78}, {%r257, %r256, %r255, %r254, %r250, %r251, %r253, %r252}, {%r265, %r264, %r263, %r262, %r258, %r259, %r261, %r260}, {%f10, %f11, %f12, %f13, %f14, %f15, %f16, %f17};
mov.b32 %r610, %f71;
mov.b32 %r609, %f72;
mov.b32 %r612, %f73;
mov.b32 %r611, %f74;
mov.b32 %r614, %f75;
mov.b32 %r613, %f76;
mov.b32 %r616, %f77;
mov.b32 %r615, %f78;
add.s32 %r599, %r599, -32;
setp.gt.s32%p35, %r599, -32;
@%p35 bra BB0_34;
mov.b64%rd145, _ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0;
mov.u64 %rd144, %rd145;
mov.u32 %r472, %tid.x;
ld.param.f32 %f145, [%rd144+428];
ld.param.u32 %r471, [%rd144+324];
shr.u32 %r470, %r472, 5;
ld.param.u32 %r469, [%rd144+332];
mov.u32 %r468, %ctaid.x;
shl.b32 %r467, %r468, 4;
mov.u32 %r466, _ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE$__cuda_local_var_29658_57_non_const_shared_storage;
mov.u32 %r465, %ctaid.y;
shl.b32 %r464, %r465, 4;
ld.param.v2.u32 {%r438, %r439}, [%rd144+288];
ld.param.u32 %r284, [%rd144+300];
ld.param.u32 %r440, [%rd144+320];
ld.param.u32 %r285, [%rd144+336];
ld.param.f32 %f18, [%rd144+424];
ld.param.u32 %r287, [%rd144+304];
ld.param.u64 %rd100, [%rd144+312];
cvta.to.global.u64 %rd101, %rd100;
and.b32 %r441, %r472, 3;
shl.b32 %r288, %r441, 2;
shr.u32 %r289, %r472, 2;
add.s32 %r290, %r289, %r464;
add.s32 %r291, %r288, %r467;
mad.lo.s32 %r442, %r290, %r440, %r291;
cvt.s64.s32%rd102, %r442;
mul.wide.s32 %rd103, %r442, 4;
add.s64 %rd27, %rd101, %rd103;
sub.s32 %r292, %r469, %r290;
shl.b32 %r293, %r470, 8;
shl.b32 %r443, %r289, 4;
add.s32 %r294, %r443, %r288;
shl.b32 %r444, %r294, 2;
add.s32 %r295, %r466, %r444;
cvt.s64.s32%rd104, %r471;
add.s64 %rd105, %rd102, %rd104;
shl.b64 %rd106, %rd105, 2;
add.s64 %rd28, %rd101, %rd106;
setp.eq.f32%p36, %f145, 0f00000000;
@%p36 bra BB0_48;
setp.lt.s32%p51, %r293, 256;
@%p51 bra BB0_50;
setp.lt.s32%p52, %r294, 256;
@%p52 bra BB0_52;
ld.param.u32 %r474, [_ZN7cutlass4gemm11gemm_kernelINS0_4GemmINS0_14WmmaGemmTraitsILNS_12MatrixLayout4KindE1ELS5_0ENS_5ShapeILi32ELi16ELi16ELi1EEEfNS0_13LinearScalingIfNS0_19FragmentMultiplyAddIfEEEEfS7_NS6_ILi16ELi16ELi16ELi1EEELi8ELi8EiNS0_20WmmaGemmTraitsHelperILS5_1ELS5_0ES7_ffSB_S7_SC_Li8ELi8EiEEEEEEEEvNT_6ParamsE_param_0];
mul.wide.s32 %rd142, %r293, 4;
add.s64 %rd32, %rd66, %rd142;
setp.lt.s32%p1, %r291, %r474;
bar.sync 0;
mov.b32 %f121, %r610;
mov.b32 %f122, %r609;
mov.b32 %f123, %r612;
mov.b32 %f124, %r611;
mov.b32 %f125, %r614;
mov.b32 %f126, %r613;
mov.b32 %f127, %r616;
mov.b32 %f128, %r615;
mov.u32 %r458, 16;
wmma.store.d.sync.col.m16n16k16.f32 [%rd32], {%f121, %f122, %f123, %f124, %f125, %f126, %f127, %f128}, %r458;
bar.sync 0;
ld.shared.v4.u32 {%r459, %r460, %r461, %r462}, [%r295+512];
setp.gt.s32%p53, %r292, 0;
and.pred %p54, %p1, %p53;
@!%p54 bra BB0_54;
bra.uni BB0_53;
ld.shared.v4.f32 {%f129, %f130, %f131, %f132}, [%r295];
mul.f32 %f137, %f18, %f132;
mul.f32 %f138, %f18, %f131;
mul.f32 %f139, %f18, %f130;
mul.f32 %f140, %f18, %f129;
st.global.v4.f32 [%rd27], {%f140, %f139, %f138, %f137};
mov.b32 %f141, %r459;
mul.f32 %f43, %f18, %f141;
mov.b32 %f142, %r460;
mul.f32 %f44, %f18, %f142;
mov.b32 %f143, %r461;
mul.f32 %f45, %f18, %f143;
mov.b32 %f144, %r462;
mul.f32 %f46, %f18, %f144;
sub.s32 %r463, %r292, %r285;
setp.gt.s32%p56, %r463, 0;
and.pred %p57, %p1, %p56;
@!%p57 bra BB0_56;
bra.uni BB0_55;
st.global.v4.f32 [%rd28], {%f43, %f44, %f45, %f46};
bra.uni BB0_56;
ret;