Last active
May 12, 2025 20:39
-
-
Save mlazos/c87b19997dc029316980c506234193d4 to your computer and use it in GitHub Desktop.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| using Accum = cutlass::epilogue::fusion::Sm90AccFetch; | |
| using ElementD = cutlass::half_t; | |
| using StrideD = cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>; | |
| using Bias = cutlass::epilogue::fusion::Sm90ColBroadcast< | |
| 0 /*Stages*/, typename EpilogueDescriptor::TileShape, cutlass::half_t, cutlass::half_t, | |
| cute::Stride<cute::Int<1>, cute::Int<0>, cute::Int<0>> | |
| >; | |
| using Compute0 = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::plus, float, float, | |
| cutlass::FloatRoundStyle::round_to_nearest | |
| >; | |
| using EVTCompute0 = cutlass::epilogue::fusion::Sm90EVT< | |
| Compute0, | |
| EVTD, | |
| Bias>; | |
| using EDescriptor = cutlass::epilogue::collective::detail::AuxStoreDescriptor< | |
| EpilogueDescriptor, cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>, cutlass::half_t | |
| >; | |
| using E = cutlass::epilogue::fusion::Sm90AuxStore< | |
| EDescriptor::Stages, typename EDescriptor::EpilogueTile, cutlass::half_t, | |
| cutlass::FloatRoundStyle::round_to_nearest, cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>, typename EDescriptor::SmemLayoutAtom, | |
| typename EDescriptor::CopyOpR2S | |
| >; | |
| using EVTE = cutlass::epilogue::fusion::Sm90EVT< | |
| E, | |
| EVTCompute0>; | |
| using ElementC = void; | |
| using StrideC = StrideD; | |
| Accum |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| using Accum = cutlass::epilogue::fusion::Sm90AccFetch; | |
| using EDescriptor = cutlass::epilogue::collective::detail::AuxStoreDescriptor< | |
| EpilogueDescriptor, cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>, cutlass::half_t | |
| >; | |
| using E = cutlass::epilogue::fusion::Sm90AuxStore< | |
| EDescriptor::Stages, typename EDescriptor::EpilogueTile, cutlass::half_t, | |
| cutlass::FloatRoundStyle::round_to_nearest, cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>, typename EDescriptor::SmemLayoutAtom, | |
| typename EDescriptor::CopyOpR2S | |
| >; | |
| using EVTE = cutlass::epilogue::fusion::Sm90EVT< | |
| E, | |
| Accum>; | |
| using Bias = cutlass::epilogue::fusion::Sm90ColBroadcast< | |
| 0 /*Stages*/, typename EpilogueDescriptor::TileShape, cutlass::half_t, cutlass::half_t, | |
| cute::Stride<cute::Int<1>, cute::Int<0>, cute::Int<0>> | |
| >; | |
| using Compute0 = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::plus, cutlass::half_t, float, | |
| cutlass::FloatRoundStyle::round_to_nearest | |
| >; | |
| using EVTCompute0 = cutlass::epilogue::fusion::Sm90EVT< | |
| Compute0, | |
| EVTE, | |
| Bias>; | |
| using ElementD = cutlass::half_t; | |
| using StrideD = cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>; | |
| using ElementC = void; | |
| using StrideC = StrideD; | |
| EVTCompute0 |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment