Merge pull request #715 from DrizztDoUrden/fix-rocm17rc-hip-md-version

Detect MD version for ROCm 1.7 RCx && HIP backend.
ROCm · Dec 18, 2017 · a9949e3 · a9949e3
2 parents f532530 + b4c2af3
commit a9949e3
Show file tree

Hide file tree

Showing 4 changed files with 26 additions and 14 deletions.
diff --git 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
@@ -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
@@ -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
@@ -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;