From 5c31eeeddb665fb8dba5aa75777079f628ba7c15 Mon Sep 17 00:00:00 2001 From: Emily Martins <65371150+ecamartins@users.noreply.github.com> Date: Fri, 6 Feb 2026 17:26:57 -0700 Subject: [PATCH] [CK_TILE] Fix MMA concepts compiler error (#4381) ## Motivation CK Tile is required to support certain older OSs; on these OSs, cpp 20 is not fully supported. For ROCm 7.2, compiler errors occur on one of these older OSs. An example of this error is as follows: ```bash /composable_kernel/include/ck_tile/core/arch/mma/amdgcn_mma.hpp:34:28: error: expected concept name with optional arguments 34 | { MmaOp::kAMBlock } -> std::convertible_to; | ``` The goal of this PR is to resolve these compiler errors. ## Technical Details The existing guards around the mma concepts only check if the concepts language feature is supported, as follows: ```cpp #if defined(__cpp_concepts) && __cpp_concepts >= 201907L // ... template concept CtrlFlagsGfx9I = requires(CtrlFlags ctrlFlags) { // Flag members for Gfx9 MFMA instructions { CtrlFlags::Cbsz } -> std::convertible_to; { CtrlFlags::Abid } -> std::convertible_to; { CtrlFlags::Blgp } -> std::convertible_to; }; #endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L ``` That said, in cases where functionality from the `` header is used (e.g., `std::convertible_to`), this guard fails to check whether the `` header is available. This change adds an additional check to the concepts that make use of functionality from the `` header to ensure the header is available. ## Test Plan I tested the changes on the relevant docker for gfx90a, gfx950, and gfx942 and the compiler issue is not present. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --- include/ck_tile/core/arch/arch.hpp | 12 ++++++++++++ include/ck_tile/core/arch/mma/amdgcn_mma.hpp | 6 ++++-- include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp | 5 +++-- include/ck_tile/core/arch/mma/mma_selector.hpp | 6 ++++-- include/ck_tile/core/arch/mma/mma_traits.hpp | 6 ++++-- include/ck_tile/core/arch/mma/mma_transforms.hpp | 6 ++++-- 6 files changed, 31 insertions(+), 10 deletions(-) diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index 8892862c0e..541e1da6f9 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -14,6 +14,18 @@ #include "ck_tile/core/arch/amd_buffer_addressing.hpp" #include "ck_tile/core/utility/ignore.hpp" +#if __has_include() +#define CK_TILE_CONCEPTS_HEADER 1 +#else +#define CK_TILE_CONCEPTS_HEADER 0 +#endif //__has_include() + +#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#define CK_TILE_CONCEPTS 1 +#else +#define CK_TILE_CONCEPTS 0 +#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L + #define CK_TILE_S_CNT_MAX 0b1100'1111'0111'1111 #define CK_TILE_VMCNT(cnt) \ ([]() { static_assert(!((cnt) >> 6), "VMCNT only has 6 bits"); }(), \ diff --git a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp index 1eef5819bc..7801e8ed3c 100644 --- a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp +++ b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp @@ -3,6 +3,7 @@ #pragma once +#include "ck_tile/core/arch/arch.hpp" #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/vector_type.hpp" #include "ck_tile/core/utility/ignore.hpp" @@ -18,7 +19,8 @@ namespace ck_tile::core::arch::mma { */ struct Unsupported; -#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER +#include /** * @concept MmaOpI * @brief Expresses the meta-data interface required for each MmaOp policy. @@ -52,7 +54,7 @@ concept MmaOpI = requires(MmaOp op) { } -> std::convertible_to; }; -#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L +#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER /** * @class amdgcn_mma 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 c7375c5e12..7db76b1919 100644 --- a/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp +++ b/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp @@ -30,7 +30,8 @@ struct DefaultMfmaCtrlFlags static constexpr uint32_t Blgp = 0; // BLGP flag, default 0 }; -#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER +#include /** * @concept CtrlFlagsGfx9I @@ -44,7 +45,7 @@ concept CtrlFlagsGfx9I = requires(CtrlFlags ctrlFlags) { { CtrlFlags::Blgp } -> std::convertible_to; }; -#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L +#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER /** * @struct amdgcn_mma diff --git a/include/ck_tile/core/arch/mma/mma_selector.hpp b/include/ck_tile/core/arch/mma/mma_selector.hpp index 070189373d..eae0f705df 100644 --- a/include/ck_tile/core/arch/mma/mma_selector.hpp +++ b/include/ck_tile/core/arch/mma/mma_selector.hpp @@ -2,6 +2,8 @@ // SPDX-License-Identifier: MIT #pragma once +#include "ck_tile/core/arch/arch.hpp" + namespace ck_tile::core::arch::mma { /** @@ -42,7 +44,7 @@ struct MmaDefaultSelector amdgcn_mma>; }; -#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#if CK_TILE_CONCEPTS /** * @concept MmaSelectorI @@ -54,7 +56,7 @@ concept MmaSelectorI = requires(MmaSelector op) { typename MmaSelector::SelectedOp; }; -#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L +#endif // CK_TILE_CONCEPTS } // namespace ck_tile::core::arch::mma diff --git a/include/ck_tile/core/arch/mma/mma_traits.hpp b/include/ck_tile/core/arch/mma/mma_traits.hpp index 8a9092f9cb..7bcf95ac55 100644 --- a/include/ck_tile/core/arch/mma/mma_traits.hpp +++ b/include/ck_tile/core/arch/mma/mma_traits.hpp @@ -2,6 +2,7 @@ // SPDX-License-Identifier: MIT #pragma once #include "amdgcn_mma.hpp" +#include "ck_tile/core/arch/arch.hpp" #include "mfma/mfma_traits.hpp" #include "wmma/wmma_traits.hpp" @@ -49,7 +50,8 @@ static constexpr bool is_mma_op_supported_v = is_mma_op_supported::value; template struct MmaOpParams; -#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER +#include /** * @concept MmaOpParamsI @@ -69,7 +71,7 @@ concept MmaOpParamsI = requires(MmaOpParams op) { { MmaOpParams::GfxTargetId } -> std::convertible_to; }; -#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L +#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER /** * @struct MmaOpParams diff --git a/include/ck_tile/core/arch/mma/mma_transforms.hpp b/include/ck_tile/core/arch/mma/mma_transforms.hpp index 4131daaced..811df04364 100644 --- a/include/ck_tile/core/arch/mma/mma_transforms.hpp +++ b/include/ck_tile/core/arch/mma/mma_transforms.hpp @@ -1,6 +1,8 @@ // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT #pragma once +#include "ck_tile/core/arch/arch.hpp" + namespace ck_tile::core::arch::mma { /** @@ -27,7 +29,7 @@ template // TODO: c++20 template struct MmaTransformsDefaultSelector; -#if defined(__cpp_concepts) && __cpp_concepts >= 201907L +#if CK_TILE_CONCEPTS /** * @concept MmaTransformsI @@ -42,6 +44,6 @@ concept MmaTransformsI = requires(MmaTransforms transforms) { typename MmaTransforms::DTransform; }; -#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L +#endif // CK_TILE_CONCEPTS } // namespace ck_tile::core::arch::mma