disabled the asm_dir_BwdWrW3x3 kernel on Fiji GPUs due to unsupported…

… ISA (#2035) * disable_asm_dir_BwdWrW3x3_on_fiji(1): disable the conv_asm_dir_BwdWrW3x3 kernel on fiji due to unsupported isa * disable conv3x3 on gfx8 * disable_asm_dir_BwdWrW3x3_on_fiji(3): disable all asm kernels on fiji * disable_asm_dir_BwdWrW3x3_on_fiji(4): disabled all binaries on fiji
ROCm · Aug 31, 2019 · ce51a4c · ce51a4c
1 parent a5c6e1a
commit ce51a4c
Show file tree

Hide file tree

Showing 3 changed files with 8 additions and 2 deletions.
diff --git a/src/mlo_dir_conv.cpp b/src/mlo_dir_conv.cpp
@@ -383,4 +383,10 @@ 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;
+    }
 }
diff --git a/src/solver/conv_asm_3x3u.cpp b/src/solver/conv_asm_3x3u.cpp
@@ -185,7 +185,7 @@ bool ConvAsm3x3U::IsApplicable(const ConvolutionContext& params) const
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
-    if(name.find("gfx8") == std::string::npos && name.find("gfx9") == std::string::npos)
+    if(name.find("gfx9") == std::string::npos)
     {
         return false;
     }

diff --git a/src/solver/conv_asm_dir_BwdWrW3x3.cpp b/src/solver/conv_asm_dir_BwdWrW3x3.cpp
@@ -343,7 +343,7 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ConvolutionContext& params) const
         return false;
 
     const std::string name = params.GetStream().GetDeviceName();
-    if(name.find("gfx8") == std::string::npos && name.find("gfx9") == std::string::npos)
+    if(name.find("gfx9") == std::string::npos)
     {
         return false;
     }