[gfx90a] explicity disable asm solvers (#1150)

* explicity disable asm solvers for gfx90a * revert unnecessary gfx90a restrictions * fix ConvBinWinogradRxS formatting
ROCm · Sep 30, 2021 · c9b674b · c9b674b
1 parent 5576c06
commit c9b674b
Show file tree

Hide file tree

Showing 10 changed files with 64 additions and 7 deletions.
diff --git a/src/ocl/batchnormocl.cpp b/src/ocl/batchnormocl.cpp
@@ -45,6 +45,8 @@
 #define WORKAROUND_SWDEV_253606 1
 #include <chrono>
 
+#define WORKAROUND_ISSUE_1146 1 // check asm solver applicability for gfx90a
+
 namespace miopen {
 
 /// Reusing the dummy instance of of the ConvolutionContext class
@@ -989,8 +991,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()))
                 {

diff --git a/src/solver/conv_MP_bidirectional_winograd.cpp b/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{});

diff --git a/src/solver/conv_asm_3x3u.cpp b/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;

diff --git a/src/solver/conv_asm_5x10u2v2b1.cpp b/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;

diff --git a/src/solver/conv_asm_5x10u2v2f1.cpp b/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;

diff --git a/src/solver/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp b/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"))
     {

diff --git a/src/solver/conv_asm_dir_BwdWrW3x3.cpp b/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;

diff --git a/src/solver/conv_bin_wino3x3U.cpp b/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.

diff --git a/src/solver/conv_multipass_wino3x3WrW.cpp b/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{});

diff --git a/src/solver/conv_winoRxS_f3x2.cpp b/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")))