[BUG] No appropriate MMATraits specialization for SM100_MMA_F8F6F4_2x1SM_SS

[infinitron]

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug
I'm unable to create a tiled MMA operation using SM100_MMA_F8F6F4_2x1SM_SS. The compilation fails as it is unable to find the appropriate MMATraits specialization
Steps/Code to reproduce bug

TiledMMA tiled_mma_mma = make_tiled_mma(
      SM100_MMA_F8F6F4_2x1SM_SS_M{});

Compiling the above code fails with the following errors

include/cute/atom/mma_traits.hpp(66): error: static assertion failed with "MMA_Traits not implemented for this MMA_Operation."
    static_assert(sizeof(MMAOperation) == 0, "MMA_Traits not implemented for this MMA_Operation.");
    ^
          detected during:
            instantiation of class "cute::MMA_Traits<MMAOperation, MMAOpArgs...> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAOpArgs=<>]" at line 50 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553 of include/cute/atom/mma_atom.hpp
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(56): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeD"
    using ValTypeD = typename Traits::ValTypeD;
                                      ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(57): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeA"
    using ValTypeA = typename Traits::ValTypeA;
                                      ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(58): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeB"
    using ValTypeB = typename Traits::ValTypeB;
                                      ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(59): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeC"
    using ValTypeC = typename Traits::ValTypeC;
                                      ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(62): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "Shape_MNK"
    using Shape_MNK = typename Traits::Shape_MNK;
                                       ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(63): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ThrID"
    using ThrID = typename Traits::ThrID;
                                   ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(64): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "CLayout"
    using LayoutC_TV = typename Traits::CLayout;
                                        ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(65): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ALayout"
    using LayoutA_TV = typename Traits::ALayout;
                                        ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(66): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "BLayout"
    using LayoutB_TV = typename Traits::BLayout;
                                        ^
          detected during:
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_traits.hpp(183): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeC"
  struct FrgTypeC_or_Default { using type = typename X::ValTypeC; };
                                                        ^
          detected during:
            instantiation of class "cute::detail::FrgTypeC_or_Default<X, <unnamed>> [with X=cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>, <unnamed>=void]" at line 69 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553 of include/cute/atom/mma_atom.hpp
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_traits.hpp(173): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeA"
  struct FrgTypeA_or_Default { using type = typename X::ValTypeA; };
                                                        ^
          detected during:
            instantiation of class "cute::detail::FrgTypeA_or_Default<X, <unnamed>> [with X=cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>, <unnamed>=void]" at line 70 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553 of include/cute/atom/mma_atom.hpp
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_traits.hpp(178): error: class "cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>" has no member "ValTypeB"
  struct FrgTypeB_or_Default { using type = typename X::ValTypeB; };
                                                        ^
          detected during:
            instantiation of class "cute::detail::FrgTypeB_or_Default<X, <unnamed>> [with X=cute::MMA_Traits<cute::SM100_MMA_F8F6F4_2x1SM_SS>, <unnamed>=void]" at line 71 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<cute::MMA_Traits<MMAOperation, Args...>> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS, Args=<>]" at line 45 of include/cute/atom/mma_atom.hpp
            instantiation of class "cute::MMA_Atom<MMAOperation> [with MMAOperation=cute::SM100_MMA_F8F6F4_2x1SM_SS]" at line 553 of include/cute/atom/mma_atom.hpp
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(225): error: no instance of overloaded function "cute::tiled_product" matches the argument list
            argument types are: (<error-type>, cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>)
    using ThrLayoutVMNK = decltype(tiled_product(AtomThrID{}, AtomLayoutMNK{}));
                                   ^
include/cute/layout_composed.hpp(511): note #3327-D: candidate function template "cute::tiled_product(const cute::ComposedLayout<A, O, B> &, const Tiler &)" failed deduction
  tiled_product(ComposedLayout<A,O,B> const& a,
  ^
include/cute/layout.hpp(1714): note #3327-D: candidate function template "cute::tiled_product(const cute::Layout<LShape, LStride> &, const Tiler &)" failed deduction
  tiled_product(Layout<LShape,LStride> const& block,
  ^
          detected during:
            instantiation of class "cute::TiledMMA<MMA_Atom, AtomLayoutMNK, PermutationMNK> [with MMA_Atom=cute::MMA_Atom<cute::SM100_MMA_F8F6F4_2x1SM_SS>, AtomLayoutMNK=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, PermutationMNK=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 540
            instantiation of "auto cute::make_tiled_mma(const cute::MMA_Atom<MMA_Op> &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 553
            instantiation of "auto cute::make_tiled_mma(const MMA_Op &, const MMAThrLayout &, const Permutations &) [with MMA_Op=cute::SM100_MMA_F8F6F4_2x1SM_SS, MMAThrLayout=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, Permutations=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

include/cute/atom/mma_atom.hpp(231): error: no instance of overloaded function "cute::tiled_product" matches the argument list
            argument types are: (<error-type>, const cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>)
        thr_layout_vmnk_(tiled_product(AtomThrID{}, thr_layout_mnk)) {}
                         ^
include/cute/layout_composed.hpp(511): note #3327-D: candidate function template "cute::tiled_product(const cute::ComposedLayout<A, O, B> &, const Tiler &)" failed deduction
  tiled_product(ComposedLayout<A,O,B> const& a,
  ^
include/cute/layout.hpp(1714): note #3327-D: candidate function template "cute::tiled_product(const cute::Layout<LShape, LStride> &, const Tiler &)" failed deduction
  tiled_product(Layout<LShape,LStride> const& block,
  ^
          detected during instantiation of "cute::TiledMMA<MMA_Atom, AtomLayoutMNK, PermutationMNK>::TiledMMA(const MMA_Atom &, const AtomLayoutMNK &) [with MMA_Atom=cute::MMA_Atom<cute::SM100_MMA_F8F6F4_2x1SM_SS>, AtomLayoutMNK=cute::Layout<cute::tuple<cute::_1, cute::_1, cute::_1>, cute::tuple<cute::_0, cute::_0, cute::C<0>>>, PermutationMNK=cute::tuple<cute::Underscore, cute::Underscore, cute::Underscore>]" at line 943 of /cutlass-dev/blackwell_tma_multicast_2SM_test.cu

Expected behavior
Should compile without errors

I can compile the code if I change SM100_MMA_F8F6F4_2x1SM_SS to a template and use it in its MMATraits specialization

template <class a_type, class b_type, class c_type, int M, int N,
          UMMA::Major a_major, UMMA::Major b_major,
          UMMA::ScaleIn a_neg = UMMA::ScaleIn::One,
          UMMA::ScaleIn b_neg = UMMA::ScaleIn::One>
struct SM100_MMA_F8F6F4_2x1SM_SS{...}


template <class a_type, class b_type, class c_type, int M, int N,
          UMMA::Major a_major, UMMA::Major b_major, UMMA::ScaleIn a_neg,
          UMMA::ScaleIn b_neg>
struct MMA_Traits<SM100_MMA_F8F6F4_2x1SM_SS_M<a_type, b_type, c_type, M, N,
                                              a_major, b_major, a_neg, b_neg>>{...}

[hwu36]

Thank you for filing the issue. I think you are right. Do you want to file a PR to add templates to

cutlass/include/cute/arch/mma_sm100_umma.hpp

Line 1455 in ae6bccf

struct SM100_MMA_F8F6F4_2x1SM_SS

?

It should look like https://github.com/NVIDIA/cutlass/blame/ae6bccf341fb4410241f696ba06873023d5ce4ed/include/cute/atom/mma_traits_sm100.hpp#L3819

cc += @ANIKET-SHIVAM

[infinitron]

Sure, here's the PR (#3209).