diff --git a/source/lib/rocprofiler-sdk/hsa/details/CMakeLists.txt b/source/lib/rocprofiler-sdk/hsa/details/CMakeLists.txt index a5be818157..be503cd2dd 100644 --- a/source/lib/rocprofiler-sdk/hsa/details/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/hsa/details/CMakeLists.txt @@ -2,7 +2,7 @@ # # set(ROCPROFILER_LIB_HSA_DETAILS_SOURCES) -set(ROCPROFILER_LIB_HSA_DETAILS_HEADERS ostream.hpp) +set(ROCPROFILER_LIB_HSA_DETAILS_HEADERS ostream.hpp fmt.hpp) target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_HSA_DETAILS_SOURCES} ${ROCPROFILER_LIB_HSA_DETAILS_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp b/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp new file mode 100644 index 0000000000..e982dad9aa --- /dev/null +++ b/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp @@ -0,0 +1,171 @@ +// MIT License +// +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + 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/hsa/queue.hpp" + +namespace fmt +{ +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_ext_amd_aql_pm4_packet_t const& pkt, Ctx& ctx) const + { + return fmt::format_to( + ctx.out(), + "[AQL_PM4_PKT, header={}, pm4_commands=[{:x}], completion_signal={}]", + pkt.header, + fmt::join(std::string_view((const char*) pkt.pm4_command, sizeof(pkt.pm4_command)), + " "), + pkt.completion_signal.handle); + } +}; + +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_kernel_dispatch_packet_t const& pkt, Ctx& ctx) const + { + return fmt::format_to(ctx.out(), + "[KERNEL_DISPATCH, header={}, dim={}, workgroup_size=[{}, {}, {}], " + "grid_size=[{}, {}, {}], private_size={}, group_size={}, " + "kernel_object={:x}, kern_arg={}, completion_signal={}]", + pkt.header, + pkt.setup, + pkt.workgroup_size_x, + pkt.workgroup_size_y, + pkt.workgroup_size_z, + pkt.grid_size_x, + pkt.grid_size_y, + pkt.grid_size_z, + pkt.private_segment_size, + pkt.group_segment_size, + pkt.kernel_object, + pkt.kernarg_address, + pkt.completion_signal.handle); + } +}; + +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_barrier_and_packet_t const& pkt, Ctx& ctx) const + { + return fmt::format_to( + ctx.out(), + "[BARRIER_AND, header={}, dep_signals=[{},{},{},{},{}], completion_signal={}]", + pkt.header, + pkt.dep_signal[0].handle, + pkt.dep_signal[1].handle, + pkt.dep_signal[2].handle, + pkt.dep_signal[3].handle, + pkt.dep_signal[4].handle, + pkt.completion_signal.handle); + } +}; + +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_barrier_or_packet_t const& pkt, Ctx& ctx) const + { + return fmt::format_to( + ctx.out(), + "[BARRIER_OR, header={}, dep_signals=[{},{},{},{},{}], completion_signal={}]", + pkt.header, + pkt.dep_signal[0].handle, + pkt.dep_signal[1].handle, + pkt.dep_signal[2].handle, + pkt.dep_signal[3].handle, + pkt.dep_signal[4].handle, + pkt.completion_signal.handle); + } +}; + +// fmt::format support for rocprofiler_packet +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(rocprofiler::hsa::rocprofiler_packet const& pkt, Ctx& ctx) const + { + static const char* type_names[] = {"HSA_PACKET_TYPE_VENDOR_SPECIFIC", + "HSA_PACKET_TYPE_INVALID", + "HSA_PACKET_TYPE_KERNEL_DISPATCH", + "HSA_PACKET_TYPE_BARRIER_AND", + "HSA_PACKET_TYPE_AGENT_DISPATCH", + "HSA_PACKET_TYPE_BARRIER_OR"}; + uint8_t t = ((pkt.ext_amd_aql_pm4.header >> HSA_PACKET_HEADER_TYPE) & + ((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)); + switch(t) + { + case 0: + // PM4 packet + return fmt::format_to(ctx.out(), "{}", pkt.ext_amd_aql_pm4); + case 2: + // Kernel dispatch + return fmt::format_to(ctx.out(), "{}", pkt.kernel_dispatch); + case 3: + // Barrier AND packet + return fmt::format_to(ctx.out(), "{}", pkt.barrier_and); + case 5: + // Barrier OR packet + return fmt::format_to(ctx.out(), "{}", pkt.barrier_or); + default: + return fmt::format_to(ctx.out(), "[Unprintable Packet of type {}]", type_names[t]); + } + } +}; +} // namespace fmt diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index 5af1f9d225..e3883c6f90 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -25,6 +25,7 @@ #include "lib/rocprofiler-sdk/buffer.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" @@ -404,6 +405,10 @@ WriteInterceptor(const void* packets, .extern_corr_ids = extern_corr_ids}); } + // Command is only executed if GLOG_v=2 or higher, otherwise it is a no-op + ROCP_TRACE << fmt::format( + "QueueID {}: {}", queue.get_id().handle, fmt::join(transformed_packets, fmt::format(" "))); + writer(transformed_packets.data(), transformed_packets.size()); } } // namespace @@ -428,7 +433,6 @@ Queue::Queue(const AgentCache& agent, : _core_api(core_api) , _ext_api(ext_api) , _agent(agent) - { LOG_IF(FATAL, _ext_api.hsa_amd_queue_intercept_create_fn(_agent.get_hsa_agent(), diff --git a/source/lib/rocprofiler-sdk/hsa/queue.hpp b/source/lib/rocprofiler-sdk/hsa/queue.hpp index 183e979a84..cd8960cad9 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -91,7 +91,7 @@ union rocprofiler_packet rocprofiler_packet(const rocprofiler_packet&) = default; rocprofiler_packet(rocprofiler_packet&&) noexcept = default; - rocprofiler_packet& operator=(const rocprofiler_packet&) = default; + rocprofiler_packet& operator=(const rocprofiler_packet&) = default; rocprofiler_packet& operator=(rocprofiler_packet&&) noexcept = default; }; enum class queue_state @@ -236,4 +236,4 @@ Queue::lock_queue(FuncT&& func) } } // namespace hsa -} // namespace rocprofiler +} // namespace rocprofiler \ No newline at end of file