diff options
Diffstat (limited to 'cutlass-example/executionFlow')
| -rw-r--r-- | cutlass-example/executionFlow | 262 |
1 files changed, 262 insertions, 0 deletions
diff --git a/cutlass-example/executionFlow b/cutlass-example/executionFlow new file mode 100644 index 0000000..356afed --- /dev/null +++ b/cutlass-example/executionFlow @@ -0,0 +1,262 @@ +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; |