From 2b24dd5681f668689a869ce451da258f95ecdcf1 Mon Sep 17 00:00:00 2001
From: Artem Tamazov <artem.tamazov@gmail.com>
Date: Tue, 30 Apr 2024 00:51:16 +0300
Subject: [PATCH] Adjustments for the latest assembler (e.g. latest changes in
 the upstream clang) (#2891)

* gcnasm-noxnack-etc(01) Remove -mxnack/mno-xnack from COMgr assembler

* gcnasm-noxnack-etc(02) Added WORKAROUND_ROCMCOMPILERSUPPORT_ISSUE_67 for the "-nogpulib" warning during assembly via COMgr

* gcnasm-noxnack-etc(03) Removed "-mno-xnack" from the offline (clang) amdgcn assembly path.
---
 src/comgr.cpp               |  8 ++++++--
 src/hip/handlehip.cpp       | 14 ++++++++++++--
 src/hip/hip_build_utils.cpp |  1 -
 src/ocl/gcn_asm_utils.cpp   |  2 --
 4 files changed, 18 insertions(+), 7 deletions(-)

diff --git a/src/comgr.cpp b/src/comgr.cpp
index 4a743aee52..4a9c985273 100644
--- a/src/comgr.cpp
+++ b/src/comgr.cpp
@@ -56,6 +56,9 @@
 /// More info at https://github.com/ROCm/MIOpen/issues/1257.
 #define WORKAROUND_ISSUE_1257 (HIP_PACKAGE_VERSION_FLAT >= 4003021331ULL)
 
+/// https://github.com/ROCm/ROCm-CompilerSupport/issues/67 about unused -nogpulib.
+#define WORKAROUND_ROCMCOMPILERSUPPORT_ISSUE_67 1
+
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_LOG_CALLS)
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMGR_LOG_SOURCE_NAMES)
 
@@ -966,9 +969,10 @@ void BuildAsm(const std::string& name,
         SetIsaName(action, target);
         action.SetLogging(true);
         auto optAsm = miopen::SplitSpaceSeparated(options);
-        if(target.Xnack() && !*target.Xnack())
-            optAsm.emplace_back("-mno-xnack");
         compiler::lc::gcnasm::RemoveOptionsUnwanted(optAsm);
+#if WORKAROUND_ROCMCOMPILERSUPPORT_ISSUE_67
+        optAsm.push_back("--rocm-path=.");
+#endif
         action.SetOptionList(optAsm);
 
         const Dataset relocatable;
diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp
index a43e102c5c..d569ec8ae3 100644
--- a/src/hip/handlehip.cpp
+++ b/src/hip/handlehip.cpp
@@ -31,6 +31,7 @@
 #include <miopen/env.hpp>
 #include <miopen/errors.hpp>
 #include <miopen/handle_lock.hpp>
+#include <miopen/hip_build_utils.hpp>
 #include <miopen/invoker.hpp>
 #include <miopen/kernel_cache.hpp>
 #include <miopen/logger.hpp>
@@ -493,8 +494,17 @@ Program Handle::LoadProgram(const std::string& program_name,
 
     std::string orig_params = params; // make a copy for target ID fallback
 
-    if(!miopen::EndsWith(program_name, ".mlir"))
-        params = params + " -mcpu=" + this->GetTargetProperties().Name();
+    if(miopen::EndsWith(program_name, ".mlir"))
+    { // no -mcpu
+    }
+    else if(miopen::EndsWith(program_name, ".s"))
+    {
+        params += " -mcpu=" + LcOptionTargetStrings{this->GetTargetProperties()}.targetId;
+    }
+    else
+    {
+        params += " -mcpu=" + this->GetTargetProperties().Name();
+    }
 
     auto hsaco = miopen::LoadBinary(
         this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
diff --git a/src/hip/hip_build_utils.cpp b/src/hip/hip_build_utils.cpp
index d0fcacb84c..f11c4f5351 100644
--- a/src/hip/hip_build_utils.cpp
+++ b/src/hip/hip_build_utils.cpp
@@ -131,7 +131,6 @@ static fs::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
     src += "\nint main() {}\n";
     WriteFile(src, tmp_dir->path / filename);
 
-    // cppcheck-suppress unreadVariable
     const LcOptionTargetStrings lots(target);
 
     if(params.find("-std=") == std::string::npos)
diff --git a/src/ocl/gcn_asm_utils.cpp b/src/ocl/gcn_asm_utils.cpp
index 42269fc911..f31321ec7c 100644
--- a/src/ocl/gcn_asm_utils.cpp
+++ b/src/ocl/gcn_asm_utils.cpp
@@ -182,8 +182,6 @@ std::string AmdgcnAssemble(const std::string& source,
 
     std::ostringstream options;
     options << " -x assembler -target amdgcn--amdhsa";
-    if(target.Xnack() && !*target.Xnack())
-        options << " -mno-xnack";
     /// \todo Hacky way to find out which CO version we need to assemble for.
     if(params.find("ROCM_METADATA_VERSION=5", 0) == std::string::npos) // Assume that !COv3 == COv2.
         if(GcnAssemblerSupportsNoCOv3()) // If assembling for COv2, then disable COv3.