Allow targetid for mcpu kernel arguments

projects/miopen/src/comgr.cpp

@@ -640,10 +640,6 @@ void BuildAsm(const std::string& name,
         SetIsaName(action, target);
         action.SetLogging(true);
         auto optAsm = miopen::SplitSpaceSeparated(options);
-#if WORKAROUND_ISSUE_3001
-        if(target.Xnack() && !*target.Xnack())
-            optAsm.emplace_back("-mno-xnack");
-#endif
         compiler::lc::gcnasm::RemoveOptionsUnwanted(optAsm);
 #if WORKAROUND_ROCMCOMPILERSUPPORT_ISSUE_67
         optAsm.push_back("--rocm-path=.");

projects/miopen/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>
@@ -544,10 +545,6 @@ Program Handle::LoadProgram(const fs::path& program_name,
 
     std::string orig_params = params; // make a copy for target ID fallback
 
-#if WORKAROUND_ISSUE_3001
-    if(program_name.extension() != ".mlir")
-        params = params + " -mcpu=" + this->GetTargetProperties().Name();
-#else
     if(program_name.extension() == ".mlir")
     { // no -mcpu
     }
@@ -559,7 +556,6 @@ Program Handle::LoadProgram(const fs::path& program_name,
     {
         params += " -mcpu=" + this->GetTargetProperties().Name();
     }
-#endif
 
     auto hsaco = miopen::LoadBinary(
         this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);

projects/miopen/src/include/miopen/binary_cache.hpp

@@ -41,17 +41,7 @@ GetCacheFile(const std::string& device, const fs::path& name, const std::string&
 
 MIOPEN_INTERNALS_EXPORT fs::path GetCachePath(bool is_system);
 
-#if !MIOPEN_ENABLE_SQLITE_KERN_CACHE
-fs::path LoadBinary(const TargetProperties& target,
-                    std::size_t num_cu,
-                    const fs::path& name,
-                    const std::string& args);
-
-fs::path SaveBinary(const fs::path& binary_path,
-                    const TargetProperties& target,
-                    const fs::path& name,
-                    const std::string& args);
-#else
+#if MIOPEN_ENABLE_SQLITE_KERN_CACHE
 std::vector<char> LoadBinary(const TargetProperties& target,
                              std::size_t num_cu,
                              const fs::path& name,
@@ -62,6 +52,16 @@ void SaveBinary(const std::vector<char>& hsaco,
                 std::size_t num_cu,
                 const fs::path& name,
                 const std::string& args);
+#else
+fs::path LoadBinary(const TargetProperties& target,
+                    std::size_t num_cu,
+                    const fs::path& name,
+                    const std::string& args);
+
+fs::path SaveBinary(const fs::path& binary_path,
+                    const TargetProperties& target,
+                    const fs::path& name,
+                    const std::string& args);
 #endif
 
 } // namespace miopen

projects/miopen/src/include/miopen/target_properties.hpp

@@ -29,8 +29,6 @@
 #include <boost/optional.hpp>
 #include <string>
 
-#define WORKAROUND_ISSUE_3001 1
-
 namespace miopen {
 
 struct Handle;

projects/miopen/src/ocl/gcn_asm_utils.cpp

@@ -178,10 +178,6 @@ std::string AmdgcnAssemble(std::string_view source,
 
     std::ostringstream options;
     options << " -x assembler -target amdgcn--amdhsa";
-#if WORKAROUND_ISSUE_3001
-    if(target.Xnack() && !*target.Xnack())
-        options << " -mno-xnack";
-#endif
     /// \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.

projects/miopen/test/gtest/db_sync.cpp

@@ -35,6 +35,7 @@
 #include <miopen/execution_context.hpp>
 
 #include <miopen/find_db.hpp>
+#include <miopen/hip_build_utils.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/conv/problem_description.hpp>
 #include <miopen/conv_algo_name.hpp>
@@ -434,6 +435,7 @@ auto LoadKDBObjects(const fs::path& filename)
             const auto kernel_name = stmt.ColumnText(0);
             const auto kernel_args = stmt.ColumnText(1);
             kdb_cache.emplace(KDBKey{kernel_name, kernel_args});
+            MIOPEN_LOG_I(kernel_name + ": " + kernel_args);
         }
         else if(rc == SQLITE_DONE)
             break;
@@ -520,11 +522,43 @@ void SetupPaths(fs::path& fdb_file_path,
         << "Db file does not exist" << kdb_file_path;
 }
 
-TEST(CPU_DBSync_NONE, KDBTargetID)
+namespace miopen {
+struct TestHandle : Handle
+{
+    TestHandle(size_t _num_cu) : Handle(), num_cu(_num_cu) {}
+
+    std::size_t GetMaxComputeUnits() const override
+    {
+        if(num_cu == 0)
+            return Handle::GetMaxComputeUnits();
+        return num_cu;
+    }
+
+    size_t num_cu = 0;
+};
+} // namespace miopen
+
+static inline miopen::TestHandle& get_test_handle(size_t num_cu)
+{
+    // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables)
+    static miopen::TestHandle h{num_cu};
+    static const std::thread::id id = std::this_thread::get_id();
+    if(std::this_thread::get_id() != id)
+    {
+        std::cout << "Cannot use handle across multiple threads\n";
+        std::abort();
+    }
+    return h;
+}
+
+void KDBTargetID(const std::string& arch, const size_t num_cu)
 {
     // Skip this test for gfx11 and gfx12 to avoid test failure (we don't have databases for those
     // devices yet)
-    const auto& handle = get_handle();
+    auto& handle = get_test_handle(num_cu);
+    if(handle.GetDeviceName() != arch)
+        GTEST_SKIP();
+    handle.num_cu = num_cu;
     if(miopen::StartsWith(handle.GetDeviceName(), "gfx11") ||
        miopen::StartsWith(handle.GetDeviceName(), "gfx12"))
     {
@@ -535,7 +569,7 @@ TEST(CPU_DBSync_NONE, KDBTargetID)
 #if WORKAROUND_ISSUE_2492
     ScopedEnvironment<std::string> issue_2492_env(MIOPEN_DEBUG_WORKAROUND_ISSUE_2492, "0");
 #endif
-    SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, get_handle());
+    SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, handle);
     std::ignore = fdb_file_path;
     std::ignore = pdb_file_path;
     EXPECT_TRUE(miopen::CheckKDBJournalMode(kdb_file_path));
@@ -565,10 +599,33 @@ void BuildKernel(const fs::path& program_file,
     try
     {
         auto p = handle.LoadProgram(program_file, program_args, "");
+
+        std::string compile_options     = program_args;
+        std::string compile_options_tid = program_args;
+        compile_options += " -mcpu=" + handle.GetDeviceName();
+        if(program_file.extension() == ".mlir") {}
+        else if(program_file.extension() == ".s")
+            compile_options_tid +=
+                " -mcpu=" + miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId;
+        else
+            compile_options_tid += " -mcpu=" + handle.GetDeviceName();
+        auto hsaco = miopen::LoadBinary(handle.GetTargetProperties(),
+                                        handle.GetMaxComputeUnits(),
+                                        program_file,
+                                        compile_options_tid);
+
+        if(!hsaco.empty())
+        {
+            miopen::SaveBinary(hsaco,
+                               handle.GetTargetProperties(),
+                               handle.GetMaxComputeUnits(),
+                               program_file,
+                               compile_options);
+        }
     }
-    catch(std::exception&)
+    catch(std::exception& e)
     {
-        MIOPEN_LOG_W("Exception thrown while building kernel");
+        MIOPEN_LOG_W("Exception thrown while building kernel, " << e.what());
     }
 #endif
 }
@@ -698,6 +755,7 @@ void CheckFDBEntry(size_t thread_index,
                    std::atomic<size_t>& counter)
 {
     fs::path fdb_file_path, pdb_file_path, kdb_file_path;
+    auto& handle = _ctx.GetStream();
     SetupPaths(fdb_file_path, pdb_file_path, kdb_file_path, _ctx.GetStream());
     std::unordered_set<KDBKey> checked_kdbs;
     const auto data_size = data.size();
@@ -842,10 +900,8 @@ void CheckFDBEntry(size_t thread_index,
                         std::string compile_options = kern.comp_options;
                         auto program_file = miopen::make_object_file_name(kern.kernel_file);
                         if(kern.kernel_file.extension() != ".mlir")
-                        {
-                            auto& handle = ctx.GetStream();
                             compile_options += " -mcpu=" + handle.GetDeviceName();
-                        }
+
                         auto search           = checked_kdbs.find({program_file, compile_options});
                         bool reported_already = search != checked_kdbs.end();
                         if(!reported_already) // we have reported this object before, no need to
@@ -866,7 +922,7 @@ void CheckFDBEntry(size_t thread_index,
                                 << compile_options; // for fdb key, solver id, solver pdb entry and
                                                     // kdb file and args
                         if(!reported_already)
-                            BuildKernel(kern.kernel_file, kern.comp_options, ctx.GetStream());
+                            BuildKernel(kern.kernel_file, kern.comp_options, handle);
                     }
                 }
             }
@@ -881,41 +937,6 @@ void CheckFDBEntry(size_t thread_index,
         counter.fetch_add(1, std::memory_order_relaxed);
     }
 }
-namespace miopen {
-struct TestHandle : Handle
-{
-    TestHandle(size_t _num_cu) : Handle(), num_cu(_num_cu) {}
-
-// Probably, according to the idea of the author of this test, the number of CUs should have been
-// substituted with the value passed to the constructor (which in fact did not happen). After
-// https://github.com/ROCm/MIOpen/pull/3175, the method became virtual, the substitution actually
-// happened, and the test broke. I disabled that part (since it doesn't work as intended anyway) to
-// keep its behavior the same.
-#if 1
-    std::size_t GetMaxComputeUnits() const override
-    {
-        if(num_cu == 0)
-            return Handle::GetMaxComputeUnits();
-        return num_cu;
-    }
-#endif
-
-    size_t num_cu = 0;
-};
-} // namespace miopen
-
-static inline miopen::TestHandle& get_test_handle(size_t num_cu)
-{
-    // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables)
-    static miopen::TestHandle h{num_cu};
-    static const std::thread::id id = std::this_thread::get_id();
-    if(std::this_thread::get_id() != id)
-    {
-        std::cout << "Cannot use handle across multiple threads\n";
-        std::abort();
-    }
-    return h;
-}
 
 void StaticFDBSync(const std::string& arch, const size_t num_cu)
 {
@@ -974,6 +995,14 @@ struct CPU_DBSync_NONE : testing::TestWithParam<std::pair<std::string, size_t>>
 {
 };
 
+TEST_P(CPU_DBSync_NONE, KDBTargetID)
+{
+    std::string arch;
+    size_t num_cu;
+    std::tie(arch, num_cu) = GetParam();
+    KDBTargetID(arch, num_cu);
+}
+
 TEST_P(CPU_DBSync_NONE, StaticFDBSync)
 {
     std::string arch;