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