[MIOpen] Add initial MIOpen support for gfx1250

projects/miopen/src/kernels/batchnorm_functions.h

@@ -177,11 +177,14 @@
 // TODO: Spaghetti code!!!
 // MIOPEN_USE_AMDGCN may be defined before this header.
 #ifndef MIOPEN_USE_AMDGCN
+// clang-format off
 #if defined(__AMDGCN__) &&                           \
     !((defined(MIO_BN_GFX103X) && MIO_BN_GFX103X) || \
       (defined(MIO_BN_GFX110X) && MIO_BN_GFX110X) || \
       (defined(MIO_BN_GFX115X) && MIO_BN_GFX115X) || \
-      (defined(MIO_BN_GFX120X) && MIO_BN_GFX120X))
+      (defined(MIO_BN_GFX120X) && MIO_BN_GFX120X) || \
+      (defined(MIO_BN_GFX125X) && MIO_BN_GFX125X))
+// clang-format on
 #define MIOPEN_USE_AMDGCN 1
 #else
 #define MIOPEN_USE_AMDGCN 0
@@ -221,6 +224,10 @@
 #define MIO_BN_GFX115X 0
 #endif
 
+#ifndef MIO_BN_GFX125X
+#define MIO_BN_GFX125X 0
+#endif
+
 #ifndef MIO_BN_VECTORIZE
 #define MIO_BN_VECTORIZE 0
 #endif

projects/miopen/src/kernels/default_configurations.hpp

@@ -145,7 +145,8 @@
 #endif
 
 #ifndef MIO_BN_LDSGCN_SIZE
-#define MIO_BN_LDSGCN_SIZE 16
+// 4 SIMD with up to 16 wave each => at most 64 waves
+#define MIO_BN_LDSGCN_SIZE 64
 #endif
 
 #ifndef MIO_BN_LDS_SIZE

projects/miopen/src/kernels/reduction_functions.hpp

@@ -175,33 +175,11 @@ __forceinline__ __device__ void lds_reduce2_2d(FloatAccumC& x,
     y = static_cast<FloatAccumC>(lcl_data[xlid * 2 + 1] * scale);
 }
 
-template <typename FloatAccum>
-__forceinline__ __device__ void dpp_interleaved_reduction(FloatAccum& temp_sum1,
-                                                          FloatAccum& temp_sum2)
-{
-    __asm__ volatile("s_nop 4\n"
-                     "v_add_f32 %0 %0 %0 row_shr:1 bound_ctrl:0\n"
-                     "v_add_f32 %1 %1 %1 row_shr:1 bound_ctrl:0\n"
-                     "s_nop 0\n"
-                     "v_add_f32 %0 %0 %0 row_shr:2 bound_ctrl:0\n"
-                     "v_add_f32 %1 %1 %1 row_shr:2 bound_ctrl:0\n"
-                     "s_nop 0\n"
-                     "v_add_f32 %0 %0 %0 row_shr:4 bank_mask:0xe\n"
-                     "v_add_f32 %1 %1 %1 row_shr:4 bank_mask:0xe\n"
-                     "s_nop 0\n"
-                     "v_add_f32 %0 %0 %0 row_shr:8 bank_mask:0xc\n"
-                     "v_add_f32 %1 %1 %1 row_shr:8 bank_mask:0xc\n"
-                     "s_nop 0\n"
-                     "v_add_f32 %0 %0 %0 row_bcast:15 row_mask:0xa\n"
-                     "v_add_f32 %1 %1 %1 row_bcast:15 row_mask:0xa\n"
-                     "s_nop 0\n"
-                     "v_add_f32 %0 %0 %0 row_bcast:31 row_mask:0xc\n"
-                     "v_add_f32 %1 %1 %1 row_bcast:31 row_mask:0xc\n"
-                     "s_nop 0"
-                     : "=v"(temp_sum1), "=v"(temp_sum2)
-                     : "0"(temp_sum1), "1"(temp_sum2));
-}
-
+// Caller must ensure: SizeLclData >= (blockDim.x * blockDim.y * blockDim.z + warpSize - 1) /
+// warpSize
+// @warning Undefined behavior if SizeLclData is too small
+// Caller must ensure: All lanes must be active
+// @warning Undefined behavior if lanes are masked
 template <typename FloatAccum, unsigned int SizeLclData>
 __forceinline__ __device__ void gcn_reduce2(FloatAccum& x,
                                             FloatAccum& y,
@@ -210,10 +188,12 @@ __forceinline__ __device__ void gcn_reduce2(FloatAccum& x,
                                             FloatAccum (&lcl_data_y)[SizeLclData],
                                             unsigned int lid)
 {
-    const unsigned int ldsidx = lid >> 6;
-    dpp_interleaved_reduction(x, y);
+    const unsigned int ldsidx         = lid / warpSize;
+    constexpr unsigned long long mask = 0xFFFFFFFFFFFFFFFFull;
+    x                                 = __reduce_add_sync(mask, x);
+    y                                 = __reduce_add_sync(mask, y);
     // Last thread
-    if((lid % 64) == 63)
+    if((lid % warpSize) == warpSize - 1)
     {
         lcl_data_x[ldsidx] = x;
         lcl_data_y[ldsidx] = y;

projects/miopen/src/legacy_composable_kernel/composable_kernel/include/utility/config.hpp

@@ -22,7 +22,7 @@
       defined(CK_AMD_GPU_GFX1036) || defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || \
       defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103) || defined(CK_AMD_GPU_GFX1150) || \
       defined(CK_AMD_GPU_GFX1151) || defined(CK_AMD_GPU_GFX1152) || defined(CK_AMD_GPU_GFX1153) || \
-      defined(CK_AMD_GPU_GFX1200) || defined(CK_AMD_GPU_GFX1201))
+      defined(CK_AMD_GPU_GFX1200) || defined(CK_AMD_GPU_GFX1201) || defined(CK_AMD_GPU_GFX1250))
 #error No CK_AMD_GPU_GFX* macro defined. Exactly one target must be defined.
 #endif
 
@@ -43,7 +43,7 @@
     defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) ||   \
     defined(CK_AMD_GPU_GFX1103) || defined(CK_AMD_GPU_GFX1150) || defined(CK_AMD_GPU_GFX1151) ||   \
     defined(CK_AMD_GPU_GFX1152) || defined(CK_AMD_GPU_GFX1153) || defined(CK_AMD_GPU_GFX1200) ||   \
-    defined(CK_AMD_GPU_GFX1201)
+    defined(CK_AMD_GPU_GFX1201) || defined(CK_AMD_GPU_GFX1250)
 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
 #endif
 
@@ -55,7 +55,7 @@
     defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) || \
     defined(CK_AMD_GPU_GFX1103) || defined(CK_AMD_GPU_GFX1150) || defined(CK_AMD_GPU_GFX1151) || \
     defined(CK_AMD_GPU_GFX1152) || defined(CK_AMD_GPU_GFX1153) || defined(CK_AMD_GPU_GFX1200) || \
-    defined(CK_AMD_GPU_GFX1201)
+    defined(CK_AMD_GPU_GFX1201) || defined(CK_AMD_GPU_GFX1250)
 #define CK_USE_AMD_V_FMAC_F32
 #define CK_USE_AMD_V_DOT2_F32_F16
 #define CK_USE_AMD_V_DOT4_I32_I8

projects/miopen/src/solver/batchnorm/backward_per_activation.cpp

@@ -121,6 +121,7 @@ BnBwdTrainingPerActivation::GetSolution(const ExecutionContext& context,
             {"MIO_BN_GFX110X", (StartsWith(handle.GetDeviceName(), "gfx110") ? "1" : "0")},
             {"MIO_BN_GFX115X", (StartsWith(handle.GetDeviceName(), "gfx115") ? "1" : "0")},
             {"MIO_BN_GFX120X", (StartsWith(handle.GetDeviceName(), "gfx120") ? "1" : "0")},
+            {"MIO_BN_GFX125X", (StartsWith(handle.GetDeviceName(), "gfx125") ? "1" : "0")},
         };
 
         kernel.comp_options = build_params.GenerateFor(kbp::HIP{});

projects/miopen/src/solver/batchnorm/backward_per_activation_fused.cpp

@@ -135,7 +135,9 @@ ConvSolution BnBwdTrgActivationFused::GetSolution(const FusionContext& context,
 
         kernel.g_wk = {xgridsize, ygridsize, zgridsize};
 
-        unsigned int ldsgcn   = xlocalsize / 64;
+        auto const waveSize = handle.GetWavefrontWidth();
+
+        unsigned int ldsgcn   = xlocalsize / waveSize;
         unsigned int ldsnogcn = xlocalsize;
 
         int variant = 0;
@@ -156,7 +158,7 @@ ConvSolution BnBwdTrgActivationFused::GetSolution(const FusionContext& context,
 
         const auto& activ_op =
             dynamic_cast<ActivBwdFusionOpDescriptor&>(*problem.fusion_plan_desc->op_map[1]);
-        const auto build_params = KernelBuildParameters{
+        auto build_params = KernelBuildParameters{
             {"MIO_BN_N", static_cast<int>(n)},
             {"MIO_BN_NCHW", static_cast<int>(n * c * h * w)},
             {"MIO_BN_NHW", static_cast<int>(n * h * w)},
@@ -172,12 +174,18 @@ ConvSolution BnBwdTrgActivationFused::GetSolution(const FusionContext& context,
             {"MIO_BN_GFX110X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx110"))},
             {"MIO_BN_GFX115X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx115"))},
             {"MIO_BN_GFX120X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx120"))},
+            {"MIO_BN_GFX125X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx125"))},
             {"MIOPEN_NRN_OP_ID", static_cast<int>(activ_op.activMode)},
             {"MIOPEN_USE_FP16", static_cast<int>(dtype == miopenHalf)},
             {"MIOPEN_USE_FP32", static_cast<int>(dtype == miopenFloat)},
             {"DATA_TYPE", data_type}};
         kernel.comp_options = build_params.GenerateFor(kbp::HIP{});
 
+        if(mode == miopenBNSpatial)
+        {
+            build_params.Define("HIP_ENABLE_EXTRA_WARP_SYNC_TYPES");
+        }
+
         result.construction_params.push_back(kernel);
     }
 

projects/miopen/src/solver/batchnorm/backward_spatial.cpp

@@ -340,8 +340,11 @@ ConvSolution BnBwdTrainingSpatial::GetSolution(const ExecutionContext& context,
                 {"MIO_BN_GFX110X", (StartsWith(handle.GetDeviceName(), "gfx110") ? "1" : "0")},
                 {"MIO_BN_GFX115X", (StartsWith(handle.GetDeviceName(), "gfx115") ? "1" : "0")},
                 {"MIO_BN_GFX120X", (StartsWith(handle.GetDeviceName(), "gfx120") ? "1" : "0")},
+                {"MIO_BN_GFX125X", (StartsWith(handle.GetDeviceName(), "gfx125") ? "1" : "0")},
             };
 
+            build_params.Define("HIP_ENABLE_EXTRA_WARP_SYNC_TYPES");
+
             kernel.comp_options = build_params.GenerateFor(kbp::HIP());
 
             kernel.l_wk.push_back(xlocalsize);

projects/miopen/src/solver/batchnorm/forward_inference.cpp

@@ -162,6 +162,7 @@ ConvSolution BnFwdInference::GetSolution(const ExecutionContext& context,
             {"MIO_BN_GFX110X", (StartsWith(handle.GetDeviceName(), "gfx110") ? "1" : "0")},
             {"MIO_BN_GFX115X", (StartsWith(handle.GetDeviceName(), "gfx115") ? "1" : "0")},
             {"MIO_BN_GFX120X", (StartsWith(handle.GetDeviceName(), "gfx120") ? "1" : "0")},
+            {"MIO_BN_GFX125X", (StartsWith(handle.GetDeviceName(), "gfx125") ? "1" : "0")},
             {"MIO_LAYOUT_NHWC", static_cast<int>(problem.IsLayoutNHWC())},
             {"MIO_BN_VECTORIZE", static_cast<int>(vectorsize > 1)},
             {"MIO_BN_VEC_SIZE", vectorsize},

projects/miopen/src/solver/batchnorm/forward_per_activation.cpp

@@ -120,6 +120,7 @@ BnFwdTrainingPerActivation::GetSolution(const ExecutionContext& context,
             {"MIO_BN_GFX110X", (StartsWith(handle.GetDeviceName(), "gfx110") ? "1" : "0")},
             {"MIO_BN_GFX115X", (StartsWith(handle.GetDeviceName(), "gfx115") ? "1" : "0")},
             {"MIO_BN_GFX120X", (StartsWith(handle.GetDeviceName(), "gfx120") ? "1" : "0")},
+            {"MIO_BN_GFX125X", (StartsWith(handle.GetDeviceName(), "gfx125") ? "1" : "0")},
         };
 
         auto kernel = KernelInfo{};

projects/miopen/src/solver/batchnorm/forward_per_activation_fused.cpp

@@ -175,6 +175,7 @@ ConvSolution BnFwdTrgActivationFused::GetSolution(const FusionContext& context,
             {"MIO_BN_GFX110X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx110"))},
             {"MIO_BN_GFX115X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx115"))},
             {"MIO_BN_GFX120X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx120"))},
+            {"MIO_BN_GFX125X", static_cast<int>(StartsWith(handle.GetDeviceName(), "gfx125"))},
             {"MIOPEN_YES_ACTIV", static_cast<int>(1)},
             {"MIOPEN_NRN_OP_ID", static_cast<int>(activ_op.activMode)},
             {"MIOPEN_USE_FP16", static_cast<int>(input_desc.GetType() == miopenHalf)},

projects/miopen/src/solver/batchnorm/forward_spatial.cpp

@@ -227,6 +227,8 @@ ConvSolution BnFwdTrainingSpatial::GetSolution(const ExecutionContext& context,
     int stash_method = 0;
     size_t nelements = 1;
 
+    auto const waveSize = handle.GetWavefrontWidth();
+
     GetVariantFromKernelId(
         config.kernel_id, variant, vectorsize, xlocalsize, ylocalsize, zlocalsize, nelements);
 
@@ -240,7 +242,7 @@ ConvSolution BnFwdTrainingSpatial::GetSolution(const ExecutionContext& context,
             xlocalsize = 256;
         }
         xgridsize = c * xlocalsize;
-        ldsgcn    = xlocalsize / 64;
+        ldsgcn    = xlocalsize / waveSize;
         ldsnogcn  = xlocalsize;
     }
     else
@@ -282,7 +284,7 @@ ConvSolution BnFwdTrainingSpatial::GetSolution(const ExecutionContext& context,
                 (xlocalsize * ylocalsize * zlocalsize) / xlocalsize_final / zlocalsize_final;
         }
         ldsnogcn = xlocalsize * ylocalsize * zlocalsize;
-        ldsgcn   = xlocalsize * ylocalsize * zlocalsize / 64;
+        ldsgcn   = xlocalsize * ylocalsize * zlocalsize / waveSize;
     }
 
     auto result = ConvSolution{miopenStatusSuccess};
@@ -316,6 +318,7 @@ ConvSolution BnFwdTrainingSpatial::GetSolution(const ExecutionContext& context,
             {"MIO_BN_GFX110X", (StartsWith(handle.GetDeviceName(), "gfx110") ? "1" : "0")},
             {"MIO_BN_GFX115X", (StartsWith(handle.GetDeviceName(), "gfx115") ? "1" : "0")},
             {"MIO_BN_GFX120X", (StartsWith(handle.GetDeviceName(), "gfx120") ? "1" : "0")},
+            {"MIO_BN_GFX125X", (StartsWith(handle.GetDeviceName(), "gfx125") ? "1" : "0")},
             {"MIO_LAYOUT_NHWC", static_cast<int>(problem.IsLayoutNHWC())},
             {"MIO_BN_VECTORIZE", static_cast<int>(vectorsize > 1)},
             {"MIO_BN_VEC_SIZE", vectorsize},
@@ -328,6 +331,8 @@ ConvSolution BnFwdTrainingSpatial::GetSolution(const ExecutionContext& context,
         build_params.Define("MIO_BN_CHW", in_nstride);
         build_params.Define("MIO_BN_NCHW", in_nchw);
 
+        build_params.Define("HIP_ENABLE_EXTRA_WARP_SYNC_TYPES");
+
         kernel.kernel_file      = "MIOpenBatchNormFwdTrainSpatial.cpp";
         std::string kernel_name = "MIOpenBatchNormFwdTrainSpatial";
         kernel.comp_options     = build_params.GenerateFor(kbp::HIP{});

projects/miopen/src/solver/conv/conv_winoRxS.cpp

@@ -694,12 +694,12 @@ static bool IsApplicableBase(const ExecutionContext& ctx, const ProblemDescripti
 
     const auto name = ctx.GetStream().GetDeviceName();
     if(!(StartsWith(name, "gfx9") || StartsWith(name, "gfx10") || StartsWith(name, "gfx11") ||
-         StartsWith(name, "gfx12")))
+         StartsWith(name, "gfx120")))
         return false;
     if(problem.IsFp16() &&
        !(name == "gfx906" || name == "gfx908" || name == "gfx90a" || name == "gfx942" ||
          StartsWith(name, "gfx95") || name == "gfx1011" || name == "gfx1012" ||
-         StartsWith(name, "gfx103") || StartsWith(name, "gfx11") || StartsWith(name, "gfx12")))
+         StartsWith(name, "gfx103") || StartsWith(name, "gfx11") || StartsWith(name, "gfx120")))
         return false;
 
     if(name == "gfx90a" && problem.IsGfx90aFp16altRequired())
@@ -868,7 +868,7 @@ ConvSolution ConvBinWinoRxS<Winodata, Winofilter>::GetSolution(
     const auto is_gfx9  = StartsWith(name, "gfx9");
     const auto is_gfx10 = StartsWith(name, "gfx10");
     const auto is_gfx11 = StartsWith(name, "gfx11");
-    const auto is_gfx12 = StartsWith(name, "gfx12");
+    const auto is_gfx12 = StartsWith(name, "gfx120");
     const auto is_v21   = IsWinogradV21Preferred<Winodata, Winofilter>(name, problem);
     size_t wg_size      = is_gfx9 ? 512 : 256;
 

projects/miopen/src/solver/conv/conv_wino_fury_RxS.cpp

@@ -326,7 +326,7 @@ class ShaderModelFactory
         {
             return std::make_unique<ShaderModelV2>(args, cu_count, n_groups, reduced_vgpr_mem);
         }
-        else if(StartsWith(dev_name, "gfx12"))
+        else if(StartsWith(dev_name, "gfx120"))
         {
             return std::make_unique<ShaderModelV4>(args, cu_count, n_groups, reduced_vgpr_mem);
         }
@@ -389,8 +389,8 @@ bool ConvWinoFuryRxSCommon<Winodata, Winofilter>::IsApplicable(const ExecutionCo
         return false;
 
     const auto dev_name = ctx.GetStream().GetDeviceName();
-    // All gfx11/gfx12 ASICs are supported
-    if(!(StartsWith(dev_name, "gfx11") || StartsWith(dev_name, "gfx12")))
+    // All gfx11/gfx120x ASICs are supported
+    if(!(StartsWith(dev_name, "gfx11") || StartsWith(dev_name, "gfx120")))
         return false;
 #if WORKAROUND_ISSUE_3044
     if(dev_name == "gfx1103")
@@ -520,7 +520,7 @@ ConvWinoFuryRxSCommon<Winodata, Winofilter>::GetSolution(const ExecutionContext&
     std::string kernel_arch    = "_gfx11";
 
     const bool is_gfx11 = StartsWith(dev_name, "gfx11");
-    const bool is_gfx12 = StartsWith(dev_name, "gfx12");
+    const bool is_gfx12 = StartsWith(dev_name, "gfx120");
 
     if(!is_gfx11 && !is_gfx12)
         MIOPEN_THROW(miopenStatusInternalError);

projects/miopen/src/solver/conv_winoRxS_fused.cpp

@@ -155,13 +155,13 @@ bool ConvBinWinogradRxSf2x3g1Fused::IsApplicable(const FusionContext& context,
 
     const std::string name = conv_ctx.GetStream().GetDeviceName();
     if(!(StartsWith(name, "gfx9") || StartsWith(name, "gfx10") || StartsWith(name, "gfx11") ||
-         StartsWith(name, "gfx12")))
+         StartsWith(name, "gfx120")))
         return false;
 
     if(conv_problem.IsFp16() &&
        !(StartsWith(name, "gfx906") || StartsWith(name, "gfx908") || StartsWith(name, "gfx90a") ||
          StartsWith(name, "gfx942") || StartsWith(name, "gfx1011") || StartsWith(name, "gfx1012") ||
-         StartsWith(name, "gfx103") || StartsWith(name, "gfx11") || StartsWith(name, "gfx12")))
+         StartsWith(name, "gfx103") || StartsWith(name, "gfx11") || StartsWith(name, "gfx120")))
         return false;
 
     // clang-format off
@@ -208,7 +208,7 @@ ConvSolution ConvBinWinogradRxSf2x3g1Fused::GetSolution(const FusionContext& con
     const auto is_gfx9  = StartsWith(name, "gfx9");
     const auto is_gfx10 = StartsWith(name, "gfx10");
     const auto is_gfx11 = StartsWith(name, "gfx11");
-    const auto is_gfx12 = StartsWith(name, "gfx12");
+    const auto is_gfx12 = StartsWith(name, "gfx120");
     const auto is_v21   = IsWinogradV21Preferred<2, 3>(name, conv_problem);
     size_t wg_size      = is_gfx9 ? 512 : 256;
     kernel.g_wk.push_back(wg_size * n_groups);

projects/miopen/src/solver/mha/mha_common.hpp

@@ -40,6 +40,7 @@
 #include <half/half.hpp>
 #if MIOPEN_ROCBLAS_VERSION_FLAT < 2045000
 #include <rocblas.h>
+#define USE_ROCBLAS_EX3 0
 #else
 #include <rocblas/rocblas.h>
 /// rocblas_gemm_ex3 supports F8 datatypes.

projects/miopen/test/CMakeLists.txt

@@ -51,6 +51,7 @@ option( MIOPEN_TEST_GFX103X "Test on Navi21/22 (gfx1030/31)" OFF )
 option( MIOPEN_TEST_GFX110X "Test on Navi31/32/33/Phoenix (gfx1100/03)" OFF )
 option( MIOPEN_TEST_GFX115X "Test on Navi35 strix halo (gfx1150/51/52/53)" OFF )
 option( MIOPEN_TEST_GFX120X "Test on gfx1200/01" OFF )
+option( MIOPEN_TEST_GFX125X "Test on gfx1250" OFF )
 option( MIOPEN_TEST_GPU_XNACK_ENABLED "Test as if XNACK mode is enabled" OFF )
 option( MIOPEN_TEST_CONV "" OFF )
 option( MIOPEN_TEST_DEEPBENCH "" OFF )
@@ -109,7 +110,8 @@ endif()
 # Also we do not detect GPU when target GPU for testing is specified explicitly.
 set(MIOPEN_TEST_GPU_DETECTION_FAILED FALSE)
 set(MIOPEN_NO_GPU FALSE)
-if (NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX94X OR MIOPEN_TEST_GFX95X OR MIOPEN_TEST_GFX103X OR MIOPEN_TEST_GFX110X OR MIOPEN_TEST_GFX115X OR MIOPEN_TEST_GFX120X OR MIOPEN_TEST_HIP_NOGPU))
+
+if(NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX94X OR MIOPEN_TEST_GFX95X OR MIOPEN_TEST_GFX103X OR MIOPEN_TEST_GFX110X OR MIOPEN_TEST_GFX115X OR MIOPEN_TEST_GFX120X OR MIOPEN_TEST_GFX125X OR MIOPEN_TEST_HIP_NOGPU))
     # Try to find rocminfo first, then hipinfo as fallback (for Windows)
     find_program(ROCMINFO
         NAMES rocminfo hipinfo
@@ -142,6 +144,8 @@ if (NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPE
             set(MIOPEN_TEST_GFX115X ON)
         elseif(ROCMINFO_OUTPUT MATCHES "gfx120.")
             set(MIOPEN_TEST_GFX120X ON)
+        elseif(ROCMINFO_OUTPUT MATCHES "gfx125.")
+            set(MIOPEN_TEST_GFX125X ON)
         elseif(ROCMINFO_OUTPUT MATCHES "gfx900")
             set(MIOPEN_TEST_GFX900 ON)
         elseif(ROCMINFO_OUTPUT MATCHES "gfx906")
@@ -196,6 +200,7 @@ message(STATUS "MIOPEN_TEST_GFX103X ${MIOPEN_TEST_GFX103X}")
 message(STATUS "MIOPEN_TEST_GFX110X ${MIOPEN_TEST_GFX110X}")
 message(STATUS "MIOPEN_TEST_GFX115X ${MIOPEN_TEST_GFX115X}")
 message(STATUS "MIOPEN_TEST_GFX120X ${MIOPEN_TEST_GFX120X}")
+message(STATUS "MIOPEN_TEST_GFX125X ${MIOPEN_TEST_GFX125X}")
 message(STATUS "MIOPEN_TEST_GPU_XNACK_ENABLED ${MIOPEN_TEST_GPU_XNACK_ENABLED}")
 message(STATUS "MIOPEN_TEST_GPU_DETECTION_FAILED ${MIOPEN_TEST_GPU_DETECTION_FAILED}")
 message(STATUS "MIOPEN_TEST_WITH_MIOPENDRIVER ${MIOPEN_TEST_WITH_MIOPENDRIVER}")