From 733aa8e43816cb93d00aab50b19012f35bb4c163 Mon Sep 17 00:00:00 2001 From: Vladimir Indic <139573562+vlaindic@users.noreply.github.com> Date: Thu, 25 Apr 2024 21:03:04 +0200 Subject: [PATCH] Restructure code object source code (#826) * public codeobj info * Restructure code object source code file layout * Update get_unloaded_code_objects + add iterate_loaded_code_objects * Remove get_unloaded_code_objects from visible internal API - iterate_loaded_code_objects + functor which filters on the hsa_executable_t effectively reproduces this behavior * Whitespace removal --------- Co-authored-by: Jonathan R. Madsen --- source/lib/rocprofiler-sdk/CMakeLists.txt | 1 + .../lib/rocprofiler-sdk/callback_tracing.cpp | 6 +- .../code_object/CMakeLists.txt | 11 + .../{hsa => code_object}/code_object.cpp | 197 +++++------------- .../{hsa => code_object}/code_object.hpp | 23 +- .../code_object/hip/CMakeLists.txt | 9 + .../code_object/hsa/CMakeLists.txt | 9 + .../code_object/hsa/code_object.cpp | 76 +++++++ .../code_object/hsa/code_object.hpp | 83 ++++++++ .../code_object/hsa/kernel_symbol.cpp | 76 +++++++ .../code_object/hsa/kernel_symbol.hpp | 75 +++++++ source/lib/rocprofiler-sdk/hsa/CMakeLists.txt | 2 - source/lib/rocprofiler-sdk/hsa/queue.cpp | 4 +- source/lib/rocprofiler-sdk/registration.cpp | 6 +- 14 files changed, 419 insertions(+), 159 deletions(-) create mode 100644 source/lib/rocprofiler-sdk/code_object/CMakeLists.txt rename source/lib/rocprofiler-sdk/{hsa => code_object}/code_object.cpp (87%) rename source/lib/rocprofiler-sdk/{hsa => code_object}/code_object.hpp (78%) create mode 100644 source/lib/rocprofiler-sdk/code_object/hip/CMakeLists.txt create mode 100644 source/lib/rocprofiler-sdk/code_object/hsa/CMakeLists.txt create mode 100644 source/lib/rocprofiler-sdk/code_object/hsa/code_object.cpp create mode 100644 source/lib/rocprofiler-sdk/code_object/hsa/code_object.hpp create mode 100644 source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.cpp create mode 100644 source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp diff --git a/source/lib/rocprofiler-sdk/CMakeLists.txt b/source/lib/rocprofiler-sdk/CMakeLists.txt index 2fc078a9c5..ffb104320a 100644 --- a/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -37,6 +37,7 @@ target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_SOURCES} add_subdirectory(hsa) add_subdirectory(hip) +add_subdirectory(code_object) add_subdirectory(context) add_subdirectory(counters) add_subdirectory(aql) diff --git a/source/lib/rocprofiler-sdk/callback_tracing.cpp b/source/lib/rocprofiler-sdk/callback_tracing.cpp index 348aeca609..fbdcdae962 100644 --- a/source/lib/rocprofiler-sdk/callback_tracing.cpp +++ b/source/lib/rocprofiler-sdk/callback_tracing.cpp @@ -27,11 +27,11 @@ #include #include +#include "lib/rocprofiler-sdk/code_object/code_object.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" #include "lib/rocprofiler-sdk/hip/hip.hpp" #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" -#include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/scratch_memory.hpp" #include "lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp" @@ -214,7 +214,7 @@ rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_trac } case ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT: { - val = rocprofiler::hsa::code_object::name_by_id(operation); + val = rocprofiler::code_object::name_by_id(operation); break; } case ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH: @@ -323,7 +323,7 @@ rocprofiler_iterate_callback_tracing_kind_operations( } case ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT: { - ops = rocprofiler::hsa::code_object::get_ids(); + ops = rocprofiler::code_object::get_ids(); break; } case ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH: diff --git a/source/lib/rocprofiler-sdk/code_object/CMakeLists.txt b/source/lib/rocprofiler-sdk/code_object/CMakeLists.txt new file mode 100644 index 0000000000..ce96785fa4 --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/CMakeLists.txt @@ -0,0 +1,11 @@ +# +# code object +# +set(ROCPROFILER_LIB_CODE_OBJECT_SOURCES code_object.cpp) +set(ROCPROFILER_LIB_CODE_OBJECT_HEADERS code_object.hpp) + +target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_CODE_OBJECT_SOURCES} + ${ROCPROFILER_LIB_CODE_OBJECT_HEADERS}) + +add_subdirectory(hsa) +add_subdirectory(hip) diff --git a/source/lib/rocprofiler-sdk/hsa/code_object.cpp b/source/lib/rocprofiler-sdk/code_object/code_object.cpp similarity index 87% rename from source/lib/rocprofiler-sdk/hsa/code_object.cpp rename to source/lib/rocprofiler-sdk/code_object/code_object.cpp index 990491aecb..2e2aabc1b6 100644 --- a/source/lib/rocprofiler-sdk/hsa/code_object.cpp +++ b/source/lib/rocprofiler-sdk/code_object/code_object.cpp @@ -20,21 +20,23 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. -#include "lib/rocprofiler-sdk/hsa/code_object.hpp" +#include "lib/rocprofiler-sdk/code_object/code_object.hpp" #include "lib/common/scope_destructor.hpp" #include "lib/common/static_object.hpp" #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" +#include "lib/rocprofiler-sdk/code_object/hsa/code_object.hpp" +#include "lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" -#include #include #include #include #include +#include #include #include @@ -43,6 +45,7 @@ #include #include #include +#include #include #if defined(ROCPROFILER_CI) @@ -55,8 +58,6 @@ namespace rocprofiler { -namespace hsa -{ namespace code_object { namespace @@ -158,7 +159,6 @@ get_names() get_names(_data, std::make_index_sequence{}); return _data; } -} // namespace code_object namespace { @@ -370,117 +370,6 @@ get_kernel_descriptor(uint64_t kernel_object) return reinterpret_cast(kernel_object); } -struct kernel_symbol -{ - using kernel_symbol_data_t = - rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; - - kernel_symbol() = default; - ~kernel_symbol() = default; - - kernel_symbol(const kernel_symbol&) = delete; - kernel_symbol(kernel_symbol&&) noexcept; - - kernel_symbol& operator=(const kernel_symbol&) = delete; - kernel_symbol& operator =(kernel_symbol&&) noexcept; - - bool beg_notified = false; - bool end_notified = false; - std::string* name = {}; - hsa_executable_t hsa_executable = {}; - hsa_agent_t hsa_agent = {}; - hsa_executable_symbol_t hsa_symbol = {}; - kernel_symbol_data_t rocp_data = common::init_public_api_struct(kernel_symbol_data_t{}); - context_user_data_map_t user_data = {}; -}; - -kernel_symbol::kernel_symbol(kernel_symbol&& rhs) noexcept { operator=(std::move(rhs)); } - -kernel_symbol& -kernel_symbol::operator=(kernel_symbol&& rhs) noexcept -{ - if(this != &rhs) - { - beg_notified = rhs.beg_notified; - end_notified = rhs.end_notified; - name = rhs.name; - hsa_executable = rhs.hsa_executable; - hsa_agent = rhs.hsa_agent; - hsa_symbol = rhs.hsa_symbol; - rocp_data = rhs.rocp_data; - user_data = std::move(rhs.user_data); - rocp_data.kernel_name = (name) ? name->c_str() : nullptr; - } - - return *this; -} - -bool -operator==(const kernel_symbol& lhs, const kernel_symbol& rhs) -{ - return std::tie(lhs.hsa_executable.handle, lhs.hsa_agent.handle, lhs.hsa_symbol.handle) == - std::tie(rhs.hsa_executable.handle, rhs.hsa_agent.handle, rhs.hsa_symbol.handle); -} - -struct code_object -{ - using code_object_data_t = rocprofiler_callback_tracing_code_object_load_data_t; - using symbol_array_t = std::vector>; - - code_object() = default; - ~code_object() = default; - - code_object(const code_object&) = delete; - code_object(code_object&&) noexcept; - - code_object& operator=(const code_object&) = delete; - code_object& operator =(code_object&&) noexcept; - - bool beg_notified = false; - bool end_notified = false; - std::string* uri = {}; - hsa_executable_t hsa_executable = {}; - hsa_loaded_code_object_t hsa_code_object = {}; - code_object_data_t rocp_data = common::init_public_api_struct(code_object_data_t{}); - symbol_array_t symbols = {}; - context_array_t contexts = {}; - context_user_data_map_t user_data = {}; -}; - -code_object::code_object(code_object&& rhs) noexcept { operator=(std::move(rhs)); } - -code_object& -code_object::operator=(code_object&& rhs) noexcept -{ - if(this != &rhs) - { - beg_notified = rhs.beg_notified; - end_notified = rhs.end_notified; - uri = rhs.uri; - hsa_executable = rhs.hsa_executable; - hsa_code_object = rhs.hsa_code_object; - rocp_data = rhs.rocp_data; - user_data = std::move(rhs.user_data); - rocp_data.uri = (uri) ? uri->c_str() : nullptr; - symbols = std::move(rhs.symbols); - } - - return *this; -} - -bool -operator==(const code_object& lhs, const code_object& rhs) -{ - return std::tie(lhs.hsa_executable.handle, lhs.hsa_code_object.handle) == - std::tie(rhs.hsa_executable.handle, rhs.hsa_code_object.handle); -} - -struct code_object_unload -{ - code_object* object = nullptr; - std::vector symbols = {}; -}; - auto& get_code_object_id() { @@ -495,12 +384,11 @@ get_kernel_symbol_id() return _v; } -using code_object_array_t = std::vector>; using kernel_object_map_t = std::unordered_map; using executable_array_t = std::vector; -using code_object_unload_array_t = std::vector; +using code_object_unload_array_t = std::vector; -std::vector +std::vector shutdown(hsa_executable_t executable); bool is_shutdown = false; @@ -517,7 +405,7 @@ get_code_objects() { static auto*& _v = common::static_object>::construct(); - static auto _dtor = common::scope_destructor{[]() { code_object_shutdown(); }}; + static auto _dtor = common::scope_destructor{[]() { finalize(); }}; return _v; } @@ -544,9 +432,9 @@ executable_iterate_agent_symbols_load_callback(hsa_executable_t executabl if(_status != HSA_STATUS_SUCCESS) return _status; \ } - auto& core_table = *get_table().core_; - auto* code_obj_v = static_cast(args); - auto symbol_v = kernel_symbol{}; + auto& core_table = *::rocprofiler::hsa::get_core_table(); + auto* code_obj_v = static_cast(args); + auto symbol_v = hsa::kernel_symbol{}; auto& data = symbol_v.rocp_data; symbol_v.hsa_executable = executable; @@ -626,7 +514,7 @@ executable_iterate_agent_symbols_load_callback(hsa_executable_t executabl data.kernel_object, data.kernel_id); - code_obj_v->symbols.emplace_back(std::make_unique(std::move(symbol_v))); + code_obj_v->symbols.emplace_back(std::make_unique(std::move(symbol_v))); return HSA_STATUS_SUCCESS; @@ -639,12 +527,12 @@ executable_iterate_agent_symbols_unload_callback(hsa_executable_t executa hsa_executable_symbol_t symbol, void* args) { - auto symbol_v = kernel_symbol{}; + auto symbol_v = hsa::kernel_symbol{}; symbol_v.hsa_executable = executable; symbol_v.hsa_agent = agent; symbol_v.hsa_symbol = symbol; - auto* code_obj_v = static_cast(args); + auto* code_obj_v = static_cast(args); CHECK_NOTNULL(code_obj_v); CHECK_NOTNULL(code_obj_v->object); @@ -672,7 +560,7 @@ code_object_load_callback(hsa_executable_t executable, } auto& loader_table = get_loader_table(); - auto code_obj_v = code_object{}; + auto code_obj_v = hsa::code_object{}; auto& data = code_obj_v.rocp_data; uint32_t _storage_type = ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_NONE; @@ -695,7 +583,7 @@ code_object_load_callback(hsa_executable_t executable, { if(itr && *itr == code_obj_v) { - get_table().core_->hsa_executable_iterate_agent_symbols_fn( + ::rocprofiler::hsa::get_core_table()->hsa_executable_iterate_agent_symbols_fn( executable, data.hsa_agent, executable_iterate_agent_symbols_load_callback, @@ -782,12 +670,12 @@ code_object_load_callback(hsa_executable_t executable, // generate a unique code object id data.code_object_id = ++get_code_object_id(); - auto _status = get_table().core_->hsa_executable_iterate_agent_symbols_fn( + auto _status = ::rocprofiler::hsa::get_core_table()->hsa_executable_iterate_agent_symbols_fn( executable, data.hsa_agent, executable_iterate_agent_symbols_load_callback, &code_obj_v); if(_status == HSA_STATUS_SUCCESS) { - code_obj_vec->emplace_back(std::make_unique(std::move(code_obj_v))); + code_obj_vec->emplace_back(std::make_unique(std::move(code_obj_v))); } else { @@ -804,7 +692,7 @@ code_object_unload_callback(hsa_executable_t executable, hsa_loaded_code_object_t loaded_code_object, void* args) { - auto code_obj_v = code_object{}; + auto code_obj_v = hsa::code_object{}; code_obj_v.hsa_executable = executable; code_obj_v.hsa_code_object = loaded_code_object; @@ -825,10 +713,11 @@ code_object_unload_callback(hsa_executable_t executable, itr->hsa_code_object.handle == loaded_code_object.handle) // if(itr && *itr == code_obj_v) { - auto& _last = code_obj_arr->emplace_back(code_object_unload{.object = itr.get()}); + auto& _last = + code_obj_arr->emplace_back(hsa::code_object_unload{.object = itr.get()}); auto agent = itr->rocp_data.hsa_agent; - get_table().core_->hsa_executable_iterate_agent_symbols_fn( + ::rocprofiler::hsa::get_core_table()->hsa_executable_iterate_agent_symbols_fn( executable, agent, executable_iterate_agent_symbols_unload_callback, &_last); } } @@ -837,6 +726,18 @@ code_object_unload_callback(hsa_executable_t executable, return HSA_STATUS_SUCCESS; } +std::vector +get_unloaded_code_objects(hsa_executable_t executable) +{ + auto _unloaded = std::vector{}; + + if(!is_shutdown && get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects) + get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects( + executable, code_object_unload_callback, &_unloaded); + + return _unloaded; +} + auto& get_freeze_function() { @@ -865,7 +766,7 @@ executable_freeze(hsa_executable_t executable, const char* options) auto* code_obj_vec = get_code_objects(); CHECK_NOTNULL(code_obj_vec)->wlock([executable](code_object_array_t& _vec) { - hsa::get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects( + get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects( executable, code_object_load_callback, &_vec); }); @@ -1005,14 +906,12 @@ executable_destroy(hsa_executable_t executable) return CHECK_NOTNULL(get_destroy_function())(executable); } -std::vector +std::vector shutdown(hsa_executable_t executable) { ROCP_INFO << "running " << __FUNCTION__ << " (executable=" << executable.handle << ")..."; - auto _unloaded = std::vector{}; - hsa::get_loader_table().hsa_ven_amd_loader_executable_iterate_loaded_code_objects( - executable, code_object_unload_callback, &_unloaded); + auto _unloaded = code_object::get_unloaded_code_objects(executable); constexpr auto CODE_OBJECT_KIND = ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT; constexpr auto CODE_OBJECT_LOAD = ROCPROFILER_CODE_OBJECT_LOAD; @@ -1085,7 +984,7 @@ shutdown(hsa_executable_t executable) } // namespace void -code_object_init(HsaApiTable* table) +initialize(HsaApiTable* table) { auto& core_table = *table->core_; @@ -1123,7 +1022,7 @@ get_kernel_id(uint64_t kernel_object) } void -code_object_shutdown() +finalize() { if(is_shutdown || !get_executables() || !get_code_objects()) return; @@ -1138,5 +1037,21 @@ code_object_shutdown() is_shutdown = true; } -} // namespace hsa + +void +iterate_loaded_code_objects(code_object_iterator_t&& func) +{ + if(is_shutdown || !get_executables() || !get_code_objects()) return; + + CHECK_NOTNULL(get_code_objects()) + ->rlock( + [](const code_object_array_t& data, code_object_iterator_t&& func_v) { + for(const auto& itr : data) + { + if(itr) func_v(*itr); + } + }, + std::move(func)); +} +} // namespace code_object } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/hsa/code_object.hpp b/source/lib/rocprofiler-sdk/code_object/code_object.hpp similarity index 78% rename from source/lib/rocprofiler-sdk/hsa/code_object.hpp rename to source/lib/rocprofiler-sdk/code_object/code_object.hpp index 5cd187aab2..0b31039ec7 100644 --- a/source/lib/rocprofiler-sdk/hsa/code_object.hpp +++ b/source/lib/rocprofiler-sdk/code_object/code_object.hpp @@ -22,17 +22,22 @@ #pragma once +#include "lib/rocprofiler-sdk/code_object/hsa/code_object.hpp" +#include "lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp" + #include #include +#include #include namespace rocprofiler { -namespace hsa -{ namespace code_object { +using code_object_array_t = std::vector>; +using code_object_iterator_t = std::function; + const char* name_by_id(uint32_t id); @@ -44,15 +49,17 @@ get_names(); std::vector get_ids(); -} // namespace code_object - -void -code_object_init(HsaApiTable* table); uint64_t get_kernel_id(uint64_t kernel_object); void -code_object_shutdown(); -} // namespace hsa +iterate_loaded_code_objects(code_object_iterator_t&& func); + +void +initialize(HsaApiTable* table); + +void +finalize(); +} // namespace code_object } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/code_object/hip/CMakeLists.txt b/source/lib/rocprofiler-sdk/code_object/hip/CMakeLists.txt new file mode 100644 index 0000000000..acdfb9215c --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hip/CMakeLists.txt @@ -0,0 +1,9 @@ +# +# code object data structures from HIP +# +set(ROCPROFILER_LIB_CODE_OBJECT_HIP_SOURCES) +set(ROCPROFILER_LIB_CODE_OBJECT_HIP_HEADERS) + +target_sources( + rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_CODE_OBJECT_HIP_SOURCES} + ${ROCPROFILER_LIB_CODE_OBJECT_HIP_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/code_object/hsa/CMakeLists.txt b/source/lib/rocprofiler-sdk/code_object/hsa/CMakeLists.txt new file mode 100644 index 0000000000..e50cfba5a6 --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hsa/CMakeLists.txt @@ -0,0 +1,9 @@ +# +# code object data structures from HSA +# +set(ROCPROFILER_LIB_CODE_OBJECT_HSA_SOURCES code_object.cpp kernel_symbol.cpp) +set(ROCPROFILER_LIB_CODE_OBJECT_HSA_HEADERS code_object.hpp kernel_symbol.hpp) + +target_sources( + rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_CODE_OBJECT_HSA_SOURCES} + ${ROCPROFILER_LIB_CODE_OBJECT_HSA_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/code_object/hsa/code_object.cpp b/source/lib/rocprofiler-sdk/code_object/hsa/code_object.cpp new file mode 100644 index 0000000000..39edd9601f --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hsa/code_object.cpp @@ -0,0 +1,76 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include "lib/rocprofiler-sdk/code_object/hsa/code_object.hpp" + +#include +#include +#include +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace code_object +{ +namespace hsa +{ +code_object::code_object(code_object&& rhs) noexcept { operator=(std::move(rhs)); } + +code_object& +code_object::operator=(code_object&& rhs) noexcept +{ + if(this != &rhs) + { + beg_notified = rhs.beg_notified; + end_notified = rhs.end_notified; + uri = rhs.uri; + hsa_executable = rhs.hsa_executable; + hsa_code_object = rhs.hsa_code_object; + rocp_data = rhs.rocp_data; + user_data = std::move(rhs.user_data); + rocp_data.uri = (uri) ? uri->c_str() : nullptr; + symbols = std::move(rhs.symbols); + } + + return *this; +} + +bool +operator==(const code_object& lhs, const code_object& rhs) +{ + return std::tie(lhs.hsa_executable.handle, lhs.hsa_code_object.handle) == + std::tie(rhs.hsa_executable.handle, rhs.hsa_code_object.handle); +} +} // namespace hsa +} // namespace code_object +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/code_object/hsa/code_object.hpp b/source/lib/rocprofiler-sdk/code_object/hsa/code_object.hpp new file mode 100644 index 0000000000..87e5348f92 --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hsa/code_object.hpp @@ -0,0 +1,83 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp" +#include "lib/rocprofiler-sdk/context/context.hpp" + +#include +#include + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace code_object +{ +namespace hsa +{ +using context_t = context::context; +using user_data_t = rocprofiler_user_data_t; +using context_user_data_map_t = std::unordered_map; +using context_array_t = context::context_array_t; +using context_user_data_map_t = std::unordered_map; + +struct code_object +{ + using code_object_data_t = rocprofiler_callback_tracing_code_object_load_data_t; + using symbol_array_t = std::vector>; + + code_object() = default; + ~code_object() = default; + + code_object(const code_object&) = delete; + code_object(code_object&&) noexcept; + + code_object& operator=(const code_object&) = delete; + code_object& operator =(code_object&&) noexcept; + + bool beg_notified = false; + bool end_notified = false; + std::string* uri = {}; + hsa_executable_t hsa_executable = {}; + hsa_loaded_code_object_t hsa_code_object = {}; + code_object_data_t rocp_data = common::init_public_api_struct(code_object_data_t{}); + symbol_array_t symbols = {}; + context_array_t contexts = {}; + context_user_data_map_t user_data = {}; +}; + +struct code_object_unload +{ + code_object* object = nullptr; + std::vector symbols = {}; +}; + +bool +operator==(const code_object& lhs, const code_object& rhs); +} // namespace hsa +} // namespace code_object +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.cpp b/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.cpp new file mode 100644 index 0000000000..18a0819f07 --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.cpp @@ -0,0 +1,76 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include "lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp" + +#include +#include +#include +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace code_object +{ +namespace hsa +{ +kernel_symbol::kernel_symbol(kernel_symbol&& rhs) noexcept { operator=(std::move(rhs)); } + +kernel_symbol& +kernel_symbol::operator=(kernel_symbol&& rhs) noexcept +{ + if(this != &rhs) + { + beg_notified = rhs.beg_notified; + end_notified = rhs.end_notified; + name = rhs.name; + hsa_executable = rhs.hsa_executable; + hsa_agent = rhs.hsa_agent; + hsa_symbol = rhs.hsa_symbol; + rocp_data = rhs.rocp_data; + user_data = std::move(rhs.user_data); + rocp_data.kernel_name = (name) ? name->c_str() : nullptr; + } + + return *this; +} + +bool +operator==(const kernel_symbol& lhs, const kernel_symbol& rhs) +{ + return std::tie(lhs.hsa_executable.handle, lhs.hsa_agent.handle, lhs.hsa_symbol.handle) == + std::tie(rhs.hsa_executable.handle, rhs.hsa_agent.handle, rhs.hsa_symbol.handle); +} +} // namespace hsa +} // namespace code_object +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp b/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp new file mode 100644 index 0000000000..c6c9947bb1 --- /dev/null +++ b/source/lib/rocprofiler-sdk/code_object/hsa/kernel_symbol.hpp @@ -0,0 +1,75 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/rocprofiler-sdk/context/context.hpp" + +#include +#include + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace code_object +{ +namespace hsa +{ +using context_t = context::context; +using user_data_t = rocprofiler_user_data_t; +using context_user_data_map_t = std::unordered_map; +using context_array_t = context::context_array_t; +using context_user_data_map_t = std::unordered_map; + +struct kernel_symbol +{ + using kernel_symbol_data_t = + rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; + + kernel_symbol() = default; + ~kernel_symbol() = default; + + kernel_symbol(const kernel_symbol&) = delete; + kernel_symbol(kernel_symbol&&) noexcept; + + kernel_symbol& operator=(const kernel_symbol&) = delete; + kernel_symbol& operator =(kernel_symbol&&) noexcept; + + bool beg_notified = false; + bool end_notified = false; + std::string* name = {}; + hsa_executable_t hsa_executable = {}; + hsa_agent_t hsa_agent = {}; + hsa_executable_symbol_t hsa_symbol = {}; + kernel_symbol_data_t rocp_data = common::init_public_api_struct(kernel_symbol_data_t{}); + context_user_data_map_t user_data = {}; +}; + +bool +operator==(const kernel_symbol& lhs, const kernel_symbol& rhs); +} // namespace hsa +} // namespace code_object +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt b/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt index dcaacd9b3c..d41c6021cb 100644 --- a/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt @@ -3,7 +3,6 @@ set(ROCPROFILER_LIB_HSA_SOURCES agent_cache.cpp aql_packet.cpp async_copy.cpp - code_object.cpp hsa_barrier.cpp hsa.cpp pc_sampling.hpp @@ -16,7 +15,6 @@ set(ROCPROFILER_LIB_HSA_HEADERS agent_cache.hpp aql_packet.hpp async_copy.hpp - code_object.hpp defines.hpp hsa_barrier.hpp hsa.hpp diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index a329ee6925..5c4774e291 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -24,8 +24,8 @@ #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" +#include "lib/rocprofiler-sdk/code_object/code_object.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/details/fmt.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" @@ -270,7 +270,7 @@ WriteInterceptor(const void* packets, queue.async_started(); // Copy kernel pkt, copy is to allow for signal to be modified rocprofiler_packet kernel_pkt = packets_arr[i]; - uint64_t kernel_id = get_kernel_id(kernel_pkt.kernel_dispatch.kernel_object); + uint64_t kernel_id = code_object::get_kernel_id(kernel_pkt.kernel_dispatch.kernel_object); queue.create_signal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &kernel_pkt.ext_amd_aql_pm4.completion_signal); diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index ca19b1db6c..4b07124a95 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -27,10 +27,10 @@ #include "lib/common/logging.hpp" #include "lib/common/static_object.hpp" #include "lib/rocprofiler-sdk/agent.hpp" +#include "lib/rocprofiler-sdk/code_object/code_object.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hip/hip.hpp" #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" -#include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" @@ -602,7 +602,7 @@ finalize() hsa::async_copy_fini(); hsa::queue_controller_fini(); page_migration::finalize(); - hsa::code_object_shutdown(); + code_object::finalize(); if(get_init_status() > 0) { invoke_client_finalizers(); @@ -749,8 +749,8 @@ rocprofiler_set_api_table(const char* name, // need to construct agent mappings before initializing the queue controller rocprofiler::agent::construct_agent_cache(hsa_api_table); rocprofiler::hsa::queue_controller_init(hsa_api_table); - rocprofiler::hsa::code_object_init(hsa_api_table); rocprofiler::hsa::async_copy_init(hsa_api_table, lib_instance); + rocprofiler::code_object::initialize(hsa_api_table); // install rocprofiler API wrappers rocprofiler::hsa::update_table(hsa_api_table->core_, lib_instance);