[BUG] Instantiated tile sizes for sm120 fail to meet kernel's layout assertion during cutlass_profiler compile

Open

#3,646 opened on Sep 20, 2026 1 comment 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 4 days

@LiRunGuo is already working on this.

Since Sep 22, 2026.

#3659 by @LiRunGuo — open

Assessment

Difficulty: 4/5
Estimated time: 3-5 days
Newbie friendliness: 50/100

Issue type: Bug
Clarity: Mostly clear
Activity status: Active
Tech stack: cpp
Domain: build-system, hpc

Research direction

Reproduce the failure with make cutlass_profiler -j8, then trace the generated SM120 GEMM file into include/cutlass/epilogue/collective/builders/sm120_builder.inl and include/cutlass/gemm/collective/builders/sm90_common.inl. Compare the instantiated tile and layout parameters shown in the log; done means cutlass_profiler compiles without the layout assertion and follow-on tile_to_shape errors.

Written by the indexing model from the issue text.

Description

? - Needs Triage bug CUTLASS C++ inactive-30d

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug

make cutlass_profiler -j8 fails with the following log.
The log suggests the generated tile sizes for the kernel cannot meet a layout assertion.

Error log:

~/cutlass-4.8.0dev/include/cutlass/gemm/collective/builders/sm90_common.inl(321): error: static assertion failed with "BLK_K0 must be a multiple of size<1>(GMMA::Layout_K_INTER_Atom<ElementType>{})"
        static_assert(BLK_K0 % size<1>(GMMA::Layout_K_INTER_Atom<ElementType>{}) == 0,
        ^
          detected during:
            instantiation of "auto cutlass::gemm::collective::detail::ss_smem_selector<major,ElementType,BLK_MN,BLK_K>() [with major=cute::SM90::GMMA::Major::K, ElementType=cutlass::float_e5m2_t, BLK_MN=cute::_128, BLK_K=cute::_8]" at line 138 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/builders/sm120_builder.inl
            instantiation of "auto cutlass::epilogue::collective::detail::sm120_get_epilogue_smem_swizzle_layout_atom<GmemStrideType,Element_,EpilogueTile_MN>() [with GmemStrideType=cute::tuple<int64_t, cute::C<1>, int64_t>, Element_=cutlass::float_e5m2_t, EpilogueTile_MN=cute::tuple<cute::_128, cute::_8>]" at line 376 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/builders/sm120_builder.inl
            instantiation of class "cutlass::epilogue::collective::detail::Sm120TmaBuilderImpl<TileShape_MNK, EpilogueTile_MN, ElementAccumulator, ElementCompute, ElementC_, GmemLayoutTagC_, AlignmentC, ElementD_, GmemLayoutTagD, AlignmentD, FusionOpOrCallbacks, DispatchPolicy> [with TileShape_MNK=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_MN=cute::tuple<cute::_128, cute::_8>, ElementAccumulator=float, ElementCompute=float, ElementC_=void, GmemLayoutTagC_=cutlass::layout::RowMajor, AlignmentC=0, ElementD_=cutlass::float_e5m2_t, GmemLayoutTagD=cutlass::layout::RowMajor, AlignmentD=16, FusionOpOrCallbacks=cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, DispatchPolicy=cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>]" at line 514 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/builders/sm120_builder.inl
            instantiation of class "cutlass::epilogue::collective::CollectiveBuilder<cutlass::arch::Sm120, OpClass, TileShape_MNK, ClusterShape_MNK, EpilogueTileType, ElementAccumulator, ElementCompute, ElementC, GmemLayoutTagC, AlignmentC, ElementD, GmemLayoutTagD, AlignmentD, Schedule, FusionOperation, std::enable_if_t<<expression>, void>> [with OpClass=cutlass::arch::OpClassBlockScaledTensorOp, TileShape_MNK=cute::tuple<cute::_128, cute::_8, cute::_128>, ClusterShape_MNK=cute::tuple<cute::_1, cute::_1, cute::_1>, EpilogueTileType=cutlass::epilogue::collective::EpilogueTileAuto, ElementAccumulator=float, ElementCompute=float, ElementC=void, GmemLayoutTagC=cutlass::layout::RowMajor, AlignmentC=0, ElementD=cutlass::float_e5m2_t, GmemLayoutTagD=cutlass::layout::RowMajor, AlignmentD=16, Schedule=cutlass::epilogue::collective::EpilogueScheduleAuto, FusionOperation=cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>]" at line 46 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp(161): error: no instance of overloaded function "tile_to_shape" matches the argument list
            argument types are: (void, cute::tuple<cute::_128, cute::_8, cute::C<1>>, cute::tuple<cute::_1, cute::_2, cute::_3>)
    using SmemLayoutC = decltype(tile_to_shape(
                                 ^
~/cutlass-4.8.0dev/include/cute/layout.hpp(1818): note #3327-D: candidate function template "cute::tile_to_shape(const cute::Layout<Shape, Stride> &, const TrgShape &, const ModeOrder &)" failed deduction
  tile_to_shape(Layout<Shape,Stride> const& block,
  ^
~/cutlass-4.8.0dev/include/cute/layout_composed.hpp(548): note #3327-D: candidate function template "cute::tile_to_shape(const cute::ComposedLayout<A, O, B> &, const Shape &, const ModeOrder &)" failed deduction
  tile_to_shape(ComposedLayout<A,O,B> const& layout,
  ^
          detected during instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp(165): error: no instance of overloaded function "tile_to_shape" matches the argument list
            argument types are: (void, cute::tuple<cute::_128, cute::_8, cute::C<1>>, cute::tuple<cute::_1, cute::_2, cute::_3>)
    using SmemLayoutD = decltype(tile_to_shape(
                                 ^
~/cutlass-4.8.0dev/include/cute/layout.hpp(1818): note #3327-D: candidate function template "cute::tile_to_shape(const cute::Layout<Shape, Stride> &, const TrgShape &, const ModeOrder &)" failed deduction
  tile_to_shape(Layout<Shape,Stride> const& block,
  ^
~/cutlass-4.8.0dev/include/cute/layout_composed.hpp(548): note #3327-D: candidate function template "cute::tile_to_shape(const cute::ComposedLayout<A, O, B> &, const Shape &, const ModeOrder &)" failed deduction
  tile_to_shape(ComposedLayout<A,O,B> const& layout,
  ^
          detected during instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp(171): error: no instance of overloaded function "cosize" matches the argument list
            argument types are: (<error-type>)
                                              && cosize(take<0,2>(SmemLayoutC{})) == cosize(take<0,2>(SmemLayoutD{}));
                                                 ^
~/cutlass-4.8.0dev/include/cute/layout.hpp(664): note #3327-D: candidate function template "cute::cosize<Is...,Shape,Stride>(const cute::Layout<Shape, Stride> &)" failed deduction
  cosize(Layout<Shape,Stride> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/layout_composed.hpp(274): note #3327-D: candidate function template "cute::cosize<Is...,A,O,B>(const cute::ComposedLayout<A, O, B> &)" failed deduction
  cosize(ComposedLayout<A,O,B> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/swizzle_layout.hpp(172): note #3327-D: candidate function template "cute::cosize<Is...,B,M,S,Offset,LayoutB>(const cute::ComposedLayout<cute::Swizzle<B, M, S>, Offset, LayoutB> &)" failed deduction
  cosize(ComposedLayout<Swizzle<B,M,S>,Offset,LayoutB> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/tensor_zip.hpp(298): note #3327-D: candidate function template "cute::cosize<Is...,Layouts...>(const cute::ZipLayout<Layouts...> &)" failed deduction
  cosize(ZipLayout<Layouts...> const& layouts)
  ^
          detected during:
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 172
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp(171): error: no instance of overloaded function "cosize" matches the argument list
            argument types are: (<error-type>)
                                              && cosize(take<0,2>(SmemLayoutC{})) == cosize(take<0,2>(SmemLayoutD{}));
                                                                                     ^
~/cutlass-4.8.0dev/include/cute/layout.hpp(664): note #3327-D: candidate function template "cute::cosize<Is...,Shape,Stride>(const cute::Layout<Shape, Stride> &)" failed deduction
  cosize(Layout<Shape,Stride> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/layout_composed.hpp(274): note #3327-D: candidate function template "cute::cosize<Is...,A,O,B>(const cute::ComposedLayout<A, O, B> &)" failed deduction
  cosize(ComposedLayout<A,O,B> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/swizzle_layout.hpp(172): note #3327-D: candidate function template "cute::cosize<Is...,B,M,S,Offset,LayoutB>(const cute::ComposedLayout<cute::Swizzle<B, M, S>, Offset, LayoutB> &)" failed deduction
  cosize(ComposedLayout<Swizzle<B,M,S>,Offset,LayoutB> const& layout)
  ^
~/cutlass-4.8.0dev/include/cute/tensor_zip.hpp(298): note #3327-D: candidate function template "cute::cosize<Is...,Layouts...>(const cute::ZipLayout<Layouts...> &)" failed deduction
  cosize(ZipLayout<Layouts...> const& layouts)
  ^
          detected during:
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 172
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cute/layout.hpp(670): error: no instance of function template "std::declval" matches the argument list
  using cosize_t = decltype(cosize(declval<Layout>()));
                                   ^
/usr/include/c++/15/type_traits(2612): note #3323-D: substituting explicit template arguments "<<error-type>>" for function template "std::declval" failed
      auto declval() noexcept -> decltype(__declval<_Tp>(0))
           ^
          detected during:
            instantiation of type "cute::cosize_t<<error-type>>" at line 673
            instantiation of "const auto cute::cosize_v [with Layout=<error-type>]" at line 178 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cute/layout.hpp(670): error: no instance of overloaded function "cute::cosize" matches the argument list
            argument types are: (<error-type>)
  using cosize_t = decltype(cosize(declval<Layout>()));
                            ^
~/cutlass-4.8.0dev/include/cute/layout.hpp(664): note #3327-D: candidate function template "cute::cosize<Is...,Shape,Stride>(const cute::Layout<Shape, Stride> &)" failed deduction
  cosize(Layout<Shape,Stride> const& layout)
  ^
          detected during:
            instantiation of type "cute::cosize_t<<error-type>>" at line 673
            instantiation of "const auto cute::cosize_v [with Layout=<error-type>]" at line 178 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

~/cutlass-4.8.0dev/include/cute/layout.hpp(673): error: name followed by "::" must be a class or namespace name
  static constexpr auto cosize_v = cosize_t<Layout>::value;
                                   ^
          detected during:
            instantiation of "const auto cute::cosize_v [with Layout=<error-type>]" at line 178 of ~/cutlass-4.8.0dev/include/cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp
            instantiation of class "cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<StagesC_, StagesD_, FragmentSize_, ReuseSmemC_, DelayTmaStore_>, CtaTileMNK_, EpilogueTile_, ElementC_, StrideC_, ElementD_, StrideD_, FusionCallbacks_, CopyOpG2S_, SmemLayoutAtomC_, CopyOpS2R_, CopyOpS2G_, SmemLayoutAtomD_, CopyOpR2S_, CopyAtomC_, CopyOpR2R_> [with StagesC_=1, StagesD_=1, FragmentSize_=4, ReuseSmemC_=false, DelayTmaStore_=true, CtaTileMNK_=cute::tuple<cute::_128, cute::_8, cute::_128>, EpilogueTile_=cute::tuple<cute::_128, cute::_8>, ElementC_=void, StrideC_=cute::tuple<int64_t, cute::C<1>, int64_t>, ElementD_=cutlass::float_e5m2_t, StrideD_=cute::tuple<int64_t, cute::C<1>, int64_t>, FusionCallbacks_=cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<1, 1, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::float_e5m2_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::_128, cute::_8, cute::_128>, cute::tuple<cute::_128, cute::_8>>, CopyOpG2S_=cute::SM90_TMA_LOAD, SmemLayoutAtomC_=void, CopyOpS2R_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyOpS2G_=cute::SM90_TMA_STORE, SmemLayoutAtomD_=void, CopyOpR2S_=cute::AutoVectorizingCopyWithAssumedAlignment<128>, CopyAtomC_=cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, CopyOpR2R_=void]" at line 59 of ~/cutlass-4.8.0dev/build/tools/library/generated/gemm/120/void_gemm_e4m3_e2m3/cutlass3x_sm120_bstensorop_gemm_ue8m0xe4m3_ue8m0xe2m3_f32_void_e5m2_128x8x128_1x1x1_0_tnt_align128_cooperative_q.cu

8 errors detected in the compilation of "~/cutlass-4.8.0dev/build/tools/library/cutlass_library_gemm_sm120_void_gemm_e4m3_e2m3_objs.unity.2da1110daef1.cu".
make[3]: *** [tools/library/CMakeFiles/cutlass_library_gemm_sm120_void_gemm_e4m3_e2m3_objs.dir/build.make:80: tools/library/CMakeFiles/cutlass_library_gemm_sm120_void_gemm_e4m3_e2m3_objs.dir/cutlass_library_gemm_sm120_void_gemm_e4m3_e2m3_objs.unity.2da1110daef1.cu.o] Error 2
make[2]: *** [CMakeFiles/Makefile2:24347: tools/library/CMakeFiles/cutlass_library_gemm_sm120_void_gemm_e4m3_e2m3_objs.dir/all] Error 2
make[2]: *** Waiting for unfinished jobs....
make[1]: *** [CMakeFiles/Makefile2:41983: tools/profiler/CMakeFiles/cutlass_profiler.dir/rule] Error 2
make: *** [Makefile:14227: cutlass_profiler] Error 2

Steps/Code to reproduce bug

cd cutlass-4.8.0dev/build
cmake .. -DCUTLASS_NVCC_ARCHS=120a -DCUTLASS_ENABLE_TESTS=OFF -DCUTLASS_UNITY_BUILD_ENABLED=ON
make cutlass_profiler -j8

Expected behavior
Successful compile without error.

Environment details (please complete the following information):
GPU: RTX 5090
Driver Version: 595.58.03 CUDA Version: 13.2

Dominant language: C++
Stars: 10.7k
Forks: 2.2k
Avg merge: 3d 20h
Merged PRs (30d): 4

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

First steps

Read the whole issue, then the project's contributing guide.
Comment on the issue to say you are picking it up — it saves two people doing the same work.
Fork the repository and make your change on a branch.
Open a pull request that references the issue number.

More from NVIDIA/cutlass

[Bug][CuTeDSL][Fuzz] `utils.gemm.sm100.epilogue` releases the accumulator before `tcgen05.ld` completes: no `tcgen05.wait::ld` before the arrive that lets the next MMA overwrite TMEM
Possibly taken @0z5a claimed this 4 days ago. Open
- Is something wrong?
bug CuTe DSL

Difficulty 2/5 1-3 hours Newbie friendliness 78/100

NVIDIA/cutlass#3805 · Opened Oct 31, 2026

Maintainers usually reply within 4 days
[QST] Remove unused Python protobuf dependency from nvidia-cutlass-dsl
Open
- Is something wrong?
? - Needs Triage question

Difficulty 2/5 1-3 hours Newbie friendliness 70/100

NVIDIA/cutlass#3796 · 1 comment · 1 reaction · Opened Oct 27, 2026

Maintainers usually reply within 4 days
[DOC] IKET quick start omits the GEMM GPU requirement and references a missing example
Possibly taken @0z5a claimed this 15 days ago. Open
- Is something wrong?
Difficulty 1/5 Under an hour Newbie friendliness 82/100

NVIDIA/cutlass#3761 · 1 comment · Opened Oct 20, 2026

Maintainers usually reply within 4 days
[BUG] CuTe DSL from_dlpack rejects float8_e4m3fn tensors (dtype_map lists Float8E5M2 only)
Possibly taken @Ammar-Alnagar claimed this 24 days ago. Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 74/100

NVIDIA/cutlass#3716 · 1 comment · Opened Oct 7, 2026

Maintainers usually reply within 4 days
EllGemm ColumnMajor-output workspace sizing uses the unswapped problem shape and undersizes split-K semaphores
Possibly taken @XFDG claimed this 52 days ago. Open
- Is something wrong?
CUTLASS C++

Difficulty 2/5 1-3 hours Newbie friendliness 88/100

NVIDIA/cutlass#3539 · 2 comments · Opened Aug 25, 2026

Maintainers usually reply within 4 days

All issues in NVIDIA/cutlass

Similar issues

[BUG] `cudf::unique_count` returns 0 for an all-null column with `null_policy::INCLUDE`
Open
- Is something wrong?
bug libcudf

Difficulty 2/5 1-3 hours Newbie friendliness 78/100

NVIDIA/cudf#25093 · Opened Nov 4, 2026

Maintainers usually reply within 2 days
5.x wodles unit tests run on Python 3.10 while the runtime embeds 3.12
Open
- Is something wrong?
level/task module/cloud monitoring type/maintenance

Difficulty 2/5 1-3 hours Newbie friendliness 78/100

wazuh/wazuh#41148 · Opened Nov 4, 2026

Maintainers usually reply within 1 day
WebGL: a canvas with width or height 0 gives a drawing buffer of 0 instead of 1 (Plotly 3D charts throw "gl-fbo: Framebuffer incomplete attachment")
Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 78/100

LadybirdBrowser/ladybird#13809 · Opened Nov 4, 2026

Maintainers usually reply within 1 day
nodes tests: test_pgvector_search fails on pgvector 0.8.7, which the test docstrings' image now ships
Open
- Is something wrong?
bug module:nodes tests

Difficulty 2/5 1-3 hours Newbie friendliness 78/100

rocketride-org/rocketride-server#3020 · Opened Nov 4, 2026

Maintainers usually reply within 2 days
fmt::formatted_size does not work with fmt::wformat_string
Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 72/100

fmtlib/fmt#5013 · Opened Nov 4, 2026

Maintainers usually reply within 9 days

More C++ issues

[BUG] Instantiated tile sizes for sm120 fail to meet kernel's layout assertion during cutlass_profiler compile

Is something wrong?

Assessment

Research direction

Description

Which component has the problem?

Bug Report

Getting set up

First steps

More from NVIDIA/cutlass

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Similar issues

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Get new issues in your inbox