[gfx90a] explicity disable asm solvers

src/ocl/batchnormocl.cpp

@@ -44,6 +44,8 @@
 
 #include <chrono>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 namespace miopen {
 
 /// Reusing the dummy instance of of the ConvolutionContext class
@@ -596,8 +598,11 @@ void BatchNormBackward(Handle& handle,
                 if((n > 64) && (n % 2 == 0) && (variant == 3) && (bfpmixparm) && (useSaved) &&
                    ctx.use_asm_kernels && ctx.rmv.IsV2orV3() &&
                    (StartsWith(handle.GetDeviceName(), "gfx8") ||
-                    (StartsWith(handle.GetDeviceName(), "gfx9") &&
-                     (handle.GetDeviceName() != "gfx90a"))) &&
+                    (StartsWith(handle.GetDeviceName(), "gfx9")
+#if WORKAROUND_ISSUE_1146
+                     && (handle.GetDeviceName() != "gfx90a")
+#endif
+                         )) &&
                    (!handle.GetTargetProperties().Xnack() ||
                     !*handle.GetTargetProperties().Xnack()))
                 {

src/solver/conv_MP_bidirectional_winograd.cpp

@@ -42,6 +42,7 @@
 #if MIOPEN_BACKEND_HIP
 
 #define WORKAROUND_SWDEV_203031 1 // See also issues #2075, #2067
+#define WORKAROUND_ISSUE_1146 1   // check asm solver applicability for gfx90a
 #endif
 
 #define WORKAROUND_SWDEV_257202 1 // For SSD convergence issue.
@@ -193,8 +194,12 @@ inline bool IsApplicableTransform(const ConvolutionContext& params)
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
-    if(!StartsWith(name, "gfx9") || name == "gfx90a")
+    if(!StartsWith(name, "gfx9"))
         return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
 
     {
         std::size_t limit = miopen::Value(MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_WORKSPACE_MAX{});

src/solver/conv_asm_3x3u.cpp

@@ -40,6 +40,8 @@
 #include <cassert>
 #include <tuple>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U_PERF_VALS)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_3X3U)
 
@@ -191,8 +193,12 @@ bool ConvAsm3x3U::IsApplicable(const ConvolutionContext& params) const
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
-    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")) || name == "gfx90a")
+    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")))
+        return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
         return false;
+#endif
     if(!params.IsLayoutDefault())
     {
         return false;

src/solver/conv_asm_5x10u2v2b1.cpp

@@ -29,6 +29,8 @@
 #include <miopen/env.hpp>
 #include <miopen/conv/invokers/gen_x_w_y_pad.hpp>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2)
 
 namespace miopen {
@@ -55,6 +57,10 @@ bool ConvAsm5x10u2v2b1::IsApplicable(const ConvolutionContext& params) const
     const bool device_is_gfx8_9_no_xnack =
         (name == "gfx800" || name == "gfx802" || name == "gfx803" || name == "gfx804" ||
          name == "gfx900" || name == "gfx904" || name == "gfx906" || name == "gfx908");
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
     if(!device_is_gfx8_9_no_xnack)
     {
         return false;

src/solver/conv_asm_5x10u2v2f1.cpp

@@ -30,6 +30,8 @@
 #include <miopen/env.hpp>
 #include <miopen/conv/invokers/gen_x_w_y_pad.hpp>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_5X10U2V2)
 
 namespace miopen {
@@ -56,6 +58,10 @@ bool ConvAsm5x10u2v2f1::IsApplicable(const ConvolutionContext& params) const
     const bool device_is_gfx8_9_no_xnack =
         (name == "gfx800" || name == "gfx802" || name == "gfx803" || name == "gfx804" ||
          name == "gfx900" || name == "gfx904" || name == "gfx906" || name == "gfx908");
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
     if(!device_is_gfx8_9_no_xnack)
     {
         return false;

src/solver/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp

@@ -30,6 +30,8 @@
 #include <miopen/env.hpp>
 #include <miopen/conv/invokers/gen_x_w_y_pad.hpp>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_7X7C3H224W224)
 
 namespace miopen {
@@ -53,6 +55,10 @@ bool ConvAsm7x7c3h224w224k64u2v2p3q3f1::IsApplicable(const ConvolutionContext& p
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
     if(!(name == "gfx800" || name == "gfx802" || name == "gfx803" || name == "gfx804" ||
          name == "gfx900" || name == "gfx904" || name == "gfx906" || name == "gfx908"))
     {

src/solver/conv_asm_dir_BwdWrW3x3.cpp

@@ -38,6 +38,7 @@
 #include <miopen/generic_search.hpp>
 
 #define WORKAROUND_ISSUE_532 1 // ConvAsmBwdWrW3x3 has precision issues with some PerformanceConfigs
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
 #define MIOPEN_GCN_ASM_DIRECT_3X3WRW_SEARCH_LWC_FIXED 0
 
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW3X3_PERF_VALS)
@@ -357,8 +358,12 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ConvolutionContext& params) const
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
-    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")) || name == "gfx90a")
+    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")))
         return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
     if(!params.IsLayoutDefault())
     {
         return false;

src/solver/conv_bin_wino3x3U.cpp

@@ -35,6 +35,8 @@
 
 #include <boost/any.hpp>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_WINOGRAD_3X3)
 
 namespace miopen {
@@ -58,6 +60,10 @@ bool ConvBinWinograd3x3U::IsApplicable(const ConvolutionContext& params) const
     const auto name = params.GetStream().GetDeviceName();
     if(!(name == "gfx803" || name == "gfx900" || name == "gfx906" || name == "gfx908"))
         return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
 
     // Check if kernel is suitable for the problem description
     // and able to correctly run with given parameters.

src/solver/conv_multipass_wino3x3WrW.cpp

@@ -37,6 +37,7 @@
 #if(MIOPEN_BACKEND_HIP && (MIOPEN_USE_ROCBLAS || MIOPEN_USE_MIOPENTENSILE))
 #define WORKAROUND_SWDEV_203031 1 // See also issues #2075, #2067
 #define WORKAROUND_SWDEV_234193 1
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
 #endif
 
 namespace miopen {
@@ -440,8 +441,12 @@ bool ConvWinograd3x3MultipassWrW<WinoDataH, WinoFilterH, WinoDataW, WinoFilterW>
          FilterTransform<WinoDataH, WinoFilterH, WinoDataW, WinoFilterW>::IsApplicable(params)))
         return false;
 
-    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")) || name == "gfx90a")
+    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")))
         return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
+        return false;
+#endif
 
     {
         std::size_t limit = miopen::Value(MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_WORKSPACE_MAX{});

src/solver/conv_winoRxS_f3x2.cpp

@@ -38,6 +38,8 @@
 
 #include <boost/any.hpp>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F3X2_PERF_VALS)
 
@@ -323,8 +325,13 @@ bool ConvBinWinogradRxSf3x2::IsApplicable(const ConvolutionContext& params) cons
         return false;
 
     const auto name = params.GetStream().GetDeviceName();
-    if(!(StartsWith(name, "gfx9") || StartsWith(name, "gfx10")) || name == "gfx90a")
+    if(!(StartsWith(name, "gfx9") || StartsWith(name, "gfx10")))
+        return false;
+#if WORKAROUND_ISSUE_1146
+    if(name == "gfx90a")
         return false;
+#endif
+
     if(params.IsFp16() &&
        !(StartsWith(name, "gfx906") || StartsWith(name, "gfx908") || StartsWith(name, "gfx1011") ||
          StartsWith(name, "gfx1012") || StartsWith(name, "gfx103")))