From b4c2af35e15b545ea43860da2085790f25fc50cc Mon Sep 17 00:00:00 2001
From: Artem Tamazov <artem.tamazov@gmail.com>
Date: Tue, 19 Dec 2017 00:28:42 +0300
Subject: [PATCH] fix-rocm17rc-hip-md-version(1) Initial implementation.

---
 src/include/miopen/gcn_asm_utils.hpp |  1 +
 src/include/miopen/mlo_internal.hpp  |  2 +-
 src/mlo_dir_conv.cpp                 | 33 +++++++++++++++++++---------
 src/ocl/gcn_asm_utils.cpp            |  4 +---
 4 files changed, 26 insertions(+), 14 deletions(-)

diff --git a/src/include/miopen/gcn_asm_utils.hpp b/src/include/miopen/gcn_asm_utils.hpp
index 9bde6bc10b..3de87c462e 100644
--- a/src/include/miopen/gcn_asm_utils.hpp
+++ b/src/include/miopen/gcn_asm_utils.hpp
@@ -33,6 +33,7 @@
 std::string GetGcnAssemblerPath();
 bool ValidateGcnAssembler();
 void AmdgcnAssemble(std::string& source, const std::string& params);
+bool GcnAssemblerHasBug34765();
 
 template <typename TValue>
 void GenerateClangDefsym(std::ostream& stream, const std::string& name, TValue value)
diff --git a/src/include/miopen/mlo_internal.hpp b/src/include/miopen/mlo_internal.hpp
index a2e3b9553f..1a6817291b 100644
--- a/src/include/miopen/mlo_internal.hpp
+++ b/src/include/miopen/mlo_internal.hpp
@@ -853,7 +853,7 @@ struct mlo_construct_direct2D
 
     std::string db_path() const { return _db_path ? _db_path : _search_params.GetPerfDbPath(); }
 
-    bool mloIsAmdRocmOpencl(rocm_meta_version& rmv) const;
+    bool mloIsAmdRocm(rocm_meta_version& rmv) const;
 
     int mloConstructBwd() { return (0); }
     int mloConstructFwd() { return (0); }
diff --git a/src/mlo_dir_conv.cpp b/src/mlo_dir_conv.cpp
index 6f310d26f0..97b5dfe8a7 100644
--- a/src/mlo_dir_conv.cpp
+++ b/src/mlo_dir_conv.cpp
@@ -63,7 +63,7 @@ void mlo_construct_direct2D::setupRocm()
     _search_params.use_binaries        = false;
     _search_params.assembler_available = false;
     _search_params.rmv                 = rocm_meta_version::Default;
-    if(mloIsAmdRocmOpencl(_search_params.rmv))
+    if(mloIsAmdRocm(_search_params.rmv))
     {
         _search_params.assembler_available =
             !miopen::IsDisabled(MIOPEN_DEBUG_GCN_ASM_KERNELS{}) && ValidateGcnAssembler();
@@ -201,6 +201,7 @@ static bool IsAmdRocmOpencl(const miopen::ConvolutionContext& context)
     const char* delimiters    = " (),*";                    // Specific for ROCm OCL driver version.
     return IsTokenWithin(driver_version, delimiters, "LC"); // Lightning Compiler.
 }
+#endif // MIOPEN_BACKEND_OPENCL
 
 static std::ostream& operator<<(std::ostream& os, const rocm_meta_version& rmv)
 {
@@ -215,8 +216,9 @@ static std::ostream& operator<<(std::ostream& os, const rocm_meta_version& rmv)
     return os << "<Error>";
 }
 
-static rocm_meta_version DetectAmdRocmOpenclVersion(const miopen::ConvolutionContext& context)
+static rocm_meta_version DetectAmdRocmMetadataVersion(const miopen::ConvolutionContext& context)
 {
+#if MIOPEN_BACKEND_OPENCL
     const auto dev                     = miopen::GetDevice(context.GetStream().GetStream());
     const auto platform                = miopen::GetDeviceInfo<CL_DEVICE_PLATFORM>(dev);
     const std::string platform_version = miopen::GetPlatformInfo<CL_PLATFORM_VERSION>(
@@ -235,25 +237,36 @@ static rocm_meta_version DetectAmdRocmOpenclVersion(const miopen::ConvolutionCon
         else
             rmv = rocm_meta_version::AMDHSA_1_0;
     }
+#else
+    /// \todo Rework this using clang-ocl.
+    (void)context;
+    rocm_meta_version rmv = rocm_meta_version::Default;
+    // Assembler is always available for HIP backend.
+    // ROCm 1.7, which uses AMDHSA_1_0 metadata, does not have bug 34765 in
+    // the assembler. Previous ROCm versions have this bug.
+    if(!GcnAssemblerHasBug34765())
+    {
+        rmv = rocm_meta_version::AMDHSA_1_0;
+    }
+#endif // MIOPEN_BACKEND_OPENCL
     MIOPEN_LOG_I(rmv);
     return rmv;
 }
-#endif // MIOPEN_BACKEND_OPENCL
 
-bool mlo_construct_direct2D::mloIsAmdRocmOpencl(rocm_meta_version& rmv) const
+bool mlo_construct_direct2D::mloIsAmdRocm(rocm_meta_version& rmv) const
 {
+    static const bool ret_bool
 #if MIOPEN_BACKEND_OPENCL
-    static const bool ret_bool = IsAmdRocmOpencl(_search_params);
+        = IsAmdRocmOpencl(_search_params);
+#else
+        = true;
+#endif // MIOPEN_BACKEND_OPENCL
     if(ret_bool)
     {
-        static const rocm_meta_version ret_rmv = DetectAmdRocmOpenclVersion(_search_params);
+        static const rocm_meta_version ret_rmv = DetectAmdRocmMetadataVersion(_search_params);
         rmv                                    = ret_rmv;
     }
     return ret_bool;
-#else
-    (void)rmv; // We don't care about metada version
-    return true;
-#endif // MIOPEN_BACKEND_OPENCL
 }
 
 bool mlo_construct_BwdWrW2D::mloIsCompilerWorkarounds() const
diff --git a/src/ocl/gcn_asm_utils.cpp b/src/ocl/gcn_asm_utils.cpp
index 43fd6dc1b6..f8bc8f6af9 100644
--- a/src/ocl/gcn_asm_utils.cpp
+++ b/src/ocl/gcn_asm_utils.cpp
@@ -247,8 +247,6 @@ static std::string CleanupPath(const char* p)
     return path;
 }
 
-static bool GcnAssemblerHasBug34765();
-
 /*
  * Temporary function which emulates online assembly feature of OpenCL-on-ROCm being developed.
  * Not intended to be used in production code, so error handling is very straghtforward,
@@ -344,7 +342,7 @@ static bool GcnAssemblerHasBug34765Impl()
     }
 }
 
-static bool GcnAssemblerHasBug34765()
+bool GcnAssemblerHasBug34765()
 {
     const static bool b = GcnAssemblerHasBug34765Impl();
     return b;