From 6dcebd21ef85abb2f49dea96fc04ed699f757c37 Mon Sep 17 00:00:00 2001
From: Karthik Reddy <karthiks@caltech.edu>
Date: Thu, 7 May 2026 12:18:21 -0700
Subject: [PATCH] fix: Add missing template specialization for F8F6F4 MMA Op
 (#3207)

---
 include/cute/arch/mma_sm100_umma.hpp   |  6 ++++--
 include/cute/atom/mma_traits_sm100.hpp | 20 ++++++++------------
 2 files changed, 12 insertions(+), 14 deletions(-)

diff --git a/include/cute/arch/mma_sm100_umma.hpp b/include/cute/arch/mma_sm100_umma.hpp
index d7dfb71596..aaaa423c03 100644
--- a/include/cute/arch/mma_sm100_umma.hpp
+++ b/include/cute/arch/mma_sm100_umma.hpp
@@ -1207,8 +1207,10 @@ struct SM100_MMA_S8_2x1SM_SS_SPARSE
   }
 };
 
-struct SM100_MMA_F8F6F4_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 SM100_MMA_F8F6F4_SS {
   using DRegisters = void;
   using ARegisters = uint64_t[1];
   using BRegisters = uint64_t[1];
diff --git a/include/cute/atom/mma_traits_sm100.hpp b/include/cute/atom/mma_traits_sm100.hpp
index 5b6af4218c..0f0fa3b885 100644
--- a/include/cute/atom/mma_traits_sm100.hpp
+++ b/include/cute/atom/mma_traits_sm100.hpp
@@ -3324,16 +3324,11 @@ struct MMA_Traits<SM100_MMA_S8_2x1SM_SS_SPARSE<a_type, b_type, c_type,
   }
 };
 
-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_SS, a_type, b_type, c_type,
-                  cute::C<M>, cute::C<N>,
-                  cute::integral_constant<UMMA::Major, a_major>,
-                  cute::integral_constant<UMMA::Major, b_major>,
-                  cute::integral_constant<UMMA::ScaleIn, a_neg>,
-                  cute::integral_constant<UMMA::ScaleIn, b_neg>>
-{
+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_SS<a_type, b_type, c_type, M, N, a_major,
+                                      b_major, a_neg, b_neg>> {
   using ValTypeD = c_type;
   using ValTypeA = a_type;
   using ValTypeB = b_type;
@@ -3390,11 +3385,12 @@ struct MMA_Traits<SM100_MMA_F8F6F4_SS, a_type, b_type, c_type,
     uint32_t tmem_c = raw_pointer_cast(D.data());
     uint64_t idesc = UMMA::make_runtime_instr_desc<>(traits.idesc_);
 
-    SM100_MMA_F8F6F4_SS::fma(desc_a, desc_b, tmem_c, uint32_t(traits.accumulate_), idesc);
+    SM100_MMA_F8F6F4_SS<a_type, b_type, c_type, M, N, a_major, b_major, a_neg,
+                        b_neg>::fma(desc_a, desc_b, tmem_c,
+                                    uint32_t(traits.accumulate_), idesc);
   }
 };
 
-
 template <class a_type, class b_type, class c_type, class sf_type,
           int M, int N, UMMA::Major a_major, UMMA::Major b_major,
           UMMA::ScaleIn a_neg, UMMA::ScaleIn b_neg>