ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 19:44:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#7850 (commit e8f2756)

Browse Source 
=?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
This commit is contained in:
Yung-sheng Tu
2026-06-29 18:51:17 +00:00
committed by assistant-librarian[bot]
parent 3719bf05c2
commit 604c56bc0e
35 changed files with 573 additions and 1141 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
example/ck_tile/51_tile_distr_enc_reg_map/example_tile_distr_enc_calc.cpp
View File

@@ -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;

18
include/ck_tile/core/arch/mma/amdgcn_mma.hpp
View File

@@ -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

8
include/ck_tile/core/arch/mma/mfma/mfma.hpp
View File

@@ -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"

670
include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp
View File

File diff suppressed because it is too large Load Diff

7
include/ck_tile/core/arch/mma/mfma/mfma_selector.hpp
View File

@@ -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 {

1
include/ck_tile/core/arch/mma/mfma/mfma_traits.hpp
View File

@@ -4,6 +4,7 @@
#pragma once
#include "ck_tile/core/config.hpp"
#include <type_traits>
namespace ck_tile::core::arch::mma {

3
include/ck_tile/core/arch/mma/mfma/mfma_transforms.hpp
View File

@@ -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 {
/**

12
include/ck_tile/core/arch/mma/mma.hpp
View File

@@ -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"

23
include/ck_tile/core/arch/mma/mma_pipeline.hpp
View File

@@ -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

10
include/ck_tile/core/arch/mma/mma_selector.hpp
View File

@@ -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"

5
include/ck_tile/core/arch/mma/mma_traits.hpp
View File

@@ -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>

4
include/ck_tile/core/arch/mma/mma_transforms.hpp
View File

@@ -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 {

22
include/ck_tile/core/arch/mma/mma_wavewise.hpp
View File

@@ -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 {

377
include/ck_tile/core/arch/mma/scale/mfma/scale_gfx9.hpp
View File

@@ -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

13
include/ck_tile/core/arch/mma/scale/mfma/selector.hpp
View File

@@ -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>

10
include/ck_tile/core/arch/mma/scale/scale_mma_pipeline.hpp
View File

@@ -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 {

3
include/ck_tile/core/arch/mma/scale/scale_traits.hpp
View File

@@ -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 {

7
include/ck_tile/core/arch/mma/sparse/mfma/selector.hpp
View File

@@ -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 {

300
include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp
View File

@@ -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

8
include/ck_tile/core/arch/mma/sparse/sparse.hpp
View File

@@ -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"

17
include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp
View File

@@ -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 {

9
include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp
View File

@@ -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 {

7
include/ck_tile/core/arch/mma/sparse/wmma/selector.hpp
View File

@@ -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 {

5
include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp
View File

@@ -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 {

11
include/ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp
View File

@@ -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.

5
include/ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp
View File

@@ -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"

5
include/ck_tile/core/arch/mma/wmma/wmma_gfx12.hpp
View File

@@ -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"

6
include/ck_tile/core/arch/mma/wmma/wmma_selector.hpp
View File

@@ -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 {

6
include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp
View File

@@ -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 {

29
include/ck_tile/core/tensor/tile_distribution.hpp
View File

@@ -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

2
include/ck_tile/ops/gemm.hpp
View File

@@ -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"

65
test/ck_tile/core/arch/mma/pipeline/test_amdgcn_scale_mma.cpp
View File

@@ -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()
// {

24
test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp
View File

@@ -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;

7
test/ck_tile/core/arch/mma/test_amdgcn_mma.cpp
View File

@@ -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;

6
test/ck_tile/core/arch/mma/test_amdgcn_mma_layout.inc
View File

@@ -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"