[rocm-libraries] ROCm/rocm-libraries#4381 (commit 5df3343)

[CK_TILE] Fix MMA concepts compiler error

## 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<unsigned int>;
      |
```

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 <typename CtrlFlags>
concept CtrlFlagsGfx9I = requires(CtrlFlags ctrlFlags) {
    // Flag members for Gfx9 MFMA instructions
    { CtrlFlags::Cbsz } -> std::convertible_to<int>;
    { CtrlFlags::Abid } -> std::convertible_to<int>;
    { CtrlFlags::Blgp } -> std::convertible_to<int>;
};

#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L
```
That said, in cases where functionality from the `<concepts>` header is
used (e.g., `std::convertible_to`), this guard fails to check whether
the `<concepts>` header is available.

This change adds an additional check to the concepts that make use of
functionality from the `<concepts>` 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

@@ -14,6 +14,18 @@
 #include "ck_tile/core/arch/amd_buffer_addressing.hpp"
 #include "ck_tile/core/utility/ignore.hpp"
 
+#if __has_include(<concepts>)
+#define CK_TILE_CONCEPTS_HEADER 1
+#else
+#define CK_TILE_CONCEPTS_HEADER 0
+#endif //__has_include(<concepts>)
+
+#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"); }(), \

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 <concepts>
 /**
  * @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<typename MmaOp::CVecType>;
 };
 
-#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L
+#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
 /**
  *  @class  amdgcn_mma

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 <concepts>
 
 /**
  * @concept CtrlFlagsGfx9I
@@ -44,7 +45,7 @@ concept CtrlFlagsGfx9I = requires(CtrlFlags ctrlFlags) {
     { CtrlFlags::Blgp } -> std::convertible_to<int>;
 };
 
-#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L
+#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
 /**
  * @struct amdgcn_mma

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<ADataType, BDataType, CDataType, FragM, FragN, FragK, void, amdgcn_target<>>;
 };
 
-#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
 

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<MmaOp>::value;
 template <typename MmaOp>
 struct MmaOpParams;
 
-#if defined(__cpp_concepts) && __cpp_concepts >= 201907L
+#if CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
+#include <concepts>
 
 /**
  *  @concept MmaOpParamsI
@@ -69,7 +71,7 @@ concept MmaOpParamsI = requires(MmaOpParams op) {
     { MmaOpParams::GfxTargetId } -> std::convertible_to<amdgcn_target_arch_id>;
 };
 
-#endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L
+#endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
 /**
  * @struct MmaOpParams

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 <typename MmaOp, typename CompilerTarget, typename Enable = void>
 // TODO: c++20 template <MmaOpI MmaOp, amdgcn_target_arch_id CompilerTarget, typename Enable = void>
 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