[QST] Gemm got 'incomplete type is not allowed' when use Sm90

When I change https://github.com/efeslab/Nanoflow/blob/main/pipeline/include/cutlassGemmWrapperImpl.cuh#L89 SmArch to Sm90, got

/root/Nanoflow/3rdparty/cutlass/include/cutlass/gemm/device/gemm.h(264): error: incomplete type is not allowed                                                                         
    using GemmKernel = typename kernel::DefaultGemm<                                                                                                                                                       
                                ^                                                                                                                                                                          
          detected during:                                                                                                                                                                                 
            instantiation of class "cutlass::gemm::device::Gemm<ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_
, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, GatherA, GatherB, ScatterD, PermuteDLayout> [with ElementA_=BaseGEMMWrapper::ElementI
nputA, LayoutA_=cutlass::layout::ColumnMajor, ElementB_=BaseGEMMWrapper::ElementInputB, LayoutB_=cutlass::layout::ColumnMajor, ElementC_=BaseGEMMWrapper::ElementOutput, LayoutC_=cutlass::layout::ColumnMa
jor, ElementAccumulator_=BaseGEMMWrapper::ElementAccumulator, OperatorClass_=cutlass::arch::OpClassTensorOp, ArchTag_=cutlass::arch::Sm90, ThreadblockShape_=cutlass::gemm::GemmShape<64, 128, 32>, WarpSha
pe_=cutlass::gemm::GemmShape<32, 64, 32>, InstructionShape_=cutlass::gemm::GemmShape<16, 8, 16>, EpilogueOutputOp_=cutlass::epilogue::thread::LinearCombination<BaseGEMMWrapper::ElementOutput, 8, BaseGEMM
Wrapper::ElementAccumulator, BaseGEMMWrapper::ElementAccumulator, cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, BaseGEMMWrapper::ElementOutput>, ThreadblockSw
izzle_=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<1>, Stages=2, AlignmentA=8, AlignmentB=8, SplitKSerial=false, Operator_=<error-type>, GatherA=false, GatherB=false, ScatterD=false, Permu
teDLayout=cutlass::layout::NoPermute]" at line 151 of /root/Nanoflow/pipeline/include/cutlassGemmWrapperImpl.cuh                                                                       
            instantiation of class "CutlassGEMMWrapper<cta_m, cta_n, cta_k, warp_m, warp_n, warp_k, split_k, stages, LayoutInputA_, LayoutInputB_, LayoutOutput_> [with cta_m=64, cta_n=128, cta_k=32, warp
_m=32, warp_n=64, warp_k=32, split_k=1, stages=2, LayoutInputA_=cutlass::layout::ColumnMajor, LayoutInputB_=cutlass::layout::ColumnMajor, LayoutOutput_=cutlass::layout::ColumnMajor]" at line 387 of /root/Nanoflow/pipeline/src/generate-gemm/cutlassGemmFactory_12.cu                                                                                                                          
                                                                                                                                                                                                           
/root/Nanoflow/3rdparty/cutlass/include/cutlass/gemm/device/gemm.h(352): error: expected a ";"                                                                                         
    typename GemmKernel::Params params_;                                                                                                                                                                   
                                ^                                                                                                                                                                          
          detected during:                                                                                                                                                                                 
            instantiation of class "cutlass::gemm::device::Gemm<ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_
, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, GatherA, GatherB, ScatterD, PermuteDLayout> [with ElementA_=BaseGEMMWrapper::ElementI
nputA, LayoutA_=cutlass::layout::ColumnMajor, ElementB_=BaseGEMMWrapper::ElementInputB, LayoutB_=cutlass::layout::ColumnMajor, ElementC_=BaseGEMMWrapper::ElementOutput, LayoutC_=cutlass::layout::ColumnMa
jor, ElementAccumulator_=BaseGEMMWrapper::ElementAccumulator, OperatorClass_=cutlass::arch::OpClassTensorOp, ArchTag_=cutlass::arch::Sm90, ThreadblockShape_=cutlass::gemm::GemmShape<64, 128, 32>, WarpSha
pe_=cutlass::gemm::GemmShape<32, 64, 32>, InstructionShape_=cutlass::gemm::GemmShape<16, 8, 16>, EpilogueOutputOp_=cutlass::epilogue::thread::LinearCombination<BaseGEMMWrapper::ElementOutput, 8, BaseGEMM
Wrapper::ElementAccumulator, BaseGEMMWrapper::ElementAccumulator, cutlass::epilogue::thread::ScaleType::Default, cutlass::FloatRoundStyle::round_to_nearest, BaseGEMMWrapper::ElementOutput>, ThreadblockSw
izzle_=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<1>, Stages=2, AlignmentA=8, AlignmentB=8, SplitKSerial=false, Operator_=<error-type>, GatherA=false, GatherB=false, ScatterD=false, Permu
teDLayout=cutlass::layout::NoPermute]" at line 151 of /root/Nanoflow/pipeline/include/cutlassGemmWrapperImpl.cuh
            instantiation of class "CutlassGEMMWrapper<cta_m, cta_n, cta_k, warp_m, warp_n, warp_k, split_k, stages, LayoutInputA_, LayoutInputB_, LayoutOutput_> [with cta_m=64, cta_n=128, cta_k=32, warp
_m=32, warp_n=64, warp_k=32, split_k=1, stages=2, LayoutInputA_=cutlass::layout::ColumnMajor, LayoutInputB_=cutlass::layout::ColumnMajor, LayoutOutput_=cutlass::layout::ColumnMajor]" at line 387 of /root/Nanoflow/pipeline/src/generate-gemm/cutlassGemmFactory_12.cu

However, there is no problem when using Sm80

Start with pipeline/include/cutlassGemmWrapperImpl.cuh at the SmArch selection and pipeline/src/generate-gemm/cutlassGemmFactory_12.cu at the reported instantiation. Compare the Sm80 and Sm90 paths, then inspect cutlass/gemm/device/gemm.h around lines 264 and 352 and reproduce the build. Done means the Sm90 configuration instantiates successfully without the incomplete-type errors while Sm80 remains working.