From 273aaba758b4ffef95f86cf58cc38ca55340f565 Mon Sep 17 00:00:00 2001 From: Artem Tamazov Date: Fri, 5 Jan 2024 08:25:55 +0300 Subject: [PATCH] [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 . 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 (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 --- src/comgr.cpp | 5 ++--- src/hip/hip_build_utils.cpp | 27 ++++++++++++++++++++++----- src/kernel_warnings.cpp | 13 +++++++++++-- src/kernels/miopen_limits.hpp | 4 ++-- src/kernels/miopen_type_traits.hpp | 2 +- test/handle_test.cpp | 7 +++---- 6 files changed, 41 insertions(+), 17 deletions(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index 971a60391e..4f99fd62d0 100644 --- a/src/comgr.cpp +++ b/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 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)) diff --git a/src/hip/hip_build_utils.cpp b/src/hip/hip_build_utils.cpp index 552618f1ae..2e52be8e6d 100644 --- a/src/hip/hip_build_utils.cpp +++ b/src/hip/hip_build_utils.cpp @@ -75,6 +75,20 @@ static boost::filesystem::path HipBuildImpl(boost::optional& 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& 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 diff --git a/src/kernel_warnings.cpp b/src/kernel_warnings.cpp index 5a9c67b37b..8da966e82c 100644 --- a/src/kernel_warnings.cpp +++ b/src/kernel_warnings.cpp @@ -23,10 +23,11 @@ * SOFTWARE. * *******************************************************************************/ -#include #include #include #include + +#include #include #include @@ -48,6 +49,9 @@ static std::vector 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 OclKernelWarnings() static std::vector HipKernelWarnings() { - return { + std::vector rv = { "-Weverything", "-Wno-c++98-compat", "-Wno-c++98-compat-pedantic", @@ -83,7 +87,12 @@ static std::vector 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& kernel_warnings, diff --git a/src/kernels/miopen_limits.hpp b/src/kernels/miopen_limits.hpp index 2a8f5e6178..bd1e27415b 100644 --- a/src/kernels/miopen_limits.hpp +++ b/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 diff --git a/src/kernels/miopen_type_traits.hpp b/src/kernels/miopen_type_traits.hpp index b9cf11cd63..f8d15f0467 100644 --- a/src/kernels/miopen_type_traits.hpp +++ b/src/kernels/miopen_type_traits.hpp @@ -76,7 +76,7 @@ struct remove_cv typedef typename remove_volatile::type>::type type; }; -#if HIP_PACKAGE_VERSION_FLAT >= 6000024000ULL +#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL && HIP_PACKAGE_VERSION_FLAT < 6001024000ULL template struct integral_constant { diff --git a/test/handle_test.cpp b/test/handle_test.cpp index ba143d7ad1..16ce78f043 100644 --- a/test/handle_test.cpp +++ b/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 #include