Fix EVT for cutlass::gemm::kernel::DefaultGemmWithVisitor's behavior when constructing GemmUniversalAdapter

[Xinyu302]

The phenomenon I encountered while using the epilogue visitor tree in C++ is quite strange...
In examples/47_ampere_gemm_universal_streamk/ampere_gemm_universal_streamk_broadcast.cu:

using EVTKernelStreamK =
    typename cutlass::gemm::kernel::DefaultGemmWithVisitor<
    ElementA, LayoutA, cutlass::ComplexTransform::kNone, AlignmentA,
    ElementB, LayoutB, cutlass::ComplexTransform::kNone, AlignmentB,
    ElementC, LayoutC, AlignmentC,
    ElementAccumulator,
    ElementCompute,
    cutlass::arch::OpClassTensorOp,
    cutlass::arch::Sm80,
    ThreadblockShape,
    WarpShape,
    InstructionShape,
    EVTD,
    cutlass::gemm::threadblock::ThreadblockSwizzleStreamK,
    NumStages,
    cutlass::arch::OpMultiplyAdd,
    EVTEpilogueStages
>::GemmKernel;

using DeviceGemmStreamK = cutlass::gemm::device::GemmUniversalAdapter<EVTKernelStreamK>;

I didn't want to use cutlass::gemm::threadblock::ThreadblockSwizzleStreamK, so I replaced it with cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<1>.
When I construct DeviceGemmStreamK::Arguments, it errors like this:

type "cutlass::gemm::kernel::GemmUniversal<Mma_, Epilogue_, ThreadblockSwizzle_, void, std::enable_if_t<<expression>, void>>::WarpShape ...
/workspace/cutlass/python/cutlass_library/../../include/cutlass/gemm/kernel/gemm_universal.h(94): here is inaccessible

I think it is because when if we don't use ThreadblockSwizzleStreamK, DefaultGemmWithVisitor will choose GemmWithEpilogueVisitor as the type of GemmKernel. So EVTKernelStreamK will be a subclass of GemmWithEpilogueVisitor in include/cutlass/gemm/kernel/gemm_universal_with_visitor.h.

The GemmWithEpilogueVisitor is a subclass of GemmUniversal, GemmUniversal has all the types we need to help us construct GemmUniversalAdapter, but when GemmWithEpilogueVisitor inherits from GemmUniversal, it doesn't use public inheritance, which means that the required types like Mma, WarpShape are not accessible in GemmWithEpilogueVisitor.

So, I changed the inheritance relationship to public to ensure the consistency of DefaultGemmWithVisitor behavior. Besides this, I also need to modify the GemmWithEpilogueVisitor class.

I would like to hear your suggestions on this issue.

[Xinyu302]

@apuaaChen Could you please help to take a look at this problem?

[gau-nernst]

Thank you for the fix. I tried this locally and it works wonder! Much faster than ThreadblockSwizzleStreamK. Hope the maintainers can merge this.

[Xinyu302]

@hwu36 Could you please take a look at this PR?

[github-actions]

This PR has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this PR if it is no longer required. Otherwise, please respond with a comment indicating any updates. This PR will be labeled inactive-90d if there is no activity in the next 60 days.