1 files changed, 0 insertions, 262 deletions
diff --git a/cutlass-example/executionFlow b/cutlass-example/executionFlow
deleted file mode 100644
index 356afed..0000000
--- a/cutlass-example/executionFlow
+++ /dev/null
@@ -1,262 +0,0 @@
-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;