From 1998dd74d39a2be184dcc4b50a479fee540b858c Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 23 Apr 2025 00:15:27 +0000 Subject: [PATCH 1/7] db_sync add targetid --- src/include/miopen/target_properties.hpp | 2 +- test/gtest/db_sync.cpp | 28 ++++++++++++++++++++---- 2 files changed, 25 insertions(+), 5 deletions(-) diff --git a/src/include/miopen/target_properties.hpp b/src/include/miopen/target_properties.hpp index 2b8b04b413..256b8d934b 100644 --- a/src/include/miopen/target_properties.hpp +++ b/src/include/miopen/target_properties.hpp @@ -29,7 +29,7 @@ #include #include -#define WORKAROUND_ISSUE_3001 1 +#define WORKAROUND_ISSUE_3001 0 namespace miopen { diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index b1edf3b899..4b78dd8561 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -35,6 +35,7 @@ #include #include +#include #include #include #include @@ -626,7 +627,16 @@ void CheckDynamicFDBEntry(size_t thread_index, auto program_file = miopen::make_object_file_name(kern.kernel_file); ASSERT_TRUE(kern.kernel_file.extension() != ".mlir") << "MLIR detected in dynamic solvers"; - compile_options += " -mcpu=" + handle.GetDeviceName(); + compile_options += " -mcpu="; + if(miopen::EndsWith(kern.kernel_file, ".s")) + { + compile_options += + miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId; + } + else + { + compile_options += handle.GetDeviceName(); + } auto search = checked_kdbs.find({program_file, compile_options}); if(search != checked_kdbs @@ -698,6 +708,7 @@ void CheckFDBEntry(size_t thread_index, std::atomic& 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 checked_kdbs; const auto data_size = data.size(); @@ -828,8 +839,17 @@ void CheckFDBEntry(size_t thread_index, 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(); + compile_options += " -mcpu=" + if(miopen::EndsWith(kern.kernel_file, ".s")) + { + compile_options += + miopen::LcOptionTargetStrings{handle.GetTargetProperties()} + .targetId; + } + else + { + compile_options += handle.GetDeviceName(); + } } auto search = checked_kdbs.find({program_file, compile_options}); bool reported_already = search != checked_kdbs.end(); @@ -851,7 +871,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); } } } From 706c4869fe369aee34e48e93c45165cf0944b0d7 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 24 Apr 2025 20:41:29 +0000 Subject: [PATCH 2/7] db_sync saves -mcpu=arch_name copy of kernel argument list --- src/binary_cache.cpp | 2 +- src/hip/handlehip.cpp | 3 +- src/include/miopen/binary_cache.hpp | 22 ++++++------- test/gtest/db_sync.cpp | 50 ++++++++++++++--------------- 4 files changed, 39 insertions(+), 38 deletions(-) diff --git a/src/binary_cache.cpp b/src/binary_cache.cpp index 3faf994621..a2942ccb3a 100644 --- a/src/binary_cache.cpp +++ b/src/binary_cache.cpp @@ -134,7 +134,7 @@ KDb GetDb(const TargetProperties& target, size_t num_cu) if(user_dir.empty()) user_path = user_dir; if(!fs::exists(sys_path)) - sys_path = sys_dir / (target.DbId() + ".kdb"); + sys_path = sys_dir / (target.Name() + ".kdb"); #if !MIOPEN_EMBED_DB if(!fs::exists(sys_path)) sys_path = fs::path{}; diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp index c56bf27bfb..f5ee7b2ed1 100644 --- a/src/hip/handlehip.cpp +++ b/src/hip/handlehip.cpp @@ -31,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -622,7 +623,7 @@ Program Handle::LoadProgram(const fs::path& program_name, else boost::filesystem::copy_file(p.GetCodeObjectPathname(), path); cache_path = miopen::SaveBinary( - path, this->GetTargetProperties(), program_name, params, is_kernel_str); + path, this->GetTargetProperties(), program_name, params); } if(force_attach_binary && p.IsCodeObjectInTempFile()) diff --git a/src/include/miopen/binary_cache.hpp b/src/include/miopen/binary_cache.hpp index 7ef9bd297f..78c9e74e7b 100644 --- a/src/include/miopen/binary_cache.hpp +++ b/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 LoadBinary(const TargetProperties& target, std::size_t num_cu, const fs::path& name, @@ -62,6 +52,16 @@ void SaveBinary(const std::vector& 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 diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index 4b78dd8561..bb3efb3565 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -435,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; @@ -566,10 +567,29 @@ void BuildKernel(const fs::path& program_file, try { auto p = handle.LoadProgram(program_file, program_args, ""); + + if(program_file.extension() == ".s") + { + std::string compile_options = program_args; + std::string compile_options_tid = program_args; + compile_options += " -mcpu=" + handle.GetDeviceName(); + compile_options_tid += " -mcpu=" + miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId; + 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 } @@ -627,16 +647,7 @@ void CheckDynamicFDBEntry(size_t thread_index, auto program_file = miopen::make_object_file_name(kern.kernel_file); ASSERT_TRUE(kern.kernel_file.extension() != ".mlir") << "MLIR detected in dynamic solvers"; - compile_options += " -mcpu="; - if(miopen::EndsWith(kern.kernel_file, ".s")) - { - compile_options += - miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId; - } - else - { - compile_options += handle.GetDeviceName(); - } + compile_options += " -mcpu=" + handle.GetDeviceName(); auto search = checked_kdbs.find({program_file, compile_options}); if(search != checked_kdbs @@ -838,19 +849,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") - { - compile_options += " -mcpu=" - if(miopen::EndsWith(kern.kernel_file, ".s")) - { - compile_options += - miopen::LcOptionTargetStrings{handle.GetTargetProperties()} - .targetId; - } - else - { - compile_options += handle.GetDeviceName(); - } - } + 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 From 722cffaec834b34e8c31097a224c8f31a87603ec Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 24 Apr 2025 21:37:18 +0000 Subject: [PATCH 3/7] db_sync saves all kernels --- test/gtest/db_sync.cpp | 31 ++++++++++++++++--------------- 1 file changed, 16 insertions(+), 15 deletions(-) diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index bb3efb3565..d665340c5b 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -568,23 +568,24 @@ void BuildKernel(const fs::path& program_file, { auto p = handle.LoadProgram(program_file, program_args, ""); - if(program_file.extension() == ".s") - { - std::string compile_options = program_args; - std::string compile_options_tid = program_args; - compile_options += " -mcpu=" + handle.GetDeviceName(); + 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; - auto hsaco = miopen::LoadBinary( - handle.GetTargetProperties(), handle.GetMaxComputeUnits(), program_file, compile_options_tid); + 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); - } + if(!hsaco.empty()) + { + miopen::SaveBinary(hsaco, + handle.GetTargetProperties(), + handle.GetMaxComputeUnits(), + program_file, + compile_options); } } catch(std::exception& e) From 7d6b2b6b358dc1ec85b62281454d1fdebc692bb0 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 25 Apr 2025 16:44:29 +0000 Subject: [PATCH 4/7] cleanup --- src/binary_cache.cpp | 2 +- src/hip/handlehip.cpp | 4 ++-- test/gtest/db_sync.cpp | 21 ++++++++++++--------- 3 files changed, 15 insertions(+), 12 deletions(-) diff --git a/src/binary_cache.cpp b/src/binary_cache.cpp index a2942ccb3a..3faf994621 100644 --- a/src/binary_cache.cpp +++ b/src/binary_cache.cpp @@ -134,7 +134,7 @@ KDb GetDb(const TargetProperties& target, size_t num_cu) if(user_dir.empty()) user_path = user_dir; if(!fs::exists(sys_path)) - sys_path = sys_dir / (target.Name() + ".kdb"); + sys_path = sys_dir / (target.DbId() + ".kdb"); #if !MIOPEN_EMBED_DB if(!fs::exists(sys_path)) sys_path = fs::path{}; diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp index f5ee7b2ed1..5c3bc88f56 100644 --- a/src/hip/handlehip.cpp +++ b/src/hip/handlehip.cpp @@ -622,8 +622,8 @@ Program Handle::LoadProgram(const fs::path& program_name, miopen::WriteFile(p.GetCodeObjectBlob(), path); else boost::filesystem::copy_file(p.GetCodeObjectPathname(), path); - cache_path = miopen::SaveBinary( - path, this->GetTargetProperties(), program_name, params); + cache_path = + miopen::SaveBinary(path, this->GetTargetProperties(), program_name, params); } if(force_attach_binary && p.IsCodeObjectInTempFile()) diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index d665340c5b..5a675f5976 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -568,24 +568,27 @@ void BuildKernel(const fs::path& program_file, { auto p = handle.LoadProgram(program_file, program_args, ""); - std::string compile_options = 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"){} + if(program_file.extension() == ".mlir") {} else if(program_file.extension() == ".s") - compile_options_tid += " -mcpu=" + miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId; + 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); + 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); + handle.GetTargetProperties(), + handle.GetMaxComputeUnits(), + program_file, + compile_options); } } catch(std::exception& e) From a376a45faad9aa180d1c76f27984b9e31cb13191 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 25 Apr 2025 19:56:34 +0000 Subject: [PATCH 5/7] update KDBTargetID test --- test/gtest/db_sync.cpp | 88 ++++++++++++++++++++++++------------------ 1 file changed, 50 insertions(+), 38 deletions(-) diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index 5a675f5976..8d4f8f0ca6 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -522,11 +522,50 @@ 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) {} + +// 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 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")) { @@ -537,7 +576,7 @@ TEST(CPU_DBSync_NONE, KDBTargetID) #if WORKAROUND_ISSUE_2492 ScopedEnvironment 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)); @@ -890,41 +929,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) { @@ -983,6 +987,14 @@ struct CPU_DBSync_NONE : testing::TestWithParam> { }; +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; From e5ee46461e3120af9c5cce2cfe43feb005fe0a9a Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 5 May 2025 15:11:28 -0500 Subject: [PATCH 6/7] remove old comments and #def --- test/gtest/db_sync.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index 8d4f8f0ca6..1826f60d31 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -527,19 +527,12 @@ 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; }; From a41e911cc8c49ee891ab40b8a859eae6d34d8e56 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 24 Jul 2025 14:31:25 -0500 Subject: [PATCH 7/7] remove WORKAROUND_ISSUE_3001 --- src/comgr.cpp | 4 ---- src/hip/handlehip.cpp | 5 ----- src/include/miopen/target_properties.hpp | 2 -- src/ocl/gcn_asm_utils.cpp | 4 ---- 4 files changed, 15 deletions(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index 4105d58eae..ad1fa7c926 100644 --- a/src/comgr.cpp +++ b/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=."); diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp index 7e2e963ba0..b177683f64 100644 --- a/src/hip/handlehip.cpp +++ b/src/hip/handlehip.cpp @@ -545,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 } @@ -560,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); diff --git a/src/include/miopen/target_properties.hpp b/src/include/miopen/target_properties.hpp index 256b8d934b..4b0d091550 100644 --- a/src/include/miopen/target_properties.hpp +++ b/src/include/miopen/target_properties.hpp @@ -29,8 +29,6 @@ #include #include -#define WORKAROUND_ISSUE_3001 0 - namespace miopen { struct Handle; diff --git a/src/ocl/gcn_asm_utils.cpp b/src/ocl/gcn_asm_utils.cpp index 2d941dd04d..500f058426 100644 --- a/src/ocl/gcn_asm_utils.cpp +++ b/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.