From 039562eed836751bea9d22300dfd1f112e8df14b Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 14 May 2025 09:24:17 -0500 Subject: [PATCH 1/9] Adding headers for CTF --- source/lib/output/CMakeLists.txt | 2 + source/lib/output/generateCTF.cpp | 78 ++++++++++++++++++++++++ source/lib/output/generateCTF.hpp | 53 ++++++++++++++++ source/lib/output/output_config.cpp | 3 +- source/lib/output/output_config.hpp | 1 + source/lib/rocprofiler-sdk-tool/tool.cpp | 1 + 6 files changed, 137 insertions(+), 1 deletion(-) create mode 100644 source/lib/output/generateCTF.cpp create mode 100644 source/lib/output/generateCTF.hpp diff --git a/source/lib/output/CMakeLists.txt b/source/lib/output/CMakeLists.txt index 2a3505a562..754b5ad563 100644 --- a/source/lib/output/CMakeLists.txt +++ b/source/lib/output/CMakeLists.txt @@ -18,6 +18,7 @@ set(TOOL_OUTPUT_HEADERS generatePerfetto.hpp generateStats.hpp generateRocpd.hpp + generateCTF.hpp generator.hpp kernel_symbol_info.hpp host_symbol_info.hpp @@ -43,6 +44,7 @@ set(TOOL_OUTPUT_SOURCES generatePerfetto.cpp generateStats.cpp generateRocpd.cpp + generateCTF.cpp metadata.cpp node_info.cpp output_config.cpp diff --git a/source/lib/output/generateCTF.cpp b/source/lib/output/generateCTF.cpp new file mode 100644 index 0000000000..b3aed02acd --- /dev/null +++ b/source/lib/output/generateCTF.cpp @@ -0,0 +1,78 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "generateJSON.hpp" +#include "output_stream.hpp" +#include "statistics.hpp" +#include "timestamps.hpp" + +#include "lib/common/filesystem.hpp" +#include "lib/common/string_entry.hpp" +#include "lib/common/utility.hpp" + +#include +#include + +#include + +namespace rocprofiler +{ +namespace tool +{ +void +setup(const output_config& cfg) +{ + namespace fs = common::filesystem; + + auto _filename = get_output_filename(cfg, "results", std::string_view{}); + auto _filepath = fs::path{_filename}; + auto _name = _filepath.filename().string(); + auto _path = _filepath.parent_path().string(); + + if(fs::exists(_filepath)) fs::remove_all(_filepath); + + fs::create_directories(_filepath); + + ROCP_ERROR << "Opened result file: " << _filename; +} + +void +write_ctf(const output_config& cfg, + const metadata& tool_metadata, + uint64_t pid, + const std::vector& agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* scratch_memory_data, + std::deque* rccl_api_data, + std::deque* memory_allocation_data, + std::deque* rocdecode_api_data, + std::deque* rocjpeg_api_data) +{ + setup(cfg); +} + +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/output/generateCTF.hpp b/source/lib/output/generateCTF.hpp new file mode 100644 index 0000000000..bea8f911c5 --- /dev/null +++ b/source/lib/output/generateCTF.hpp @@ -0,0 +1,53 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "agent_info.hpp" +#include "buffered_output.hpp" +#include "metadata.hpp" +#include "output_config.hpp" +#include "output_stream.hpp" +#include "statistics.hpp" + +namespace rocprofiler +{ +namespace tool +{ +void +write_ctf(const output_config& cfg, + const metadata& tool_metadata, + uint64_t pid, + const std::vector& agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* scratch_memory_data, + std::deque* rccl_api_data, + std::deque* memory_allocation_data, + std::deque* rocdecode_api_data, + std::deque* rocjpeg_api_data); + +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/output/output_config.cpp b/source/lib/output/output_config.cpp index 806337fe9f..7f30f901b6 100644 --- a/source/lib/output/output_config.cpp +++ b/source/lib/output/output_config.cpp @@ -79,9 +79,10 @@ output_config::parse_env() pftrace_output = entries.count("PFTRACE") > 0; otf2_output = entries.count("OTF2") > 0; rocpd_output = entries.count("ROCPD") > 0 || entries.empty(); + ctf_output = entries.count("CTF") > 0; const auto supported_formats = - std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD"}; + std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD", "CTF"}; for(const auto& itr : entries) { LOG_IF(FATAL, supported_formats.count(itr) == 0) diff --git a/source/lib/output/output_config.hpp b/source/lib/output/output_config.hpp index 4284f9bdff..b112fb8dac 100644 --- a/source/lib/output/output_config.hpp +++ b/source/lib/output/output_config.hpp @@ -71,6 +71,7 @@ struct output_config bool summary_output = false; bool kernel_rename = false; bool group_by_queue = false; + bool ctf_output = false; uint64_t stats_summary_unit_value = 1; size_t perfetto_shmem_size_hint = defaults::perfetto_shmem_size_hint_kb; size_t perfetto_buffer_size = defaults::perfetto_buffer_size_kb; diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index d1a789c571..b01656f9cd 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -46,6 +46,7 @@ #include "lib/output/csv_output_file.hpp" #include "lib/output/domain_type.hpp" #include "lib/output/generateCSV.hpp" +#include "lib/output/generateCTF.hpp" #include "lib/output/generateJSON.hpp" #include "lib/output/generateOTF2.hpp" #include "lib/output/generatePerfetto.hpp" From f14a5d9f10aff785f25f69436d0a20640b847f65 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 14 May 2025 13:06:14 -0500 Subject: [PATCH 2/9] Updates --- source/lib/output/CMakeLists.txt | 5 +- source/lib/output/generateCTF.cpp | 346 +++++++++++++++++++++++++++++- source/lib/output/generateCTF.hpp | 31 ++- 3 files changed, 375 insertions(+), 7 deletions(-) diff --git a/source/lib/output/CMakeLists.txt b/source/lib/output/CMakeLists.txt index 754b5ad563..3fee93296d 100644 --- a/source/lib/output/CMakeLists.txt +++ b/source/lib/output/CMakeLists.txt @@ -54,6 +54,8 @@ set(TOOL_OUTPUT_SOURCES tmp_file_buffer.cpp tmp_file.cpp) +find_library(Babeltrace2 babeltrace2 REQUIRED HINTS "/usr/lib/x86_64-linux-gnu") + add_library(rocprofiler-sdk-output-library STATIC) add_library(rocprofiler-sdk::rocprofiler-sdk-output-library ALIAS rocprofiler-sdk-output-library) @@ -72,7 +74,8 @@ target_link_libraries( rocprofiler-sdk::rocprofiler-sdk-amd-comgr rocprofiler-sdk::rocprofiler-sdk-dw rocprofiler-sdk::rocprofiler-sdk-elf - rocprofiler-sdk::rocprofiler-sdk-sqlite3) + rocprofiler-sdk::rocprofiler-sdk-sqlite3 + babeltrace2) target_compile_definitions(rocprofiler-sdk-output-library PRIVATE PROJECT_BINARY_DIR="${PROJECT_BINARY_DIR}") diff --git a/source/lib/output/generateCTF.cpp b/source/lib/output/generateCTF.cpp index b3aed02acd..12963cb9bc 100644 --- a/source/lib/output/generateCTF.cpp +++ b/source/lib/output/generateCTF.cpp @@ -20,7 +20,7 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. -#include "generateJSON.hpp" +#include "generateCTF.hpp" #include "output_stream.hpp" #include "statistics.hpp" #include "timestamps.hpp" @@ -38,11 +38,287 @@ namespace rocprofiler { namespace tool { +namespace fs = common::filesystem; + +ctf_output::ctf_output(const output_config& cfg) +{ + auto _filename = get_output_filename(cfg, "results", std::string_view{}); + auto _filepath = fs::path{_filename}; + + // Initialize Babeltrace 2 logging + bt_logging_set_global_level(BT_LOGGING_LEVEL_DEBUG); + + // Create a Babeltrace 2 component graph + graph = bt_graph_create(0); // Pass 0 for default MIP version + if(!graph) + { + fprintf(stderr, "Failed to create Babeltrace 2 graph\n"); + return; + } + + // Prepare parameters for the CTF sink + bt_value* params = bt_value_map_create(); + if(!params) + { + fprintf(stderr, "Failed to create parameters map\n"); + bt_graph_put_ref(graph); + return; + } + bt_value* out_path = bt_value_string_create_init(_filepath.c_str()); + bt_value_map_insert_entry(params, "path", out_path); + + // Add the CTF writer sink component from the ctf plugin + bt_component* ctf_writer_comp = nullptr; + // Discover the plugin and component class + const bt_plugin_set* plugin_set = nullptr; + bt_plugin_find_all_status find_status = bt_plugin_find_all(BT_TRUE, // find_in_std_env_var + BT_TRUE, // find_in_user_dir + BT_TRUE, // find_in_system_dir + BT_TRUE, // find_in_static + BT_TRUE, // find_in_static_linked + &plugin_set); + + if(find_status != BT_PLUGIN_FIND_ALL_STATUS_OK || !plugin_set) + { + fprintf(stderr, "Failed to find plugins\n"); + bt_graph_put_ref(graph); + return; + } + + const bt_plugin* ctf_plugin = nullptr; + uint64_t plugin_count = bt_plugin_set_get_plugin_count(plugin_set); + for(uint64_t i = 0; i < plugin_count; ++i) + { + const bt_plugin* plugin = bt_plugin_set_borrow_plugin_by_index_const(plugin_set, i); + if(plugin && strcmp(bt_plugin_get_name(plugin), "ctf") == 0) + { + ctf_plugin = plugin; + break; + } + } + if(!ctf_plugin) + { + fprintf(stderr, "Failed to find 'ctf' plugin\n"); + bt_graph_put_ref(graph); + return; + } + + const bt_component_class_sink* ctf_sink_class = + bt_plugin_borrow_sink_component_class_by_name_const(ctf_plugin, "fs"); + if(!ctf_sink_class) + { + fprintf(stderr, "Failed to find 'fs' sink class in 'ctf' plugin\n"); + bt_graph_put_ref(graph); + return; + } + + // Now add the sink component + bt_graph_add_component_status status = + bt_graph_add_sink_component(graph, + ctf_sink_class, + "ctf_writer", + params, + BT_LOGGING_LEVEL_NONE, + (const bt_component_sink**) &ctf_writer_comp); + bt_plugin_set_put_ref(plugin_set); + bt_value_put_ref(params); + + if(status != BT_GRAPH_ADD_COMPONENT_STATUS_OK || !ctf_writer_comp) + { + fprintf(stderr, "Failed to add CTF writer sink component to the graph\n"); + bt_graph_put_ref(graph); + return; + } +} + +ctf_output::~ctf_output() +{ + close(); +} + +void ctf_output::close() { + // Cleanup + bt_graph_put_ref(graph); +} + +template +void ctf_output::write_event(const record_type& event) { + write_event_impl(event, ctf_event_tag{}); +} + +// ---- HIP API EXT ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_hip_api_ext_record_t& event, + ctf_event_tag) +{ + // Example: Extract fields + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + // event.args, event.retval + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- HSA API ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_hsa_api_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- Kernel Dispatch ---- +void ctf_output::write_event_impl( + const tool_buffer_tracing_kernel_dispatch_ext_record_t& event, + ctf_event_tag) +{ + auto& base = static_cast(event); + auto kind = base.kind; + auto op = base.operation; + auto corr = base.correlation_id; + auto tid = base.thread_id; + auto start = base.start_timestamp; + auto end = base.end_timestamp; + auto dispatch_info = base.dispatch_info; + auto stream_id = event.stream_id; + auto kernel_rename_val = event.kernel_rename_val; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- Memory Copy ---- +void ctf_output::write_event_impl( + const tool_buffer_tracing_memory_copy_ext_record_t& event, + ctf_event_tag) +{ + auto& base = static_cast(event); + auto kind = base.kind; + auto op = base.operation; + auto corr = base.correlation_id; + auto tid = base.thread_id; + auto start = base.start_timestamp; + auto end = base.end_timestamp; + auto dst_agent_id = base.dst_agent_id; + auto src_agent_id = base.src_agent_id; + auto bytes = base.bytes; + auto dst_addr = base.dst_address; + auto src_addr = base.src_address; + auto stream_id = event.stream_id; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- Marker API ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_marker_api_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- Scratch Memory ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_scratch_memory_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto agent_id = event.agent_id; + auto queue_id = event.queue_id; + auto tid = event.thread_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto flags = event.flags; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- RCCL API ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_rccl_api_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- Memory Allocation ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_memory_allocation_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto tid = event.thread_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto agent_id = event.agent_id; + auto address = event.address; + auto alloc_size = event.allocation_size; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- rocDecode API EXT ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_rocdecode_api_ext_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + // event.args, event.retval + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + +// ---- rocJPEG API ---- +void ctf_output::write_event_impl( + const rocprofiler_buffer_tracing_rocjpeg_api_record_t& event, + ctf_event_tag) +{ + auto kind = event.kind; + auto op = event.operation; + auto corr = event.correlation_id; + auto start = event.start_timestamp; + auto end = event.end_timestamp; + auto tid = event.thread_id; + + // TODO: Use Babeltrace 2 API to create a message/event and set these fields +} + void setup(const output_config& cfg) { - namespace fs = common::filesystem; - auto _filename = get_output_filename(cfg, "results", std::string_view{}); auto _filepath = fs::path{_filename}; auto _name = _filepath.filename().string(); @@ -55,8 +331,18 @@ setup(const output_config& cfg) ROCP_ERROR << "Opened result file: " << _filename; } +ctf_output open_ctf_stream(const output_config& cfg) { + setup(cfg); + return ctf_output{cfg}; +} + +void close_ctf_stream(ctf_output& ctf_out) { + ctf_out.close(); +} + void -write_ctf(const output_config& cfg, +write_ctf(ctf_output& ctf_out, + const output_config& cfg, const metadata& tool_metadata, uint64_t pid, const std::vector& agent_data, @@ -71,7 +357,57 @@ write_ctf(const output_config& cfg std::deque* rocdecode_api_data, std::deque* rocjpeg_api_data) { - setup(cfg); + // Loop over each deque and call write_event for each record + if (hip_api_data) { + for (const auto& rec : *hip_api_data) { + ctf_out.write_event(rec); + } + } + if (hsa_api_data) { + for (const auto& rec : *hsa_api_data) { + ctf_out.write_event(rec); + } + } + if (kernel_dispatch_data) { + for (const auto& rec : *kernel_dispatch_data) { + ctf_out.write_event(rec); + } + } + if (memory_copy_data) { + for (const auto& rec : *memory_copy_data) { + ctf_out.write_event(rec); + } + } + if (marker_api_data) { + for (const auto& rec : *marker_api_data) { + ctf_out.write_event(rec); + } + } + if (scratch_memory_data) { + for (const auto& rec : *scratch_memory_data) { + ctf_out.write_event(rec); + } + } + if (rccl_api_data) { + for (const auto& rec : *rccl_api_data) { + ctf_out.write_event(rec); + } + } + if (memory_allocation_data) { + for (const auto& rec : *memory_allocation_data) { + ctf_out.write_event(rec); + } + } + if (rocdecode_api_data) { + for (const auto& rec : *rocdecode_api_data) { + ctf_out.write_event(rec); + } + } + if (rocjpeg_api_data) { + for (const auto& rec : *rocjpeg_api_data) { + ctf_out.write_event(rec); + } + } } } // namespace tool diff --git a/source/lib/output/generateCTF.hpp b/source/lib/output/generateCTF.hpp index bea8f911c5..58944528a9 100644 --- a/source/lib/output/generateCTF.hpp +++ b/source/lib/output/generateCTF.hpp @@ -20,6 +20,11 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#include +#include +#include +#include + #pragma once #include "agent_info.hpp" @@ -33,8 +38,32 @@ namespace rocprofiler { namespace tool { +struct ctf_output +{ + ctf_output(const output_config& cfg); + ~ctf_output(); + + template + void write_event(const record_type& event); + void close(); + +private: + bt_graph* graph; + bt_component* ctf_writer_comp; + bt_message_iterator* msg_iter; + bt_event_class* event_class; + + // Helper: Tag dispatch for each record type + template struct ctf_event_tag {}; +}; + +ctf_output open_ctf_stream(const output_config& cfg); + +void close_ctf_stream(ctf_output& ctf_out); + void -write_ctf(const output_config& cfg, +write_ctf(ctf_output& ctf_out, + const output_config& cfg, const metadata& tool_metadata, uint64_t pid, const std::vector& agent_data, From ae2676360789085041898928c22b0d3fa5b651bf Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Thu, 15 May 2025 10:34:30 -0500 Subject: [PATCH 3/9] Addind HIP API EXT Data for prototype --- source/lib/output/ctf/hip_api_ext.hpp | 127 +++++++++ source/lib/output/generateCTF.cpp | 364 ++++++++++---------------- source/lib/output/generateCTF.hpp | 32 ++- 3 files changed, 293 insertions(+), 230 deletions(-) create mode 100644 source/lib/output/ctf/hip_api_ext.hpp diff --git a/source/lib/output/ctf/hip_api_ext.hpp b/source/lib/output/ctf/hip_api_ext.hpp new file mode 100644 index 0000000000..6b01757251 --- /dev/null +++ b/source/lib/output/ctf/hip_api_ext.hpp @@ -0,0 +1,127 @@ +#include +#include +#include +#include +#include + +#include +#include + +#include "lib/common/filesystem.hpp" +#include "lib/common/string_entry.hpp" +#include "lib/common/utility.hpp" + +#include +#include + +#include + +// State for the source component +struct hip_api_source_state { + std::deque* hip_api_data; + bt_stream *stream; + bt_event_class *event_class; +}; + +// --- Component methods --- + +// Initialization method: create trace/stream/event classes and store state +static +bt_component_class_initialize_method_status hip_api_source_init( + bt_self_component_source *self_comp_src, + bt_self_component_source_configuration * /*config*/, + const bt_value * /*params*/, + void * init_data) +{ + std::deque* hip_api_data = (std::deque*)init_data; + + // Create trace class + bt_self_component *self_comp = bt_self_component_source_as_self_component(self_comp_src); + bt_trace_class *trace_class = bt_trace_class_create(self_comp); + + // Create stream class + bt_stream_class *stream_class = bt_stream_class_create(trace_class); + + // Create event class + bt_event_class *event_class = bt_event_class_create(stream_class); + + // Create payload field class + bt_field_class *payload_fc = bt_field_class_structure_create(trace_class); + bt_field_class_structure_append_member(payload_fc, "kind", bt_field_class_integer_unsigned_create(trace_class)); + bt_field_class_structure_append_member(payload_fc, "operation", bt_field_class_integer_unsigned_create(trace_class)); + bt_field_class_structure_append_member(payload_fc, "correlation_id", bt_field_class_integer_unsigned_create(trace_class)); + bt_field_class_structure_append_member(payload_fc, "thread_id", bt_field_class_integer_unsigned_create(trace_class)); + bt_event_class_set_payload_field_class(event_class, payload_fc); + + // Create trace and stream + bt_trace *trace = bt_trace_create(trace_class); + bt_stream *stream = bt_stream_create(stream_class, trace); + + // Store state + hip_api_source_state* state = static_cast(malloc(sizeof(hip_api_source_state))); + state->hip_api_data = hip_api_data; + state->stream = stream; + state->event_class = event_class; + bt_self_component_set_data(bt_self_component_source_as_self_component(self_comp_src), state); + + return BT_COMPONENT_CLASS_INITIALIZE_METHOD_STATUS_OK; +} + +// Next method: emit events from the queue +static +bt_message_iterator_class_next_method_status hip_api_source_next( + bt_self_message_iterator *self_message_iterator, + bt_message_array_const msgs, uint64_t capacity, + uint64_t *count) +{ + // Get the component's user data (state) + bt_self_component_source *self_comp_src = + (bt_self_component_source *) bt_self_message_iterator_borrow_component(self_message_iterator); + hip_api_source_state* state = (hip_api_source_state*) + bt_self_component_get_data(bt_self_component_source_as_self_component(self_comp_src)); + + uint64_t produced = 0; + + while (produced < capacity && !state->hip_api_data->empty()) { + rocprofiler_buffer_tracing_hip_api_ext_record_t* rec = &state->hip_api_data->back(); + + // Create event message + bt_message *msg = bt_message_event_create( + self_message_iterator, + state->event_class, + state->stream + ); + if (!msg) break; + + // Set event payload fields + bt_event *evt = bt_message_event_borrow_event(msg); + bt_field *payload = bt_event_borrow_payload_field(evt); + + bt_field *kind_field = bt_field_structure_borrow_member_field_by_name(payload, "kind"); + bt_field_integer_unsigned_set_value(kind_field, rec->kind); + + bt_field *op_field = bt_field_structure_borrow_member_field_by_name(payload, "operation"); + bt_field_integer_signed_set_value(op_field, rec->operation); + + bt_field *corr_field = bt_field_structure_borrow_member_field_by_name(payload, "correlation_id"); + bt_field_integer_unsigned_set_value(corr_field, rec->correlation_id.internal); + + bt_field *tid_field = bt_field_structure_borrow_member_field_by_name(payload, "thread_id"); + bt_field_integer_unsigned_set_value(tid_field, rec->thread_id); + + msgs[produced++] = msg; + state->hip_api_data->pop_back(); + } + + *count = produced; + return BT_MESSAGE_ITERATOR_CLASS_NEXT_METHOD_STATUS_OK; +} + +// Finalize method: cleanup +static +void hip_api_source_finalize(bt_self_component_source *self_comp_src) +{ + hip_api_source_state* state = (hip_api_source_state*)bt_self_component_get_data( + bt_self_component_source_as_self_component(self_comp_src)); + free(state); +} \ No newline at end of file diff --git a/source/lib/output/generateCTF.cpp b/source/lib/output/generateCTF.cpp index 12963cb9bc..3a2c2dd496 100644 --- a/source/lib/output/generateCTF.cpp +++ b/source/lib/output/generateCTF.cpp @@ -25,6 +25,8 @@ #include "statistics.hpp" #include "timestamps.hpp" +#include "ctf/hip_api_ext.hpp" + #include "lib/common/filesystem.hpp" #include "lib/common/string_entry.hpp" #include "lib/common/utility.hpp" @@ -131,189 +133,60 @@ ctf_output::ctf_output(const output_config& cfg) } } -ctf_output::~ctf_output() -{ - close(); -} +ctf_output::~ctf_output() { close(); } -void ctf_output::close() { +void +ctf_output::close() +{ // Cleanup bt_graph_put_ref(graph); } -template -void ctf_output::write_event(const record_type& event) { - write_event_impl(event, ctf_event_tag{}); -} - -// ---- HIP API EXT ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_hip_api_ext_record_t& event, - ctf_event_tag) -{ - // Example: Extract fields - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - // event.args, event.retval - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- HSA API ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_hsa_api_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- Kernel Dispatch ---- -void ctf_output::write_event_impl( - const tool_buffer_tracing_kernel_dispatch_ext_record_t& event, - ctf_event_tag) -{ - auto& base = static_cast(event); - auto kind = base.kind; - auto op = base.operation; - auto corr = base.correlation_id; - auto tid = base.thread_id; - auto start = base.start_timestamp; - auto end = base.end_timestamp; - auto dispatch_info = base.dispatch_info; - auto stream_id = event.stream_id; - auto kernel_rename_val = event.kernel_rename_val; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- Memory Copy ---- -void ctf_output::write_event_impl( - const tool_buffer_tracing_memory_copy_ext_record_t& event, - ctf_event_tag) -{ - auto& base = static_cast(event); - auto kind = base.kind; - auto op = base.operation; - auto corr = base.correlation_id; - auto tid = base.thread_id; - auto start = base.start_timestamp; - auto end = base.end_timestamp; - auto dst_agent_id = base.dst_agent_id; - auto src_agent_id = base.src_agent_id; - auto bytes = base.bytes; - auto dst_addr = base.dst_address; - auto src_addr = base.src_address; - auto stream_id = event.stream_id; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- Marker API ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_marker_api_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- Scratch Memory ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_scratch_memory_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto agent_id = event.agent_id; - auto queue_id = event.queue_id; - auto tid = event.thread_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto flags = event.flags; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- RCCL API ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_rccl_api_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- Memory Allocation ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_memory_allocation_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto tid = event.thread_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto agent_id = event.agent_id; - auto address = event.address; - auto alloc_size = event.allocation_size; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- rocDecode API EXT ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_rocdecode_api_ext_record_t& event, - ctf_event_tag) -{ - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - // event.args, event.retval - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields -} - -// ---- rocJPEG API ---- -void ctf_output::write_event_impl( - const rocprofiler_buffer_tracing_rocjpeg_api_record_t& event, - ctf_event_tag) +void +ctf_output::write_event_source_component( + const char* source_name, + bt_component_class_initialize_method_status (*init_method)( + bt_self_component_source*, + bt_self_component_source_configuration*, + const bt_value*, + void*), + void (*finalize_method)(bt_self_component_source*), + bt_message_iterator_class_next_method_status (*next_method)( + bt_self_message_iterator* self_message_iterator, + bt_message_array_const msgs, + uint64_t capacity, + uint64_t* count), + void* data) { - auto kind = event.kind; - auto op = event.operation; - auto corr = event.correlation_id; - auto start = event.start_timestamp; - auto end = event.end_timestamp; - auto tid = event.thread_id; - - // TODO: Use Babeltrace 2 API to create a message/event and set these fields + bt_message_iterator_class* msg_iter_cls = bt_message_iterator_class_create(next_method); + + // 2. Optionally set msg_iter init/finalize + // bt_message_iterator_class_set_initialize_method(msg_iter_cls, hip_api_source_msg_iter_init); + // bt_message_iterator_class_set_finalize_method(msg_iter_cls, + // hip_api_source_msg_iter_finalize); + + // 3. Create source component class + bt_component_class_source* src_class = + bt_component_class_source_create(source_name, msg_iter_cls); + + // 4. Set component class init/finalize + bt_component_class_source_set_initialize_method(src_class, init_method); + bt_component_class_source_set_finalize_method(src_class, finalize_method); + + // 5. Add the source component to the graph + const bt_component_source* src_comp = nullptr; + bt_graph_add_source_component_with_initialize_method_data( + graph, src_class, source_name, NULL, data, BT_LOGGING_LEVEL_NONE, &src_comp); + + bt_port_output* src_out_port = + (bt_port_output*) bt_self_component_source_borrow_output_port_by_index( + (bt_self_component_source*) src_comp, 0); + bt_port_input* sink_in_port = + (bt_port_input*) bt_self_component_sink_borrow_input_port_by_index( + (bt_self_component_sink*) ctf_writer_comp, 0); + bt_graph_connect_ports(graph, src_out_port, sink_in_port, NULL); + + bt_graph_run(graph); } void @@ -331,21 +204,25 @@ setup(const output_config& cfg) ROCP_ERROR << "Opened result file: " << _filename; } -ctf_output open_ctf_stream(const output_config& cfg) { +ctf_output +open_ctf_stream(const output_config& cfg) +{ setup(cfg); return ctf_output{cfg}; } -void close_ctf_stream(ctf_output& ctf_out) { +void +close_ctf_stream(ctf_output& ctf_out) +{ ctf_out.close(); } void -write_ctf(ctf_output& ctf_out, - const output_config& cfg, - const metadata& tool_metadata, - uint64_t pid, - const std::vector& agent_data, +write_ctf(ctf_output& ctf_out, + const output_config& /*cfg*/, + const metadata& /*tool_metadata*/, + uint64_t /*pid*/, + const std::vector& /*agent_data*/, std::deque* hip_api_data, std::deque* hsa_api_data, std::deque* kernel_dispatch_data, @@ -358,55 +235,96 @@ write_ctf(ctf_output& ctf std::deque* rocjpeg_api_data) { // Loop over each deque and call write_event for each record - if (hip_api_data) { - for (const auto& rec : *hip_api_data) { - ctf_out.write_event(rec); + if(hip_api_data) + { + if(!ctf_out.hip_api_ext_initialized) + { + ctf_out.hip_api_data = + new std::deque(); + ctf_out.write_event_source_component("hip_api_source", + hip_api_source_init, + hip_api_source_finalize, + hip_api_source_next, + (void*) ctf_out.hip_api_data); + ctf_out.hip_api_ext_initialized = true; } - } - if (hsa_api_data) { - for (const auto& rec : *hsa_api_data) { - ctf_out.write_event(rec); + + while(!hip_api_data->empty()) { + ctf_out.hip_api_data->push_front(hip_api_data->back()); + hip_api_data->pop_back(); } } - if (kernel_dispatch_data) { - for (const auto& rec : *kernel_dispatch_data) { - ctf_out.write_event(rec); - } + if(hsa_api_data) + { + // ctf_out.write_event_source_component("hsa_api_source", + // hsa_api_source_init, + // hsa_api_source_finalize, + // hsa_api_source_next, + // hsa_api_data); } - if (memory_copy_data) { - for (const auto& rec : *memory_copy_data) { - ctf_out.write_event(rec); - } + if(kernel_dispatch_data) + { + // ctf_out.write_event_source_component("kernel_dispatch_source", + // kernel_dispatch_source_init, + // kernel_dispatch_source_finalize, + // kernel_dispatch_source_next, + // kernel_dispatch_data); } - if (marker_api_data) { - for (const auto& rec : *marker_api_data) { - ctf_out.write_event(rec); - } + if(memory_copy_data) + { + // ctf_out.write_event_source_component("memory_copy_source", + // memory_copy_source_init, + // memory_copy_source_finalize, + // memory_copy_source_next, + // memory_copy_data); } - if (scratch_memory_data) { - for (const auto& rec : *scratch_memory_data) { - ctf_out.write_event(rec); - } + if(marker_api_data) + { + // ctf_out.write_event_source_component("marker_api_source", + // marker_api_source_init, + // marker_api_source_finalize, + // marker_api_source_next, + // marker_api_data); } - if (rccl_api_data) { - for (const auto& rec : *rccl_api_data) { - ctf_out.write_event(rec); - } + if(scratch_memory_data) + { + // ctf_out.write_event_source_component("scratch_memory_source", + // scratch_memory_source_init, + // scratch_memory_source_finalize, + // scratch_memory_source_next, + // scratch_memory_data); } - if (memory_allocation_data) { - for (const auto& rec : *memory_allocation_data) { - ctf_out.write_event(rec); - } + if(rccl_api_data) + { + // ctf_out.write_event_source_component("rccl_api_source", + // rccl_api_source_init, + // rccl_api_source_finalize, + // rccl_api_source_next, + // rccl_api_data); } - if (rocdecode_api_data) { - for (const auto& rec : *rocdecode_api_data) { - ctf_out.write_event(rec); - } + if(memory_allocation_data) + { + // ctf_out.write_event_source_component("memory_allocation_source", + // memory_allocation_source_init, + // memory_allocation_source_finalize, + // memory_allocation_source_next, + // memory_allocation_data); } - if (rocjpeg_api_data) { - for (const auto& rec : *rocjpeg_api_data) { - ctf_out.write_event(rec); - } + if(rocdecode_api_data) + { + // ctf_out.write_event_source_component("rocdecode_api_source", + // rocdecode_api_source_init, + // rocdecode_api_source_finalize, + // rocdecode_api_source_next, + // rocdecode_api_data); + } + if(rocjpeg_api_data) + { + // ctf_out.write_event_source_component("rocjpeg_api_source", + // rocjpeg_api_source_init, + // rocjpeg_api_source_finalize, + // rocjpeg_api_source_next, + // rocjpeg_api_data); } } diff --git a/source/lib/output/generateCTF.hpp b/source/lib/output/generateCTF.hpp index 58944528a9..be95aa4dfe 100644 --- a/source/lib/output/generateCTF.hpp +++ b/source/lib/output/generateCTF.hpp @@ -38,28 +38,46 @@ namespace rocprofiler { namespace tool { + + enum record_type_t { + HIP_API_EXT = 0, + + }; struct ctf_output { ctf_output(const output_config& cfg); ~ctf_output(); - template - void write_event(const record_type& event); + void write_event_source_component(const char* source_name, + bt_component_class_initialize_method_status (*init_method)( + bt_self_component_source*, + bt_self_component_source_configuration*, + const bt_value*, + void*), + void (*finalize_method)(bt_self_component_source*), + bt_message_iterator_class_next_method_status (*next_method)( + bt_self_message_iterator* self_message_iterator, + bt_message_array_const msgs, + uint64_t capacity, + uint64_t* count), void* data); void close(); + bool hip_api_ext_initialized{false}; + + std::deque* hip_api_data = nullptr; + private: bt_graph* graph; bt_component* ctf_writer_comp; bt_message_iterator* msg_iter; bt_event_class* event_class; - - // Helper: Tag dispatch for each record type - template struct ctf_event_tag {}; }; -ctf_output open_ctf_stream(const output_config& cfg); +ctf_output +open_ctf_stream(const output_config& cfg); -void close_ctf_stream(ctf_output& ctf_out); +void +close_ctf_stream(ctf_output& ctf_out); void write_ctf(ctf_output& ctf_out, From fa5618158b37819bdf1d8104a6105d8ee3416646 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Mon, 2 Jun 2025 13:42:53 -0500 Subject: [PATCH 4/9] Replacing the usage of babeltrace2 with lttng-ust --- source/bin/rocprofv3.py | 4 +- .../comparing-with-legacy-tools.rst | 2 +- source/lib/output/CMakeLists.txt | 9 +- source/lib/output/ctf/hip_api_ext.hpp | 127 ------- source/lib/output/generateCTF.cpp | 332 ------------------ source/lib/output/generateCTF.hpp | 100 ------ source/lib/output/generateLTTng.cpp | 209 +++++++++++ source/lib/output/generateLTTng.hpp | 53 +++ source/lib/output/lttng/CMakeLists.txt | 9 + .../lttng/rocprofiler_sdk_trace_provider.cpp | 4 + .../lttng/rocprofiler_sdk_trace_provider.hpp | 224 ++++++++++++ source/lib/output/output_config.cpp | 4 +- source/lib/output/output_config.hpp | 2 +- source/lib/rocprofiler-sdk-tool/tool.cpp | 32 +- 14 files changed, 540 insertions(+), 571 deletions(-) delete mode 100644 source/lib/output/ctf/hip_api_ext.hpp delete mode 100644 source/lib/output/generateCTF.cpp delete mode 100644 source/lib/output/generateCTF.hpp create mode 100644 source/lib/output/generateLTTng.cpp create mode 100644 source/lib/output/generateLTTng.hpp create mode 100644 source/lib/output/lttng/CMakeLists.txt create mode 100644 source/lib/output/lttng/rocprofiler_sdk_trace_provider.cpp create mode 100644 source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp diff --git a/source/bin/rocprofv3.py b/source/bin/rocprofv3.py index 0a807eb446..29f2f853dd 100755 --- a/source/bin/rocprofv3.py +++ b/source/bin/rocprofv3.py @@ -275,10 +275,10 @@ def add_parser_bool_argument(gparser, *args, **kwargs): io_options.add_argument( "-f", "--output-format", - help="For adding output format (supported formats: csv, json, pftrace, otf2, rocpd)", + help="For adding output format (supported formats: csv, json, pftrace, otf2, rocpd, lttng)", nargs="+", default=None, - choices=("csv", "json", "pftrace", "otf2", "rocpd"), + choices=("csv", "json", "pftrace", "otf2", "rocpd", "lttng"), type=str.lower, ) add_parser_bool_argument( diff --git a/source/docs/conceptual/comparing-with-legacy-tools.rst b/source/docs/conceptual/comparing-with-legacy-tools.rst index 195a2af5df..214cebadda 100644 --- a/source/docs/conceptual/comparing-with-legacy-tools.rst +++ b/source/docs/conceptual/comparing-with-legacy-tools.rst @@ -328,7 +328,7 @@ ROCprofiler-SDK introduces a new command-line tool, `rocprofv3`, which is a more * - I/O options - Output Formats - CSV, JSON (Chrome-Tracing format) - - CSV, JSON (Chrome-Tracing format), Perfetto, CTF + - CSV, JSON (Chrome-Tracing format), Perfetto, LTTng - CSV, JSON (custom schema), Perfetto, OTF2 - | # Multiple output formats can be supported in single run. | # OTF2 can visualize larger trace files compared to perfetto. diff --git a/source/lib/output/CMakeLists.txt b/source/lib/output/CMakeLists.txt index 3fee93296d..b3c17f7f7f 100644 --- a/source/lib/output/CMakeLists.txt +++ b/source/lib/output/CMakeLists.txt @@ -18,7 +18,7 @@ set(TOOL_OUTPUT_HEADERS generatePerfetto.hpp generateStats.hpp generateRocpd.hpp - generateCTF.hpp + generateLTTng.hpp generator.hpp kernel_symbol_info.hpp host_symbol_info.hpp @@ -44,7 +44,7 @@ set(TOOL_OUTPUT_SOURCES generatePerfetto.cpp generateStats.cpp generateRocpd.cpp - generateCTF.cpp + generateLTTng.cpp metadata.cpp node_info.cpp output_config.cpp @@ -54,8 +54,6 @@ set(TOOL_OUTPUT_SOURCES tmp_file_buffer.cpp tmp_file.cpp) -find_library(Babeltrace2 babeltrace2 REQUIRED HINTS "/usr/lib/x86_64-linux-gnu") - add_library(rocprofiler-sdk-output-library STATIC) add_library(rocprofiler-sdk::rocprofiler-sdk-output-library ALIAS rocprofiler-sdk-output-library) @@ -75,9 +73,10 @@ target_link_libraries( rocprofiler-sdk::rocprofiler-sdk-dw rocprofiler-sdk::rocprofiler-sdk-elf rocprofiler-sdk::rocprofiler-sdk-sqlite3 - babeltrace2) + rocprofiler-sdk-tool-output-lttng) target_compile_definitions(rocprofiler-sdk-output-library PRIVATE PROJECT_BINARY_DIR="${PROJECT_BINARY_DIR}") add_subdirectory(sql) +add_subdirectory(lttng) diff --git a/source/lib/output/ctf/hip_api_ext.hpp b/source/lib/output/ctf/hip_api_ext.hpp deleted file mode 100644 index 6b01757251..0000000000 --- a/source/lib/output/ctf/hip_api_ext.hpp +++ /dev/null @@ -1,127 +0,0 @@ -#include -#include -#include -#include -#include - -#include -#include - -#include "lib/common/filesystem.hpp" -#include "lib/common/string_entry.hpp" -#include "lib/common/utility.hpp" - -#include -#include - -#include - -// State for the source component -struct hip_api_source_state { - std::deque* hip_api_data; - bt_stream *stream; - bt_event_class *event_class; -}; - -// --- Component methods --- - -// Initialization method: create trace/stream/event classes and store state -static -bt_component_class_initialize_method_status hip_api_source_init( - bt_self_component_source *self_comp_src, - bt_self_component_source_configuration * /*config*/, - const bt_value * /*params*/, - void * init_data) -{ - std::deque* hip_api_data = (std::deque*)init_data; - - // Create trace class - bt_self_component *self_comp = bt_self_component_source_as_self_component(self_comp_src); - bt_trace_class *trace_class = bt_trace_class_create(self_comp); - - // Create stream class - bt_stream_class *stream_class = bt_stream_class_create(trace_class); - - // Create event class - bt_event_class *event_class = bt_event_class_create(stream_class); - - // Create payload field class - bt_field_class *payload_fc = bt_field_class_structure_create(trace_class); - bt_field_class_structure_append_member(payload_fc, "kind", bt_field_class_integer_unsigned_create(trace_class)); - bt_field_class_structure_append_member(payload_fc, "operation", bt_field_class_integer_unsigned_create(trace_class)); - bt_field_class_structure_append_member(payload_fc, "correlation_id", bt_field_class_integer_unsigned_create(trace_class)); - bt_field_class_structure_append_member(payload_fc, "thread_id", bt_field_class_integer_unsigned_create(trace_class)); - bt_event_class_set_payload_field_class(event_class, payload_fc); - - // Create trace and stream - bt_trace *trace = bt_trace_create(trace_class); - bt_stream *stream = bt_stream_create(stream_class, trace); - - // Store state - hip_api_source_state* state = static_cast(malloc(sizeof(hip_api_source_state))); - state->hip_api_data = hip_api_data; - state->stream = stream; - state->event_class = event_class; - bt_self_component_set_data(bt_self_component_source_as_self_component(self_comp_src), state); - - return BT_COMPONENT_CLASS_INITIALIZE_METHOD_STATUS_OK; -} - -// Next method: emit events from the queue -static -bt_message_iterator_class_next_method_status hip_api_source_next( - bt_self_message_iterator *self_message_iterator, - bt_message_array_const msgs, uint64_t capacity, - uint64_t *count) -{ - // Get the component's user data (state) - bt_self_component_source *self_comp_src = - (bt_self_component_source *) bt_self_message_iterator_borrow_component(self_message_iterator); - hip_api_source_state* state = (hip_api_source_state*) - bt_self_component_get_data(bt_self_component_source_as_self_component(self_comp_src)); - - uint64_t produced = 0; - - while (produced < capacity && !state->hip_api_data->empty()) { - rocprofiler_buffer_tracing_hip_api_ext_record_t* rec = &state->hip_api_data->back(); - - // Create event message - bt_message *msg = bt_message_event_create( - self_message_iterator, - state->event_class, - state->stream - ); - if (!msg) break; - - // Set event payload fields - bt_event *evt = bt_message_event_borrow_event(msg); - bt_field *payload = bt_event_borrow_payload_field(evt); - - bt_field *kind_field = bt_field_structure_borrow_member_field_by_name(payload, "kind"); - bt_field_integer_unsigned_set_value(kind_field, rec->kind); - - bt_field *op_field = bt_field_structure_borrow_member_field_by_name(payload, "operation"); - bt_field_integer_signed_set_value(op_field, rec->operation); - - bt_field *corr_field = bt_field_structure_borrow_member_field_by_name(payload, "correlation_id"); - bt_field_integer_unsigned_set_value(corr_field, rec->correlation_id.internal); - - bt_field *tid_field = bt_field_structure_borrow_member_field_by_name(payload, "thread_id"); - bt_field_integer_unsigned_set_value(tid_field, rec->thread_id); - - msgs[produced++] = msg; - state->hip_api_data->pop_back(); - } - - *count = produced; - return BT_MESSAGE_ITERATOR_CLASS_NEXT_METHOD_STATUS_OK; -} - -// Finalize method: cleanup -static -void hip_api_source_finalize(bt_self_component_source *self_comp_src) -{ - hip_api_source_state* state = (hip_api_source_state*)bt_self_component_get_data( - bt_self_component_source_as_self_component(self_comp_src)); - free(state); -} \ No newline at end of file diff --git a/source/lib/output/generateCTF.cpp b/source/lib/output/generateCTF.cpp deleted file mode 100644 index 3a2c2dd496..0000000000 --- a/source/lib/output/generateCTF.cpp +++ /dev/null @@ -1,332 +0,0 @@ -// MIT License -// -// Copyright (c) 2023-2025 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 "generateCTF.hpp" -#include "output_stream.hpp" -#include "statistics.hpp" -#include "timestamps.hpp" - -#include "ctf/hip_api_ext.hpp" - -#include "lib/common/filesystem.hpp" -#include "lib/common/string_entry.hpp" -#include "lib/common/utility.hpp" - -#include -#include - -#include - -namespace rocprofiler -{ -namespace tool -{ -namespace fs = common::filesystem; - -ctf_output::ctf_output(const output_config& cfg) -{ - auto _filename = get_output_filename(cfg, "results", std::string_view{}); - auto _filepath = fs::path{_filename}; - - // Initialize Babeltrace 2 logging - bt_logging_set_global_level(BT_LOGGING_LEVEL_DEBUG); - - // Create a Babeltrace 2 component graph - graph = bt_graph_create(0); // Pass 0 for default MIP version - if(!graph) - { - fprintf(stderr, "Failed to create Babeltrace 2 graph\n"); - return; - } - - // Prepare parameters for the CTF sink - bt_value* params = bt_value_map_create(); - if(!params) - { - fprintf(stderr, "Failed to create parameters map\n"); - bt_graph_put_ref(graph); - return; - } - bt_value* out_path = bt_value_string_create_init(_filepath.c_str()); - bt_value_map_insert_entry(params, "path", out_path); - - // Add the CTF writer sink component from the ctf plugin - bt_component* ctf_writer_comp = nullptr; - // Discover the plugin and component class - const bt_plugin_set* plugin_set = nullptr; - bt_plugin_find_all_status find_status = bt_plugin_find_all(BT_TRUE, // find_in_std_env_var - BT_TRUE, // find_in_user_dir - BT_TRUE, // find_in_system_dir - BT_TRUE, // find_in_static - BT_TRUE, // find_in_static_linked - &plugin_set); - - if(find_status != BT_PLUGIN_FIND_ALL_STATUS_OK || !plugin_set) - { - fprintf(stderr, "Failed to find plugins\n"); - bt_graph_put_ref(graph); - return; - } - - const bt_plugin* ctf_plugin = nullptr; - uint64_t plugin_count = bt_plugin_set_get_plugin_count(plugin_set); - for(uint64_t i = 0; i < plugin_count; ++i) - { - const bt_plugin* plugin = bt_plugin_set_borrow_plugin_by_index_const(plugin_set, i); - if(plugin && strcmp(bt_plugin_get_name(plugin), "ctf") == 0) - { - ctf_plugin = plugin; - break; - } - } - if(!ctf_plugin) - { - fprintf(stderr, "Failed to find 'ctf' plugin\n"); - bt_graph_put_ref(graph); - return; - } - - const bt_component_class_sink* ctf_sink_class = - bt_plugin_borrow_sink_component_class_by_name_const(ctf_plugin, "fs"); - if(!ctf_sink_class) - { - fprintf(stderr, "Failed to find 'fs' sink class in 'ctf' plugin\n"); - bt_graph_put_ref(graph); - return; - } - - // Now add the sink component - bt_graph_add_component_status status = - bt_graph_add_sink_component(graph, - ctf_sink_class, - "ctf_writer", - params, - BT_LOGGING_LEVEL_NONE, - (const bt_component_sink**) &ctf_writer_comp); - bt_plugin_set_put_ref(plugin_set); - bt_value_put_ref(params); - - if(status != BT_GRAPH_ADD_COMPONENT_STATUS_OK || !ctf_writer_comp) - { - fprintf(stderr, "Failed to add CTF writer sink component to the graph\n"); - bt_graph_put_ref(graph); - return; - } -} - -ctf_output::~ctf_output() { close(); } - -void -ctf_output::close() -{ - // Cleanup - bt_graph_put_ref(graph); -} - -void -ctf_output::write_event_source_component( - const char* source_name, - bt_component_class_initialize_method_status (*init_method)( - bt_self_component_source*, - bt_self_component_source_configuration*, - const bt_value*, - void*), - void (*finalize_method)(bt_self_component_source*), - bt_message_iterator_class_next_method_status (*next_method)( - bt_self_message_iterator* self_message_iterator, - bt_message_array_const msgs, - uint64_t capacity, - uint64_t* count), - void* data) -{ - bt_message_iterator_class* msg_iter_cls = bt_message_iterator_class_create(next_method); - - // 2. Optionally set msg_iter init/finalize - // bt_message_iterator_class_set_initialize_method(msg_iter_cls, hip_api_source_msg_iter_init); - // bt_message_iterator_class_set_finalize_method(msg_iter_cls, - // hip_api_source_msg_iter_finalize); - - // 3. Create source component class - bt_component_class_source* src_class = - bt_component_class_source_create(source_name, msg_iter_cls); - - // 4. Set component class init/finalize - bt_component_class_source_set_initialize_method(src_class, init_method); - bt_component_class_source_set_finalize_method(src_class, finalize_method); - - // 5. Add the source component to the graph - const bt_component_source* src_comp = nullptr; - bt_graph_add_source_component_with_initialize_method_data( - graph, src_class, source_name, NULL, data, BT_LOGGING_LEVEL_NONE, &src_comp); - - bt_port_output* src_out_port = - (bt_port_output*) bt_self_component_source_borrow_output_port_by_index( - (bt_self_component_source*) src_comp, 0); - bt_port_input* sink_in_port = - (bt_port_input*) bt_self_component_sink_borrow_input_port_by_index( - (bt_self_component_sink*) ctf_writer_comp, 0); - bt_graph_connect_ports(graph, src_out_port, sink_in_port, NULL); - - bt_graph_run(graph); -} - -void -setup(const output_config& cfg) -{ - auto _filename = get_output_filename(cfg, "results", std::string_view{}); - auto _filepath = fs::path{_filename}; - auto _name = _filepath.filename().string(); - auto _path = _filepath.parent_path().string(); - - if(fs::exists(_filepath)) fs::remove_all(_filepath); - - fs::create_directories(_filepath); - - ROCP_ERROR << "Opened result file: " << _filename; -} - -ctf_output -open_ctf_stream(const output_config& cfg) -{ - setup(cfg); - return ctf_output{cfg}; -} - -void -close_ctf_stream(ctf_output& ctf_out) -{ - ctf_out.close(); -} - -void -write_ctf(ctf_output& ctf_out, - const output_config& /*cfg*/, - const metadata& /*tool_metadata*/, - uint64_t /*pid*/, - const std::vector& /*agent_data*/, - std::deque* hip_api_data, - std::deque* hsa_api_data, - std::deque* kernel_dispatch_data, - std::deque* memory_copy_data, - std::deque* marker_api_data, - std::deque* scratch_memory_data, - std::deque* rccl_api_data, - std::deque* memory_allocation_data, - std::deque* rocdecode_api_data, - std::deque* rocjpeg_api_data) -{ - // Loop over each deque and call write_event for each record - if(hip_api_data) - { - if(!ctf_out.hip_api_ext_initialized) - { - ctf_out.hip_api_data = - new std::deque(); - ctf_out.write_event_source_component("hip_api_source", - hip_api_source_init, - hip_api_source_finalize, - hip_api_source_next, - (void*) ctf_out.hip_api_data); - ctf_out.hip_api_ext_initialized = true; - } - - while(!hip_api_data->empty()) { - ctf_out.hip_api_data->push_front(hip_api_data->back()); - hip_api_data->pop_back(); - } - } - if(hsa_api_data) - { - // ctf_out.write_event_source_component("hsa_api_source", - // hsa_api_source_init, - // hsa_api_source_finalize, - // hsa_api_source_next, - // hsa_api_data); - } - if(kernel_dispatch_data) - { - // ctf_out.write_event_source_component("kernel_dispatch_source", - // kernel_dispatch_source_init, - // kernel_dispatch_source_finalize, - // kernel_dispatch_source_next, - // kernel_dispatch_data); - } - if(memory_copy_data) - { - // ctf_out.write_event_source_component("memory_copy_source", - // memory_copy_source_init, - // memory_copy_source_finalize, - // memory_copy_source_next, - // memory_copy_data); - } - if(marker_api_data) - { - // ctf_out.write_event_source_component("marker_api_source", - // marker_api_source_init, - // marker_api_source_finalize, - // marker_api_source_next, - // marker_api_data); - } - if(scratch_memory_data) - { - // ctf_out.write_event_source_component("scratch_memory_source", - // scratch_memory_source_init, - // scratch_memory_source_finalize, - // scratch_memory_source_next, - // scratch_memory_data); - } - if(rccl_api_data) - { - // ctf_out.write_event_source_component("rccl_api_source", - // rccl_api_source_init, - // rccl_api_source_finalize, - // rccl_api_source_next, - // rccl_api_data); - } - if(memory_allocation_data) - { - // ctf_out.write_event_source_component("memory_allocation_source", - // memory_allocation_source_init, - // memory_allocation_source_finalize, - // memory_allocation_source_next, - // memory_allocation_data); - } - if(rocdecode_api_data) - { - // ctf_out.write_event_source_component("rocdecode_api_source", - // rocdecode_api_source_init, - // rocdecode_api_source_finalize, - // rocdecode_api_source_next, - // rocdecode_api_data); - } - if(rocjpeg_api_data) - { - // ctf_out.write_event_source_component("rocjpeg_api_source", - // rocjpeg_api_source_init, - // rocjpeg_api_source_finalize, - // rocjpeg_api_source_next, - // rocjpeg_api_data); - } -} - -} // namespace tool -} // namespace rocprofiler diff --git a/source/lib/output/generateCTF.hpp b/source/lib/output/generateCTF.hpp deleted file mode 100644 index be95aa4dfe..0000000000 --- a/source/lib/output/generateCTF.hpp +++ /dev/null @@ -1,100 +0,0 @@ -// MIT License -// -// Copyright (c) 2023-2025 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 -#include -#include -#include - -#pragma once - -#include "agent_info.hpp" -#include "buffered_output.hpp" -#include "metadata.hpp" -#include "output_config.hpp" -#include "output_stream.hpp" -#include "statistics.hpp" - -namespace rocprofiler -{ -namespace tool -{ - - enum record_type_t { - HIP_API_EXT = 0, - - }; -struct ctf_output -{ - ctf_output(const output_config& cfg); - ~ctf_output(); - - void write_event_source_component(const char* source_name, - bt_component_class_initialize_method_status (*init_method)( - bt_self_component_source*, - bt_self_component_source_configuration*, - const bt_value*, - void*), - void (*finalize_method)(bt_self_component_source*), - bt_message_iterator_class_next_method_status (*next_method)( - bt_self_message_iterator* self_message_iterator, - bt_message_array_const msgs, - uint64_t capacity, - uint64_t* count), void* data); - void close(); - - bool hip_api_ext_initialized{false}; - - std::deque* hip_api_data = nullptr; - -private: - bt_graph* graph; - bt_component* ctf_writer_comp; - bt_message_iterator* msg_iter; - bt_event_class* event_class; -}; - -ctf_output -open_ctf_stream(const output_config& cfg); - -void -close_ctf_stream(ctf_output& ctf_out); - -void -write_ctf(ctf_output& ctf_out, - const output_config& cfg, - const metadata& tool_metadata, - uint64_t pid, - const std::vector& agent_data, - std::deque* hip_api_data, - std::deque* hsa_api_data, - std::deque* kernel_dispatch_data, - std::deque* memory_copy_data, - std::deque* marker_api_data, - std::deque* scratch_memory_data, - std::deque* rccl_api_data, - std::deque* memory_allocation_data, - std::deque* rocdecode_api_data, - std::deque* rocjpeg_api_data); - -} // namespace tool -} // namespace rocprofiler diff --git a/source/lib/output/generateLTTng.cpp b/source/lib/output/generateLTTng.cpp new file mode 100644 index 0000000000..2ec4f2b08d --- /dev/null +++ b/source/lib/output/generateLTTng.cpp @@ -0,0 +1,209 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "generateLTTng.hpp" +#include "output_stream.hpp" +#include "statistics.hpp" +#include "timestamps.hpp" + +#include "lttng/rocprofiler_sdk_trace_provider.hpp" + +#include "lib/common/string_entry.hpp" +#include "lib/common/utility.hpp" + +#include +#include + +#include + +namespace rocprofiler +{ +namespace tool +{ +void +write_lttng(const output_config& /*cfg*/, + const metadata& tool_metadata, + uint64_t pid, + const std::vector& agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* scratch_memory_data, + std::deque* rccl_api_data, + std::deque* memory_allocation_data, + std::deque* rocdecode_api_data, + std::deque* rocjpeg_api_data) +{ + auto buffer_names = sdk::get_buffer_tracing_names(); + auto callbk_name_info = sdk::get_callback_tracing_names(); + + for(auto& _agent_info : agent_data) + { + tracepoint(rocprofv3_trace, + agents_info, + pid, + _agent_info.type, + _agent_info.node_id, + _agent_info.logical_node_id, + _agent_info.logical_node_type_id, + _agent_info.gpu_id, + _agent_info.name, + _agent_info.vendor_name, + _agent_info.product_name, + _agent_info.model_name); + } + // Loop over each deque and call write_event for each record + if(hip_api_data) + { + for (auto hip_api_record : *hip_api_data) + { + // --- Prepare raw args data --- + const uint8_t* raw_args_ptr = reinterpret_cast(&hip_api_record.args); + uint64_t raw_args_size_val = static_cast( + sizeof(hip_api_record.args)); // Or sizeof(rocprofiler_hip_api_args_t) + + auto api_name = buffer_names.at(hip_api_record.kind, hip_api_record.operation); + + tracepoint(rocprofv3_trace, + hip_api, + pid, + api_name.data(), + hip_api_record.correlation_id.ancestor, + hip_api_record.correlation_id.internal, + hip_api_record.start_timestamp, + hip_api_record.end_timestamp, + hip_api_record.thread_id, + raw_args_ptr, + raw_args_size_val // Sending args as bytes + ); + } + } + if(hsa_api_data) + { + for (auto hsa_api_record : *hsa_api_data) + { + auto api_name = buffer_names.at(hsa_api_record.kind, hsa_api_record.operation); + + tracepoint(rocprofv3_trace, + hsa_api, + pid, + api_name.data(), + hsa_api_record.correlation_id.ancestor, + hsa_api_record.correlation_id.internal, + hsa_api_record.start_timestamp, + hsa_api_record.end_timestamp, + hsa_api_record.thread_id + ); + } + } + if(kernel_dispatch_data) + { + for(auto kernel_dispatch_record : *kernel_dispatch_data) { + auto name = + tool_metadata.get_kernel_name(kernel_dispatch_record.dispatch_info.kernel_id, kernel_dispatch_record.correlation_id.external.value); + + tracepoint(rocprofv3_trace, + kernel_dispatch, + pid, + kernel_dispatch_record.correlation_id.internal, + kernel_dispatch_record.start_timestamp, + kernel_dispatch_record.end_timestamp, + kernel_dispatch_record.thread_id, + kernel_dispatch_record.dispatch_info.agent_id.handle, + kernel_dispatch_record.dispatch_info.queue_id.handle, + kernel_dispatch_record.stream_id.handle, + name.data() + ); + } + } + if(memory_copy_data) + { + for(auto memory_copy_record : *memory_copy_data) { + tracepoint(rocprofv3_trace, + memory_copy, + pid, + memory_copy_record.operation, + memory_copy_record.correlation_id.internal, + memory_copy_record.start_timestamp, + memory_copy_record.end_timestamp, + memory_copy_record.thread_id, + memory_copy_record.src_agent_id.handle, + memory_copy_record.dst_agent_id.handle, + memory_copy_record.stream_id.handle, + memory_copy_record.bytes + ); + } + } + if(marker_api_data) + { + // ctf_out.write_event_source_component("marker_api_source", + // marker_api_source_init, + // marker_api_source_finalize, + // marker_api_source_next, + // marker_api_data); + } + if(scratch_memory_data) + { + // ctf_out.write_event_source_component("scratch_memory_source", + // scratch_memory_source_init, + // scratch_memory_source_finalize, + // scratch_memory_source_next, + // scratch_memory_data); + } + if(rccl_api_data) + { + // ctf_out.write_event_source_component("rccl_api_source", + // rccl_api_source_init, + // rccl_api_source_finalize, + // rccl_api_source_next, + // rccl_api_data); + } + if(memory_allocation_data) + { + // ctf_out.write_event_source_component("memory_allocation_source", + // memory_allocation_source_init, + // memory_allocation_source_finalize, + // memory_allocation_source_next, + // memory_allocation_data); + } + if(rocdecode_api_data) + { + // ctf_out.write_event_source_component("rocdecode_api_source", + // rocdecode_api_source_init, + // rocdecode_api_source_finalize, + // rocdecode_api_source_next, + // rocdecode_api_data); + } + if(rocjpeg_api_data) + { + // ctf_out.write_event_source_component("rocjpeg_api_source", + // rocjpeg_api_source_init, + // rocjpeg_api_source_finalize, + // rocjpeg_api_source_next, + // rocjpeg_api_data); + } +} + +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/output/generateLTTng.hpp b/source/lib/output/generateLTTng.hpp new file mode 100644 index 0000000000..6b27bb6114 --- /dev/null +++ b/source/lib/output/generateLTTng.hpp @@ -0,0 +1,53 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 "agent_info.hpp" +#include "buffered_output.hpp" +#include "metadata.hpp" +#include "output_config.hpp" +#include "output_stream.hpp" +#include "statistics.hpp" + +namespace rocprofiler +{ +namespace tool +{ +void +write_lttng(const output_config& cfg, + const metadata& tool_metadata, + uint64_t pid, + const std::vector& agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* scratch_memory_data, + std::deque* rccl_api_data, + std::deque* memory_allocation_data, + std::deque* rocdecode_api_data, + std::deque* rocjpeg_api_data); + +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/output/lttng/CMakeLists.txt b/source/lib/output/lttng/CMakeLists.txt new file mode 100644 index 0000000000..8fa7921eee --- /dev/null +++ b/source/lib/output/lttng/CMakeLists.txt @@ -0,0 +1,9 @@ +find_package(LTTngUST REQUIRED) + +add_library(rocprofiler-sdk-tool-output-lttng STATIC rocprofiler_sdk_trace_provider.cpp + rocprofiler_sdk_trace_provider.hpp) + +target_include_directories(rocprofiler-sdk-tool-output-lttng + PUBLIC ${LTTNGUST_INCLUDE_DIRS} ${CMAKE_CURRENT_SOURCE_DIR}) + +target_link_libraries(rocprofiler-sdk-tool-output-lttng ${LTTNGUST_LIBRARIES} dl) diff --git a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.cpp b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.cpp new file mode 100644 index 0000000000..1b03167342 --- /dev/null +++ b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.cpp @@ -0,0 +1,4 @@ +// rocprofiler_sdk_trace_provider.cpp +#define TRACEPOINT_CREATE_PROBES +#define TRACEPOINT_DEFINE +#include "rocprofiler_sdk_trace_provider.hpp" \ No newline at end of file diff --git a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp new file mode 100644 index 0000000000..f9f260ffb5 --- /dev/null +++ b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp @@ -0,0 +1,224 @@ +// clang-format off +// rocprofiler_sdk_trace_provider.hpp + +#undef TRACEPOINT_PROVIDER +#define TRACEPOINT_PROVIDER rocprofv3_trace // APP Provider Name + +#undef TRACEPOINT_INCLUDE +#define TRACEPOINT_INCLUDE "./rocprofiler_sdk_trace_provider.hpp" // Must point to itself + +#if !defined(_ROCPROFILER_SDK_TRACE_PROVIDER_H) || defined(TRACEPOINT_HEADER_MULTI_READ) +#define _ROCPROFILER_SDK_TRACE_PROVIDER_H + +#include + +// Define the tracepoint event for recording the agents +TRACEPOINT_EVENT( + rocprofv3_trace, + agents_info, + TP_ARGS( + uint64_t, rec_pid, + + uint32_t, rec_type, + + uint32_t, rec_node_id, + int32_t, rec_logical_node_id, + int32_t, rec_logical_node_type_id, + uint64_t, rec_gpu_id, + + const char*, rec_name, + const char*, rec_vendor_name, + const char*, rec_product_name, + const char*, rec_model_name + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint32_t, type, rec_type) + + ctf_integer(uint32_t, node_id, rec_node_id) + ctf_integer(int32_t, logical_node_id, rec_logical_node_id) + ctf_integer(int32_t, logical_node_type_id, rec_logical_node_type_id) + ctf_integer(uint64_t, gpu_id, rec_gpu_id) + + ctf_string(name, rec_name) + ctf_string(vendor_name, rec_vendor_name) + ctf_string(product_name, rec_product_name) + ctf_string(model_name, rec_model_name) + ) +) + +// Define the tracepoint event for recording HIP API calls +TRACEPOINT_EVENT( + rocprofv3_trace, + hip_api, + TP_ARGS( + uint64_t, rec_pid, + + const char*, rec_api_name, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + + /* Raw dump of the rocprofiler_hip_api_args_t union */ + const uint8_t*, raw_args_data, // Pointer to the start of the args union + uint64_t, raw_args_size // Size of the args union (sizeof) + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_string(api_name, rec_api_name) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + + /* Raw dump of the args union */ + ctf_sequence(uint8_t, args_raw_payload, raw_args_data, uint64_t, raw_args_size) + ) +) + +// Define the tracepoint event for recording HSA API calls +TRACEPOINT_EVENT( + rocprofv3_trace, + hsa_api, + TP_ARGS( + uint64_t, rec_pid, + + const char*, rec_api_name, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_string(api_name, rec_api_name) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + ) +) + +// Define the tracepoint event for recording Kernel Dispatchs +TRACEPOINT_EVENT( + rocprofv3_trace, + kernel_dispatch, + TP_ARGS( + uint64_t, rec_pid, + + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + uint64_t, rec_agent_id, + uint64_t, rec_queue_id, + uint64_t, rec_stream_id, + + // uint32_t, rec_private_segment_size, + // uint32_t, rec_group_segment_size, + + // uint32_t, rec_workgroup_size_x, + // uint32_t, rec_workgroup_size_y, + // uint32_t, rec_workgroup_size_z, + + // uint32_t, rec_grid_size_x, + // uint32_t, rec_grid_size_y, + // uint32_t, rec_grid_size_z, + + const char*, rec_kernel_name + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, queue_id, rec_queue_id) + ctf_integer(uint64_t, stream_id, rec_stream_id) + + // ctf_integer(uint32_t, private_segment_size, rec_private_segment_size) + // ctf_integer(uint32_t, group_segment_size, rec_group_segment_size) + + // ctf_integer(uint32_t, workgroup_size_x, rec_workgroup_size_x) + // ctf_integer(uint32_t, workgroup_size_y, rec_workgroup_size_y) + // ctf_integer(uint32_t, workgroup_size_z, rec_workgroup_size_z) + + // ctf_integer(uint32_t, grid_size_x, rec_grid_size_x) + // ctf_integer(uint32_t, grid_size_y, rec_grid_size_y) + // ctf_integer(uint32_t, grid_size_z, rec_grid_size_z) + + ctf_string(kernel_name, rec_kernel_name) + ) +) + +// Define the tracepoint event for recording Memory Copy Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + memory_copy, + TP_ARGS( + uint64_t, rec_pid, + + uint32_t, rec_operation, + + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + uint64_t, rec_src_agent_id, + uint64_t, rec_dst_agent_id, + uint64_t, rec_stream_id, + + uint64_t, rec_size + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint32_t, operation, rec_operation) + + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + ctf_integer(uint64_t, src_agent_id, rec_src_agent_id) + ctf_integer(uint64_t, dst_agent_id, rec_dst_agent_id) + ctf_integer(uint64_t, stream_id, rec_stream_id) + + ctf_integer(uint64_t, size, rec_size) + ) +) + +#endif /* _ROCPROFILER_SDK_TRACE_PROVIDER_H */ + +#include // Must be last + +// clang-format on \ No newline at end of file diff --git a/source/lib/output/output_config.cpp b/source/lib/output/output_config.cpp index 7f30f901b6..7a99a78007 100644 --- a/source/lib/output/output_config.cpp +++ b/source/lib/output/output_config.cpp @@ -79,10 +79,10 @@ output_config::parse_env() pftrace_output = entries.count("PFTRACE") > 0; otf2_output = entries.count("OTF2") > 0; rocpd_output = entries.count("ROCPD") > 0 || entries.empty(); - ctf_output = entries.count("CTF") > 0; + lttng_output = entries.count("LTTNG") > 0; const auto supported_formats = - std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD", "CTF"}; + std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD", "LTTNG"}; for(const auto& itr : entries) { LOG_IF(FATAL, supported_formats.count(itr) == 0) diff --git a/source/lib/output/output_config.hpp b/source/lib/output/output_config.hpp index b112fb8dac..d50ef256c0 100644 --- a/source/lib/output/output_config.hpp +++ b/source/lib/output/output_config.hpp @@ -71,7 +71,7 @@ struct output_config bool summary_output = false; bool kernel_rename = false; bool group_by_queue = false; - bool ctf_output = false; + bool lttng_output = false; uint64_t stats_summary_unit_value = 1; size_t perfetto_shmem_size_hint = defaults::perfetto_shmem_size_hint_kb; size_t perfetto_buffer_size = defaults::perfetto_buffer_size_kb; diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index b01656f9cd..4ed888e721 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -46,8 +46,8 @@ #include "lib/output/csv_output_file.hpp" #include "lib/output/domain_type.hpp" #include "lib/output/generateCSV.hpp" -#include "lib/output/generateCTF.hpp" #include "lib/output/generateJSON.hpp" +#include "lib/output/generateLTTng.hpp" #include "lib/output/generateOTF2.hpp" #include "lib/output/generatePerfetto.hpp" #include "lib/output/generateRocpd.hpp" @@ -2464,6 +2464,36 @@ tool_fini(void* /*tool_data*/) counters_output.get_generator()); } + if(tool::get_config().lttng_output && outdata.num_output > 0 && + outdata.num_bytes >= tool::get_config().minimum_output_bytes) + { + auto hip_elem_data = hip_output.load_all(); + auto hsa_elem_data = hsa_output.load_all(); + auto kernel_dispatch_elem_data = kernel_dispatch_output.load_all(); + auto memory_copy_elem_data = memory_copy_output.load_all(); + auto marker_elem_data = marker_output.load_all(); + auto scratch_memory_elem_data = scratch_memory_output.load_all(); + auto rccl_elem_data = rccl_output.load_all(); + auto memory_allocation_elem_data = memory_allocation_output.load_all(); + auto rocdecode_elem_data = rocdecode_output.load_all(); + auto rocjpeg_elem_data = rocjpeg_output.load_all(); + + tool::write_lttng(tool::get_config(), + *tool_metadata, + getpid(), + agents_output, + &hip_elem_data, + &hsa_elem_data, + &kernel_dispatch_elem_data, + &memory_copy_elem_data, + &marker_elem_data, + &scratch_memory_elem_data, + &rccl_elem_data, + &memory_allocation_elem_data, + &rocdecode_elem_data, + &rocjpeg_elem_data); + } + if(tool::get_config().otf2_output && outdata.num_output > 0 && outdata.num_bytes >= tool::get_config().minimum_output_bytes) { From b44e45146442026b534735cebeb3a0baf735b85b Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 4 Jun 2025 11:27:13 -0500 Subject: [PATCH 5/9] Adding support for Marker, RCCL, ROCDecode, ROCJPEG APIs and Memory Allocation Trace Support for LTTng output --- source/lib/output/generateLTTng.cpp | 153 ++++++++---- .../lttng/rocprofiler_sdk_trace_provider.hpp | 226 ++++++++++++++++++ 2 files changed, 337 insertions(+), 42 deletions(-) diff --git a/source/lib/output/generateLTTng.cpp b/source/lib/output/generateLTTng.cpp index 2ec4f2b08d..1cc1518e8f 100644 --- a/source/lib/output/generateLTTng.cpp +++ b/source/lib/output/generateLTTng.cpp @@ -40,7 +40,7 @@ namespace rocprofiler namespace tool { void -write_lttng(const output_config& /*cfg*/, +write_lttng(const output_config& /*cfg*/, const metadata& tool_metadata, uint64_t pid, const std::vector& agent_data, @@ -76,7 +76,7 @@ write_lttng(const output_config& /*cf // Loop over each deque and call write_event for each record if(hip_api_data) { - for (auto hip_api_record : *hip_api_data) + for(auto hip_api_record : *hip_api_data) { // --- Prepare raw args data --- const uint8_t* raw_args_ptr = reinterpret_cast(&hip_api_record.args); @@ -101,7 +101,7 @@ write_lttng(const output_config& /*cf } if(hsa_api_data) { - for (auto hsa_api_record : *hsa_api_data) + for(auto hsa_api_record : *hsa_api_data) { auto api_name = buffer_names.at(hsa_api_record.kind, hsa_api_record.operation); @@ -113,15 +113,16 @@ write_lttng(const output_config& /*cf hsa_api_record.correlation_id.internal, hsa_api_record.start_timestamp, hsa_api_record.end_timestamp, - hsa_api_record.thread_id - ); + hsa_api_record.thread_id); } } if(kernel_dispatch_data) { - for(auto kernel_dispatch_record : *kernel_dispatch_data) { + for(auto kernel_dispatch_record : *kernel_dispatch_data) + { auto name = - tool_metadata.get_kernel_name(kernel_dispatch_record.dispatch_info.kernel_id, kernel_dispatch_record.correlation_id.external.value); + tool_metadata.get_kernel_name(kernel_dispatch_record.dispatch_info.kernel_id, + kernel_dispatch_record.correlation_id.external.value); tracepoint(rocprofv3_trace, kernel_dispatch, @@ -133,13 +134,13 @@ write_lttng(const output_config& /*cf kernel_dispatch_record.dispatch_info.agent_id.handle, kernel_dispatch_record.dispatch_info.queue_id.handle, kernel_dispatch_record.stream_id.handle, - name.data() - ); + name.data()); } } if(memory_copy_data) { - for(auto memory_copy_record : *memory_copy_data) { + for(auto memory_copy_record : *memory_copy_data) + { tracepoint(rocprofv3_trace, memory_copy, pid, @@ -151,57 +152,125 @@ write_lttng(const output_config& /*cf memory_copy_record.src_agent_id.handle, memory_copy_record.dst_agent_id.handle, memory_copy_record.stream_id.handle, - memory_copy_record.bytes - ); + memory_copy_record.bytes); } } if(marker_api_data) { - // ctf_out.write_event_source_component("marker_api_source", - // marker_api_source_init, - // marker_api_source_finalize, - // marker_api_source_next, - // marker_api_data); + for(auto marker_api_record : *marker_api_data) + { + auto name = + (marker_api_record.kind == ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API && + marker_api_record.operation != ROCPROFILER_MARKER_CORE_API_ID_roctxGetThreadId) + ? tool_metadata.get_marker_message(marker_api_record.correlation_id.internal) + : buffer_names.at(marker_api_record.kind, marker_api_record.operation); + tracepoint(rocprofv3_trace, + marker_api, + pid, + marker_api_record.operation, + marker_api_record.correlation_id.ancestor, + marker_api_record.correlation_id.internal, + marker_api_record.start_timestamp, + marker_api_record.end_timestamp, + marker_api_record.thread_id, + name.data()); + } } if(scratch_memory_data) { - // ctf_out.write_event_source_component("scratch_memory_source", - // scratch_memory_source_init, - // scratch_memory_source_finalize, - // scratch_memory_source_next, - // scratch_memory_data); + for(auto scratch_memory_record : *scratch_memory_data) + { + tracepoint(rocprofv3_trace, + scratch_memory, + pid, + scratch_memory_record.operation, + scratch_memory_record.correlation_id.ancestor, + scratch_memory_record.correlation_id.internal, + scratch_memory_record.start_timestamp, + scratch_memory_record.end_timestamp, + scratch_memory_record.thread_id, + scratch_memory_record.agent_id.handle, + scratch_memory_record.queue_id.handle, + scratch_memory_record.flags); + } } if(rccl_api_data) { - // ctf_out.write_event_source_component("rccl_api_source", - // rccl_api_source_init, - // rccl_api_source_finalize, - // rccl_api_source_next, - // rccl_api_data); + for(auto rccl_api_record : *rccl_api_data) + { + auto name = buffer_names.at(rccl_api_record.kind, rccl_api_record.operation); + tracepoint(rocprofv3_trace, + rccl_api, + pid, + rccl_api_record.correlation_id.ancestor, + rccl_api_record.correlation_id.internal, + rccl_api_record.start_timestamp, + rccl_api_record.end_timestamp, + rccl_api_record.thread_id, + name.data()); + } } if(memory_allocation_data) { - // ctf_out.write_event_source_component("memory_allocation_source", - // memory_allocation_source_init, - // memory_allocation_source_finalize, - // memory_allocation_source_next, - // memory_allocation_data); + for(auto memory_allocation_record : *memory_allocation_data) + { + tracepoint(rocprofv3_trace, + memory_allocation, + pid, + memory_allocation_record.operation, + memory_allocation_record.correlation_id.internal, + memory_allocation_record.start_timestamp, + memory_allocation_record.end_timestamp, + memory_allocation_record.thread_id, + memory_allocation_record.agent_id.handle, + memory_allocation_record.stream_id.handle, + memory_allocation_record.address.value, + memory_allocation_record.allocation_size); + } } if(rocdecode_api_data) { - // ctf_out.write_event_source_component("rocdecode_api_source", - // rocdecode_api_source_init, - // rocdecode_api_source_finalize, - // rocdecode_api_source_next, - // rocdecode_api_data); + for(auto rocdecode_api_record : *rocdecode_api_data) + { + auto name = buffer_names.at(rocdecode_api_record.kind, rocdecode_api_record.operation); + + // TODO(aelwazir): Uncomment if sending the data as bytes didn't work + // auto rocdecode_args = sdk::serialization::get_buffer_tracing_args(itr); + + // --- Prepare raw args data --- + const uint8_t* raw_args_ptr = + reinterpret_cast(&rocdecode_api_record.args); + uint64_t raw_args_size_val = static_cast(sizeof(rocdecode_api_record.args)); + + tracepoint(rocprofv3_trace, + rocdecode_api, + pid, + rocdecode_api_record.correlation_id.ancestor, + rocdecode_api_record.correlation_id.internal, + rocdecode_api_record.start_timestamp, + rocdecode_api_record.end_timestamp, + rocdecode_api_record.thread_id, + name.data(), + raw_args_ptr, + raw_args_size_val); + } } if(rocjpeg_api_data) { - // ctf_out.write_event_source_component("rocjpeg_api_source", - // rocjpeg_api_source_init, - // rocjpeg_api_source_finalize, - // rocjpeg_api_source_next, - // rocjpeg_api_data); + for(auto rocjpeg_api_record : *rocjpeg_api_data) + { + auto name = buffer_names.at(rocjpeg_api_record.kind, rocjpeg_api_record.operation); + + tracepoint(rocprofv3_trace, + rocjpeg_api, + pid, + rocjpeg_api_record.correlation_id.ancestor, + rocjpeg_api_record.correlation_id.internal, + rocjpeg_api_record.start_timestamp, + rocjpeg_api_record.end_timestamp, + rocjpeg_api_record.thread_id, + name.data()); + } } } diff --git a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp index f9f260ffb5..533f97301b 100644 --- a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp +++ b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp @@ -217,6 +217,232 @@ TRACEPOINT_EVENT( ) ) +// Define the tracepoint event for recording Marker API Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + marker_api, + TP_ARGS( + uint64_t, rec_pid, + + uint32_t, rec_operation, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + + const char*, rec_name + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint32_t, operation, rec_operation) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + + ctf_string(name, rec_name) + ) +) + +// Define the tracepoint event for recording Scratch Memory Copy Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + scratch_memory, + TP_ARGS( + uint64_t, rec_pid, + + uint32_t, rec_operation, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + uint64_t, rec_agent_id, + uint64_t, rec_queue_id, + + uint32_t, rec_flags + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint32_t, operation, rec_operation) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, dst_queue_id, rec_queue_id) + + ctf_integer(uint64_t, flags, rec_flags) + ) +) + +// Define the tracepoint event for recording RCCL API Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + rccl_api, + TP_ARGS( + uint64_t, rec_pid, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + + const char*, rec_name + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + + ctf_string(name, rec_name) + ) +) + +// Define the tracepoint event for recording Memory Allocation Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + memory_allocation, + TP_ARGS( + uint64_t, rec_pid, + + uint64_t, rec_operation, + + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + uint64_t, rec_agent_id, + uint64_t, rec_stream_id, + + // uint64_t, rec_ptr_address, + uint64_t, rec_ptr_value, + uint64_t, rec_allocation_size + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint64_t, operation, rec_operation) + + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, stream_id, rec_stream_id) + + // ctf_integer(uint64_t, ptr_address, rec_ptr_address) + ctf_integer(uint64_t, ptr_value, rec_ptr_value) + ctf_integer(uint64_t, allocation_size, rec_allocation_size) + ) +) + +// Define the tracepoint event for recording ROCDecode API Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + rocdecode_api, + TP_ARGS( + uint64_t, rec_pid, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + + const char*, rec_name, + + /* Raw dump of the rocprofiler_hip_api_args_t union */ + const uint8_t*, raw_args_data, // Pointer to the start of the args union + uint64_t, raw_args_size // Size of the args union (sizeof) + + // uint64_t, rec_retval + // const char*, rec_retval_str + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + + ctf_string(name, rec_name) + + /* Raw dump of the args union */ + ctf_sequence(uint8_t, args_raw_payload, raw_args_data, uint64_t, raw_args_size) + + // ctf_integer(uint64_t, retval, rec_retval) + // ctf_string(retval_str, rec_retval_str) + ) +) + +// Define the tracepoint event for recording ROCJPEG API Traces +TRACEPOINT_EVENT( + rocprofv3_trace, + rocjpeg_api, + TP_ARGS( + uint64_t, rec_pid, + + uint64_t, rec_ancestor_correlation_id, + uint64_t, rec_internal_correlation_id, + + uint64_t, rec_start_ts, + uint64_t, rec_end_ts, + + uint64_t, rec_thread_id, + + const char*, rec_name + ), + TP_FIELDS( + ctf_integer(uint64_t, pid, rec_pid) + + ctf_integer(uint64_t, ancestor_correlation_id, rec_ancestor_correlation_id) + ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) + + ctf_integer(uint64_t, start_timestamp, rec_start_ts) + ctf_integer(uint64_t, end_timestamp, rec_end_ts) + + ctf_integer(uint64_t, thread_id, rec_thread_id) + + ctf_string(name, rec_name) + ) +) + #endif /* _ROCPROFILER_SDK_TRACE_PROVIDER_H */ #include // Must be last From db7db71c2738efbcd07b417b86580b90b5553031 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 4 Jun 2025 13:20:11 -0500 Subject: [PATCH 6/9] Adding LTTng-UST as external Project --- .github/workflows/continuous_integration.yml | 2 +- .gitmodules | 3 + cmake/rocprofiler_interfaces.cmake | 2 + cmake/rocprofiler_options.cmake | 2 + external/CMakeLists.txt | 58 ++++++++++++++++++++ external/lttng | 1 + source/lib/output/CMakeLists.txt | 3 +- source/lib/output/lttng/CMakeLists.txt | 10 ++-- 8 files changed, 74 insertions(+), 7 deletions(-) create mode 160000 external/lttng diff --git a/.github/workflows/continuous_integration.yml b/.github/workflows/continuous_integration.yml index b191454815..ca5a7ef87a 100644 --- a/.github/workflows/continuous_integration.yml +++ b/.github/workflows/continuous_integration.yml @@ -70,7 +70,7 @@ jobs: run: | git config --global --add safe.directory '*' apt-get update - apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev libsqlite3-dev rccl-dev rccl-unittests rocjpeg-dev rocjpeg-test rocdecode-dev rocdecode-test + apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev libsqlite3-dev rccl-dev rccl-unittests rocjpeg-dev rocjpeg-test rocdecode-dev rocdecode-test liblttng-ust-dev lttng-tools update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11 --slave /usr/bin/gcov gcov /usr/bin/gcov-11 update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12 --slave /usr/bin/gcov gcov /usr/bin/gcov-12 python3 -m pip install -U --user -r requirements.txt diff --git a/.gitmodules b/.gitmodules index 7bbab5a350..8ff1394edf 100644 --- a/.gitmodules +++ b/.gitmodules @@ -40,3 +40,6 @@ [submodule "external/gotcha"] path = external/gotcha url = https://jrmadsen@github.com/jrmadsen/GOTCHA +[submodule "external/lttng"] + path = external/lttng + url = https://github.com/lttng/lttng-ust.git diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index df604b80ff..49b6f88334 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -24,6 +24,8 @@ rocprofiler_add_interface_library(rocprofiler-sdk-cereal "Enables Cereal support INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-sqlite3 "Enables SQLite3 support" INTERNAL) +rocprofiler_add_interface_library(rocprofiler-sdk-lttng "Enables LTTng support" + INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-pybind11 "Enables PyBind11 support" INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-gotcha "Enables GOTCHA support" diff --git a/cmake/rocprofiler_options.cmake b/cmake/rocprofiler_options.cmake index cbc17a961d..5cd5fb7d33 100644 --- a/cmake/rocprofiler_options.cmake +++ b/cmake/rocprofiler_options.cmake @@ -61,6 +61,8 @@ rocprofiler_add_option(ROCPROFILER_BUILD_GLOG "Enable building glog (Google logging) library internally" ON) rocprofiler_add_option(ROCPROFILER_BUILD_SQLITE3 "Enable building sqlite3 library internally" OFF) +rocprofiler_add_option(ROCPROFILER_BUILD_LTTNG + "Enable building lttng-ust library internally" OFF) rocprofiler_add_option(ROCPROFILER_BUILD_PYBIND11 "Enable building pybind11 library internally" ON) rocprofiler_add_option(ROCPROFILER_BUILD_GOTCHA diff --git a/external/CMakeLists.txt b/external/CMakeLists.txt index 4756b81774..519e55f323 100644 --- a/external/CMakeLists.txt +++ b/external/CMakeLists.txt @@ -278,6 +278,64 @@ else() target_link_libraries(rocprofiler-sdk-sqlite3 INTERFACE SQLite::SQLite3) endif() +# +# LTTng-UST +# +if(ROCPROFILER_BUILD_LTTNG) + # checkout submodule if not already checked out or clone repo if no .gitmodules file + rocprofiler_checkout_git_submodule( + RECURSIVE + RELATIVE_PATH external/lttng + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + TEST_FILE bootstrap + REPO_URL https://github.com/lttng/lttng-ust + REPO_BRANCH "stable-2.13") + + find_program( + MAKE_COMMAND + NAMES make gmake + PATH_SUFFIXES bin REQUIRED) + + execute_process( + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}/external/lttng + COMMAND ./bootstrap + RESULT_VARIABLE RET) + + if(RET GREATER 0) + message(STATUS "Bootstrap command for LTTng-UST failed.") + message(FATAL_ERROR "Command: ./bootstrap") + endif() + + include(ExternalProject) + externalproject_add( + rocprofiler-sdk-lttng-build + PREFIX ${PROJECT_BINARY_DIR}/external/lttng/build + SOURCE_DIR ${PROJECT_SOURCE_DIR}/external/lttng + BUILD_IN_SOURCE 0 + CONFIGURE_COMMAND + /configure --prefix=${PROJECT_BINARY_DIR}/external/lttng/install + --libdir=${PROJECT_BINARY_DIR}/external/lttng/install/lib + --disable-man-pages --with-pic CFLAGS=-O3\ -g1 + BUILD_COMMAND ${MAKE_COMMAND} -j ${CMAKE_BUILD_PARALLEL_LEVEL} install -s + INSTALL_COMMAND "") + + target_link_libraries( + rocprofiler-sdk-lttng + INTERFACE + $ + ) + target_include_directories( + rocprofiler-sdk-lttng SYSTEM + INTERFACE $ + ) + add_dependencies(rocprofiler-sdk-lttng rocprofiler-sdk-lttng-build) +else() + find_package(LTTngUST REQUIRED) + target_include_directories(rocprofiler-sdk-lttng + SYSTEM INTERFACE ${LTTNGUST_INCLUDE_DIRS}) + target_link_libraries(rocprofiler-sdk-lttng INTERFACE ${LTTNGUST_LIBRARIES}) +endif() + # # PyBind11 # diff --git a/external/lttng b/external/lttng new file mode 160000 index 0000000000..e8306a0e0e --- /dev/null +++ b/external/lttng @@ -0,0 +1 @@ +Subproject commit e8306a0e0e8c694251418dc43bf8df92d50c3de5 diff --git a/source/lib/output/CMakeLists.txt b/source/lib/output/CMakeLists.txt index b3c17f7f7f..e419fd4e19 100644 --- a/source/lib/output/CMakeLists.txt +++ b/source/lib/output/CMakeLists.txt @@ -72,8 +72,7 @@ target_link_libraries( rocprofiler-sdk::rocprofiler-sdk-amd-comgr rocprofiler-sdk::rocprofiler-sdk-dw rocprofiler-sdk::rocprofiler-sdk-elf - rocprofiler-sdk::rocprofiler-sdk-sqlite3 - rocprofiler-sdk-tool-output-lttng) + rocprofiler-sdk::rocprofiler-sdk-sqlite3) target_compile_definitions(rocprofiler-sdk-output-library PRIVATE PROJECT_BINARY_DIR="${PROJECT_BINARY_DIR}") diff --git a/source/lib/output/lttng/CMakeLists.txt b/source/lib/output/lttng/CMakeLists.txt index 8fa7921eee..461a1fe616 100644 --- a/source/lib/output/lttng/CMakeLists.txt +++ b/source/lib/output/lttng/CMakeLists.txt @@ -1,9 +1,11 @@ -find_package(LTTngUST REQUIRED) - add_library(rocprofiler-sdk-tool-output-lttng STATIC rocprofiler_sdk_trace_provider.cpp rocprofiler_sdk_trace_provider.hpp) target_include_directories(rocprofiler-sdk-tool-output-lttng - PUBLIC ${LTTNGUST_INCLUDE_DIRS} ${CMAKE_CURRENT_SOURCE_DIR}) + PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) + +target_link_libraries(rocprofiler-sdk-tool-output-lttng rocprofiler-sdk::rocprofiler-sdk-lttng dl) -target_link_libraries(rocprofiler-sdk-tool-output-lttng ${LTTNGUST_LIBRARIES} dl) +target_link_libraries( + rocprofiler-sdk-output-library + PRIVATE rocprofiler-sdk-tool-output-lttng) From 07e04cf0ed939761171785d36f5bb4439722f84f Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 4 Jun 2025 19:07:15 -0500 Subject: [PATCH 7/9] Fixing structures and sending readable data --- cmake/rocprofiler_interfaces.cmake | 3 +- external/CMakeLists.txt | 11 +- source/lib/output/generateLTTng.cpp | 105 ++++++++++++++---- source/lib/output/lttng/CMakeLists.txt | 14 ++- .../lttng/rocprofiler_sdk_trace_provider.hpp | 66 ++++++----- 5 files changed, 130 insertions(+), 69 deletions(-) diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index 49b6f88334..37fb844e56 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -24,8 +24,7 @@ rocprofiler_add_interface_library(rocprofiler-sdk-cereal "Enables Cereal support INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-sqlite3 "Enables SQLite3 support" INTERNAL) -rocprofiler_add_interface_library(rocprofiler-sdk-lttng "Enables LTTng support" - INTERNAL) +rocprofiler_add_interface_library(rocprofiler-sdk-lttng "Enables LTTng support" INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-pybind11 "Enables PyBind11 support" INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-gotcha "Enables GOTCHA support" diff --git a/external/CMakeLists.txt b/external/CMakeLists.txt index 519e55f323..7e14b7570d 100644 --- a/external/CMakeLists.txt +++ b/external/CMakeLists.txt @@ -314,8 +314,8 @@ if(ROCPROFILER_BUILD_LTTNG) BUILD_IN_SOURCE 0 CONFIGURE_COMMAND /configure --prefix=${PROJECT_BINARY_DIR}/external/lttng/install - --libdir=${PROJECT_BINARY_DIR}/external/lttng/install/lib - --disable-man-pages --with-pic CFLAGS=-O3\ -g1 + --libdir=${PROJECT_BINARY_DIR}/external/lttng/install/lib --disable-man-pages + --with-pic CFLAGS=-O3\ -g1 BUILD_COMMAND ${MAKE_COMMAND} -j ${CMAKE_BUILD_PARALLEL_LEVEL} install -s INSTALL_COMMAND "") @@ -326,13 +326,12 @@ if(ROCPROFILER_BUILD_LTTNG) ) target_include_directories( rocprofiler-sdk-lttng SYSTEM - INTERFACE $ - ) + INTERFACE $) add_dependencies(rocprofiler-sdk-lttng rocprofiler-sdk-lttng-build) else() find_package(LTTngUST REQUIRED) - target_include_directories(rocprofiler-sdk-lttng - SYSTEM INTERFACE ${LTTNGUST_INCLUDE_DIRS}) + target_include_directories(rocprofiler-sdk-lttng SYSTEM + INTERFACE ${LTTNGUST_INCLUDE_DIRS}) target_link_libraries(rocprofiler-sdk-lttng INTERFACE ${LTTNGUST_LIBRARIES}) endif() diff --git a/source/lib/output/generateLTTng.cpp b/source/lib/output/generateLTTng.cpp index 1cc1518e8f..183b60a139 100644 --- a/source/lib/output/generateLTTng.cpp +++ b/source/lib/output/generateLTTng.cpp @@ -39,6 +39,30 @@ namespace rocprofiler { namespace tool { +struct args_info +{ + std::string type = {}; + std::string value = {}; +}; +int +iterate_args_callback(rocprofiler_buffer_tracing_kind_t /*kind*/, + rocprofiler_tracing_operation_t /*operation*/, + uint32_t /*arg_number*/, + const void* const /*arg_value_addr*/, + int32_t /*arg_indirection_count*/, + const char* arg_type, + const char* arg_name, + const char* arg_value_str, + void* data) +{ + ROCP_FATAL_IF(data == nullptr) << "nullptr to data for iterate_args_callback"; + + auto* _data = static_cast*>(data); + if(arg_type && arg_name && arg_value_str) + _data->emplace_back(args_info{arg_name, arg_value_str}); + return 0; +} + void write_lttng(const output_config& /*cfg*/, const metadata& tool_metadata, @@ -58,7 +82,7 @@ write_lttng(const output_config& /*cfg*/, auto buffer_names = sdk::get_buffer_tracing_names(); auto callbk_name_info = sdk::get_callback_tracing_names(); - for(auto& _agent_info : agent_data) + for(const auto& _agent_info : agent_data) { tracepoint(rocprofv3_trace, agents_info, @@ -78,13 +102,27 @@ write_lttng(const output_config& /*cfg*/, { for(auto hip_api_record : *hip_api_data) { - // --- Prepare raw args data --- - const uint8_t* raw_args_ptr = reinterpret_cast(&hip_api_record.args); - uint64_t raw_args_size_val = static_cast( - sizeof(hip_api_record.args)); // Or sizeof(rocprofiler_hip_api_args_t) - auto api_name = buffer_names.at(hip_api_record.kind, hip_api_record.operation); + std::vector args_types = {}; + std::vector args_values = {}; + std::vector args = {}; + { + auto _record = rocprofiler_record_header_t{ + .hash = rocprofiler_record_header_compute_hash( + ROCPROFILER_BUFFER_CATEGORY_TRACING, hip_api_record.kind), + .payload = &hip_api_record}; + + rocprofiler_iterate_buffer_tracing_record_args( + _record, iterate_args_callback, &args); + + for(const auto& arg : args) + { + args_types.push_back(arg.type.c_str()); + args_values.push_back(arg.value.c_str()); + } + } + tracepoint(rocprofv3_trace, hip_api, pid, @@ -94,9 +132,9 @@ write_lttng(const output_config& /*cfg*/, hip_api_record.start_timestamp, hip_api_record.end_timestamp, hip_api_record.thread_id, - raw_args_ptr, - raw_args_size_val // Sending args as bytes - ); + args_types.data(), + args_values.data(), + args.size()); } } if(hsa_api_data) @@ -123,6 +161,8 @@ write_lttng(const output_config& /*cfg*/, auto name = tool_metadata.get_kernel_name(kernel_dispatch_record.dispatch_info.kernel_id, kernel_dispatch_record.correlation_id.external.value); + const auto& agent = + tool_metadata.get_agent(kernel_dispatch_record.dispatch_info.agent_id); tracepoint(rocprofv3_trace, kernel_dispatch, @@ -131,7 +171,7 @@ write_lttng(const output_config& /*cfg*/, kernel_dispatch_record.start_timestamp, kernel_dispatch_record.end_timestamp, kernel_dispatch_record.thread_id, - kernel_dispatch_record.dispatch_info.agent_id.handle, + agent->node_id, kernel_dispatch_record.dispatch_info.queue_id.handle, kernel_dispatch_record.stream_id.handle, name.data()); @@ -141,6 +181,9 @@ write_lttng(const output_config& /*cfg*/, { for(auto memory_copy_record : *memory_copy_data) { + const auto& src_agent = tool_metadata.get_agent(memory_copy_record.src_agent_id); + const auto& dst_agent = tool_metadata.get_agent(memory_copy_record.dst_agent_id); + tracepoint(rocprofv3_trace, memory_copy, pid, @@ -149,8 +192,8 @@ write_lttng(const output_config& /*cfg*/, memory_copy_record.start_timestamp, memory_copy_record.end_timestamp, memory_copy_record.thread_id, - memory_copy_record.src_agent_id.handle, - memory_copy_record.dst_agent_id.handle, + src_agent->node_id, + dst_agent->node_id, memory_copy_record.stream_id.handle, memory_copy_record.bytes); } @@ -180,6 +223,8 @@ write_lttng(const output_config& /*cfg*/, { for(auto scratch_memory_record : *scratch_memory_data) { + const auto& agent = tool_metadata.get_agent(scratch_memory_record.agent_id); + tracepoint(rocprofv3_trace, scratch_memory, pid, @@ -189,7 +234,7 @@ write_lttng(const output_config& /*cfg*/, scratch_memory_record.start_timestamp, scratch_memory_record.end_timestamp, scratch_memory_record.thread_id, - scratch_memory_record.agent_id.handle, + agent->node_id, scratch_memory_record.queue_id.handle, scratch_memory_record.flags); } @@ -214,6 +259,8 @@ write_lttng(const output_config& /*cfg*/, { for(auto memory_allocation_record : *memory_allocation_data) { + const auto& agent = tool_metadata.get_agent(memory_allocation_record.agent_id); + tracepoint(rocprofv3_trace, memory_allocation, pid, @@ -222,9 +269,9 @@ write_lttng(const output_config& /*cfg*/, memory_allocation_record.start_timestamp, memory_allocation_record.end_timestamp, memory_allocation_record.thread_id, - memory_allocation_record.agent_id.handle, + agent->node_id, memory_allocation_record.stream_id.handle, - memory_allocation_record.address.value, + memory_allocation_record.address, memory_allocation_record.allocation_size); } } @@ -234,13 +281,24 @@ write_lttng(const output_config& /*cfg*/, { auto name = buffer_names.at(rocdecode_api_record.kind, rocdecode_api_record.operation); - // TODO(aelwazir): Uncomment if sending the data as bytes didn't work - // auto rocdecode_args = sdk::serialization::get_buffer_tracing_args(itr); + std::vector args_types = {}; + std::vector args_values = {}; + std::vector args = {}; + { + auto _record = rocprofiler_record_header_t{ + .hash = rocprofiler_record_header_compute_hash( + ROCPROFILER_BUFFER_CATEGORY_TRACING, rocdecode_api_record.kind), + .payload = &rocdecode_api_record}; + + rocprofiler_iterate_buffer_tracing_record_args( + _record, iterate_args_callback, &args); - // --- Prepare raw args data --- - const uint8_t* raw_args_ptr = - reinterpret_cast(&rocdecode_api_record.args); - uint64_t raw_args_size_val = static_cast(sizeof(rocdecode_api_record.args)); + for(const auto& arg : args) + { + args_types.push_back(arg.type.c_str()); + args_values.push_back(arg.value.c_str()); + } + } tracepoint(rocprofv3_trace, rocdecode_api, @@ -251,8 +309,9 @@ write_lttng(const output_config& /*cfg*/, rocdecode_api_record.end_timestamp, rocdecode_api_record.thread_id, name.data(), - raw_args_ptr, - raw_args_size_val); + args_types.data(), + args_values.data(), + args.size()); } } if(rocjpeg_api_data) diff --git a/source/lib/output/lttng/CMakeLists.txt b/source/lib/output/lttng/CMakeLists.txt index 461a1fe616..d3c652b239 100644 --- a/source/lib/output/lttng/CMakeLists.txt +++ b/source/lib/output/lttng/CMakeLists.txt @@ -4,8 +4,14 @@ add_library(rocprofiler-sdk-tool-output-lttng STATIC rocprofiler_sdk_trace_provi target_include_directories(rocprofiler-sdk-tool-output-lttng PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) -target_link_libraries(rocprofiler-sdk-tool-output-lttng rocprofiler-sdk::rocprofiler-sdk-lttng dl) - target_link_libraries( - rocprofiler-sdk-output-library - PRIVATE rocprofiler-sdk-tool-output-lttng) + rocprofiler-sdk-tool-output-lttng + PRIVATE rocprofiler-sdk::rocprofiler-sdk-headers + rocprofiler-sdk::rocprofiler-sdk-build-flags + rocprofiler-sdk::rocprofiler-sdk-memcheck + rocprofiler-sdk::rocprofiler-sdk-common-library + rocprofiler-sdk::rocprofiler-sdk-lttng + dl) + +target_link_libraries(rocprofiler-sdk-output-library + PRIVATE rocprofiler-sdk-tool-output-lttng) diff --git a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp index 533f97301b..bae7a97e1d 100644 --- a/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp +++ b/source/lib/output/lttng/rocprofiler_sdk_trace_provider.hpp @@ -10,6 +10,8 @@ #if !defined(_ROCPROFILER_SDK_TRACE_PROVIDER_H) || defined(TRACEPOINT_HEADER_MULTI_READ) #define _ROCPROFILER_SDK_TRACE_PROVIDER_H +#include + #include // Define the tracepoint event for recording the agents @@ -65,9 +67,9 @@ TRACEPOINT_EVENT( uint64_t, rec_thread_id, - /* Raw dump of the rocprofiler_hip_api_args_t union */ - const uint8_t*, raw_args_data, // Pointer to the start of the args union - uint64_t, raw_args_size // Size of the args union (sizeof) + const char**, rec_args_types, + const char**, rec_args_values, + size_t, rec_args_len ), TP_FIELDS( ctf_integer(uint64_t, pid, rec_pid) @@ -78,12 +80,13 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) - /* Raw dump of the args union */ - ctf_sequence(uint8_t, args_raw_payload, raw_args_data, uint64_t, raw_args_size) + ctf_sequence(uint8_t, args_types_payload, (const uint8_t*)&rec_args_types, uint64_t, rec_args_len*sizeof(const char*)) + ctf_sequence(uint8_t, args_values_payload, (const uint8_t*)&rec_args_values, uint64_t, rec_args_len*sizeof(const char*)) + ctf_integer(size_t, args_count, rec_args_len) ) ) @@ -113,7 +116,7 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) ) @@ -155,10 +158,10 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) - ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, agent_node_id, rec_agent_id) ctf_integer(uint64_t, queue_id, rec_queue_id) ctf_integer(uint64_t, stream_id, rec_stream_id) @@ -206,11 +209,11 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) - ctf_integer(uint64_t, src_agent_id, rec_src_agent_id) - ctf_integer(uint64_t, dst_agent_id, rec_dst_agent_id) + ctf_integer(uint64_t, src_agent_node_id, rec_src_agent_id) + ctf_integer(uint64_t, dst_agent_node_id, rec_dst_agent_id) ctf_integer(uint64_t, stream_id, rec_stream_id) ctf_integer(uint64_t, size, rec_size) @@ -245,7 +248,7 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) @@ -283,10 +286,10 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) - ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, agent_node_id, rec_agent_id) ctf_integer(uint64_t, dst_queue_id, rec_queue_id) ctf_integer(uint64_t, flags, rec_flags) @@ -317,7 +320,7 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) @@ -344,7 +347,7 @@ TRACEPOINT_EVENT( uint64_t, rec_stream_id, // uint64_t, rec_ptr_address, - uint64_t, rec_ptr_value, + rocprofiler_address_t, rec_address, uint64_t, rec_allocation_size ), TP_FIELDS( @@ -355,14 +358,14 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) - ctf_integer(uint64_t, agent_id, rec_agent_id) + ctf_integer(uint64_t, agent_node_id, rec_agent_id) ctf_integer(uint64_t, stream_id, rec_stream_id) - // ctf_integer(uint64_t, ptr_address, rec_ptr_address) - ctf_integer(uint64_t, ptr_value, rec_ptr_value) + ctf_integer(uint64_t, ptr_value, rec_address.value) + ctf_sequence(uint8_t, ptr, (const uint8_t*)(rec_address.ptr), uint64_t, sizeof(rec_address.ptr)) ctf_integer(uint64_t, allocation_size, rec_allocation_size) ) ) @@ -384,12 +387,9 @@ TRACEPOINT_EVENT( const char*, rec_name, - /* Raw dump of the rocprofiler_hip_api_args_t union */ - const uint8_t*, raw_args_data, // Pointer to the start of the args union - uint64_t, raw_args_size // Size of the args union (sizeof) - - // uint64_t, rec_retval - // const char*, rec_retval_str + const char**, rec_args_types, + const char**, rec_args_values, + size_t, rec_args_len ), TP_FIELDS( ctf_integer(uint64_t, pid, rec_pid) @@ -397,17 +397,15 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) ctf_string(name, rec_name) - /* Raw dump of the args union */ - ctf_sequence(uint8_t, args_raw_payload, raw_args_data, uint64_t, raw_args_size) - - // ctf_integer(uint64_t, retval, rec_retval) - // ctf_string(retval_str, rec_retval_str) + ctf_sequence(uint8_t, args_types_payload, (const uint8_t*)&rec_args_types, uint64_t, rec_args_len*sizeof(const char*)) + ctf_sequence(uint8_t, args_values_payload, (const uint8_t*)&rec_args_values, uint64_t, rec_args_len*sizeof(const char*)) + ctf_integer(size_t, args_count, rec_args_len) ) ) @@ -435,7 +433,7 @@ TRACEPOINT_EVENT( ctf_integer(uint64_t, internal_correlation_id, rec_internal_correlation_id) ctf_integer(uint64_t, start_timestamp, rec_start_ts) - ctf_integer(uint64_t, end_timestamp, rec_end_ts) + ctf_integer(uint64_t, duration, rec_end_ts-rec_start_ts) ctf_integer(uint64_t, thread_id, rec_thread_id) From 42c9eeef4635946a93b64368c36488ae5f2c13ad Mon Sep 17 00:00:00 2001 From: "Elwazir, Ammar" Date: Wed, 4 Jun 2025 19:15:04 -0500 Subject: [PATCH 8/9] Update source/docs/conceptual/comparing-with-legacy-tools.rst --- source/docs/conceptual/comparing-with-legacy-tools.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/docs/conceptual/comparing-with-legacy-tools.rst b/source/docs/conceptual/comparing-with-legacy-tools.rst index 214cebadda..195a2af5df 100644 --- a/source/docs/conceptual/comparing-with-legacy-tools.rst +++ b/source/docs/conceptual/comparing-with-legacy-tools.rst @@ -328,7 +328,7 @@ ROCprofiler-SDK introduces a new command-line tool, `rocprofv3`, which is a more * - I/O options - Output Formats - CSV, JSON (Chrome-Tracing format) - - CSV, JSON (Chrome-Tracing format), Perfetto, LTTng + - CSV, JSON (Chrome-Tracing format), Perfetto, CTF - CSV, JSON (custom schema), Perfetto, OTF2 - | # Multiple output formats can be supported in single run. | # OTF2 can visualize larger trace files compared to perfetto. From aea1fa3bb7d8a2d4baadb48973c6d35884bdf471 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Wed, 18 Jun 2025 17:27:33 -0500 Subject: [PATCH 9/9] Adding Post Processing CTF2 output using barectf --- source/lib/output/output_config.cpp | 3 +- source/lib/output/output_config.hpp | 2 + source/lib/python/rocpd/README.md | 2 +- source/lib/python/rocpd/__init__.py | 24 + source/lib/python/rocpd/__main__.py | 9 +- source/lib/python/rocpd/ctf2.py | 120 ++ source/lib/python/rocpd/libpyrocpd.cpp | 153 +++ source/lib/python/rocpd/output_config.py | 2 +- source/lib/python/rocpd/source/CMakeLists.txt | 9 +- .../rocpd/source/ctf/barectf-bitfield.h | 144 +++ source/lib/python/rocpd/source/ctf/barectf.c | 1104 +++++++++++++++++ source/lib/python/rocpd/source/ctf/barectf.h | 182 +++ .../lib/python/rocpd/source/ctf/barectf.yml | 128 ++ source/lib/python/rocpd/source/ctf/metadata | 296 +++++ source/lib/python/rocpd/source/ctf2.cpp | 382 ++++++ source/lib/python/rocpd/source/ctf2.hpp | 111 ++ source/lib/python/utilities.cmake | 1 + 17 files changed, 2664 insertions(+), 8 deletions(-) create mode 100644 source/lib/python/rocpd/ctf2.py create mode 100644 source/lib/python/rocpd/source/ctf/barectf-bitfield.h create mode 100644 source/lib/python/rocpd/source/ctf/barectf.c create mode 100644 source/lib/python/rocpd/source/ctf/barectf.h create mode 100644 source/lib/python/rocpd/source/ctf/barectf.yml create mode 100644 source/lib/python/rocpd/source/ctf/metadata create mode 100644 source/lib/python/rocpd/source/ctf2.cpp create mode 100644 source/lib/python/rocpd/source/ctf2.hpp diff --git a/source/lib/output/output_config.cpp b/source/lib/output/output_config.cpp index 7a99a78007..64e49e8b7f 100644 --- a/source/lib/output/output_config.cpp +++ b/source/lib/output/output_config.cpp @@ -78,11 +78,12 @@ output_config::parse_env() json_output = entries.count("JSON") > 0; pftrace_output = entries.count("PFTRACE") > 0; otf2_output = entries.count("OTF2") > 0; + ctf2_output = entries.count("CTF2") > 0; rocpd_output = entries.count("ROCPD") > 0 || entries.empty(); lttng_output = entries.count("LTTNG") > 0; const auto supported_formats = - std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD", "LTTNG"}; + std::set{"CSV", "JSON", "PFTRACE", "OTF2", "ROCPD", "LTTNG", "CTF2"}; for(const auto& itr : entries) { LOG_IF(FATAL, supported_formats.count(itr) == 0) diff --git a/source/lib/output/output_config.hpp b/source/lib/output/output_config.hpp index d50ef256c0..c904ea76de 100644 --- a/source/lib/output/output_config.hpp +++ b/source/lib/output/output_config.hpp @@ -67,6 +67,7 @@ struct output_config bool json_output = false; bool pftrace_output = false; bool otf2_output = false; + bool ctf2_output = false; bool rocpd_output = false; bool summary_output = false; bool kernel_rename = false; @@ -130,6 +131,7 @@ output_config::save(ArchiveT& ar) const CFG_SERIALIZE_MEMBER(json_output); CFG_SERIALIZE_MEMBER(pftrace_output); CFG_SERIALIZE_MEMBER(otf2_output); + CFG_SERIALIZE_MEMBER(ctf2_output); CFG_SERIALIZE_MEMBER(summary_output); CFG_SERIALIZE_MEMBER(rocpd_output); CFG_SERIALIZE_MEMBER(kernel_rename); diff --git a/source/lib/python/rocpd/README.md b/source/lib/python/rocpd/README.md index fe97b2575d..8ad0f2ffd4 100644 --- a/source/lib/python/rocpd/README.md +++ b/source/lib/python/rocpd/README.md @@ -6,7 +6,7 @@ collected with the ROCm profiling tools suite. ## Background In the past, the ROCm profiling tools (e.g. rocprofv3, rocprofiler-systems, etc.) have directly written data to -various output formats such as CSV, JSON, Perfetto, OTF2, etc. This approach has a significant number of flaws: +various output formats such as CSV, JSON, Perfetto, OTF2, CTF2, etc. This approach has a significant number of flaws: ### No standardization in the CSV and JSON output formats diff --git a/source/lib/python/rocpd/__init__.py b/source/lib/python/rocpd/__init__.py index bd8b171833..ccc5783fc3 100644 --- a/source/lib/python/rocpd/__init__.py +++ b/source/lib/python/rocpd/__init__.py @@ -45,6 +45,7 @@ "write_perfetto", "write_csv", "write_otf2", + "write_ctf2", "RocpdImportData", ] @@ -147,3 +148,26 @@ def write_otf2(connection, config=None, **kwargs): ) return libpyrocpd.write_otf2(connection, config) + +def write_ctf2(connection, config=None, **kwargs): + """ + Write CTF@ output file + + Args: + connection (rocpd.RocpdImportData): + rocPD instance of database connection(s) + config (rocpd.output_config.output_config): + Output specification + + Returns: + bool: returns True if successful + """ + from . import output_config + + config = ( + output_config.output_config(**kwargs) + if config is None + else config.update(**kwargs) + ) + + return libpyrocpd.write_ctf2(connection, config) diff --git a/source/lib/python/rocpd/__main__.py b/source/lib/python/rocpd/__main__.py index a056d497dd..8ae54634ed 100644 --- a/source/lib/python/rocpd/__main__.py +++ b/source/lib/python/rocpd/__main__.py @@ -43,6 +43,7 @@ def main(argv=None, config=None): from . import pftrace from . import csv from . import otf2 + from . import ctf2 from .importer import RocpdImportData convert_examples = """ @@ -94,10 +95,10 @@ def get_output_type(val): required_params.add_argument( "-f", "--output-format", - help="For adding output format (supported formats: csv, pftrace, otf2)", + help="For adding output format (supported formats: csv, pftrace, otf2, ctf2)", nargs="+", default=None, - choices=("csv", "pftrace", "otf2"), + choices=("csv", "pftrace", "otf2", "ctf2"), type=get_output_type, required=True, ) @@ -108,6 +109,7 @@ def get_output_type(val): valid_pftrace_args = pftrace.add_args(converter) valid_csv_args = csv.add_args(converter) valid_otf2_args = otf2.add_args(converter) + valid_ctf2_args = ctf2.add_args(converter) valid_time_window_args = time_window.add_args(converter) # parse the command line arguments @@ -119,6 +121,7 @@ def get_output_type(val): pftrace_args = pftrace.process_args(args, valid_pftrace_args) csv_args = csv.process_args(args, valid_csv_args) otf2_args = otf2.process_args(args, valid_otf2_args) + ctf2_args = ctf2.process_args(args, valid_ctf2_args) window_args = time_window.process_args(args, valid_time_window_args) # now start processing the data. Import the data and merge the views @@ -134,6 +137,7 @@ def get_output_type(val): **pftrace_args, **csv_args, **otf2_args, + **ctf2_args, } # setup the config args config = ( @@ -147,6 +151,7 @@ def get_output_type(val): "pftrace": pftrace.write_pftrace, "csv": csv.write_csv, "otf2": otf2.write_otf2, + "ctf2": ctf2.write_ctf2, } for out_format in args.output_format: diff --git a/source/lib/python/rocpd/ctf2.py b/source/lib/python/rocpd/ctf2.py new file mode 100644 index 0000000000..d09b890330 --- /dev/null +++ b/source/lib/python/rocpd/ctf2.py @@ -0,0 +1,120 @@ +#!/usr/bin/env python3 +############################################################################### +# MIT License +# +# Copyright (c) 2023 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. +############################################################################### + +from .importer import RocpdImportData +from .time_window import apply_time_window +from . import output_config +from . import libpyrocpd + + +def write_ctf2(importData, config): + return libpyrocpd.write_ctf2(importData, config) + + +def execute(input, config=None, window_args=None, **kwargs): + + importData = RocpdImportData(input) + + apply_time_window(importData, **window_args) + + config = ( + output_config.output_config(**kwargs) + if config is None + else config.update(**kwargs) + ) + + write_ctf2(importData, config) + + +def add_args(parser): + """Add ctf2 arguments.""" + + # Currently, no ctf2 specific args + + # ctf2_options = parser.add_argument_group("CTF2 options") + + # ctf2_options.add_argument( + # "--kernel-rename", + # help="Use kernel names from debugging symbols if available", + # action="store_true", + # default=False, + # ) + + return [] + + +def process_args(args, valid_args): + + ret = {} + for itr in valid_args: + if hasattr(args, itr): + val = getattr(args, itr) + if val is not None: + ret[itr] = val + return ret + + +def main(argv=None): + import argparse + from .time_window import add_args as add_args_time_window + from .time_window import process_args as process_args_time_window + from .output_config import add_args as add_args_output_config + from .output_config import process_args as process_args_output_config + from .output_config import add_generic_args, process_generic_args + + parser = argparse.ArgumentParser( + description="Convert rocPD to CTF2 format", allow_abbrev=False + ) + + required_params = parser.add_argument_group("Required arguments") + + required_params.add_argument( + "-i", + "--input", + required=True, + type=output_config.check_file_exists, + nargs="+", + help="Input path and filename to one or more database(s), separated by spaces", + ) + + valid_out_config_args = add_args_output_config(parser) + valid_ctf2_args = add_args(parser) + valid_generic_args = add_generic_args(parser) + valid_time_window_args = add_args_time_window(parser) + + args = parser.parse_args(argv) + + out_cfg_args = process_args_output_config(args, valid_out_config_args) + generic_out_cfg_args = process_generic_args(args, valid_generic_args) + window_args = process_args_time_window(args, valid_time_window_args) + ctf2_args = process_args(args, valid_ctf2_args) + + all_args = {**out_cfg_args, **ctf2_args, **generic_out_cfg_args} + + execute(args.input, window_args=window_args, **all_args) + + +if __name__ == "__main__": + main() diff --git a/source/lib/python/rocpd/libpyrocpd.cpp b/source/lib/python/rocpd/libpyrocpd.cpp index 0f62325724..1bb8b703bd 100644 --- a/source/lib/python/rocpd/libpyrocpd.cpp +++ b/source/lib/python/rocpd/libpyrocpd.cpp @@ -27,6 +27,7 @@ #include "lib/python/rocpd/source/functions.hpp" #include "lib/python/rocpd/source/interop.hpp" #include "lib/python/rocpd/source/otf2.hpp" +#include "lib/python/rocpd/source/ctf2.hpp" #include "lib/python/rocpd/source/perfetto.hpp" #include "lib/python/rocpd/source/serialization/sql.hpp" #include "lib/python/rocpd/source/sql_generator.hpp" @@ -282,6 +283,7 @@ PYBIND11_MODULE(libpyrocpd, pyrocpd) .def_readwrite("csv", &tool::output_config::csv_output) .def_readwrite("pftrace", &tool::output_config::pftrace_output) .def_readwrite("otf2", &tool::output_config::otf2_output) + .def_readwrite("ctf2", &tool::output_config::ctf2_output) .def_readwrite("kernel_rename", &tool::output_config::kernel_rename) .def_readwrite("agent_index_value", &tool::output_config::agent_index_value) .def_readwrite("group_by_queue", &tool::output_config::group_by_queue) @@ -762,6 +764,157 @@ PYBIND11_MODULE(libpyrocpd, pyrocpd) }, "Write OTF2 output file from rocpd SQLite3 database"); + pyrocpd.def( + "write_ctf2", + [](rocpd::RocpdImportData& data, const tool::output_config& output_cfg) { + auto _create_agent_index = + [&output_cfg](const rocpd::types::agent& _agent) -> tool::agent_index { + auto ret_index = tool::create_agent_index( + output_cfg.agent_index_value, + _agent.node_id, // absolute index + static_cast(_agent.logical_node_id), // relative index + static_cast(_agent.logical_node_type_id), // type-relative index + std::string_view(_agent.type)); + return ret_index; + }; + + constexpr auto kernels_order_by = + "agent_abs_index ASC, stream_id ASC, queue_id ASC, start ASC, end DESC"; + + // to initialise the CTF@ session properly we need to know: + // (1) the process with the earliest start time + // (2) find the process with the longest duration + uint64_t min_start_time = std::numeric_limits::max(); + uint64_t max_fini_time = 0; + for(auto obj : {data.connection}) + { + auto* conn = rocpd::interop::get_connection(std::move(obj)); + + // min start + sqlite3_stmt* _stmt_min_start; + sqlite3_prepare_v2( + conn, "SELECT MIN(start) FROM processes;", -1, &_stmt_min_start, nullptr); + uint64_t _min_start_time = std::numeric_limits::max(); + if(sqlite3_step(_stmt_min_start) == SQLITE_ROW) + { + _min_start_time = + static_cast(sqlite3_column_int64(_stmt_min_start, 0)); + } + + sqlite3_finalize(_stmt_min_start); + if(min_start_time > _min_start_time) + { + min_start_time = _min_start_time; + } + //// max fini + sqlite3_stmt* _stmt_max_fini; + sqlite3_prepare_v2( + conn, "SELECT MAX(fini) FROM processes;", -1, &_stmt_max_fini, nullptr); + uint64_t _max_fini_time = 0; + if(sqlite3_step(_stmt_max_fini) == SQLITE_ROW) + { + _max_fini_time = static_cast(sqlite3_column_int64(_stmt_max_fini, 0)); + } + + sqlite3_finalize(_stmt_max_fini); + if(max_fini_time < _max_fini_time) + { + max_fini_time = _max_fini_time; + } + } + + auto ctf2_session = + rocpd::output::CTF2Session(output_cfg, min_start_time, max_fini_time); + + auto sqlgen_ctf2 = common::simple_timer{ + fmt::format("CTF2 generation from {} SQL database(s)", data.size())}; + + uint16_t _process_counter = 0; + for(auto obj : {data.connection}) + { + auto* conn = rocpd::interop::get_connection(std::move(obj)); + auto nodes = rocpd::read(conn); + for(const auto& nitr : nodes) + { + auto agents = rocpd::read( + conn, fmt::format("WHERE guid = '{}' AND nid = {}", nitr.guid, nitr.id)); + auto processes = rocpd::read( + conn, fmt::format("WHERE guid = '{}' AND nid = {}", nitr.guid, nitr.id)); + + // absolute_index |-> (agent, agent_index) + auto agents_map = std::unordered_map{}; + + for(const auto& itr : agents) + { + const rocprofiler::tool::agent_index new_index = _create_agent_index(itr); + const std::string labeled_name = fmt::format("{}", itr.name); + agents_map.emplace( + itr.absolute_index, + rocpd::output::extended_agent_ctf{itr, new_index, labeled_name}); + } + + for(const auto& pitr : processes) + { + ROCP_FATAL_IF(pitr.nid != nitr.id || pitr.guid != nitr.guid) + << fmt::format("Found process with a mismatched nid/guid. process: " + "{}/{} vs. node: {}/{}", + pitr.nid, + pitr.guid, + nitr.id, + nitr.guid); + + auto select_guid_nid_pid = + [&nitr, &pitr](std::string_view tbl, + std::string_view where_extra_condition = "") { + return fmt::format("SELECT * FROM {} WHERE guid = '{}' AND " + "nid = {} AND pid = {} {}", + tbl, + pitr.guid, + nitr.id, + pitr.pid, + where_extra_condition); + }; + + constexpr auto region_order_by = "start ASC, end DESC"; + + auto _sqlgen_ctf2 = common::simple_timer{fmt::format( + "CTF2 generation from SQL for process {} (total)", pitr.pid)}; + + auto kernels = rocpd::sql_generator{ + conn, select_guid_nid_pid("kernels"), kernels_order_by}; + + auto memory_allocations = + rocpd::sql_generator{ + conn, select_guid_nid_pid("memory_allocations"), region_order_by}; + + auto memory_copies = rocpd::sql_generator{ + conn, select_guid_nid_pid("memory_copies"), region_order_by}; + + auto regions = rocpd::sql_generator{ + conn, select_guid_nid_pid("regions"), region_order_by}; + + auto threads = rocpd::sql_generator{ + conn, select_guid_nid_pid("threads")}; + + ROCP_TRACE << "Starting CTF2 generation from SQL for process " << pitr.pid; + auto _sqlgen_perfw = common::simple_timer{fmt::format( + "CTF2 generation from SQL for process {} (write)", pitr.pid)}; + rocpd::output::write_ctf2(ctf2_session, + pitr, + _process_counter, + agents_map, + threads, + regions, + kernels, + memory_copies, + memory_allocations); + _process_counter++; + } + } + } + }, + "Write CTF2 output file from rocpd SQLite3 database"); + // NOLINTEND(performance-unnecessary-value-param) // reads in all the agent info from database diff --git a/source/lib/python/rocpd/output_config.py b/source/lib/python/rocpd/output_config.py index d0877384c1..c5ef3956e4 100644 --- a/source/lib/python/rocpd/output_config.py +++ b/source/lib/python/rocpd/output_config.py @@ -160,7 +160,7 @@ def add_generic_args(parser): generic_options.add_argument( "--agent-index-value", choices=("absolute", "relative", "type-relative"), - help="""Device identification format in CSV/Perfetto/OTF2 output (default: relative): + help="""Device identification format in CSV/Perfetto/OTF2/CTF2 output (default: relative): absolute: uses node_id (Agent-0, Agent-2, Agent-4) ignoring cgroups restrictions. relative: uses logical_node_id (Agent-0, Agent-1, Agent-2) considering cgroups restrictions. type-relative: uses logical_node_type_id (CPU-0, GPU-0, GPU-1) with numbering that resets for each device type.""", diff --git a/source/lib/python/rocpd/source/CMakeLists.txt b/source/lib/python/rocpd/source/CMakeLists.txt index 99a6927b43..252c736b92 100644 --- a/source/lib/python/rocpd/source/CMakeLists.txt +++ b/source/lib/python/rocpd/source/CMakeLists.txt @@ -9,12 +9,15 @@ set(libpyrocpd_source_headers perfetto.hpp csv.hpp otf2.hpp + ctf2.hpp sql_generator.hpp pysqlite_Connection.h - types.hpp) + types.hpp + ctf/barectf.h + ctf/barectf-bitfield.h) -set(libpyrocpd_source_sources csv.cpp functions.cpp interop.cpp otf2.cpp perfetto.cpp - types.cpp) +set(libpyrocpd_source_sources csv.cpp functions.cpp interop.cpp otf2.cpp ctf2.cpp perfetto.cpp + types.cpp ctf/barectf.c) foreach(_PYTHON_VERSION ${ROCPROFILER_PYTHON_VERSIONS}) rocprofiler_rocpd_python_bindings_target_sources( diff --git a/source/lib/python/rocpd/source/ctf/barectf-bitfield.h b/source/lib/python/rocpd/source/ctf/barectf-bitfield.h new file mode 100644 index 0000000000..bc77c939b3 --- /dev/null +++ b/source/lib/python/rocpd/source/ctf/barectf-bitfield.h @@ -0,0 +1,144 @@ +#ifndef _BARECTF_BITFIELD_H +#define _BARECTF_BITFIELD_H + +#include + +/* + * BabelTrace + * + * Bitfields read/write functions. + * + * Copyright (c) 2010-2020 Mathieu Desnoyers + * + * 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. + * + * - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + * + * The following code was generated by barectf v3.1.2 + * on 2025-06-18T16:33:18.853411. + * + * For more details, see . + */ + +#include + +#ifdef __cplusplus +# define _CAST_PTR(_type, _value) \ + static_cast<_type>(static_cast(_value)) +#else +# define _CAST_PTR(_type, _value) ((void *) (_value)) +#endif + +/* We can't shift a int from 32 bit, >> 32 and << 32 on int is undefined */ +#define _bt_piecewise_rshift(_vtype, _v, _shift) \ +do { \ + unsigned long ___shift = (_shift); \ + unsigned long sb = (___shift) / (sizeof(_v) * CHAR_BIT - 1); \ + unsigned long final = (___shift) % (sizeof(_v) * CHAR_BIT - 1); \ + \ + for (; sb; sb--) \ + _v >>= sizeof(_v) * CHAR_BIT - 1; \ + _v >>= final; \ +} while (0) + +/* + * bt_bitfield_write - write integer to a bitfield in native endianness + * + * Save integer to the bitfield, which starts at the "start" bit, has "len" + * bits. + * The inside of a bitfield is from high bits to low bits. + * Uses native endianness. + * For unsigned "v", pad MSB with 0 if bitfield is larger than v. + * For signed "v", sign-extend v if bitfield is larger than v. + * + * On little endian, bytes are placed from the less significant to the most + * significant. Also, consecutive bitfields are placed from lower bits to higher + * bits. + * + * On big endian, bytes are places from most significant to less significant. + * Also, consecutive bitfields are placed from higher to lower bits. + */ + +/* Trace byte order: little-endian */ + +#define _bt_bitfield_write_le(_ptr, type, _start, _length, _vtype, _v) \ +do { \ + _vtype __v = (_v); \ + type *__ptr = _CAST_PTR(type *, _ptr); \ + unsigned long __start = (_start), __length = (_length); \ + type mask, cmask; \ + unsigned long ts = sizeof(type) * CHAR_BIT; /* type size */ \ + unsigned long start_unit, end_unit, this_unit; \ + unsigned long end, cshift; /* cshift is "complement shift" */ \ + \ + if (!__length) \ + break; \ + \ + end = __start + __length; \ + start_unit = __start / ts; \ + end_unit = (end + (ts - 1)) / ts; \ + \ + /* Trim v high bits */ \ + if (__length < sizeof(__v) * CHAR_BIT) \ + __v &= ~((~(_vtype) 0) << __length); \ + \ + /* We can now append v with a simple "or", shift it piece-wise */ \ + this_unit = start_unit; \ + if (start_unit == end_unit - 1) { \ + mask = ~((~(type) 0) << (__start % ts)); \ + if (end % ts) \ + mask |= (~(type) 0) << (end % ts); \ + cmask = (type) __v << (__start % ts); \ + cmask &= ~mask; \ + __ptr[this_unit] &= mask; \ + __ptr[this_unit] |= cmask; \ + break; \ + } \ + if (__start % ts) { \ + cshift = __start % ts; \ + mask = ~((~(type) 0) << cshift); \ + cmask = (type) __v << cshift; \ + cmask &= ~mask; \ + __ptr[this_unit] &= mask; \ + __ptr[this_unit] |= cmask; \ + _bt_piecewise_rshift(_vtype, __v, ts - cshift); \ + __start += ts - cshift; \ + this_unit++; \ + } \ + for (; this_unit < end_unit - 1; this_unit++) { \ + __ptr[this_unit] = (type) __v; \ + _bt_piecewise_rshift(_vtype, __v, ts); \ + __start += ts; \ + } \ + if (end % ts) { \ + mask = (~(type) 0) << (end % ts); \ + cmask = (type) __v; \ + cmask &= ~mask; \ + __ptr[this_unit] &= mask; \ + __ptr[this_unit] |= cmask; \ + } else \ + __ptr[this_unit] = (type) __v; \ +} while (0) + +#define bt_bitfield_write_le(ptr, _start, _length, _vtype, _v) \ + _bt_bitfield_write_le(ptr, uint8_t, _start, _length, _vtype, _v) + +#endif /* _BARECTF_BITFIELD_H */ diff --git a/source/lib/python/rocpd/source/ctf/barectf.c b/source/lib/python/rocpd/source/ctf/barectf.c new file mode 100644 index 0000000000..a050b690cc --- /dev/null +++ b/source/lib/python/rocpd/source/ctf/barectf.c @@ -0,0 +1,1104 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2020 Philippe Proulx + * + * 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. + * + * - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + * + * The following code was generated by barectf v3.1.2 + * on 2025-06-18T16:33:18.853411. + * + * For more details, see . + */ + +#include +#include +#include + +#include "barectf.h" +#include "barectf-bitfield.h" + +#define _ALIGN(_at_var, _align) \ + do { \ + (_at_var) = ((_at_var) + ((_align) - 1)) & -(_align); \ + } while (0) + +#ifdef __cplusplus +# define _TO_VOID_PTR(_value) static_cast(_value) +# define _FROM_VOID_PTR(_type, _value) static_cast<_type *>(_value) +#else +# define _TO_VOID_PTR(_value) ((void *) (_value)) +# define _FROM_VOID_PTR(_type, _value) ((_type *) (_value)) +#endif + +#define _BITS_TO_BYTES(_x) ((_x) >> 3) +#define _BYTES_TO_BITS(_x) ((_x) << 3) + +union _f2u { + float f; + uint32_t u; +}; + +union _d2u { + double f; + uint64_t u; +}; + +uint32_t barectf_packet_size(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->packet_size; +} + +int barectf_packet_is_full(const void * const vctx) +{ + const struct barectf_ctx * const ctx = _FROM_VOID_PTR(const struct barectf_ctx, vctx); + + return ctx->at == ctx->packet_size; +} + +int barectf_packet_is_empty(const void * const vctx) +{ + const struct barectf_ctx * const ctx = _FROM_VOID_PTR(const struct barectf_ctx, vctx); + + return ctx->at <= ctx->off_content; +} + +uint32_t barectf_packet_events_discarded(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->events_discarded; +} + +uint32_t barectf_discarded_event_records_count(const void * const vctx) +{ + return barectf_packet_events_discarded(vctx); +} + +uint32_t barectf_packet_sequence_number(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->sequence_number; +} + +uint8_t *barectf_packet_buf(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->buf; +} + +uint8_t *barectf_packet_buf_addr(const void * const vctx) +{ + return barectf_packet_buf(vctx); +} + +uint32_t barectf_packet_buf_size(const void * const vctx) +{ + const struct barectf_ctx * const ctx = _FROM_VOID_PTR(const struct barectf_ctx, vctx); + + return _BITS_TO_BYTES(ctx->packet_size); +} + +void barectf_packet_set_buf(void * const vctx, uint8_t * const buf, + const uint32_t buf_size) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + ctx->buf = buf; + + if (ctx->at == ctx->packet_size) { + /* Keep full packet state */ + ctx->at = _BYTES_TO_BITS(buf_size); + } + + ctx->packet_size = _BYTES_TO_BITS(buf_size); +} + +int barectf_packet_is_open(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->packet_is_open; +} + +int barectf_is_in_tracing_section(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->in_tracing_section; +} + +volatile const int *barectf_is_in_tracing_section_ptr(const void * const vctx) +{ + return &_FROM_VOID_PTR(const struct barectf_ctx, vctx)->in_tracing_section; +} + +int barectf_is_tracing_enabled(const void * const vctx) +{ + return _FROM_VOID_PTR(const struct barectf_ctx, vctx)->is_tracing_enabled; +} + +void barectf_enable_tracing(void * const vctx, const int enable) +{ + _FROM_VOID_PTR(struct barectf_ctx, vctx)->is_tracing_enabled = enable; +} + +static +void _write_c_str(struct barectf_ctx * const ctx, const char * const src) +{ + const uint32_t sz = strlen(src) + 1; + + memcpy(&ctx->buf[_BITS_TO_BYTES(ctx->at)], src, sz); + ctx->at += _BYTES_TO_BITS(sz); +} + +static +int _reserve_er_space(void * const vctx, const uint32_t er_size) +{ + int ret; + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + /* Event _cannot_ fit? */ + if (er_size > (ctx->packet_size - ctx->off_content)) { + goto no_space; + } + + /* Packet is full? */ + if (barectf_packet_is_full(ctx)) { + /* Yes: is the back end full? */ + if (ctx->cbs.is_backend_full(ctx->data)) { + /* Yes: discard event record */ + goto no_space; + } + + /* Back-end is _not_ full: open new packet */ + ctx->use_cur_last_event_ts = 1; + ctx->cbs.open_packet(ctx->data); + // ctx->use_cur_last_event_ts = 0; + } + + /* Event fits the current packet? */ + if (er_size > (ctx->packet_size - ctx->at)) { + /* No: close packet now */ + ctx->use_cur_last_event_ts = 1; + ctx->cbs.close_packet(ctx->data); + // ctx->use_cur_last_event_ts = 0; + + /* Is the back end full? */ + if (ctx->cbs.is_backend_full(ctx->data)) { + /* Yes: discard event record */ + goto no_space; + } + + /* Back-end is _not_ full: open new packet */ + ctx->use_cur_last_event_ts = 1; + ctx->cbs.open_packet(ctx->data); + // ctx->use_cur_last_event_ts = 0; + assert(er_size <= (ctx->packet_size - ctx->at)); + } + + ret = 1; + goto end; + +no_space: + ctx->events_discarded++; + ret = 0; + +end: + return ret; +} + +static +void _commit_er(void * const vctx) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + /* Is the packet full? */ + if (barectf_packet_is_full(ctx)) { + /* Yes: close it now */ + ctx->cbs.close_packet(ctx->data); + } +} + +/* Initialize context */ +void barectf_init(void *vctx, + uint8_t * const buf, const uint32_t buf_size, + const struct barectf_platform_callbacks cbs, void * const data) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + ctx->cbs = cbs; + ctx->data = data; + ctx->buf = buf; + ctx->packet_size = _BYTES_TO_BITS(buf_size); + ctx->at = 0; + ctx->events_discarded = 0; + ctx->sequence_number = 0; + ctx->packet_is_open = 0; + ctx->in_tracing_section = 0; + ctx->is_tracing_enabled = 1; + ctx->use_cur_last_event_ts = 1; +} + +/* Open packet for data stream type `default` */ +void barectf_default_open_packet( + struct barectf_default_ctx * const sctx) +{ + struct barectf_ctx * const ctx = &sctx->parent; + const int saved_in_tracing_section = ctx->in_tracing_section; + + /* + * This function is either called by a tracing function, or + * directly by the platform. + * + * If it's called by a tracing function, then + * `ctx->in_tracing_section` is 1, so it's safe to open + * the packet here (alter the packet), even if tracing was + * disabled in the meantime because we're already in a tracing + * section (which finishes at the end of the tracing function + * call). + * + * If it's called directly by the platform, then if tracing is + * disabled, we don't want to alter the packet, and return + * immediately. + */ + if (!ctx->is_tracing_enabled && !saved_in_tracing_section) { + ctx->in_tracing_section = 0; + goto end; + } + + /* We can alter the packet */ + ctx->in_tracing_section = 1; + + /* Do not open a packet that is already open */ + if (ctx->packet_is_open) { + ctx->in_tracing_section = saved_in_tracing_section; + goto end; + } + + ctx->at = 0; + + /* Write packet header structure */ + { + /* Align for packet header structure */ + _ALIGN(ctx->at, 8); + } + + /* Write packet context structure */ + { + /* Align for packet context structure */ + _ALIGN(ctx->at, 8); + + /* Align for `packet_size` field */ + _ALIGN(ctx->at, 8); + + /* Write packet total size field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) ctx->packet_size); + ctx->at += 64; + + /* Align for `content_size` field */ + _ALIGN(ctx->at, 8); + + /* Do not write `content_size` field; save its offset */ + sctx->off_pc_content_size = ctx->at; + ctx->at += 64; + + /* Align for `events_discarded` field */ + _ALIGN(ctx->at, 8); + + /* Do not write `events_discarded` field; save its offset */ + sctx->off_pc_events_discarded = ctx->at; + ctx->at += 16; + } + + /* Save content beginning's offset */ + ctx->off_content = ctx->at; + + /* Mark current packet as open */ + ctx->packet_is_open = 1; + + /* Not tracing anymore */ + ctx->in_tracing_section = saved_in_tracing_section; + +end: + return; +} + +/* Close packet for data stream type `default` */ +void barectf_default_close_packet(struct barectf_default_ctx * const sctx) +{ + struct barectf_ctx * const ctx = &sctx->parent; + const int saved_in_tracing_section = ctx->in_tracing_section; + + /* + * This function is either called by a tracing function, or + * directly by the platform. + * + * If it's called by a tracing function, then + * `ctx->in_tracing_section` is 1, so it's safe to close + * the packet here (alter the packet), even if tracing was + * disabled in the meantime, because we're already in a tracing + * section (which finishes at the end of the tracing function + * call). + * + * If it's called directly by the platform, then if tracing is + * disabled, we don't want to alter the packet, and return + * immediately. + */ + if (!ctx->is_tracing_enabled && !saved_in_tracing_section) { + ctx->in_tracing_section = 0; + goto end; + } + + /* We can alter the packet */ + ctx->in_tracing_section = 1; + + /* Do not close a packet that is not open */ + if (!ctx->packet_is_open) { + ctx->in_tracing_section = saved_in_tracing_section; + goto end; + } + + /* Save content size */ + ctx->content_size = ctx->at; + + /* Go back to `content_size` field offset */ + ctx->at = sctx->off_pc_content_size; + + /* Write `content_size` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) ctx->content_size); + ctx->at += 64; + + /* Go back to `events_discarded` field offset */ + ctx->at = sctx->off_pc_events_discarded; + + /* Write `events_discarded` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 16, + uint16_t, (uint16_t) ctx->events_discarded); + ctx->at += 16; + + /* Go back to end of packet */ + ctx->at = ctx->packet_size; + + /* Mark packet as closed */ + ctx->packet_is_open = 0; + + /* Not tracing anymore */ + ctx->in_tracing_section = saved_in_tracing_section; + +end: + return; +} + +static void _serialize_er_header_default(void * const vctx, + const uint32_t ert_id, uint64_t timestamp) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + struct barectf_default_ctx * const sctx = _FROM_VOID_PTR(struct barectf_default_ctx, vctx); + const uint32_t ts = timestamp; + + /* Write header structure */ + { + /* Align for header structure */ + _ALIGN(ctx->at, 8); + + /* Align for `id` field */ + _ALIGN(ctx->at, 8); + + /* Write event record type ID field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 8, + uint8_t, (uint8_t) ert_id); + ctx->at += 8; + + /* Align for `timestamp` field */ + _ALIGN(ctx->at, 8); + + /* Write timestamp field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) ts); + ctx->at += 64; + } +} + +static void _serialize_er_default_api(void * const vctx, + const uint8_t p_phase, + const uint64_t p_pid, + const char * const p_name, + const uint64_t p_ancestor_correlation_id, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + /* Serialize header */ + _serialize_er_header_default(ctx, 0, p_timestamp); + + /* Write payload structure */ + { + /* Align for payload structure */ + _ALIGN(ctx->at, 8); + + /* Align for `phase` field */ + _ALIGN(ctx->at, 8); + + /* Write `phase` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 8, + uint8_t, (uint8_t) p_phase); + ctx->at += 8; + + /* Align for `pid` field */ + _ALIGN(ctx->at, 8); + + /* Write `pid` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_pid); + ctx->at += 64; + + /* Align for `name` field */ + _ALIGN(ctx->at, 8); + + /* Write `name` field */ + _write_c_str(ctx, p_name); + + /* Align for `ancestor_correlation_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `ancestor_correlation_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_ancestor_correlation_id); + ctx->at += 64; + + /* Align for `internal_correlation_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `internal_correlation_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_internal_correlation_id); + ctx->at += 64; + + /* Align for `timestamp` field */ + _ALIGN(ctx->at, 8); + + /* Write `timestamp` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_timestamp); + ctx->at += 64; + + /* Align for `thread_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `thread_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_thread_id); + ctx->at += 64; + } +} + +static void _serialize_er_default_kernel_dispatch(void * const vctx, + const uint8_t p_phase, + const uint64_t p_pid, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id, + const uint64_t p_agent_node_id, + const uint64_t p_queue_id, + const uint64_t p_stream_id, + const char * const p_kernel_name) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + /* Serialize header */ + _serialize_er_header_default(ctx, 1, p_timestamp); + + /* Write payload structure */ + { + /* Align for payload structure */ + _ALIGN(ctx->at, 8); + + /* Align for `phase` field */ + _ALIGN(ctx->at, 8); + + /* Write `phase` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 8, + uint8_t, (uint8_t) p_phase); + ctx->at += 8; + + /* Align for `pid` field */ + _ALIGN(ctx->at, 8); + + /* Write `pid` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_pid); + ctx->at += 64; + + /* Align for `internal_correlation_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `internal_correlation_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_internal_correlation_id); + ctx->at += 64; + + /* Align for `timestamp` field */ + _ALIGN(ctx->at, 8); + + /* Write `timestamp` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_timestamp); + ctx->at += 64; + + /* Align for `thread_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `thread_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_thread_id); + ctx->at += 64; + + /* Align for `agent_node_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `agent_node_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_agent_node_id); + ctx->at += 64; + + /* Align for `queue_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `queue_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_queue_id); + ctx->at += 64; + + /* Align for `stream_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `stream_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_stream_id); + ctx->at += 64; + + /* Align for `kernel_name` field */ + _ALIGN(ctx->at, 8); + + /* Write `kernel_name` field */ + _write_c_str(ctx, p_kernel_name); + } +} + +static void _serialize_er_default_memory_copy(void * const vctx, + const uint8_t p_phase, + const uint64_t p_pid, + const char * const p_operation, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id, + const uint64_t p_src_agent_node_id, + const uint64_t p_dst_agent_node_id, + const uint64_t p_stream_id, + const uint64_t p_size) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + + /* Serialize header */ + _serialize_er_header_default(ctx, 2, p_timestamp); + + /* Write payload structure */ + { + /* Align for payload structure */ + _ALIGN(ctx->at, 8); + + /* Align for `phase` field */ + _ALIGN(ctx->at, 8); + + /* Write `phase` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 8, + uint8_t, (uint8_t) p_phase); + ctx->at += 8; + + /* Align for `pid` field */ + _ALIGN(ctx->at, 8); + + /* Write `pid` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_pid); + ctx->at += 64; + + /* Align for `operation` field */ + _ALIGN(ctx->at, 8); + + /* Write `operation` field */ + _write_c_str(ctx, p_operation); + + /* Align for `internal_correlation_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `internal_correlation_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_internal_correlation_id); + ctx->at += 64; + + /* Align for `timestamp` field */ + _ALIGN(ctx->at, 8); + + /* Write `timestamp` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_timestamp); + ctx->at += 64; + + /* Align for `thread_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `thread_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_thread_id); + ctx->at += 64; + + /* Align for `src_agent_node_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `src_agent_node_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_src_agent_node_id); + ctx->at += 64; + + /* Align for `dst_agent_node_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `dst_agent_node_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_dst_agent_node_id); + ctx->at += 64; + + /* Align for `stream_id` field */ + _ALIGN(ctx->at, 8); + + /* Write `stream_id` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_stream_id); + ctx->at += 64; + + /* Align for `size` field */ + _ALIGN(ctx->at, 8); + + /* Write `size` field */ + bt_bitfield_write_le(&ctx->buf[_BITS_TO_BYTES(ctx->at)], 0, 64, + uint64_t, (uint64_t) p_size); + ctx->at += 64; + } +} + +static uint32_t _er_size_default_api(void * const vctx, + const char * const p_name) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + uint32_t at = ctx->at; + + /* Add header structure size */ + { + /* Align for header structure */ + _ALIGN(at, 8); + + /* Align for `id` field */ + _ALIGN(at, 8); + + /* Add `id` bit array field's size */ + at += 8; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + } + + /* Add payload structure size */ + { + /* Align for payload structure */ + _ALIGN(at, 8); + + /* Align for `phase` field */ + _ALIGN(at, 8); + + /* Add `phase` bit array field's size */ + at += 8; + + /* Align for `pid` field */ + _ALIGN(at, 8); + + /* Add `pid` bit array field's size */ + at += 64; + + /* Align for `name` field */ + _ALIGN(at, 8); + + /* Add `name` string field's size */ + at += _BYTES_TO_BITS(strlen(p_name) + 1); + + /* Align for `ancestor_correlation_id` field */ + _ALIGN(at, 8); + + /* Add `ancestor_correlation_id` bit array field's size */ + at += 64; + + /* Align for `internal_correlation_id` field */ + _ALIGN(at, 8); + + /* Add `internal_correlation_id` bit array field's size */ + at += 64; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + + /* Align for `thread_id` field */ + _ALIGN(at, 8); + + /* Add `thread_id` bit array field's size */ + at += 64; + } + + return at - ctx->at; +} + +static uint32_t _er_size_default_kernel_dispatch(void * const vctx, + const char * const p_kernel_name) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + uint32_t at = ctx->at; + + /* Add header structure size */ + { + /* Align for header structure */ + _ALIGN(at, 8); + + /* Align for `id` field */ + _ALIGN(at, 8); + + /* Add `id` bit array field's size */ + at += 8; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + } + + /* Add payload structure size */ + { + /* Align for payload structure */ + _ALIGN(at, 8); + + /* Align for `phase` field */ + _ALIGN(at, 8); + + /* Add `phase` bit array field's size */ + at += 8; + + /* Align for `pid` field */ + _ALIGN(at, 8); + + /* Add `pid` bit array field's size */ + at += 64; + + /* Align for `internal_correlation_id` field */ + _ALIGN(at, 8); + + /* Add `internal_correlation_id` bit array field's size */ + at += 64; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + + /* Align for `thread_id` field */ + _ALIGN(at, 8); + + /* Add `thread_id` bit array field's size */ + at += 64; + + /* Align for `agent_node_id` field */ + _ALIGN(at, 8); + + /* Add `agent_node_id` bit array field's size */ + at += 64; + + /* Align for `queue_id` field */ + _ALIGN(at, 8); + + /* Add `queue_id` bit array field's size */ + at += 64; + + /* Align for `stream_id` field */ + _ALIGN(at, 8); + + /* Add `stream_id` bit array field's size */ + at += 64; + + /* Align for `kernel_name` field */ + _ALIGN(at, 8); + + /* Add `kernel_name` string field's size */ + at += _BYTES_TO_BITS(strlen(p_kernel_name) + 1); + } + + return at - ctx->at; +} + +static uint32_t _er_size_default_memory_copy(void * const vctx, + const char * const p_operation) +{ + struct barectf_ctx * const ctx = _FROM_VOID_PTR(struct barectf_ctx, vctx); + uint32_t at = ctx->at; + + /* Add header structure size */ + { + /* Align for header structure */ + _ALIGN(at, 8); + + /* Align for `id` field */ + _ALIGN(at, 8); + + /* Add `id` bit array field's size */ + at += 8; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + } + + /* Add payload structure size */ + { + /* Align for payload structure */ + _ALIGN(at, 8); + + /* Align for `phase` field */ + _ALIGN(at, 8); + + /* Add `phase` bit array field's size */ + at += 8; + + /* Align for `pid` field */ + _ALIGN(at, 8); + + /* Add `pid` bit array field's size */ + at += 64; + + /* Align for `operation` field */ + _ALIGN(at, 8); + + /* Add `operation` string field's size */ + at += _BYTES_TO_BITS(strlen(p_operation) + 1); + + /* Align for `internal_correlation_id` field */ + _ALIGN(at, 8); + + /* Add `internal_correlation_id` bit array field's size */ + at += 64; + + /* Align for `timestamp` field */ + _ALIGN(at, 8); + + /* Add `timestamp` bit array field's size */ + at += 64; + + /* Align for `thread_id` field */ + _ALIGN(at, 8); + + /* Add `thread_id` bit array field's size */ + at += 64; + + /* Align for `src_agent_node_id` field */ + _ALIGN(at, 8); + + /* Add `src_agent_node_id` bit array field's size */ + at += 64; + + /* Align for `dst_agent_node_id` field */ + _ALIGN(at, 8); + + /* Add `dst_agent_node_id` bit array field's size */ + at += 64; + + /* Align for `stream_id` field */ + _ALIGN(at, 8); + + /* Add `stream_id` bit array field's size */ + at += 64; + + /* Align for `size` field */ + _ALIGN(at, 8); + + /* Add `size` bit array field's size */ + at += 64; + } + + return at - ctx->at; +} + +/* Trace (data stream type `default`, event record type `api`) */ +void barectf_default_trace_api(struct barectf_default_ctx * const sctx, + const uint8_t p_phase, + const uint64_t p_pid, + const char * const p_name, + const uint64_t p_ancestor_correlation_id, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id) +{ + struct barectf_ctx * const ctx = &sctx->parent; + uint32_t er_size; + + /* Save timestamp */ + sctx->cur_last_event_ts = p_timestamp; + + if (!ctx->is_tracing_enabled) { + goto end; + } + + /* We can alter the packet */ + ctx->in_tracing_section = 1; + + /* Compute event record size */ + er_size = _er_size_default_api(_TO_VOID_PTR(ctx), p_name); + + /* Is there enough space to serialize? */ + if (!_reserve_er_space(_TO_VOID_PTR(ctx), er_size)) { + /* no: forget this */ + ctx->in_tracing_section = 0; + goto end; + } + + /* Serialize event record */ + _serialize_er_default_api(_TO_VOID_PTR(ctx), p_phase, p_pid, p_name, p_ancestor_correlation_id, p_internal_correlation_id, p_timestamp, p_thread_id); + + /* Commit event record */ + _commit_er(_TO_VOID_PTR(ctx)); + + /* Not tracing anymore */ + ctx->in_tracing_section = 0; + +end: + return; +} + +/* Trace (data stream type `default`, event record type `kernel_dispatch`) */ +void barectf_default_trace_kernel_dispatch(struct barectf_default_ctx * const sctx, + const uint8_t p_phase, + const uint64_t p_pid, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id, + const uint64_t p_agent_node_id, + const uint64_t p_queue_id, + const uint64_t p_stream_id, + const char * const p_kernel_name) +{ + struct barectf_ctx * const ctx = &sctx->parent; + uint32_t er_size; + + /* Save timestamp */ + sctx->cur_last_event_ts = p_timestamp; + + if (!ctx->is_tracing_enabled) { + goto end; + } + + /* We can alter the packet */ + ctx->in_tracing_section = 1; + + /* Compute event record size */ + er_size = _er_size_default_kernel_dispatch(_TO_VOID_PTR(ctx), p_kernel_name); + + /* Is there enough space to serialize? */ + if (!_reserve_er_space(_TO_VOID_PTR(ctx), er_size)) { + /* no: forget this */ + ctx->in_tracing_section = 0; + goto end; + } + + sctx->cur_last_event_ts = p_timestamp; + + /* Serialize event record */ + _serialize_er_default_kernel_dispatch(_TO_VOID_PTR(ctx), p_phase, p_pid, p_internal_correlation_id, p_timestamp, p_thread_id, p_agent_node_id, p_queue_id, p_stream_id, p_kernel_name); + + /* Commit event record */ + _commit_er(_TO_VOID_PTR(ctx)); + + /* Not tracing anymore */ + ctx->in_tracing_section = 0; + +end: + return; +} + +/* Trace (data stream type `default`, event record type `memory_copy`) */ +void barectf_default_trace_memory_copy(struct barectf_default_ctx * const sctx, + const uint8_t p_phase, + const uint64_t p_pid, + const char * const p_operation, + const uint64_t p_internal_correlation_id, + const uint64_t p_timestamp, + const uint64_t p_thread_id, + const uint64_t p_src_agent_node_id, + const uint64_t p_dst_agent_node_id, + const uint64_t p_stream_id, + const uint64_t p_size) +{ + struct barectf_ctx * const ctx = &sctx->parent; + uint32_t er_size; + + /* Save timestamp */ + sctx->cur_last_event_ts = p_timestamp; + + if (!ctx->is_tracing_enabled) { + goto end; + } + + /* We can alter the packet */ + ctx->in_tracing_section = 1; + + /* Compute event record size */ + er_size = _er_size_default_memory_copy(_TO_VOID_PTR(ctx), p_operation); + + /* Is there enough space to serialize? */ + if (!_reserve_er_space(_TO_VOID_PTR(ctx), er_size)) { + /* no: forget this */ + ctx->in_tracing_section = 0; + goto end; + } + + /* Serialize event record */ + _serialize_er_default_memory_copy(_TO_VOID_PTR(ctx), p_phase, p_pid, p_operation, p_internal_correlation_id, p_timestamp, p_thread_id, p_src_agent_node_id, p_dst_agent_node_id, p_stream_id, p_size); + + /* Commit event record */ + _commit_er(_TO_VOID_PTR(ctx)); + + /* Not tracing anymore */ + ctx->in_tracing_section = 0; + +end: + return; +} diff --git a/source/lib/python/rocpd/source/ctf/barectf.h b/source/lib/python/rocpd/source/ctf/barectf.h new file mode 100644 index 0000000000..b7cd7564ff --- /dev/null +++ b/source/lib/python/rocpd/source/ctf/barectf.h @@ -0,0 +1,182 @@ +#ifndef _BARECTF_H +#define _BARECTF_H + +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2020 Philippe Proulx + * + * 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. + * + * - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + * + * The following code was generated by barectf v3.1.2 + * on 2025-06-18T16:33:18.853411. + * + * For more details, see . + */ + +#include + +#ifdef __cplusplus +extern "C" { +#endif + + +struct barectf_ctx; + +uint32_t barectf_packet_size(const void *vctx); +int barectf_packet_is_full(const void *vctx); +int barectf_packet_is_empty(const void *vctx); +uint32_t barectf_packet_events_discarded(const void *vctx); +uint32_t barectf_discarded_event_records_count(const void * const vctx); +uint32_t barectf_packet_sequence_number(const void * const vctx); +uint8_t *barectf_packet_buf(const void *vctx); +uint8_t *barectf_packet_buf_addr(const void * const vctx); +void barectf_packet_set_buf(void *vctx, uint8_t *buf, uint32_t buf_size); +uint32_t barectf_packet_buf_size(const void *vctx); +int barectf_packet_is_open(const void *vctx); +int barectf_is_in_tracing_section(const void *vctx); +volatile const int *barectf_is_in_tracing_section_ptr(const void *vctx); +int barectf_is_tracing_enabled(const void *vctx); +void barectf_enable_tracing(void *vctx, int enable); + +/* barectf platform callbacks */ +struct barectf_platform_callbacks { + /* Clock source callbacks */ + uint64_t (*default_clock_get_value)(void *); + + /* Is the back end full? */ + int (*is_backend_full)(void *); + + /* Open packet */ + void (*open_packet)(void *); + + /* Close packet */ + void (*close_packet)(void *); +}; + +/* Common barectf context */ +struct barectf_ctx { + /* Platform callbacks */ + struct barectf_platform_callbacks cbs; + + /* Platform data (passed to callbacks) */ + void *data; + + /* Output buffer (will contain a CTF binary packet) */ + uint8_t *buf; + + /* Packet's total size (bits) */ + uint32_t packet_size; + + /* Packet's content size (bits) */ + uint32_t content_size; + + /* Current position from beginning of packet (bits) */ + uint32_t at; + + /* Size of packet header + context fields (content offset) */ + uint32_t off_content; + + /* Discarded event records counter snapshot */ + uint32_t events_discarded; + + /* Packet's sequence number */ + uint32_t sequence_number; + + /* Current packet is open? */ + int packet_is_open; + + /* In tracing code? */ + volatile int in_tracing_section; + + /* Tracing is enabled? */ + volatile int is_tracing_enabled; + + /* Use current/last event record timestamp when opening/closing packets */ + uint64_t use_cur_last_event_ts; +}; + +/* Context for data stream type `default` */ +struct barectf_default_ctx { + /* Parent */ + struct barectf_ctx parent; + + /* Config-specific members follow */ + uint32_t off_pc_packet_size; + uint32_t off_pc_content_size; + uint32_t off_pc_events_discarded; + uint64_t cur_last_event_ts; +}; + +/* Initialize context */ +void barectf_init(void *vctx, + uint8_t *buf, uint32_t buf_size, + const struct barectf_platform_callbacks cbs, void *data); + +/* Open packet for data stream type `default` */ +void barectf_default_open_packet( + struct barectf_default_ctx *sctx); + +/* Close packet for data stream type `default` */ +void barectf_default_close_packet(struct barectf_default_ctx *sctx); + +/* Trace (data stream type `default`, event record type `api`) */ +void barectf_default_trace_api(struct barectf_default_ctx *sctx, + uint8_t p_phase, + uint64_t p_pid, + const char *p_name, + uint64_t p_ancestor_correlation_id, + uint64_t p_internal_correlation_id, + uint64_t p_timestamp, + uint64_t p_thread_id); + +/* Trace (data stream type `default`, event record type `kernel_dispatch`) */ +void barectf_default_trace_kernel_dispatch(struct barectf_default_ctx *sctx, + uint8_t p_phase, + uint64_t p_pid, + uint64_t p_internal_correlation_id, + uint64_t p_timestamp, + uint64_t p_thread_id, + uint64_t p_agent_node_id, + uint64_t p_queue_id, + uint64_t p_stream_id, + const char *p_kernel_name); + +/* Trace (data stream type `default`, event record type `memory_copy`) */ +void barectf_default_trace_memory_copy(struct barectf_default_ctx *sctx, + uint8_t p_phase, + uint64_t p_pid, + const char *p_operation, + uint64_t p_internal_correlation_id, + uint64_t p_timestamp, + uint64_t p_thread_id, + uint64_t p_src_agent_node_id, + uint64_t p_dst_agent_node_id, + uint64_t p_stream_id, + uint64_t p_size); + +#ifdef __cplusplus +} +#endif + +#endif /* _BARECTF_H */ diff --git a/source/lib/python/rocpd/source/ctf/barectf.yml b/source/lib/python/rocpd/source/ctf/barectf.yml new file mode 100644 index 0000000000..a975572477 --- /dev/null +++ b/source/lib/python/rocpd/source/ctf/barectf.yml @@ -0,0 +1,128 @@ +# barectf.yml +version: "2.0" + +metadata: + trace: + byte-order: le + + clocks: + default: + freq: 1000000000 + + streams: + default: + event-header-type: + class: struct + fields: + id: + class: integer + size: 8 + timestamp: + class: integer + size: 64 + property-mappings: + - type: clock + name: default + property: value + packet-context-type: + class: struct + fields: + packet_size: + class: integer + size: 64 + content_size: + class: integer + size: 64 + events_discarded: + class: int + size: 16 + events: + api: + payload-type: + class: struct + fields: + phase: + class: integer + size: 8 + pid: + class: integer + size: 64 + name: + class: string + ancestor_correlation_id: + class: integer + size: 64 + internal_correlation_id: + class: integer + size: 64 + timestamp: + class: integer + size: 64 + thread_id: + class: integer + size: 64 + + kernel_dispatch: + payload-type: + class: struct + fields: + phase: + class: integer + size: 8 + pid: + class: integer + size: 64 + internal_correlation_id: + class: integer + size: 64 + timestamp: + class: integer + size: 64 + thread_id: + class: integer + size: 64 + agent_node_id: + class: integer + size: 64 + queue_id: + class: integer + size: 64 + stream_id: + class: integer + size: 64 + kernel_name: + class: string + + memory_copy: + payload-type: + class: struct + fields: + phase: + class: integer + size: 8 + pid: + class: integer + size: 64 + operation: + class: string + internal_correlation_id: + class: integer + size: 64 + timestamp: + class: integer + size: 64 + thread_id: + class: integer + size: 64 + src_agent_node_id: + class: integer + size: 64 + dst_agent_node_id: + class: integer + size: 64 + stream_id: + class: integer + size: 64 + size: + class: integer + size: 64 \ No newline at end of file diff --git a/source/lib/python/rocpd/source/ctf/metadata b/source/lib/python/rocpd/source/ctf/metadata new file mode 100644 index 0000000000..edf6ef89a6 --- /dev/null +++ b/source/lib/python/rocpd/source/ctf/metadata @@ -0,0 +1,296 @@ +/* CTF 1.8 */ + +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2020 Philippe Proulx + * + * 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. + * + * - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + * + * The following code was generated by barectf v3.1.2 + * on 2025-06-18T16:33:18.853411. + * + * For more details, see . + */ + +trace { + major = 1; + minor = 8; + byte_order = le; + packet.header := struct { + } align(8); +}; + +env { + domain = "bare"; + tracer_name = "barectf"; + tracer_major = 3; + tracer_minor = 1; + tracer_patch = 2; + tracer_pre = ""; + barectf_gen_date = "2025-06-18T16:33:18.853411"; +}; + +clock { + name = default; + freq = 1000000000; + precision = 0; + offset_s = 0; + offset = 0; + absolute = true; +}; + +/* Data stream type `default` */ +stream { + packet.context := struct { + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } packet_size; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } content_size; + integer { + signed = false; + size = 16; + align = 8; + byte_order = native; + base = 10; + } events_discarded; + } align(8); + event.header := struct { + integer { + signed = false; + size = 8; + align = 8; + byte_order = native; + base = 10; + } id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + map = clock.default.value; + } timestamp; + } align(8); +}; + +event { + id = 0; + name = "api"; + fields := struct { + integer { + signed = false; + size = 8; + align = 8; + byte_order = native; + base = 10; + } phase; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } pid; + string { + encoding = UTF8; + } name; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } ancestor_correlation_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } internal_correlation_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } timestamp; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } thread_id; + } align(1); +}; + +event { + id = 1; + name = "kernel_dispatch"; + fields := struct { + integer { + signed = false; + size = 8; + align = 8; + byte_order = native; + base = 10; + } phase; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } pid; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } internal_correlation_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } timestamp; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } thread_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } agent_node_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } queue_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } stream_id; + string { + encoding = UTF8; + } kernel_name; + } align(1); +}; + +event { + id = 2; + name = "memory_copy"; + fields := struct { + integer { + signed = false; + size = 8; + align = 8; + byte_order = native; + base = 10; + } phase; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } pid; + string { + encoding = UTF8; + } operation; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } internal_correlation_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } timestamp; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } thread_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } src_agent_node_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } dst_agent_node_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } stream_id; + integer { + signed = false; + size = 64; + align = 8; + byte_order = native; + base = 10; + } size; + } align(1); +}; diff --git a/source/lib/python/rocpd/source/ctf2.cpp b/source/lib/python/rocpd/source/ctf2.cpp new file mode 100644 index 0000000000..279df3f54d --- /dev/null +++ b/source/lib/python/rocpd/source/ctf2.cpp @@ -0,0 +1,382 @@ +// MIT License +// +// Copyright (c) 2023-2025 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/python/rocpd/source/ctf2.hpp" + +#include "lib/common/defines.hpp" +#include "lib/common/filesystem.hpp" +#include "lib/common/hasher.hpp" +#include "lib/common/logging.hpp" +#include "lib/common/mpl.hpp" +#include "lib/common/units.hpp" +#include "lib/common/utility.hpp" +#include "lib/output/generator.hpp" +#include "lib/output/metadata.hpp" +#include "lib/output/node_info.hpp" +#include "lib/output/output_config.hpp" +#include "lib/output/output_stream.hpp" +#include "lib/output/sql/common.hpp" +#include "lib/output/stream_info.hpp" +#include "lib/output/timestamps.hpp" +#include "lib/rocprofiler-sdk-tool/config.hpp" + +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define CTF2_CHECK(result) \ + { \ + CTF2_ErrorCode ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) = result; \ + if(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) != CTF2_SUCCESS) \ + { \ + auto _err_name = CTF2_Error_GetName(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__)); \ + auto _err_msg = \ + CTF2_Error_GetDescription(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__)); \ + ROCP_FATAL << #result << " failed with error code " << _err_name \ + << " (code=" << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \ + << ") :: " << _err_msg; \ + } \ + } + +#define BUFFER_SIZE (1024 * 10) // 10 KiB + +namespace rocpd +{ +namespace output +{ +namespace +{ +// --- Platform callbacks for barectf --- +// Simple clock_gettime for nanoseconds +static uint64_t +get_time(void* const data) +{ + (void) data; // Unused + auto _ts = rocprofiler_timestamp_t{}; + rocprofiler_get_timestamp(&_ts); + return (uint64_t) _ts; +} + +// Dummy "is backend full" - for this example, we assume it's never full +static int +is_backend_full(void* const data) +{ + (void) data; /* Optional */ + return 0; // Never full +} + +static void +write_packet(const struct my_platform_ctx* const platform_ctx) +{ + /* Append current packet to data stream file */ + const size_t nmemb = fwrite(barectf_packet_buf_addr(platform_ctx->ctx), + barectf_packet_buf_size(platform_ctx->ctx), + 1, + platform_ctx->fh); + + assert(nmemb == 1); +} + +static void +open_packet(void* const data) +{ + struct my_platform_ctx* const platform_ctx = + reinterpret_cast(data); + + barectf_default_open_packet(platform_ctx->ctx); +} +static void +close_packet(void* const data) +{ + struct my_platform_ctx* const platform_ctx = + reinterpret_cast(data); + + /* Close packet now */ + barectf_default_close_packet(platform_ctx->ctx); + + /* Write packet to file */ + write_packet(platform_ctx); +} + +} // namespace + +void +CTF2Session::add_event(const extended_event_data& event_data, const types::process& process) const +{ + if(!platform_ctx || !platform_ctx->ctx) + { + // ROCP_FATAL << "Platform context is not initialized"; + return; + } + switch(event_data.event_type) + { + case ctf2_event_type::kernel_dispatch: + { + auto kernel_dispatch = + static_cast(event_data.kernel_dispatch); + if(!kernel_dispatch) + { + ROCP_FATAL << "Kernel dispatch event data is null"; + return; + } + barectf_default_trace_kernel_dispatch( + platform_ctx->ctx, + (event_data.event_phase == ctf2_event_phase::start ? 0 : 1), + process.pid, + kernel_dispatch->corr_id, + event_data.timestamp, + kernel_dispatch->tid, + kernel_dispatch->agent_type_index, + kernel_dispatch->queue_id, + 0, // Assuming stream_id is not used in this example + event_data.name.c_str()); + break; + } + case ctf2_event_type::api: + { + auto api_region = static_cast(event_data.api_region); + if(!api_region) + { + ROCP_FATAL << "API region event data is null"; + return; + } + barectf_default_trace_api(platform_ctx->ctx, + (event_data.event_phase == ctf2_event_phase::start ? 0 : 1), + process.pid, + event_data.name.c_str(), + api_region->parent_stack_id, + api_region->corr_id, + event_data.timestamp, + api_region->tid); + break; + } + case ctf2_event_type::memory_copy: + { + auto memory_copy = static_cast(event_data.memory_copy); + if(!memory_copy) + { + ROCP_FATAL << "Memory copy event data is null"; + return; + } + barectf_default_trace_memory_copy( + platform_ctx->ctx, + (event_data.event_phase == ctf2_event_phase::start ? 0 : 1), + process.pid, + memory_copy->category.c_str(), + memory_copy->corr_id, + event_data.timestamp, + memory_copy->tid, + memory_copy->src_agent_type_index, + memory_copy->dst_agent_type_index, + 0, + memory_copy->size); + break; + } + // Add other event types as needed + default: + ROCP_WARNING << "Unsupported event type: " << static_cast(event_data.event_type); + } +} + +CTF2Session::CTF2Session(const tool::output_config& output_cfg, + uint64_t min_start, + uint64_t max_fini) +: config{output_cfg} +{ + namespace fs = rocprofiler::common::filesystem; + + auto _filename = + rocprofiler::tool::get_output_filename(output_cfg, "results", std::string_view{}); + auto _filepath = fs::path{_filename}; + + if(fs::exists(_filepath)) fs::remove_all(_filepath); + + struct barectf_platform_callbacks cbs; + + /* Set platform callback functions */ + cbs.default_clock_get_value = get_time; + cbs.is_backend_full = is_backend_full; + cbs.open_packet = open_packet; + cbs.close_packet = close_packet; + + /* Allocate platform context (which contains a barectf context) */ + platform_ctx = reinterpret_cast(malloc(sizeof(*platform_ctx))); + + if(!platform_ctx) + { + // goto error; + } + + /* Allocate packet buffer */ + ctf_buffer = reinterpret_cast(malloc(BUFFER_SIZE)); + + if(!ctf_buffer) + { + // goto error; + } + + /* Open data stream file */ + platform_ctx->fh = fopen(_filepath.c_str(), "wb"); + + if(!platform_ctx->fh) + { + // goto error; + } + + /* Initialize barectf context */ + barectf_init(platform_ctx->ctx, ctf_buffer, BUFFER_SIZE, cbs, platform_ctx); + + /* Open the first packet */ + open_packet(platform_ctx); + + std::cout << "CTF2 session initialized with output file: " << _filepath << std::endl; +} + +CTF2Session::~CTF2Session() +{ + /* Close current packet if needed */ + if(barectf_packet_is_open(platform_ctx->ctx) && !barectf_packet_is_empty(platform_ctx->ctx)) + { + close_packet(platform_ctx); + } + + /* Close data stream file */ + fclose(platform_ctx->fh); + + /* Deallocate packet buffer */ + free(ctf_buffer); + + /* Deallocate platform context */ + if(platform_ctx) free(platform_ctx); + + std::cout << "CTF2 session finalized and output file closed." << std::endl; +} + +void +write_ctf2(const CTF2Session& ctf2_session, + const types::process& process, + const uint16_t tree_node_id, + const std::unordered_map& agent_data, + const tool::generator& thread_gen, + const tool::generator& api_gen, + const tool::generator& kernel_dispatch_gen, + const tool::generator& memory_copy_gen, + const tool::generator& memory_allocation_gen) +{ + const auto& ocfg = ctf2_session.config; + + auto _app_ts = rocprofiler::tool::timestamps_t{process.start, process.fini}; + + auto _data = std::deque{}; + + for(auto ditr : kernel_dispatch_gen) + for(const auto& itr : kernel_dispatch_gen.get(ditr)) + { + auto _name = fmt::format( + "{}", (ocfg.kernel_rename && !itr.region.empty()) ? itr.region : itr.name); + + _data.emplace_back(extended_event_data{ + itr.start, ctf2_event_type::kernel_dispatch, ctf2_event_phase::start, &itr, _name}); + + _data.emplace_back(extended_event_data{ + itr.end, ctf2_event_type::kernel_dispatch, ctf2_event_phase::end, &itr, _name}); + } + + for(auto ditr : memory_copy_gen) + for(const auto& itr : memory_copy_gen.get(ditr)) + { + std::string _name = itr.name; + + _data.emplace_back(extended_event_data{itr.start, + ctf2_event_type::kernel_dispatch, + ctf2_event_phase::start, + .memory_copy = &itr, + _name}); + + _data.emplace_back(extended_event_data{itr.end, + ctf2_event_type::kernel_dispatch, + ctf2_event_phase::end, + .memory_copy = &itr, + _name}); + } + + for(auto ditr : api_gen) + for(const auto& itr : api_gen.get(ditr)) + { + std::string _name = itr.name; + + _data.emplace_back(extended_event_data{itr.start, + ctf2_event_type::api, + ctf2_event_phase::start, + .api_region = &itr, + _name}); + + _data.emplace_back(extended_event_data{ + itr.end, ctf2_event_type::api, ctf2_event_phase::end, .api_region = &itr, _name}); + } + + std::sort(_data.begin(), + _data.end(), + [](const extended_event_data& lhs, const extended_event_data& rhs) { + if(lhs.timestamp != rhs.timestamp) return (lhs.timestamp < rhs.timestamp); + return (lhs.event_phase >= rhs.event_phase); + }); + + uint64_t largest_timestamp = 0; + for(const auto& itr : _data) + { + ROCP_ERROR_IF(itr.timestamp < largest_timestamp) + << "event found with timestamp < last event timestamp by " + << (largest_timestamp - itr.timestamp) << " nsec"; + + ctf2_session.add_event(itr, process); + + largest_timestamp = itr.timestamp; + + ROCP_ERROR_IF(itr.timestamp < _app_ts.app_start_time) + << "event found with timestamp < app start time by " + << (_app_ts.app_start_time - itr.timestamp) << " nsec"; + ROCP_ERROR_IF(itr.timestamp > _app_ts.app_end_time) + << "event found with timestamp > app end time by " + << (itr.timestamp - _app_ts.app_end_time) << " nsec"; + } +} + +} // namespace output +} // namespace rocpd diff --git a/source/lib/python/rocpd/source/ctf2.hpp b/source/lib/python/rocpd/source/ctf2.hpp new file mode 100644 index 0000000000..f93dde57ec --- /dev/null +++ b/source/lib/python/rocpd/source/ctf2.hpp @@ -0,0 +1,111 @@ +// MIT License +// +// Copyright (c) 2025 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/python/rocpd/source/types.hpp" +#include "lib/python/rocpd/source/ctf/barectf.h" + +#include "lib/common/defines.hpp" +#include "lib/output/generator.hpp" +#include "lib/output/metadata.hpp" +#include "lib/output/node_info.hpp" +#include "lib/output/output_config.hpp" +#include "lib/output/sql/common.hpp" +#include "lib/output/stream_info.hpp" +#include "lib/rocprofiler-sdk-tool/config.hpp" + +#include + +namespace rocpd +{ +namespace output +{ +namespace tool = rocprofiler::tool; + +enum class ctf2_event_type : uint8_t +{ + none = 0, + kernel_dispatch, + memory_copy, + api +}; + +enum class ctf2_event_phase : uint8_t +{ + none = 0, + start, + end +}; + +struct extended_event_data +{ + uint64_t timestamp = 0; + ctf2_event_type event_type = ctf2_event_type::none; + ctf2_event_phase event_phase = ctf2_event_phase::none; + union + { + const types::kernel_dispatch* kernel_dispatch; + const types::memory_copies* memory_copy; + const types::region* api_region; + }; + std::string name = ""; +}; + +struct extended_agent_ctf +{ + const rocpd::types::agent& types_agent; + const tool::agent_index agent_index; + const std::string labeled_name; +}; + +/* Platform context */ +struct my_platform_ctx { + barectf_default_ctx* ctx = reinterpret_cast(malloc(sizeof(::barectf_default_ctx)));; + FILE *fh; +}; + +struct CTF2Session +{ + CTF2Session(const tool::output_config& output_cfg, uint64_t min_start, uint64_t max_fini); + + ~CTF2Session(); + + void add_event(const extended_event_data& event_data, const types::process& process) const; + + const tool::output_config& config; + struct my_platform_ctx *platform_ctx; + uint8_t* ctf_buffer = nullptr; +}; + +void +write_ctf2(const CTF2Session& session, + const types::process& process, + const uint16_t tree_node_id, + const std::unordered_map& agent_data, + const tool::generator& thread_gen, + const tool::generator& api_gen, + const tool::generator& kernel_dispatch_gen, + const tool::generator& memory_copy_gen, + const tool::generator& memory_allocation_gen); +} // namespace output +} // namespace rocpd diff --git a/source/lib/python/utilities.cmake b/source/lib/python/utilities.cmake index 6bae034df2..e365ad9b4c 100644 --- a/source/lib/python/utilities.cmake +++ b/source/lib/python/utilities.cmake @@ -153,6 +153,7 @@ function(rocprofiler_rocpd_python_bindings _VERSION) __main__.py output_config.py otf2.py + ctf2.py pftrace.py schema.py time_window.py)