[gfx8] Enable assembly kernels (#2334)

* asm-gfx8-enable(01) Enable gfx8 asm kernels.

* asm-gfx8-enable(02) Enable gfx8 for 3x3 Fwd/Bwd and 3x3 WrW.

Co-authored-by: Daniel Lowell <daniel.lowell@amd.com>

Author: atamazov

src/mlo_dir_conv.cpp

@@ -462,10 +462,4 @@ void miopen::ConvolutionContext::DetectRocm()
         use_binaries = !miopen::IsDisabled(MIOPEN_DEBUG_AMD_ROCM_PRECOMPILED_BINARIES{});
 #endif
     }
-
-    if(StartsWith(GetStream().GetDeviceName(), "gfx8"))
-    {
-        use_asm_kernels = false;
-        use_binaries    = false;
-    }
 }

src/solver/conv_asm_3x3u.cpp

@@ -186,12 +186,9 @@ bool ConvAsm3x3U::IsApplicable(const ConvolutionContext& params) const
         return false;
     if(!params.rmv.IsV2orV3())
         return false;
-
     const std::string name = params.GetStream().GetDeviceName();
-    if(name.find("gfx9") == std::string::npos)
-    {
+    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")))
         return false;
-    }
     assert(params.weights_layout.length() == 0); // FIXME _weights_layout is not supported yet.
     // clang-format off
     return params.pad_w == 1

src/solver/conv_asm_dir_BwdWrW3x3.cpp

@@ -344,12 +344,9 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ConvolutionContext& params) const
         return false;
     if(!params.rmv.IsV2orV3())
         return false;
-
     const std::string name = params.GetStream().GetDeviceName();
-    if(name.find("gfx9") == std::string::npos)
-    {
+    if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx9")))
         return false;
-    }
     assert(params.weights_layout.length() == 0); // _weights_layout is not supported yet
     // clang-format off
     bool ok = params.pad_w == 1           // -q  pad_w
@@ -364,16 +361,12 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ConvolutionContext& params) const
         && (params.IsFp32() || params.IsFp16())
         && params.in_layout == "NCHW";
     if(!ok)
-    {
         return false; // Early exit to speed up the check.
-    }
 
     if(params.IsFp16()
-          && (name.find("gfx8") != std::string::npos // Not supported.
+          && (StartsWith(name, "gfx8") // Not supported.
              || params.batch_sz % 2 != 0)) /// \todo Initial version.
-    {
        return false;
-    }
 
     // Check limits:
     const auto h_w     = static_cast<long>(params.out_height) * params.out_width;
@@ -401,7 +394,6 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ConvolutionContext& params) const
          && n_c_h_w < std::pow(2, 29)
          && n_k_h_w < std::pow(2, 29)
          && c_k_r_s < std::pow(2, 29);              // clang-format on
-
     return ok;
 }