From 604c56bc0e18097724d28c92a54deae7b1a3bf27 Mon Sep 17 00:00:00 2001
From: Yung-sheng Tu <112800063+yungshengtu@users.noreply.github.com>
Date: Mon, 29 Jun 2026 18:51:17 +0000
Subject: [PATCH] [rocm-libraries] ROCm/rocm-libraries#7850 (commit e8f2756)

=?UTF-8?q?style:=20[CK=20TILE]=20Unification=20Work=20?=
 =?UTF-8?q?=E2=80=93=20Unify=20format=20MFMA=20part=20(#7850)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Motivation

This PR unifies the parameter comments and simplifies the docs for
`amdgcn_mma` specialisations of `MfmaOp`.

## Technical Details

Except for the two things mentioned above, it also simplifies the sparse
traits, unifies the usages of `enable_if_target_id_t`, and cleans up the
files in
[include/ck_tile/core/arch/mma](https://github.com/ROCm/rocm-libraries/tree/users/yungshengtu/ck/unification/unify_format_mfma/projects/composablekernel/include/ck_tile/core/arch/mma).

**NOTE: The first commit is not in the scope of this PR.**

## Test Plan

Test has existed.

## Test Result

Test should pass.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

close #8907
---
 .../example_tile_distr_enc_calc.cpp           |   9 +-
 include/ck_tile/core/arch/mma/amdgcn_mma.hpp  |  18 +-
 include/ck_tile/core/arch/mma/mfma/mfma.hpp   |   8 +-
 .../ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp  | 670 ++++--------------
 .../core/arch/mma/mfma/mfma_selector.hpp      |   7 +-
 .../core/arch/mma/mfma/mfma_traits.hpp        |   1 +
 .../core/arch/mma/mfma/mfma_transforms.hpp    |   3 +
 include/ck_tile/core/arch/mma/mma.hpp         |  12 +-
 .../ck_tile/core/arch/mma/mma_pipeline.hpp    |  23 +-
 .../ck_tile/core/arch/mma/mma_selector.hpp    |  10 +-
 include/ck_tile/core/arch/mma/mma_traits.hpp  |   5 +-
 .../ck_tile/core/arch/mma/mma_transforms.hpp  |   4 +
 .../ck_tile/core/arch/mma/mma_wavewise.hpp    |  22 +-
 .../core/arch/mma/scale/mfma/scale_gfx9.hpp   | 377 +++++-----
 .../core/arch/mma/scale/mfma/selector.hpp     |  13 +-
 .../arch/mma/scale/scale_mma_pipeline.hpp     |  10 +-
 .../core/arch/mma/scale/scale_traits.hpp      |   3 +-
 .../core/arch/mma/sparse/mfma/selector.hpp    |   7 +-
 .../core/arch/mma/sparse/mfma/sparse_gfx9.hpp | 300 ++------
 .../ck_tile/core/arch/mma/sparse/sparse.hpp   |   8 +-
 .../arch/mma/sparse/sparse_mma_pipeline.hpp   |  17 +-
 .../arch/mma/sparse/sparse_transforms.hpp     |   9 +-
 .../core/arch/mma/sparse/wmma/selector.hpp    |   7 +-
 .../arch/mma/sparse/wmma/sparse_gfx12.hpp     |   5 +
 ..._distribution_encoding_register_mapper.hpp |  11 +-
 .../ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp |   5 +-
 .../ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp |   5 +-
 .../core/arch/mma/wmma/wmma_selector.hpp      |   6 +-
 .../core/arch/mma/wmma/wmma_transforms.hpp    |   6 +-
 .../ck_tile/core/tensor/tile_distribution.hpp |  29 +-
 include/ck_tile/ops/gemm.hpp                  |   2 +-
 .../mma/pipeline/test_amdgcn_scale_mma.cpp    |  65 +-
 .../mma/pipeline/test_amdgcn_sparse_mma.cpp   |  24 +-
 .../ck_tile/core/arch/mma/test_amdgcn_mma.cpp |   7 +-
 .../core/arch/mma/test_amdgcn_mma_layout.inc  |   6 +-
 35 files changed, 573 insertions(+), 1141 deletions(-)

diff --git a/example/ck_tile/51_tile_distr_enc_reg_map/example_tile_distr_enc_calc.cpp b/example/ck_tile/51_tile_distr_enc_reg_map/example_tile_distr_enc_calc.cpp
index 41b954a6de..cd98d91b44 100644
--- a/example/ck_tile/51_tile_distr_enc_reg_map/example_tile_distr_enc_calc.cpp
+++ b/example/ck_tile/51_tile_distr_enc_reg_map/example_tile_distr_enc_calc.cpp
@@ -1,15 +1,16 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 
-#include <cstdio>
-#include <type_traits>
-#include <tuple>
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/mma.hpp"
-#include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp"
 #include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_calculator.hpp"
+#include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp"
 #include "ck_tile/core/container/tuple.hpp"
 
+#include <cstdio>
+#include <tuple>
+#include <type_traits>
+
 using namespace ck_tile;
 using namespace ck_tile::core::arch;
 using namespace mma;
diff --git a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp
index 938fc1791d..1e26c2386d 100644
--- a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp
+++ b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp
@@ -10,23 +10,18 @@
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/numeric/ext_vector_base.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/numeric.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/ignore.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
 #include "ck_tile/ops/common/utils.hpp"
 
-#include <cstdint>
 #include <stdio.h>
 #include <type_traits>
 #if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 #include <concepts>
-#include <utility>
 #endif
 
-#if __clang_major__ >= 23
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
-#endif
 namespace ck_tile::core::arch::mma {
 
 /**---------------------------------------------------
@@ -314,7 +309,7 @@ struct amdgcn_mma : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 1u, 1u, 1u, 1u, 1, 1
 {
     // This is a default pass-through implementation that doesn't do anything practical.
     template <typename... Params>
-    CK_TILE_DEVICE static auto
+    CK_TILE_DEVICE static CVecType
     exec(AVecType const& regsA, BVecType const& regsB, CVecType const& regsC)
     {
         // Prints once across all thread blocks and threads.
@@ -332,9 +327,9 @@ struct amdgcn_mma : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 1u, 1u, 1u, 1u, 1, 1
 template <typename ADataType,
           typename BDataType,
           typename CDataType,
-          std::uint32_t FragM,
-          std::uint32_t FragN,
-          std::uint32_t FragK,
+          uint32_t FragM,
+          uint32_t FragN,
+          uint32_t FragK,
           typename CompilerTarget,
           MmaOpFamily OpFamily_,
           typename Enabler = void>
@@ -382,6 +377,3 @@ CK_TILE_HOST_DEVICE void print(amdgcn_mma<ADataType,
 }
 
 } // namespace ck_tile::core::arch::mma
-#if __clang_major__ >= 23
-#pragma clang diagnostic pop
-#endif
diff --git a/include/ck_tile/core/arch/mma/mfma/mfma.hpp b/include/ck_tile/core/arch/mma/mfma/mfma.hpp
index 55817b5ba8..832de70d22 100644
--- a/include/ck_tile/core/arch/mma/mfma/mfma.hpp
+++ b/include/ck_tile/core/arch/mma/mfma/mfma.hpp
@@ -4,7 +4,7 @@
 #pragma once
 
 // Include the architecture-specific MFMA implementations and traits
-#include "mfma_gfx9.hpp"
-#include "mfma_traits.hpp"
-#include "mfma_selector.hpp"
-#include "mfma_transforms.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_selector.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_traits.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_transforms.hpp"
diff --git a/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp b/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp
index ad4a055a06..48d91fbacf 100644
--- a/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp
+++ b/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp
@@ -29,17 +29,27 @@ namespace ck_tile::core::arch::mma {
 // one packed register for each input to be able to process smaller K values by padding.
 
 /**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
+ * @defgroup dense_mfma_gfx9 Dense MFMA for GFX9
+ * @brief Dense specializations of @ref amdgcn_mma for GFX9 family.
+ *
+ * Template parameters A/B/C denote input/output types,
+ * M/N/K are the fragment (MmaTile) sizes,
+ * and `enable_if_target_*` restricts the specialization to specific GPU targets.
+ *
+ * @tparam CompilerTarget Current compiler target.
+ *
+ * @sa amdgcn_mma_base for base template parameter documentation.
+ * @{
  */
+
 // TODO: c++20 template <amdgcn_target CompilerTarget>
 // TODO: c++20 requires
+
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 32u, 64u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 32u, 64u, 1u, 64u, 1, 1, 2, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -55,18 +65,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 32u, 64u, 1u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 32u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 64u, 32u, 1u, 64u, 1, 1, 1, 1, 2, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -82,18 +85,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 32u, 1u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 16u, 64u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 16u, 64u, 1u, 64u, 1, 1, 4, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -109,18 +105,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 16u, 64u, 1u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 16u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 64u, 16u, 1u, 64u, 1, 1, 1, 1, 4, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -136,18 +125,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 16u, 1u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams  |BPar |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 4u, 64u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 4u, 64u, 1u, 64u, 1, 1, 16, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -163,18 +145,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 4u, 64u, 1u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams |BPar  |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 4u, 1u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                   |WS  |AParams |BPar  |CPar |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 64u, 4u, 1u, 64u, 1, 1, 1, 1, 16, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -190,18 +165,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 64u, 4u, 1u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 32u, 32u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 32u, 32u, 2u, 64u, 1, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -217,18 +185,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 32u, 32u, 2u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp32_t, fp32_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 16u, 16u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp32_t, fp32_t, fp32_t, 16u, 16u, 4u, 64u, 1, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -244,18 +205,11 @@ struct amdgcn_mma<fp32_t, fp32_t, fp32_t, 16u, 16u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 32u, 64u, 4u, 64u, 4, 1, 2, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -270,18 +224,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 64u, 32u, 4u, 64u, 4, 1, 1, 1, 2, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -296,18 +243,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 16u, 64u, 4u, 64u, 4, 1, 4, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -322,18 +262,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 64u, 16u, 4u, 64u, 4, 1, 1, 1, 4, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -348,18 +281,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams  |BPar |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 4u, 64u, 4u, 64u, 4, 1, 16, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -374,18 +300,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams |BPar  |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                   |WS  |AParams |BPar  |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 64u, 4u, 4u, 64u, 4, 1, 1, 1, 16, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -400,18 +319,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 32u, 32u, 8u, 64u, 4, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -426,18 +338,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 16u, 16u, 16u, 64u, 4, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -452,18 +357,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes        |MNK          |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 64u, 4u, 64u, 4, 1, 2, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -479,18 +377,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes        |MNK          |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 64u, 32u, 4u, 64u, 4, 1, 1, 1, 2, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -506,18 +397,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK          |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 64u, 4u, 64u, 4, 1, 4, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -533,18 +417,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK          |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 64u, 16u, 4u, 64u, 4, 1, 1, 1, 4, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -560,18 +437,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize  |AParams  |BPar |CPar |
+//               |A B C DataTypes        |MNK         |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 4u, 64u, 4u, 64u, 4, 1, 16, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -587,18 +457,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX9
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize  |AParams |BPar  |CPar |
+//               |A B C DataTypes        |MNK         |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_family_gfx9_t<CompilerTarget>>
+//                                                    |WS  |AParams |BPar  |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 64u, 4u, 4u, 64u, 4, 1, 1, 1, 16, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -614,18 +477,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX908 and
- * GFX90a architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes        |MNK          |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 32u, 8u, 64u, 4, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -641,18 +497,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX908 and
- * GFX90a architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                      |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 16u, 16u, 64u, 4, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -668,18 +517,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 16u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 64u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 64u, 2u, 64u, 2, 1, 2, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -694,18 +536,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 64u, 2u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 32u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 32u, 2u, 64u, 2, 1, 1, 1, 2, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -720,18 +555,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 32u, 2u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 64u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 64u, 2u, 64u, 2, 1, 4, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -746,18 +574,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 64u, 2u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 16u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 16u, 2u, 64u, 2, 1, 1, 1, 4, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -772,18 +593,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 16u, 2u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams  |BPar |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 4u, 64u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 4u, 64u, 2u, 64u, 2, 1, 16, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -798,18 +612,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 4u, 64u, 2u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams |BPar  |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 4u, 2u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                   |WS  |AParams |BPar  |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 4u, 2u, 64u, 2, 1, 1, 1, 16, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -824,18 +631,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 4u, 2u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 32u, 4u, 64u, 2, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -850,18 +650,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX908 and GFX90a
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 8u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX908, amdgcn_target_id::GFX90A>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 16u, 8u, 64u, 2, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -876,18 +669,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 8u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 64u, 4u, 64u, 4, 1, 2, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -903,18 +689,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 64u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 32u, 4u, 64u, 4, 1, 1, 1, 2, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -930,18 +709,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 32u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 64u, 4u, 64u, 4, 1, 4, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -957,18 +729,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 64u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 16u, 4u, 64u, 4, 1, 1, 1, 4, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -984,18 +749,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 16u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams  |BPar |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 4u, 64u, 4u, 64u, 4, 1, 16, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1011,18 +769,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 4u, 64u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize  |AParams |BPar  |CPar |
+//               |A B C DataTypes       |MNK         |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar  |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 64u, 4u, 4u, 64u, 4, 1, 1, 1, 16, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1038,18 +789,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 64u, 4u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 32u, 8u, 64u, 4, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1065,18 +809,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 8u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 16u, 16u, 64u, 4, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1092,18 +829,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp64_t, fp64_t, fp64_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize   |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK          |
 struct amdgcn_mma<fp64_t, fp64_t, fp64_t, 16u, 16u, 4u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX90A, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp64_t, fp64_t, fp64_t, 16u, 16u, 4u, 64u, 1, 1, 1, 1, 1, 4, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1125,14 +855,6 @@ struct amdgcn_mma<fp64_t, fp64_t, fp64_t, 16u, 16u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp64_t, fp64_t, fp64_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
 //               |A B C DataTypes       |MNK         |
@@ -1157,14 +879,6 @@ struct amdgcn_mma<fp64_t, fp64_t, fp64_t, 4u, 16u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp64_t, fp64_t, fp64_t MMA operation on GFX90A, GFX942,
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
 //               |A B C DataTypes       |MNK         |
@@ -1189,18 +903,11 @@ struct amdgcn_mma<fp64_t, fp64_t, fp64_t, 16u, 4u, 4u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1216,18 +923,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 32u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1243,14 +943,6 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 16u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for tf32_t, tf32_t, float MMA operation on GFX942
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
 //               |A B C DataTypes       |MNK          |
@@ -1271,14 +963,6 @@ struct amdgcn_mma<tf32_t, tf32_t, fp32_t, 16u, 16u, 8u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for tf32_t, tf32_t, float MMA operation on GFX942
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
 //               |A B C DataTypes       |MNK          |
@@ -1299,18 +983,11 @@ struct amdgcn_mma<tf32_t, tf32_t, fp32_t, 32u, 32u, 4u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1326,18 +1003,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1353,18 +1023,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1380,18 +1043,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1407,18 +1063,11 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1434,18 +1083,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1461,18 +1103,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1488,18 +1123,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX942, GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes    | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1515,18 +1143,11 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1542,18 +1163,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1569,18 +1183,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1596,18 +1203,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar  |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1623,18 +1223,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize    |AParams  |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1650,18 +1243,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 64u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes       | MNK + WaveSize    |AParams  |BPar |CPar  |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::DENSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::DENSE>
 // clang-format on
 {
@@ -1677,4 +1263,6 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 32u, CompilerTarget, MmaOpF
     }
 };
 
+/** @} */ // dense_mfma_gfx9
+
 } // namespace ck_tile::core::arch::mma
diff --git a/include/ck_tile/core/arch/mma/mfma/mfma_selector.hpp b/include/ck_tile/core/arch/mma/mfma/mfma_selector.hpp
index 233ee77526..e760c000cd 100644
--- a/include/ck_tile/core/arch/mma/mfma/mfma_selector.hpp
+++ b/include/ck_tile/core/arch/mma/mfma/mfma_selector.hpp
@@ -4,10 +4,11 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
-#include "mfma_gfx9.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/mfma/mfma_traits.hpp b/include/ck_tile/core/arch/mma/mfma/mfma_traits.hpp
index 29340fece7..66d9d2fbac 100644
--- a/include/ck_tile/core/arch/mma/mfma/mfma_traits.hpp
+++ b/include/ck_tile/core/arch/mma/mfma/mfma_traits.hpp
@@ -4,6 +4,7 @@
 #pragma once
 
 #include "ck_tile/core/config.hpp"
+
 #include <type_traits>
 
 namespace ck_tile::core::arch::mma {
diff --git a/include/ck_tile/core/arch/mma/mfma/mfma_transforms.hpp b/include/ck_tile/core/arch/mma/mfma/mfma_transforms.hpp
index 9609ed5116..e97c46b81d 100644
--- a/include/ck_tile/core/arch/mma/mfma/mfma_transforms.hpp
+++ b/include/ck_tile/core/arch/mma/mfma/mfma_transforms.hpp
@@ -4,8 +4,11 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_transforms.hpp"
 
+#include <type_traits>
+
 namespace ck_tile::core::arch::mma {
 
 /**
diff --git a/include/ck_tile/core/arch/mma/mma.hpp b/include/ck_tile/core/arch/mma/mma.hpp
index ec38fe78e3..65c18f6e47 100644
--- a/include/ck_tile/core/arch/mma/mma.hpp
+++ b/include/ck_tile/core/arch/mma/mma.hpp
@@ -3,6 +3,12 @@
 
 #pragma once
 
-#include "wmma/wmma.hpp"
-#include "mfma/mfma.hpp"
-#include "sparse/sparse.hpp"
+#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
+#include "ck_tile/core/arch/mma/mma_selector.hpp"
+#include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/arch/mma/mma_transforms.hpp"
+#include "ck_tile/core/arch/mma/scale/scale.hpp"
+#include "ck_tile/core/arch/mma/sparse/sparse.hpp"
+#include "ck_tile/core/arch/mma/wmma/wmma.hpp"
diff --git a/include/ck_tile/core/arch/mma/mma_pipeline.hpp b/include/ck_tile/core/arch/mma/mma_pipeline.hpp
index b687c1adc9..0dc52b13f8 100644
--- a/include/ck_tile/core/arch/mma/mma_pipeline.hpp
+++ b/include/ck_tile/core/arch/mma/mma_pipeline.hpp
@@ -1,19 +1,16 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 #pragma once
+
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/numeric/vector_type.hpp"
-#include "ck_tile/ops/gemm/warp/warp_gemm_params.hpp"
+#include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/config.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 
-#include "amdgcn_mma.hpp"
-#include "mma_selector.hpp"
-#include "mma_traits.hpp"
-#include "mma_transforms.hpp"
-
-#if __clang_major__ >= 23
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+#if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
+#include <concepts>
 #endif
+
 namespace ck_tile::core::arch::mma {
 
 /**
@@ -144,8 +141,6 @@ struct MmaPipelineBase
 
 #if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
-#include <concepts>
-
 /**
  * @concept MmaPipelineI
  * @brief  Expresses the meta-data interface required for a CRTP MmaPipeline.
@@ -156,7 +151,3 @@ concept MmaPipelineInterface = std::derived_from<Derived, MmaPipelineBase<Derive
 #endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
 } // namespace ck_tile::core::arch::mma
-
-#if __clang_major__ >= 23
-#pragma clang diagnostic pop
-#endif
diff --git a/include/ck_tile/core/arch/mma/mma_selector.hpp b/include/ck_tile/core/arch/mma/mma_selector.hpp
index b8d6f31558..f79a456c73 100644
--- a/include/ck_tile/core/arch/mma/mma_selector.hpp
+++ b/include/ck_tile/core/arch/mma/mma_selector.hpp
@@ -2,9 +2,14 @@
 // SPDX-License-Identifier: MIT
 
 #pragma once
+
 #include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
@@ -141,8 +146,3 @@ struct MmaKSearchSelector<ADataType,
 };
 
 } // namespace ck_tile::core::arch::mma
-
-// Include the implementations
-#include "wmma/wmma_selector.hpp"
-#include "mfma/mfma_selector.hpp"
-#include "sparse/sparse_selector.hpp"
diff --git a/include/ck_tile/core/arch/mma/mma_traits.hpp b/include/ck_tile/core/arch/mma/mma_traits.hpp
index 863ad07bbb..53a76efc06 100644
--- a/include/ck_tile/core/arch/mma/mma_traits.hpp
+++ b/include/ck_tile/core/arch/mma/mma_traits.hpp
@@ -2,11 +2,10 @@
 // SPDX-License-Identifier: MIT
 #pragma once
 
-#include "amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/config.hpp"
-#include "mfma/mfma_traits.hpp"
-#include "wmma/wmma_traits.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 
 #include <stdio.h>
 #include <type_traits>
diff --git a/include/ck_tile/core/arch/mma/mma_transforms.hpp b/include/ck_tile/core/arch/mma/mma_transforms.hpp
index b870a343ce..d1748cdc62 100644
--- a/include/ck_tile/core/arch/mma/mma_transforms.hpp
+++ b/include/ck_tile/core/arch/mma/mma_transforms.hpp
@@ -1,7 +1,11 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 #pragma once
+
 #include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/config.hpp"
+
+#include <utility>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/mma_wavewise.hpp b/include/ck_tile/core/arch/mma/mma_wavewise.hpp
index 253457525b..6c6b9c320a 100644
--- a/include/ck_tile/core/arch/mma/mma_wavewise.hpp
+++ b/include/ck_tile/core/arch/mma/mma_wavewise.hpp
@@ -1,19 +1,23 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 #pragma once
+
 #include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_selector.hpp"
+#include "ck_tile/core/arch/mma/mfma/mfma_transforms.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
+#include "ck_tile/core/arch/mma/mma_pipeline.hpp"
+#include "ck_tile/core/arch/mma/mma_traits.hpp"
 #include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_calculator.hpp"
+#include "ck_tile/core/arch/mma/wmma/wmma_selector.hpp"
+#include "ck_tile/core/arch/mma/wmma/wmma_transforms.hpp"
+#include "ck_tile/core/container/thread_buffer.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/tensor/static_distributed_tensor.hpp"
-#include "ck_tile/core/numeric/vector_type.hpp"
+#include "ck_tile/core/tensor/tile_distribution.hpp"
+#include "ck_tile/core/utility/type_traits.hpp"
 
-#include "amdgcn_mma.hpp"
-#include "mma_pipeline.hpp"
-#include "mma_selector.hpp"
-#include "mma_transforms.hpp"
-
-#include "mfma/mfma.hpp"
-#include "wmma/wmma.hpp"
-#include <tuple>
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp b/include/ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp
index 32dd252d0d..2f77593b7d 100644
--- a/include/ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp
+++ b/include/ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp
@@ -10,6 +10,8 @@
 #include "ck_tile/core/arch/mma/scale/scale_traits.hpp"
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/numeric/float8.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/pk_f6.hpp"
 #include "ck_tile/core/numeric/pk_fp4.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/bit_cast.hpp"
@@ -18,29 +20,39 @@
 namespace ck_tile::core::arch::mma {
 
 /**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
+ * @defgroup scale_mfma_gfx9 Scale MFMA for GFX9
+ * @brief Scale specializations of @ref amdgcn_mma for GFX9 family.
  *
- * This specialization implements the Scale MFMA instruction for fp8_t A and B
- * matrices with fp32_t accumulator, with 16x16x128 block sizes.
+ * Template parameters A/B/C denote input/output types,
+ * M/N/K are the fragment (MmaTile) sizes,
+ * and `enable_if_target_*` restricts the specialization to specific GPU targets.
  *
- * @tparam CompilerTarget Current compiler target
+ * @tparam CompilerTarget Current compiler target.
+ *
+ * @sa amdgcn_mma_base for base template parameter documentation.
+ * @{
  */
+
 // TODO: c++20 template <amdgcn_target CompilerTarget>
 // TODO: c++20 requires
+
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 2, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes    | MNK + WaveSize     |AParams  |BPar |CPar |
 // clang-format on
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
@@ -56,30 +68,23 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for bf8_t A and B
- * matrices with fp32_t accumulator, with 16x16x128 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 2, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes    | MNK + WaveSize     |AParams  |BPar |CPar |
 // clang-format on
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
@@ -95,71 +100,22 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_fp4_t A and B
- * matrices with fp32_t accumulator, with 16x16x128 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
-template <typename CompilerTarget>
-// clang-format off
-struct amdgcn_mma<pk_fp4_t, pk_fp4_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
-: amdgcn_mma_base<pk_fp4_t, pk_fp4_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes          | MNK + WaveSize     |AParams  |BPar |CPar |
-// clang-format on
-{
-    static constexpr const char* instruction_name =
-        "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
-
-    template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
-    {
-        using P         = WarpGemmParamsParser<Params...>;
-        int32x4_t arg_a = bit_cast<int32x4_t>(aVec);
-        int32x4_t arg_b = bit_cast<int32x4_t>(bVec);
-
-        return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
-            int32x8_t{arg_a[0], arg_a[1], arg_a[2], arg_a[3], 0, 0, 0, 0},
-            int32x8_t{arg_b[0], arg_b[1], arg_b[2], arg_b[3], 0, 0, 0, 0},
-            cVec,
-            scale::detail::ScaleDataTypeToFlag_v<pk_fp4_t>,
-            scale::detail::ScaleDataTypeToFlag_v<pk_fp4_t>,
-            P::op_sel_a,
-            scale_A,
-            P::op_sel_b,
-            scale_B)};
-    }
-};
-
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_fp6x16_t A and B
- * matrices with fp32_t accumulator, with 16x16x128 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes                 |MNK            |
 struct amdgcn_mma<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                                |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes                | MNK + WaveSize     |AParams  |BPar |CPar |
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
@@ -176,29 +132,22 @@ struct amdgcn_mma<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 16u, 16u, 128u, CompilerTarg
 };
 // clang-format on
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_bf6x16_t A and B
- * matrices with fp32_t accumulator, with 16x16x128 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes                 |MNK            |
 struct amdgcn_mma<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                                |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes                | MNK + WaveSize     |AParams  |BPar |CPar |
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
-    
+
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
@@ -215,30 +164,58 @@ struct amdgcn_mma<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 16u, 16u, 128u, CompilerTarg
 };
 // clang-format on
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for fp8_t A and B
- * matrices with fp32_t accumulator, with 32x32x64 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes           |MNK            |
+struct amdgcn_mma<pk_fp4_t, pk_fp4_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                          |WS  |AParams  |BPar |CPar |
+: amdgcn_mma_base<pk_fp4_t, pk_fp4_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SCALE>
+// clang-format on
+{
+    static constexpr const char* instruction_name =
+        "__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4";
+
+    template <typename... Params>
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
+    {
+        using P         = WarpGemmParamsParser<Params...>;
+        int32x4_t arg_a = bit_cast<int32x4_t>(aVec);
+        int32x4_t arg_b = bit_cast<int32x4_t>(bVec);
+
+        return {__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(
+            int32x8_t{arg_a[0], arg_a[1], arg_a[2], arg_a[3], 0, 0, 0, 0},
+            int32x8_t{arg_b[0], arg_b[1], arg_b[2], arg_b[3], 0, 0, 0, 0},
+            cVec,
+            scale::detail::ScaleDataTypeToFlag_v<pk_fp4_t>,
+            scale::detail::ScaleDataTypeToFlag_v<pk_fp4_t>,
+            P::op_sel_a,
+            scale_A,
+            P::op_sel_b,
+            scale_B)};
+    }
+};
+
+template <typename CompilerTarget>
+// clang-format off
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 2, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes    | MNK + WaveSize    |AParams  |BPar |CPar  |
 // clang-format on
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
@@ -254,30 +231,23 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for bf8_t A and B
- * matrices with fp32_t accumulator, with 32x32x64 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 2, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes    | MNK + WaveSize    |AParams  |BPar |CPar  |
 // clang-format on
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
@@ -293,30 +263,87 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_fp4_t A and B
- * matrices with fp32_t accumulator, with 32x32x64 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
 template <typename CompilerTarget>
 // clang-format off
-struct amdgcn_mma<pk_fp4_t, pk_fp4_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>> 
+//               |A B C DataTypes                 |MNK           |
+struct amdgcn_mma<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                               |WS  |AParams  |BPar |CPar  |
+: amdgcn_mma_base<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
+{
+    static constexpr const char* instruction_name =
+        "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
+
+    template <typename... Params>
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
+    {
+        using P = WarpGemmParamsParser<Params...>;
+        return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
+            int32x8_t{aVec.data[0], aVec.data[1], aVec.data[2], aVec.data[3], aVec.data[4], aVec.data[5], 0, 0},
+            int32x8_t{bVec.data[0], bVec.data[1], bVec.data[2], bVec.data[3], bVec.data[4], bVec.data[5], 0, 0},
+            cVec,
+            scale::detail::ScaleDataTypeToFlag_v<pk_fp6x16_t>,
+            scale::detail::ScaleDataTypeToFlag_v<pk_fp6x16_t>,
+            P::op_sel_a,
+            scale_A,
+            P::op_sel_b,
+            scale_B)};
+    }
+};
+// clang-format on
+
+template <typename CompilerTarget>
+// clang-format off
+//               |A B C DataTypes                 |MNK           |
+struct amdgcn_mma<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                               |WS  |AParams  |BPar |CPar  |
+: amdgcn_mma_base<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
+{
+    static constexpr const char* instruction_name =
+        "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
+
+    template <typename... Params>
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
+    {
+        using P = WarpGemmParamsParser<Params...>;
+        return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
+            int32x8_t{aVec.data[0], aVec.data[1], aVec.data[2], aVec.data[3], aVec.data[4], aVec.data[5], 0, 0},
+            int32x8_t{bVec.data[0], bVec.data[1], bVec.data[2], bVec.data[3], bVec.data[4], bVec.data[5], 0, 0},
+            cVec,
+            scale::detail::ScaleDataTypeToFlag_v<pk_bf6x16_t>,
+            scale::detail::ScaleDataTypeToFlag_v<pk_bf6x16_t>,
+            P::op_sel_a,
+            scale_A,
+            P::op_sel_b,
+            scale_B)};
+    }
+};
+// clang-format on
+
+template <typename CompilerTarget>
+// clang-format off
+//               |A B C DataTypes           |MNK           |
+struct amdgcn_mma<pk_fp4_t, pk_fp4_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                         |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<pk_fp4_t, pk_fp4_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes          | MNK + WaveSize    |AParams  |BPar |CPar  |
 // clang-format on
 {
     static constexpr const char* instruction_name =
         "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
 
     template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
+    CK_TILE_DEVICE static CVecType exec(AVecType const& aVec,
+                                        BVecType const& bVec,
+                                        CVecType const& cVec,
+                                        int32_t scale_A,
+                                        int32_t scale_B)
     {
         using P         = WarpGemmParamsParser<Params...>;
         int32x4_t arg_a = bit_cast<int32x4_t>(aVec);
@@ -335,82 +362,6 @@ struct amdgcn_mma<pk_fp4_t, pk_fp4_t, fp32_t, 32u, 32u, 64u, CompilerTarget, Mma
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_fp6x16_t A and B
- * matrices with fp32_t accumulator, with 32x32x64 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
-template <typename CompilerTarget>
-// clang-format off
-struct amdgcn_mma<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
-: amdgcn_mma_base<pk_fp6x16_t, pk_fp6x16_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes                | MNK + WaveSize    |AParams  |BPar |CPar  |
-{
-    static constexpr const char* instruction_name =
-        "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
-
-    template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
-    {
-        using P = WarpGemmParamsParser<Params...>;
-        return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
-            int32x8_t{aVec.data[0], aVec.data[1], aVec.data[2], aVec.data[3], aVec.data[4], aVec.data[5], 0, 0},
-            int32x8_t{bVec.data[0], bVec.data[1], bVec.data[2], bVec.data[3], bVec.data[4], bVec.data[5], 0, 0},
-            cVec,
-            scale::detail::ScaleDataTypeToFlag_v<pk_fp6x16_t>,
-            scale::detail::ScaleDataTypeToFlag_v<pk_fp6x16_t>,
-            P::op_sel_a,
-            scale_A,
-            P::op_sel_b,
-            scale_B)};
-    }
-};
-// clang-format on
-
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Scale MFMA on GFX950 targets
- *
- * This specialization implements the Scale MFMA instruction for pk_bf6x16_t A and B
- * matrices with fp32_t accumulator, with 32x32x64 block sizes.
- *
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
-// TODO: c++20 requires
-template <typename CompilerTarget>
-// clang-format off
-struct amdgcn_mma<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SCALE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
-: amdgcn_mma_base<pk_bf6x16_t, pk_bf6x16_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SCALE>
-//               | A B C DataTypes                | MNK + WaveSize    |AParams  |BPar |CPar  |
-{
-    static constexpr const char* instruction_name =
-        "__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4";
-
-    template <typename... Params>
-    CK_TILE_DEVICE static CVecType
-    exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int scale_A, int scale_B)
-    {
-        using P = WarpGemmParamsParser<Params...>;
-        return {__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(
-            int32x8_t{aVec.data[0], aVec.data[1], aVec.data[2], aVec.data[3], aVec.data[4], aVec.data[5], 0, 0},
-            int32x8_t{bVec.data[0], bVec.data[1], bVec.data[2], bVec.data[3], bVec.data[4], bVec.data[5], 0, 0},
-            cVec,
-            scale::detail::ScaleDataTypeToFlag_v<pk_bf6x16_t>,
-            scale::detail::ScaleDataTypeToFlag_v<pk_bf6x16_t>,
-            P::op_sel_a,
-            scale_A,
-            P::op_sel_b,
-            scale_B)};
-    }
-};
-// clang-format on
+/** @} */ // scale_mfma_gfx9
 
 } // namespace ck_tile::core::arch::mma
diff --git a/include/ck_tile/core/arch/mma/scale/mfma/selector.hpp b/include/ck_tile/core/arch/mma/scale/mfma/selector.hpp
index bfdb78de09..86d81451d6 100644
--- a/include/ck_tile/core/arch/mma/scale/mfma/selector.hpp
+++ b/include/ck_tile/core/arch/mma/scale/mfma/selector.hpp
@@ -4,10 +4,13 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/utility/type_traits.hpp"
 #include "ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/utility/type_traits.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 // TODO: We do not allow M / N composition for now, since it is not used in current CK Tile.
@@ -30,9 +33,9 @@ namespace ck_tile::core::arch::mma {
 template <typename ADataType,
           typename BDataType,
           typename CDataType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK,
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK,
           typename CompilerTarget,
           MmaOpFamily OpFamily>
 // TODO: c++20 amdgcn_target_arch_id CompilerTarget>
diff --git a/include/ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp b/include/ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp
index bed2a6506b..8dc0012bd1 100644
--- a/include/ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp
+++ b/include/ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp
@@ -5,15 +5,19 @@
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_pipeline.hpp"
+#include "ck_tile/core/arch/mma/mma_traits.hpp"
 #include "ck_tile/core/arch/mma/mma_wavewise.hpp"
 #include "ck_tile/core/arch/mma/scale/scale_selector.hpp"
 #include "ck_tile/core/arch/mma/scale/scale_transforms.hpp"
+#include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_calculator.hpp"
 #include "ck_tile/core/config.hpp"
+#include "ck_tile/core/container/thread_buffer.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/tensor/static_distributed_tensor.hpp"
+#include "ck_tile/core/tensor/tile_distribution.hpp"
+#include "ck_tile/core/utility/type_traits.hpp"
 
-#include <cstdint>
-#include <tuple>
 #include <type_traits>
-#include <utility>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/scale/scale_traits.hpp b/include/ck_tile/core/arch/mma/scale/scale_traits.hpp
index ea879feebc..24f7cf967c 100644
--- a/include/ck_tile/core/arch/mma/scale/scale_traits.hpp
+++ b/include/ck_tile/core/arch/mma/scale/scale_traits.hpp
@@ -4,8 +4,9 @@
 #pragma once
 
 #include "ck_tile/core/numeric/float8.hpp"
-#include "ck_tile/core/numeric/pk_fp4.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/pk_f6.hpp"
+#include "ck_tile/core/numeric/pk_fp4.hpp"
 
 namespace ck_tile::core::arch::mma {
 namespace scale::detail {
diff --git a/include/ck_tile/core/arch/mma/sparse/mfma/selector.hpp b/include/ck_tile/core/arch/mma/sparse/mfma/selector.hpp
index 9ea90daab8..1d869fd1c6 100644
--- a/include/ck_tile/core/arch/mma/sparse/mfma/selector.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/mfma/selector.hpp
@@ -4,9 +4,12 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/utility/type_traits.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp b/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp
index bdb50cc232..1119ab8178 100644
--- a/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp
@@ -8,29 +8,38 @@
 #include "ck_tile/core/arch/mma/mfma/mfma_traits.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/config.hpp"
+#include "ck_tile/core/numeric/bfloat16.hpp"
+#include "ck_tile/core/numeric/float8.hpp"
 #include "ck_tile/core/numeric/half.hpp"
+#include "ck_tile/core/numeric/int8.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
-#include "ck_tile/core/utility/type_traits.hpp"
 #include "ck_tile/ops/gemm/warp/warp_gemm_params.hpp"
 
-#include <type_traits>
-
 namespace ck_tile::core::arch::mma {
 
 /**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for Sparse MFMA (SMFMA) on GFX942, GFX950 targets
+ * @defgroup sparse_mfma_gfx9 Sparse MFMA for GFX9
+ * @brief Sparse specializations of @ref amdgcn_mma for GFX9 family.
  *
- * This specialization implements the SMFMA instruction for fp16_t A and B
- * matrices with structured sparsity, fp32_t accumulator, with 16x16x32 fragment sizes.
+ * Template parameters A/B/C denote input/output types,
+ * M/N/K are the fragment (MmaTile) sizes,
+ * and `enable_if_target_*` restricts the specialization to specific GPU targets.
  *
- * @tparam CompilerTarget Current compiler target
+ * @tparam CompilerTarget Current compiler target.
+ *
+ * @sa amdgcn_mma_base for base template parameter documentation.
+ * @{
  */
+
 // TODO: c++20 template <amdgcn_target CompilerTarget>
+// TODO: c++20 requires
+
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -51,17 +60,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFa
     };
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -76,17 +79,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, 64u, 8, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -101,17 +98,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 32u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, 64u, 8, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -126,17 +117,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 16u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX942 and
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -151,17 +136,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 64u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX942 and
- * GFX950 architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -176,17 +155,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 32u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -202,17 +175,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -228,17 +195,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -254,17 +215,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 1, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -280,17 +235,11 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -306,17 +255,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -332,17 +275,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -358,17 +295,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX942 and GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX942, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 1, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -384,17 +315,11 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -409,17 +334,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp16_t, fp16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp16_t, fp16_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -434,17 +353,11 @@ struct amdgcn_mma<fp16_t, fp16_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 16u, 16u, 64u, 64u, 16, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -459,17 +372,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 16u, 16u, 64u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf16_t, bf16_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes       |MNK           |
 struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                     |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf16_t, bf16_t, fp32_t, 32u, 32u, 32u, 64u, 16, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -484,17 +391,11 @@ struct amdgcn_mma<bf16_t, bf16_t, fp32_t, 32u, 32u, 32u, CompilerTarget, MmaOpFa
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK            |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                       |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -509,17 +410,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 16u, 16u, 128u, CompilerTarget, MmaOp
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for int8_t, int8_t, int32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes        |MNK           |
 struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                      |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<int8_t, int8_t, int32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -534,17 +429,11 @@ struct amdgcn_mma<int8_t, int8_t, int32_t, 32u, 32u, 64u, CompilerTarget, MmaOpF
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -560,17 +449,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -586,17 +469,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -612,17 +489,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK            |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                    |WS  |AParams  |BPar |CPar |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, 64u, 32, 1, 1, 2, 1, 4, 1, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -638,17 +509,11 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 16u, 16u, 128u, CompilerTarget, MmaOpFam
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, bf8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -664,17 +529,11 @@ struct amdgcn_mma<bf8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for bf8_t, fp8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<bf8_t, fp8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -690,17 +549,11 @@ struct amdgcn_mma<bf8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, bf8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, bf8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -716,17 +569,11 @@ struct amdgcn_mma<fp8_t, bf8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
     }
 };
 
-/**
- * @struct amdgcn_mma
- * @brief Specialization of amdgcn_mma for fp8_t, fp8_t, fp32_t MMA operation on GFX950
- * architecture.
- * @tparam CompilerTarget Current compiler target
- */
-// TODO: c++20 template <amdgcn_target CompilerTarget>
 template <typename CompilerTarget>
 // clang-format off
-//               | A B C DataTypes      | MNK + WaveSize    |AParams |BPar |CPar |
+//               |A B C DataTypes     |MNK           |
 struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFamily::SPARSE, enable_if_target_id_t<CompilerTarget, amdgcn_target_id::GFX950>>
+//                                                   |WS  |AParams  |BPar |CPar  |
 : amdgcn_mma_base<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, 64u, 32, 1, 1, 2, 1, 16, 4, MfmaOp, MmaOpFamily::SPARSE>
 // clang-format on
 {
@@ -741,4 +588,7 @@ struct amdgcn_mma<fp8_t, fp8_t, fp32_t, 32u, 32u, 64u, CompilerTarget, MmaOpFami
             __builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)};
     }
 };
+
+/** @} */ // sparse_mfma_gfx9
+
 } // namespace ck_tile::core::arch::mma
diff --git a/include/ck_tile/core/arch/mma/sparse/sparse.hpp b/include/ck_tile/core/arch/mma/sparse/sparse.hpp
index 303687b1f8..fac1178ac0 100644
--- a/include/ck_tile/core/arch/mma/sparse/sparse.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/sparse.hpp
@@ -3,12 +3,8 @@
 
 #pragma once
 
-namespace ck_tile::core::arch::mma {
-
-} // namespace ck_tile::core::arch::mma
-
 // Include sparse MFMA traits and architecture-specific implementations
 #include "ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp"
-#include "ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp"
-#include "ck_tile/core/arch/mma/sparse/sparse_transforms.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_selector.hpp"
+#include "ck_tile/core/arch/mma/sparse/sparse_transforms.hpp"
+#include "ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp"
diff --git a/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp b/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp
index b39cdae770..7cf8435e1e 100644
--- a/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp
@@ -2,13 +2,28 @@
 // SPDX-License-Identifier: MIT
 #pragma once
 
+#include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_pipeline.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
 #include "ck_tile/core/arch/mma/mma_traits.hpp"
 #include "ck_tile/core/arch/mma/mma_wavewise.hpp"
+#include "ck_tile/core/arch/mma/sparse/sparse_selector.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_transforms.hpp"
+#include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_calculator.hpp"
+#include "ck_tile/core/config.hpp"
+#include "ck_tile/core/container/thread_buffer.hpp"
+#include "ck_tile/core/numeric/ext_vector_base.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
-#include <cstdint>
+#include "ck_tile/core/tensor/static_distributed_tensor.hpp"
+#include "ck_tile/core/tensor/tile_distribution.hpp"
+#include "ck_tile/core/utility/bit_cast.hpp"
+#include "ck_tile/core/utility/type_traits.hpp"
+
+#include <tuple>
+#include <type_traits>
+#include <utility>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp b/include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp
index c4390abb2f..a0d55ce468 100644
--- a/include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp
@@ -3,12 +3,17 @@
 
 #pragma once
 
-#include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_transforms.hpp"
+#include "ck_tile/core/config.hpp"
+#include "ck_tile/core/numeric/ext_vector_base.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
-#include <cstdint>
+#include "ck_tile/core/utility/bit_cast.hpp"
+#include "ck_tile/core/utility/functional.hpp"
+
+#include <tuple>
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/sparse/wmma/selector.hpp b/include/ck_tile/core/arch/mma/sparse/wmma/selector.hpp
index 87d46bd98f..ae6c456323 100644
--- a/include/ck_tile/core/arch/mma/sparse/wmma/selector.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/wmma/selector.hpp
@@ -4,10 +4,11 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
-#include "ck_tile/core/arch/mma/wmma/wmma_traits.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp b/include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp
index 1418bc909c..a5a46bf278 100644
--- a/include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp
+++ b/include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp
@@ -8,9 +8,14 @@
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/wmma/wmma_traits.hpp"
 #include "ck_tile/core/config.hpp"
+#include "ck_tile/core/numeric/bfloat16.hpp"
+#include "ck_tile/core/numeric/float8.hpp"
 #include "ck_tile/core/numeric/half.hpp"
+#include "ck_tile/core/numeric/int8.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/pk_int4.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
+#include "ck_tile/core/utility/bit_cast.hpp"
 #include "ck_tile/ops/gemm/warp/warp_gemm_params.hpp"
 
 namespace ck_tile::core::arch::mma {
diff --git a/include/ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp b/include/ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp
index eb2deef987..2f0147d5fe 100644
--- a/include/ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp
+++ b/include/ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp
@@ -21,10 +21,17 @@
 
 #pragma once
 
-#include <stdio.h>
-#include "ck_tile/core/tensor/tensor_descriptor.hpp"
+#include "ck_tile/core/config.hpp"
+#include "ck_tile/core/container/array.hpp"
+#include "ck_tile/core/container/container_helper.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/integral_constant.hpp"
+#include "ck_tile/core/numeric/math.hpp"
 #include "ck_tile/core/tensor/tile_distribution.hpp"
 
+#include <array>
+#include <stdio.h>
+
 namespace ck_tile::core::arch::mma {
 
 // Utility to calculate register mappings from a Tile Distribution Encoding.
diff --git a/include/ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp b/include/ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp
index 2c4767fde1..1410b3d090 100644
--- a/include/ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp
+++ b/include/ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp
@@ -3,17 +3,16 @@
 
 #pragma once
 
-#include "wmma_traits.hpp"
-
-#include "ck_tile/core/config.hpp"
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
+#include "ck_tile/core/arch/mma/wmma/wmma_traits.hpp"
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/numeric/bfloat16.hpp"
 #include "ck_tile/core/numeric/half.hpp"
 #include "ck_tile/core/numeric/int8.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/pk_int4.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/bit_cast.hpp"
 #include "ck_tile/ops/gemm/warp/warp_gemm_params.hpp"
diff --git a/include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp b/include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp
index 9146d6e250..a80143c0e1 100644
--- a/include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp
+++ b/include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp
@@ -3,18 +3,17 @@
 
 #pragma once
 
-#include "wmma_traits.hpp"
-
-#include "ck_tile/core/config.hpp"
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
+#include "ck_tile/core/arch/mma/wmma/wmma_traits.hpp"
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/numeric/bfloat16.hpp"
 #include "ck_tile/core/numeric/float8.hpp"
 #include "ck_tile/core/numeric/half.hpp"
 #include "ck_tile/core/numeric/int8.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/pk_int4.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/bit_cast.hpp"
 #include "ck_tile/ops/gemm/warp/warp_gemm_params.hpp"
diff --git a/include/ck_tile/core/arch/mma/wmma/wmma_selector.hpp b/include/ck_tile/core/arch/mma/wmma/wmma_selector.hpp
index 645fdc81e6..4027a7c11c 100644
--- a/include/ck_tile/core/arch/mma/wmma/wmma_selector.hpp
+++ b/include/ck_tile/core/arch/mma/wmma/wmma_selector.hpp
@@ -4,9 +4,11 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+
+#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp b/include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp
index 7a04fb4633..5f90384910 100644
--- a/include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp
+++ b/include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp
@@ -3,9 +3,13 @@
 
 #pragma once
 
-#include "ck_tile/core/config.hpp"
 #include "ck_tile/core/arch/arch.hpp"
+#include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_transforms.hpp"
+#include "ck_tile/core/config.hpp"
+
+#include <type_traits>
+#include <utility>
 
 namespace ck_tile::core::arch::mma {
 
diff --git a/include/ck_tile/core/tensor/tile_distribution.hpp b/include/ck_tile/core/tensor/tile_distribution.hpp
index 606b346e19..9374af55c3 100644
--- a/include/ck_tile/core/tensor/tile_distribution.hpp
+++ b/include/ck_tile/core/tensor/tile_distribution.hpp
@@ -3,22 +3,29 @@
 
 #pragma once
 
+#include "ck_tile/core/algorithm/coordinate_transform.hpp"
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/container/array.hpp"
-#include "ck_tile/core/container/sequence.hpp"
-#include "ck_tile/core/container/tuple.hpp"
 #include "ck_tile/core/container/container_helper.hpp"
 #include "ck_tile/core/container/meta_data_buffer.hpp"
+#include "ck_tile/core/container/multi_index.hpp"
+#include "ck_tile/core/container/sequence.hpp"
+#include "ck_tile/core/container/tuple.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
+#include "ck_tile/core/numeric/integral_constant.hpp"
+#include "ck_tile/core/numeric/math.hpp"
 #include "ck_tile/core/tensor/tensor_adaptor.hpp"
+#include "ck_tile/core/tensor/tensor_adaptor_coordinate.hpp"
+#include "ck_tile/core/tensor/tensor_descriptor.hpp"
 #include "ck_tile/core/tensor/tile_distribution_encoding.hpp"
 #include "ck_tile/core/utility/functional.hpp"
+#include "ck_tile/core/utility/to_sequence.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
 
-#if __clang_major__ >= 23
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
-#endif
+#include <stdio.h>
+#include <type_traits>
+
 namespace ck_tile {
 
 template <typename Distribution>
@@ -136,11 +143,16 @@ struct tile_distribution
     }
 
     CK_TILE_HOST_DEVICE constexpr const auto& get_ps_ys_to_xs_adaptor() const
+        [[clang::lifetimebound]]
     {
         return ps_ys_to_xs_;
     }
 
-    CK_TILE_HOST_DEVICE constexpr const auto& get_ys_to_d_descriptor() const { return ys_to_d_; }
+    CK_TILE_HOST_DEVICE constexpr const auto& get_ys_to_d_descriptor() const
+        [[clang::lifetimebound]]
+    {
+        return ys_to_d_;
+    }
 
     CK_TILE_HOST_DEVICE static constexpr auto get_static_tile_distribution_encoding()
     {
@@ -708,6 +720,3 @@ CK_TILE_HOST_DEVICE void print(const tile_distribution<PsYs2XsAdaptor_,
 }
 
 } // namespace ck_tile
-#if __clang_major__ >= 23
-#pragma clang diagnostic pop
-#endif
diff --git a/include/ck_tile/ops/gemm.hpp b/include/ck_tile/ops/gemm.hpp
index fe825e2605..6779da556e 100644
--- a/include/ck_tile/ops/gemm.hpp
+++ b/include/ck_tile/ops/gemm.hpp
@@ -63,8 +63,8 @@
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v5_default_policy.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v6_default_policy.hpp"
-#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_default_policy.hpp"
+#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_mem.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_wavelet.hpp"
diff --git a/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_scale_mma.cpp b/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_scale_mma.cpp
index 864997d5f1..7692ee6a45 100644
--- a/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_scale_mma.cpp
+++ b/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_scale_mma.cpp
@@ -4,19 +4,16 @@
 #include "pipeline_tests_helper.hpp"
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma.hpp"
-#include "ck_tile/core/arch/mma/mma_op_family.hpp"
-#include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
 #include "ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp"
 #include "ck_tile/core/numeric/float8.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/pk_fp4.hpp"
+#include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/functional.hpp"
 
 #include <gtest/gtest.h>
 
-#include <cstdint>
 #include <iostream>
 #include <type_traits>
 
@@ -29,9 +26,9 @@ using CompilerTargetGfx950 = decltype(make_amdgcn_gfx9_target<amdgcn_target_id::
 template <typename AType,
           typename BType,
           typename CType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK>
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK>
 void ScaleMfmaGfx950Specialization_impl()
 {
     using TestScaleMma = amdgcn_mma<AType,
@@ -82,9 +79,9 @@ TEST(ScaleMMATrait, ScaleMfmaGfx950Specialization)
 template <typename AType,
           typename BType,
           typename CType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK>
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK>
 void TestConceptRequirements_impl()
 {
     using TestScaleMma = amdgcn_mma<AType,
@@ -122,9 +119,9 @@ void ScaleSelector_impl()
             using Selected = typename MmaDefaultSelector<AType,
                                                          BType,
                                                          CType,
-                                                         static_cast<std::uint32_t>(i),
-                                                         static_cast<std::uint32_t>(i),
-                                                         static_cast<std::uint32_t>(k_factor * i),
+                                                         static_cast<uint32_t>(i),
+                                                         static_cast<uint32_t>(i),
+                                                         static_cast<uint32_t>(k_factor * i),
                                                          CompilerTargetGfx950,
                                                          MmaOpFamily::SCALE>::SelectedOp;
             static constexpr bool isValid =
@@ -158,9 +155,9 @@ template <typename AType,
           typename CType,
           typename ScaleAType,
           typename ScaleBType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK>
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK>
 struct ScalePipelineKernel
 {
     static constexpr int kBlockSize = mma_pipeline_test::getCMakeWaveSize();
@@ -221,9 +218,9 @@ struct ScalePipelineKernel
 template <typename AType,
           typename BType,
           typename CType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK>
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK>
 struct ScalePipelineFactory
 {
     template <typename Target>
@@ -249,13 +246,13 @@ struct ScalePipelineFactory
 template <typename AType,
           typename BType,
           typename CType,
-          std::uint32_t WaveTileM,
-          std::uint32_t WaveTileN,
-          std::uint32_t WaveTileK>
+          uint32_t WaveTileM,
+          uint32_t WaveTileN,
+          uint32_t WaveTileK>
 void MmaSelector_Scale_Real_impl()
 {
-    using ScaleAType = std::int32_t;
-    using ScaleBType = std::int32_t;
+    using ScaleAType = int32_t;
+    using ScaleBType = int32_t;
 
     const auto should_skip = [](amdgcn_target_id currentArchId) {
         bool isSupportedMfma = (currentArchId == amdgcn_target_id::GFX950);
@@ -339,9 +336,9 @@ TEST(ScaleMMATrait, MmaSelector_Scale_BF8_BF8_F32_32x32x64_Real)
 //           typename CType,
 //           typename ScaleAType,
 //           typename ScaleBType,
-//           std::uint32_t WaveTileM,
-//           std::uint32_t WaveTileN,
-//           std::uint32_t WaveTileK,
+//           uint32_t WaveTileM,
+//           uint32_t WaveTileN,
+//           uint32_t WaveTileK,
 //           MmaAccumPolicy AccumPolicy>
 // struct ScaleWaveWisePipelineKernel
 // {
@@ -402,9 +399,9 @@ TEST(ScaleMMATrait, MmaSelector_Scale_BF8_BF8_F32_32x32x64_Real)
 // template <typename AType,
 //           typename BType,
 //           typename CType,
-//           std::uint32_t WaveTileM,
-//           std::uint32_t WaveTileN,
-//           std::uint32_t WaveTileK,
+//           uint32_t WaveTileM,
+//           uint32_t WaveTileN,
+//           uint32_t WaveTileK,
 //           MmaAccumPolicy AccumPolicy>
 // struct ScaleWaveWisePipelineFactory
 // {
@@ -429,9 +426,9 @@ TEST(ScaleMMATrait, MmaSelector_Scale_BF8_BF8_F32_32x32x64_Real)
 // template <typename AType,
 //           typename BType,
 //           typename CType,
-//           std::uint32_t WaveTileM,
-//           std::uint32_t WaveTileN,
-//           std::uint32_t WaveTileK,
+//           uint32_t WaveTileM,
+//           uint32_t WaveTileN,
+//           uint32_t WaveTileK,
 //           MmaAccumPolicy AccumPolicy = MmaAccumPolicy::ROW_MAJOR>
 // void MmaSelector_Scale_WaveWise_Real_impl()
 // {
diff --git a/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp b/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp
index ff156e8692..af4e470004 100644
--- a/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp
+++ b/test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp
@@ -1,28 +1,26 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 
-#include <cstdint>
-#include <gtest/gtest.h>
-#include <iostream>
-#include <numeric>
+#include "pipeline_tests_helper.hpp"
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma.hpp"
-#include "ck_tile/core/arch/mma/mma_op_family.hpp"
-#include "ck_tile/core/arch/mma/mma_selector.hpp"
 #include "ck_tile/core/arch/mma/mma_wavewise.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp"
-#include <hip/hip_runtime.h>
-#include "ck_tile/core/numeric/bfloat16.hpp"
-#include "ck_tile/core/numeric/float8.hpp"
+#include "ck_tile/core/numeric/ext_vector_base.hpp"
 #include "ck_tile/core/numeric/half.hpp"
 #include "ck_tile/core/numeric/integer.hpp"
-#include "ck_tile/host/hip_check_error.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
+#include "ck_tile/host/hip_check_error.hpp"
+#include "ck_tile/host/kernel_launch.hpp"
+#include "ck_tile/host/stream_config.hpp"
 
-#include "pipeline_tests_helper.hpp"
+#include <gtest/gtest.h>
+#include <hip/hip_runtime.h>
+
+#include <tuple>
+#include <type_traits>
+#include <vector>
 
 using namespace ck_tile;
 using namespace ck_tile::core::arch;
diff --git a/test/ck_tile/core/arch/mma/test_amdgcn_mma.cpp b/test/ck_tile/core/arch/mma/test_amdgcn_mma.cpp
index f15c21bfe6..75cd507097 100644
--- a/test/ck_tile/core/arch/mma/test_amdgcn_mma.cpp
+++ b/test/ck_tile/core/arch/mma/test_amdgcn_mma.cpp
@@ -2,11 +2,9 @@
 // SPDX-License-Identifier: MIT
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
-#include "ck_tile/core/arch/mma/mma_op_family.hpp"
-#include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
+#include "ck_tile/core/arch/mma/mma.hpp"
 #include "ck_tile/core/config.hpp"
+#include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/host/hip_check_error.hpp"
 
@@ -15,7 +13,6 @@
 #include <gtest/gtest.h>
 #include <hip/hip_runtime.h>
 
-#include <cstdint>
 #include <type_traits>
 
 using namespace ck_tile;
diff --git a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc
index 7cd471f562..cc27210d6e 100644
--- a/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc
+++ b/test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc
@@ -4,13 +4,9 @@
 #pragma once
 
 #include "ck_tile/core/arch/arch.hpp"
-#include "ck_tile/core/arch/mma/mfma/mfma.hpp"
-#include "ck_tile/core/arch/mma/mma_traits.hpp"
-#include "ck_tile/core/arch/mma/scale/scale.hpp"
-#include "ck_tile/core/arch/mma/sparse/sparse.hpp"
+#include "ck_tile/core/arch/mma/mma.hpp"
 #include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_calculator.hpp"
 #include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp"
-#include "ck_tile/core/arch/mma/wmma/wmma.hpp"
 #include "ck_tile/core/numeric/bfloat16.hpp"
 #include "ck_tile/core/numeric/float8.hpp"
 #include "ck_tile/core/numeric/half.hpp"