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 7157c8657c..b177683f64 100644 --- a/src/hip/handlehip.cpp +++ b/src/hip/handlehip.cpp @@ -31,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -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); 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/src/include/miopen/target_properties.hpp b/src/include/miopen/target_properties.hpp index 2b8b04b413..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 1 - 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. diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp index 9727b75011..556e4b8415 100644 --- a/test/gtest/db_sync.cpp +++ b/test/gtest/db_sync.cpp @@ -35,6 +35,7 @@ #include #include +#include #include #include #include @@ -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 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& 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(); @@ -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> { }; +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;