code.cpp

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