From f230af4408374b8da331a2aefbd98db3e570ab50 Mon Sep 17 00:00:00 2001
From: BrianHarrisonAMD <169072757+BrianHarrisonAMD@users.noreply.github.com>
Date: Mon, 4 Nov 2024 10:02:13 -0700
Subject: [PATCH] Temporarily disable igemm wrw solvers for bfloat16 unless
 device is gfx94 (#3362)

---
 .../conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp     | 4 +++-
 src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp | 4 +++-
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
index f30c5db6bb..feca383423 100644
--- a/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
+++ b/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
@@ -541,7 +541,9 @@ bool ConvHipImplicitGemm3DGroupWrwXdlops::IsApplicable(
     case miopenHalf: return CheckCKApplicability<ck::half_t>(problem);
     case miopenFloat: return CheckCKApplicability<float>(problem);
     case miopenInt8: return CheckCKApplicability<int8_t>(problem);
-    case miopenBFloat16: return CheckCKApplicability<ck::bhalf_t>(problem);
+    case miopenBFloat16:
+        return StartsWith(ctx.GetStream().GetDeviceName(), "gfx94") &&
+               CheckCKApplicability<ck::bhalf_t>(problem);
     case miopenInt64:
     case miopenInt32:
     case miopenFloat8:
diff --git a/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp b/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
index ad5be773d7..d4e53d5e5d 100644
--- a/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
+++ b/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
@@ -591,7 +591,9 @@ bool ConvHipImplicitGemmGroupWrwXdlops::IsApplicable(
     case miopenHalf: return CheckCKApplicability<ck::half_t>(problem);
     case miopenFloat: return CheckCKApplicability<float>(problem);
     case miopenInt8: return CheckCKApplicability<int8_t>(problem);
-    case miopenBFloat16: return CheckCKApplicability<ck::bhalf_t>(problem);
+    case miopenBFloat16:
+        return StartsWith(ctx.GetStream().GetDeviceName(), "gfx94") &&
+               CheckCKApplicability<ck::bhalf_t>(problem);
     case miopenInt64:
     case miopenInt32:
     case miopenFloat8: