From ce51a4c474541793a8dc3d08a2416e81e3d4d9dd Mon Sep 17 00:00:00 2001 From: zjing14 Date: Sat, 31 Aug 2019 08:44:24 -0500 Subject: [PATCH] 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 --- src/mlo_dir_conv.cpp | 6 ++++++ src/solver/conv_asm_3x3u.cpp | 2 +- src/solver/conv_asm_dir_BwdWrW3x3.cpp | 2 +- 3 files changed, 8 insertions(+), 2 deletions(-) diff --git a/src/mlo_dir_conv.cpp b/src/mlo_dir_conv.cpp index 5613e23682..b8ea8ee33e 100644 --- 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 index 5b2e622618..993b38971b 100644 --- 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 index b97b29eee4..ce86295949 100644 --- 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; }