Workaround for hcc-clang version detection (#2517)

* hcc-clang-version-detect(01) Formalize workaround for SWDEV-200782 (NFC) * hcc-clang-version-detect(02) Formatting * hcc-clang-version-detect(03) W/A for issue 2514. Assume version 3.2.0 for HIP compiler that is not hcc.
ROCm · Apr 1, 2020 · 82b61f1 · 82b61f1
1 parent de8c3ce
commit 82b61f1
Show file tree

Hide file tree

Showing 2 changed files with 29 additions and 5 deletions.
diff --git a/src/hip/hip_build_utils.cpp b/src/hip/hip_build_utils.cpp
@@ -35,6 +35,8 @@
 
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_HIP_ENFORCE_COV3)
 
+#define WORKAROUND_ISSUE_2514 1
+
 namespace miopen {
 
 namespace {
@@ -172,6 +174,11 @@ static external_tool_version_t HipGetHccVersionImpl()
         if(miopen::exec::Run(path + " --version", nullptr, &out) != 0)
             break;
 
+#if WORKAROUND_ISSUE_2514
+        // If compiler is not hcc and mandatory prefix is not found,
+        // then assume hip-clang 3.2.0.
+        bool mandatory_prefix_found = false;
+#endif
         std::string line;
         while(!out.eof())
         {
@@ -181,6 +188,9 @@ static external_tool_version_t HipGetHccVersionImpl()
             if(begin == std::string::npos)
                 continue;
 
+#if WORKAROUND_ISSUE_2514
+            mandatory_prefix_found = true;
+#endif
             begin += mandatory_prefix.size();
             int v3, v2, v1 = v2 = v3 = -1;
             char c2, c1 = c2 = 'X';
@@ -198,6 +208,15 @@ static external_tool_version_t HipGetHccVersionImpl()
             }
             break;
         }
+#if WORKAROUND_ISSUE_2514
+        if(!mandatory_prefix_found && !IsHccCompiler())
+        {
+            MIOPEN_LOG_NQI2("Assuming 3.2.0 (hip-clang?)");
+            hcc_version.major = 3;
+            hcc_version.minor = 2;
+            hcc_version.patch = 0;
+        }
+#endif
     } while(false);
     MIOPEN_LOG_NQI("HCC base: " << hcc_version.major << '.' << hcc_version.minor << '.'
                                 << hcc_version.patch);

diff --git a/src/solver/implicitgemm_util.hpp b/src/solver/implicitgemm_util.hpp
@@ -13,6 +13,8 @@ MIOPEN_DECLARE_ENV_VAR(
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_IMPLICIT_GEMM_NON_XDLOPS_INLINE_ASM)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_IMPLICIT_GEMM_XDLOPS_INLINE_ASM)
 
+#define WORKAROUND_SWDEV_200782 1
+
 namespace miopen {
 namespace solver {
 
@@ -331,15 +333,18 @@ static inline bool IsXdlopsSupport(const ConvolutionContext& c)
     if(miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS_EMULATE{}))
         return true;
 
+    // disable xdlops kernels by default due to possible failures:
+    // 1) inline asm may crash
+    // 2) llvm intrin may has incorrect results
     return StartsWith(c.GetStream().GetDeviceName(), "gfx908") &&
-           // disable xdlops kernels by default due to possible failures:
-           // 1) inline asm may crash
-           // 2) llvm intrin may has incorrect results
-           /// \todo enable xdlops kernels by default after llvm intrin fix (SWDEV-200782) in
-           /// release
+#if WORKAROUND_SWDEV_200782
+           /// \todo Remove workaround when we drop suport of HCC older than 2.10.19392.
            ((miopen::HipGetHccVersion() >= external_tool_version_t{2, 10, 19392})
                 ? !miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS{})
                 : miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS{}));
+#else
+           !miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS{});
+#endif
 }
 
 inline static uint32_t GetReadWriteVectorSize(const int v)