[HOTFIX] Adapt to changes in HIP Mainline 417 (possibly future 6.1 RC) (#2652)

* fix-rocm61rc417(01) Disable new kernel build warnings. [NFC] Sort headers properly.

* fix-rocm61rc417(02) [ROCm 6.1][HIPRTC] Use custom implementations instead of standard <limits>. This fixes build issues with ROCm 6.1.

* fix-rocm61rc417(03) [ROCm 6.1][HIPRTC][Bugfix] Fixed issue in miopen_limits.h that prevented the use of custom implementations.

* fix-rocm61rc417(04) [ROCm 6.1 RC][HIPRTC] Disable some of the custom implementations from <type_traits> (like `integral_constant`) for HIP mainline 417. This fixes some build issues.

* fix-rocm61rc417(05) [ROCm 6.1 RC][offline compiler] Removed "-mcpu" from build options. This resolves kernel build issues with HIP mainline 417 (offline compiler). Improved diagnostic messages output onto console after offline build failures.

* fix-rocm61rc417(06) [tests] Disable some testcase from handle_test as #2600 still persists in Hip Mainline 417.

---------

Co-authored-by: Jun Liu <[email protected]>

Author: atamazov

src/comgr.cpp

@@ -1302,10 +1302,9 @@ void BuildHip(const std::string& name,
 #endif
         opts.push_back("-DHIP_PACKAGE_VERSION_FLAT=" + std::to_string(HIP_PACKAGE_VERSION_FLAT));
         opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS");
-        /// For now, use only standard <limits> to avoid possibility of
-        /// correctnes or performance regressions.
-        /// \todo Test and enable "custom" local implementation.
+#if HIP_PACKAGE_VERSION_FLAT < 6001024000ULL
         opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1");
+#endif
 #if WORKAROUND_ISSUE_1431
         if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) &&
            !miopen::comgr::IsWave64Enforced(opts))

src/hip/hip_build_utils.cpp

@@ -75,6 +75,20 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
     if(params.find("-std=") == std::string::npos)
         params += " --std=c++17";
 
+#if HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL
+    size_t pos = 0;
+    while((pos = params.find("-mcpu=", pos)) != std::string::npos)
+    {
+        size_t endpos = params.find(' ', pos);
+        if(endpos == std::string::npos)
+        {
+            params.erase(pos, std::string::npos);
+            break;
+        }
+        params.erase(pos, endpos - pos);
+    }
+#endif
+
 #if HIP_PACKAGE_VERSION_FLAT < 4001000000ULL
     params += " --cuda-gpu-arch=" + lots.device;
 #else
@@ -110,11 +124,14 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
     auto bin_file = tmp_dir->path / (filename + ".o");
 
     // compile
-    const std::string redirector = testing_mode ? " 1>/dev/null 2>&1" : "";
-    tmp_dir->Execute(env + std::string(" ") + MIOPEN_HIP_COMPILER,
-                     params + filename + " -o " + bin_file.string() + redirector);
-    if(!boost::filesystem::exists(bin_file))
-        MIOPEN_THROW(filename + " failed to compile");
+    {
+        const std::string redirector = testing_mode ? " 1>/dev/null 2>&1" : "";
+        const std::string cmd        = env + std::string(" ") + MIOPEN_HIP_COMPILER;
+        const std::string args       = params + filename + " -o " + bin_file.string() + redirector;
+        tmp_dir->Execute(cmd, args);
+        if(!boost::filesystem::exists(bin_file))
+            MIOPEN_THROW("Failed cmd: '" + cmd + "', args: '" + args + '\'');
+    }
 
 #if defined(MIOPEN_OFFLOADBUNDLER_BIN) && !MIOPEN_BACKEND_HIP
     // Unbundling is not required for HIP runtime && hip-clang

src/kernel_warnings.cpp

@@ -23,10 +23,11 @@
  * SOFTWARE.
  *
  *******************************************************************************/
-#include <iterator>
 #include <miopen/config.h>
 #include <miopen/kernel_warnings.hpp>
 #include <miopen/stringutils.hpp>
+
+#include <iterator>
 #include <numeric>
 #include <sstream>
 
@@ -48,6 +49,9 @@ static std::vector<std::string> OclKernelWarnings()
         "-Wno-shorten-64-to-32",
         "-Wno-sign-compare",
         "-Wno-sign-conversion",
+#if HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL
+        "-Wno-unsafe-buffer-usage",
+#endif
         "-Wno-unused-function",
         "-Wno-unused-macros",
         "-Wno-declaration-after-statement", // W/A for SWDEV-337356
@@ -58,7 +62,7 @@ static std::vector<std::string> OclKernelWarnings()
 
 static std::vector<std::string> HipKernelWarnings()
 {
-    return {
+    std::vector<std::string> rv = {
         "-Weverything",
         "-Wno-c++98-compat",
         "-Wno-c++98-compat-pedantic",
@@ -83,7 +87,12 @@ static std::vector<std::string> HipKernelWarnings()
         "-Wno-covered-switch-default",
         "-Wno-disabled-macro-expansion",
         "-Wno-undefined-reinterpret-cast",
+#if HIP_PACKAGE_VERSION_FLAT >= 6001024000ULL
+        "-Wno-unsafe-buffer-usage",
+#endif
     };
+
+    return rv;
 }
 
 static std::string MakeKernelWarningsString(const std::vector<std::string>& kernel_warnings,

src/kernels/miopen_limits.hpp

@@ -25,8 +25,8 @@
  *******************************************************************************/
 #pragma once
 
-#ifndef WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS
-#define WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS 0
+#ifndef WORKAROUND_DONT_USE_CUSTOM_LIMITS
+#define WORKAROUND_DONT_USE_CUSTOM_LIMITS 0
 #endif
 
 #if defined(MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS) && !WORKAROUND_DONT_USE_CUSTOM_LIMITS

src/kernels/miopen_type_traits.hpp

@@ -76,7 +76,7 @@ struct remove_cv
     typedef typename remove_volatile<typename remove_const<T>::type>::type type;
 };
 
-#if HIP_PACKAGE_VERSION_FLAT >= 6000024000ULL
+#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL && HIP_PACKAGE_VERSION_FLAT < 6001024000ULL
 template <class T, T v>
 struct integral_constant
 {

test/handle_test.cpp

@@ -29,12 +29,11 @@
 #define WORKAROUND_SWDEV_257056_PCH_MISSING_MACROS 1
 
 // https://gerrit-git.amd.com/c/compute/ec/clr/+/972441
-// "HIP_PACKAGE_VERSION_FLAT == 6001000000ULL" is for ROCm 6.1 RC where issue #2600 is not
-// yet fixed in the compiler. In order to test such release candidates, we have to
-// override HIP version to 6.1.0.
+// Issue #2600 is not fixed in 6.0 and still persists in 6.1 release candidates.
+// We are expecting it to be fixed in 6.1 RC after week 4 in 2024.
 #define WORKAROUND_ISSUE_2600                                                                    \
     ((HIP_PACKAGE_VERSION_FLAT >= 6000000000ULL && HIP_PACKAGE_VERSION_FLAT <= 6000999999ULL) || \
-     HIP_PACKAGE_VERSION_FLAT == 6001000000ULL)
+     (HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL && HIP_PACKAGE_VERSION_FLAT <= 6001024049ULL))
 
 #include <miopen/config.h>
 #include <miopen/handle.hpp>