diff --git a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.cpp b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.cpp
index 546148be62..c411aaa8f4 100644
--- a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.cpp
+++ b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.cpp
@@ -80,7 +80,8 @@ struct MmaLayoutTestKernel
                                     BlockM,
                                     BlockN,
                                     BlockK,
-                                    decltype(ck_tile::core::arch::get_compiler_target())>;
+                                    decltype(ck_tile::core::arch::get_compiler_target()),
+                                    mma::MmaOpFamily::DENSE>;
         using MmaOp     = typename Selector::SelectedOp;
         using MmaTraits = mma::MmaOpTraits<MmaOp>;
 
@@ -253,12 +254,30 @@ using MmaGfx90aCompilerTarget  = decltype(ck_tile::core::arch::make_amdgcn_gfx9_
 using MmaGfx1100CompilerTarget = decltype(ck_tile::core::arch::make_amdgcn_gfx11_target<
                                           ck_tile::core::arch::amdgcn_target_id::GFX1100>());
 
-using MmaGfx1201Selector = mma::
-    MmaDefaultSelector<ck::fp16_t, ck::fp16_t, ck::fp32_t, 16u, 16u, 16u, MmaGfx1201CompilerTarget>;
-using MmaGfx90aSelector = mma::
-    MmaDefaultSelector<ck::fp16_t, ck::fp16_t, ck::fp32_t, 16u, 16u, 16u, MmaGfx90aCompilerTarget>;
-using MmaGfx1100Selector = mma::
-    MmaDefaultSelector<ck::fp16_t, ck::fp16_t, ck::fp32_t, 16u, 16u, 16u, MmaGfx1100CompilerTarget>;
+using MmaGfx1201Selector = mma::MmaDefaultSelector<ck::fp16_t,
+                                                   ck::fp16_t,
+                                                   ck::fp32_t,
+                                                   16u,
+                                                   16u,
+                                                   16u,
+                                                   MmaGfx1201CompilerTarget,
+                                                   mma::MmaOpFamily::DENSE>;
+using MmaGfx90aSelector  = mma::MmaDefaultSelector<ck::fp16_t,
+                                                   ck::fp16_t,
+                                                   ck::fp32_t,
+                                                   16u,
+                                                   16u,
+                                                   16u,
+                                                   MmaGfx90aCompilerTarget,
+                                                   mma::MmaOpFamily::DENSE>;
+using MmaGfx1100Selector = mma::MmaDefaultSelector<ck::fp16_t,
+                                                   ck::fp16_t,
+                                                   ck::fp32_t,
+                                                   16u,
+                                                   16u,
+                                                   16u,
+                                                   MmaGfx1100CompilerTarget,
+                                                   mma::MmaOpFamily::DENSE>;
 
 // clang-format off
 using KernelTypes = ::testing::Types<
diff --git a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout_util.hpp b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout_util.hpp
index cb14e1676d..850435d256 100644
--- a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout_util.hpp
+++ b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout_util.hpp
@@ -68,7 +68,9 @@ struct RegisterMap
 /**
  * @brief RegisterMapTraits for GFX12 WMMA 16x16x16_F16_F16_F32_GFX12
  */
-template <typename CtrlFlags, typename CompilerTarget>
+template <typename CtrlFlags,
+          typename CompilerTarget,
+          ck_tile::core::arch::mma::MmaOpFamily OpFamily_>
 struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     ck_tile::fp16_t,
     ck_tile::fp16_t,
@@ -78,6 +80,7 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     16u,
     CtrlFlags,
     CompilerTarget,
+    OpFamily_,
     ck_tile::core::arch::enable_if_target_family_gfx12_t<CompilerTarget>>>
 {
     using MmaOp = ck_tile::core::arch::mma::amdgcn_mma<ck_tile::fp16_t,
@@ -87,7 +90,8 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
                                                        16u,
                                                        16u,
                                                        CtrlFlags,
-                                                       CompilerTarget>;
+                                                       CompilerTarget,
+                                                       OpFamily_>;
 
     using MmaTraits = ck_tile::core::arch::mma::MmaOpTraits<MmaOp>;
     static constexpr index_t WaveSize =
@@ -147,7 +151,9 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
 /**
  * @brief RegisterMapTraits for GFX9 MFMA 16x16x16_F16_F16_F32_GFX9
  */
-template <typename CtrlFlags, typename CompilerTarget>
+template <typename CtrlFlags,
+          typename CompilerTarget,
+          ck_tile::core::arch::mma::MmaOpFamily OpFamily_>
 struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     ck_tile::fp16_t,
     ck_tile::fp16_t,
@@ -157,6 +163,7 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     16u,
     CtrlFlags,
     CompilerTarget,
+    OpFamily_,
     ck_tile::core::arch::enable_if_target_family_gfx9_t<CompilerTarget>>>
 {
     using MmaOp = ck_tile::core::arch::mma::amdgcn_mma<ck_tile::fp16_t,
@@ -166,7 +173,8 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
                                                        16u,
                                                        16u,
                                                        CtrlFlags,
-                                                       CompilerTarget>;
+                                                       CompilerTarget,
+                                                       OpFamily_>;
 
     using MmaTraits = ck_tile::core::arch::mma::MmaOpTraits<MmaOp>;
     static constexpr index_t WaveSize =
@@ -212,7 +220,9 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
 /**
  * @brief RegisterMapTraits for GFX11 WMMA 16x16x16_F16_F16_F32_GFX11
  */
-template <typename CtrlFlags, typename CompilerTarget>
+template <typename CtrlFlags,
+          typename CompilerTarget,
+          ck_tile::core::arch::mma::MmaOpFamily OpFamily_>
 struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     ck_tile::fp16_t,
     ck_tile::fp16_t,
@@ -222,6 +232,7 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
     16u,
     CtrlFlags,
     CompilerTarget,
+    OpFamily_,
     ck_tile::core::arch::enable_if_target_family_gfx11_t<CompilerTarget>>>
 {
     using MmaOp = ck_tile::core::arch::mma::amdgcn_mma<ck_tile::fp16_t,
@@ -231,7 +242,8 @@ struct RegisterMapTraits<ck_tile::core::arch::mma::amdgcn_mma<
                                                        16u,
                                                        16u,
                                                        CtrlFlags,
-                                                       CompilerTarget>;
+                                                       CompilerTarget,
+                                                       OpFamily_>;
 
     using MmaTraits = ck_tile::core::arch::mma::MmaOpTraits<MmaOp>;
     static constexpr index_t WaveSize =