Adjustments for the latest assembler (e.g. latest changes in the upst…

…ream 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.
ROCm · Apr 29, 2024 · 2b24dd5 · 2b24dd5
1 parent 1e09879
commit 2b24dd5
Show file tree

Hide file tree

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