[rocm-libraries] ROCm/rocm-libraries#12203 (commit 3d0a0e8)

[CK] enable new gfx1250-strict target.

## Motivation

These changes will enable building CK for the new gfx1250-strict target
by limiting access to some of the unsupported builtins. I have verified
that the code builds successfully for the gfx1250-strict with a
rocm/compiler installation that supports it.

JIRA ID : AICK-2239

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Author: illsilin

CMakeLists.txt

@@ -196,7 +196,7 @@ function(_ck_gpu_target_string_to_id TARGET_STR OUT_VAR)
         set(${OUT_VAR} "0x1201" PARENT_SCOPE)
     elseif(_tgt MATCHES "^gfx12-generic$")
         set(${OUT_VAR} "0x12FF" PARENT_SCOPE)
-    elseif(_tgt STREQUAL "gfx1250")
+    elseif(_tgt STREQUAL "gfx1250" OR _tgt STREQUAL "gfx1250-strict")
         set(${OUT_VAR} "0x1250" PARENT_SCOPE)
     else()
         message(WARNING "_ck_gpu_target_string_to_id: unknown GPU target '${TARGET_STR}', skipping")

example/ck_tile/42_mx_gemm/run_mx_gemm.inc

@@ -188,12 +188,11 @@ int run_mx_gemm_with_layouts(int argc, char* argv[], ALayout, BLayout, CLayout)
     {
         ck_tile::preShuffleScaleBufferPermuteN_gfx950<GemmConfig::N_Warp,
                                                       GemmConfig::N_Tile,
-                                                      XdlMNThread>(
-            scale_b_aligned.mData.data(),
-            scale_b_shuffled.mData.data(),
-            N_scale_aligned,
-            scale_k_size,
-            true);
+                                                      XdlMNThread>(scale_b_aligned.mData.data(),
+                                                                   scale_b_shuffled.mData.data(),
+                                                                   N_scale_aligned,
+                                                                   scale_k_size,
+                                                                   true);
     }
     else
     {

include/ck/ck.hpp

@@ -51,13 +51,12 @@
 #endif
 
 // define general macros for various architectures
-#if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || \
-    defined(__gfx9_4_generic__)
-#define __gfx9__
-#endif
 #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__)
 #define __gfx94__
 #endif
+#if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__)
+#define __gfx9__
+#endif
 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || \
     defined(__gfx1013__) || defined(__gfx10_1_generic__)
 #define __gfx101__
@@ -72,16 +71,15 @@
     defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx11_generic__)
 #define __gfx11__
 #endif
-#if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__) || \
-    defined(__gfx1250__)
-#define __gfx12__
-#endif
 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)
 #define __gfx120__
 #endif
-#if defined(__gfx1250__)
+#if defined(__gfx1250__) || defined(__gfx1250_strict__)
 #define __gfx125__
 #endif
+#if defined(__gfx120__) || defined(__gfx125__)
+#define __gfx12__
+#endif
 // buffer resource
 #ifndef __HIP_DEVICE_COMPILE__ // for host code
 #define CK_BUFFER_RESOURCE_3RD_DWORD -1

include/ck/utility/amd_wmma.hpp

@@ -20,7 +20,7 @@ namespace ck {
 #define __gfx120__
 #endif
 
-#if defined(__gfx1250__)
+#if defined(__gfx1250__) || defined(__gfx1250_strict__)
 #define __gfx125__
 #endif
 
@@ -1426,7 +1426,7 @@ struct intrin_wmma_scale_f32_32x16x128_f4<32, 16, ScaleOpselB, ScaleTypeA, Scale
                           is_same_v<ScaleTypeB, e5m3x4_scale_t> ||
                           is_same_v<ScaleTypeB, e4m3x4_scale_t>,
                       "ScaleTypeB must be e8m0x4_bexp_t, e5m3x4_scale_t, or e4m3x4_scale_t");
-#if defined(__gfx125__)
+#if defined(__gfx1250__)
         int32x16_t arg_a = bit_cast<int32x16_t>(reg_a);
         int32x8_t arg_b  = bit_cast<int32x8_t>(reg_b);
         reg_c.template AsType<float16_t>()(Number<0>{}) =
@@ -1480,7 +1480,7 @@ struct intrin_wmma_scale16_f32_32x16x128_f4<32, 16, ScaleOpselB, ScaleTypeA, Sca
                           is_same_v<ScaleTypeB, e5m3x8_scale_t> ||
                           is_same_v<ScaleTypeB, e4m3x8_scale_t>,
                       "ScaleTypeB must be e8m0x8_bexp_t, e5m3x8_scale_t, or e4m3x8_scale_t");
-#if defined(__gfx125__)
+#if defined(__gfx1250__)
         int32x16_t arg_a = bit_cast<int32x16_t>(reg_a);
         int32x8_t arg_b  = bit_cast<int32x8_t>(reg_b);
         reg_c.template AsType<float16_t>()(Number<0>{}) =

include/ck/utility/e4m3.hpp

@@ -35,7 +35,7 @@ struct e4m3_scale_t
     }
     __host__ __device__ explicit e4m3_scale_t(float scale)
     {
-#if defined(__gfx1250__)
+#if defined(__gfx125__)
         union
         {
             float fval;
@@ -59,7 +59,7 @@ struct e4m3_scale_t
 
     __host__ __device__ explicit operator float() const
     {
-#if defined(__gfx1250__)
+#if defined(__gfx125__)
         union
         {
             unsigned int i32val;

include/ck/utility/e5m3.hpp

@@ -35,7 +35,7 @@ struct e5m3_scale_t
     }
     __host__ __device__ explicit e5m3_scale_t(float scale)
     {
-#if defined(__gfx1250__)
+#if defined(__gfx125__)
         union
         {
             float fval;
@@ -59,7 +59,7 @@ struct e5m3_scale_t
 
     __host__ __device__ explicit operator float() const
     {
-#if defined(__gfx1250__)
+#if defined(__gfx125__)
         union
         {
             unsigned int i32val;

include/ck/utility/mxf4_utils.hpp

@@ -8,7 +8,7 @@
 #include "ck/utility/mxfp_utils.hpp"
 #include "dtype_vector.hpp"
 
-#if CK_MX_ARCH_950 || CK_MX_ARCH_125
+#if CK_MX_ARCH_950 || CK_MX_ARCH_1250
 #define CK_MX_FP4_CVT_FAST_PATH 1
 #else
 #define CK_MX_FP4_CVT_FAST_PATH 0
@@ -287,7 +287,7 @@ static inline __device__ f4x8_t cast_to_f4_scaled(T x, float scale)
     return ret.vf4;
 }
 
-#elif CK_MX_ARCH_125
+#elif CK_MX_ARCH_1250
 // from f4
 template <typename T>
 static inline __device__ enable_if_t<scalar_type<T>::vector_size == 1, T>

include/ck/utility/mxf6_utils.hpp

@@ -7,7 +7,7 @@
 #include "ck/utility/numeric_limits.hpp"
 #include "ck/utility/mxfp_utils.hpp"
 
-#if CK_MX_ARCH_950 || CK_MX_ARCH_125
+#if CK_MX_ARCH_950 || CK_MX_ARCH_1250
 #define CK_MX_FP6_CVT_FAST_PATH 1
 #else
 #define CK_MX_FP6_CVT_FAST_PATH 0
@@ -721,7 +721,7 @@ inline __device__ T_F6 cast_to_f6_scaled(T x, float scale)
     }
 }
 
-#elif CK_MX_ARCH_125
+#elif CK_MX_ARCH_1250
 // from f6
 template <typename T, typename T_F6>
 inline __device__ enable_if_t<scalar_type<T>::vector_size == 1, T> cast_from_f6_scaled(T_F6 x,

include/ck/utility/mxf8_utils.hpp

@@ -4,7 +4,7 @@
 #include "ck/utility/numeric_limits.hpp"
 #include "ck/utility/mxfp_utils.hpp"
 
-#if CK_MX_ARCH_950 || CK_MX_ARCH_125
+#if CK_MX_ARCH_950 || CK_MX_ARCH_1250
 #define CK_MX_FP8_CVT_FAST_PATH 1
 #else
 #define CK_MX_FP8_CVT_FAST_PATH 0
@@ -660,7 +660,7 @@ static __device__ fp8x8_storage_t cast_to_f8_from_bf16_scaled(bhalf8_t v,
     return ret.v8f8x1;
 }
 
-#elif CK_MX_ARCH_125
+#elif CK_MX_ARCH_1250
 
 // fp8 -> float 8
 template <ck_fp8_interpretation_t interpret, typename Ts, int Opsel>
@@ -1011,7 +1011,7 @@ static __device__ bhalf2_t cast_to_bf16_from_f8_scaled(float scale, fp8x2_storag
     out.v8x1 = cast_to_bf16_from_f8_scaled<interpret>(scale, v8);
     return out.v2x4[0];
 }
-#endif // CK_MX_ARCH_125
+#endif // CK_MX_ARCH_1250
 
 #endif // CK_MX_FP8_CVT_FAST_PATH
 

include/ck/utility/mxfp_utils.hpp

@@ -16,6 +16,12 @@
 #define CK_MX_ARCH_125 0
 #endif
 
+#if defined(__gfx1250__) && !defined(__gfx1250_strict__) && __HIP_DEVICE_COMPILE__
+#define CK_MX_ARCH_1250 1
+#else
+#define CK_MX_ARCH_1250 0
+#endif
+
 #ifdef CK_CODE_GEN_RTC
 #define UINT_MAX 4294967295
 #endif