From f27f76716e641439bd74341c56a59137c7e0bdfa Mon Sep 17 00:00:00 2001 From: "Kuricheti, Mythreya" Date: Thu, 20 Mar 2025 22:52:48 -0700 Subject: [PATCH] [SDK] Add Stack IDs (#269) * Add Stack IDs * Add memcpy test * Add async corr id record * Async events use `rocprofiler_async_correlation_id_t` * Sync events use `rocprofiler_correlation_id_t` * Update ATT to use asnyc IDs * Review comments --- .../amd_detail/thread_trace_dispatch.h | 14 +- .../include/rocprofiler-sdk/buffer_tracing.h | 4 +- .../rocprofiler-sdk/cxx/serialization.hpp | 9 ++ .../dispatch_counting_service.h | 14 +- source/include/rocprofiler-sdk/fwd.h | 30 ++++ source/include/rocprofiler-sdk/pc_sampling.h | 11 +- source/lib/output/generatePerfetto.cpp | 26 +++- source/lib/rocprofiler-sdk-tool/tool.cpp | 2 +- .../context/correlation_id.cpp | 7 +- .../context/correlation_id.hpp | 4 +- .../counters/dispatch_handlers.cpp | 2 +- .../counters/sample_processing.cpp | 3 +- .../rocprofiler-sdk/counters/tests/core.cpp | 18 +-- source/lib/rocprofiler-sdk/hip/hip.cpp | 3 + source/lib/rocprofiler-sdk/hip/stream.cpp | 6 + source/lib/rocprofiler-sdk/hsa/async_copy.cpp | 4 +- source/lib/rocprofiler-sdk/hsa/hsa.cpp | 4 +- .../rocprofiler-sdk/hsa/memory_allocation.cpp | 7 +- source/lib/rocprofiler-sdk/hsa/queue.cpp | 2 + .../rocprofiler-sdk/hsa/scratch_memory.cpp | 3 + .../kernel_dispatch/tracing.cpp | 5 +- source/lib/rocprofiler-sdk/marker/marker.cpp | 3 + source/lib/rocprofiler-sdk/ompt/ompt.cpp | 7 + .../pc_sampling/parser/correlation.hpp | 12 +- .../rocprofiler-sdk/pc_sampling/parser/rocr.h | 4 +- .../pc_sampling/parser/stochastic_records.h | 8 +- source/lib/rocprofiler-sdk/rccl/rccl.cpp | 3 + .../rocprofiler-sdk/rocdecode/rocdecode.cpp | 3 + .../lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp | 3 + .../runtime_initialization.cpp | 5 +- .../rocprofiler-sdk/thread_trace/att_core.cpp | 9 +- .../thread_trace/tests/att_packet_test.cpp | 4 +- .../lib/rocprofiler-sdk/tracing/tracing.hpp | 33 ++++- tests/async-copy-tracing/validate.py | 132 ++++++++++++++++++ tests/thread-trace/multi_dispatch.cpp | 2 +- tests/thread-trace/single_dispatch.cpp | 2 +- 36 files changed, 341 insertions(+), 67 deletions(-) diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h index 772376e271..62fa3e4d7c 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -55,13 +55,13 @@ typedef enum * @param [out] userdata_shader Userdata to be passed in shader_callback */ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)( - rocprofiler_agent_id_t agent_id, - rocprofiler_queue_id_t queue_id, - rocprofiler_correlation_id_t correlation_id, - rocprofiler_kernel_id_t kernel_id, - rocprofiler_dispatch_id_t dispatch_id, - void* userdata_config, - rocprofiler_user_data_t* userdata_shader); + rocprofiler_agent_id_t agent_id, + rocprofiler_queue_id_t queue_id, + rocprofiler_async_correlation_id_t correlation_id, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, + void* userdata_config, + rocprofiler_user_data_t* userdata_shader); /** * @brief Enables the advanced thread trace service for dispatch-based tracing. diff --git a/source/include/rocprofiler-sdk/buffer_tracing.h b/source/include/rocprofiler-sdk/buffer_tracing.h index a213faeb66..e241fdf6e7 100644 --- a/source/include/rocprofiler-sdk/buffer_tracing.h +++ b/source/include/rocprofiler-sdk/buffer_tracing.h @@ -228,7 +228,7 @@ typedef struct uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; rocprofiler_memory_copy_operation_t operation; - rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record + rocprofiler_async_correlation_id_t correlation_id; ///< correlation ids for record rocprofiler_thread_id_t thread_id; ///< id for thread that triggered copy rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds @@ -276,7 +276,7 @@ typedef struct rocprofiler_buffer_tracing_kernel_dispatch_record_t uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH rocprofiler_kernel_dispatch_operation_t operation; - rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record + rocprofiler_async_correlation_id_t correlation_id; ///< correlation ids for record rocprofiler_thread_id_t thread_id; ///< id for thread that launched kernel rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds diff --git a/source/include/rocprofiler-sdk/cxx/serialization.hpp b/source/include/rocprofiler-sdk/cxx/serialization.hpp index cd159e50c6..4ab8ff6a08 100644 --- a/source/include/rocprofiler-sdk/cxx/serialization.hpp +++ b/source/include/rocprofiler-sdk/cxx/serialization.hpp @@ -131,6 +131,15 @@ save(ArchiveT& ar, rocprofiler_counter_id_t data) template void save(ArchiveT& ar, rocprofiler_correlation_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(internal); + ROCP_SDK_SAVE_DATA_VALUE("external", external.value); + ROCP_SDK_SAVE_DATA_VALUE("ancestor", ancestor); +} + +template +void +save(ArchiveT& ar, rocprofiler_async_correlation_id_t data) { ROCP_SDK_SAVE_DATA_FIELD(internal); ROCP_SDK_SAVE_DATA_VALUE("external", external.value); diff --git a/source/include/rocprofiler-sdk/dispatch_counting_service.h b/source/include/rocprofiler-sdk/dispatch_counting_service.h index f2898ae662..cd5c4cc2ea 100644 --- a/source/include/rocprofiler-sdk/dispatch_counting_service.h +++ b/source/include/rocprofiler-sdk/dispatch_counting_service.h @@ -44,7 +44,7 @@ ROCPROFILER_EXTERN_C_INIT typedef struct rocprofiler_dispatch_counting_service_data_t { uint64_t size; ///< Size of this struct - rocprofiler_correlation_id_t correlation_id; ///< Correlation ID for this dispatch + rocprofiler_async_correlation_id_t correlation_id; ///< Correlation ID for this dispatch rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds rocprofiler_kernel_dispatch_info_t dispatch_info; ///< Dispatch info @@ -57,12 +57,12 @@ typedef struct rocprofiler_dispatch_counting_service_data_t */ typedef struct rocprofiler_dispatch_counting_service_record_t { - uint64_t size; ///< Size of this struct - uint64_t num_records; ///< number of ::rocprofiler_record_counter_t records - rocprofiler_correlation_id_t correlation_id; ///< Correlation ID for this dispatch - rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds - rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds - rocprofiler_kernel_dispatch_info_t dispatch_info; ///< Contains the `dispatch_id` + uint64_t size; ///< Size of this struct + uint64_t num_records; ///< number of ::rocprofiler_record_counter_t records + rocprofiler_async_correlation_id_t correlation_id; ///< Correlation ID for this dispatch + rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds + rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds + rocprofiler_kernel_dispatch_info_t dispatch_info; ///< Contains the `dispatch_id` } rocprofiler_dispatch_counting_service_record_t; /** diff --git a/source/include/rocprofiler-sdk/fwd.h b/source/include/rocprofiler-sdk/fwd.h index 29a8cfec5c..71130ef52a 100644 --- a/source/include/rocprofiler-sdk/fwd.h +++ b/source/include/rocprofiler-sdk/fwd.h @@ -605,13 +605,43 @@ typedef struct rocprofiler_correlation_id_t { uint64_t internal; rocprofiler_user_data_t external; + uint64_t ancestor; + + /// @var internal + /// @brief A unique ID created by rocprofiler-sdk when an API call is invoked. + /// @var external + /// @brief An ID specified by tools to associate external events. + /// See include/rocprofiler-sdk/external_correlation.h + /// @var ancestor + /// @brief Stores the ::internal value of the API call that generated this + /// API call. } rocprofiler_correlation_id_t; +/** + * @brief ROCProfiler Correlation ID record for async activity. + */ +typedef struct rocprofiler_async_correlation_id_t +{ + uint64_t internal; + rocprofiler_user_data_t external; + + /// @var internal + /// @brief A unique ID created by rocprofiler-sdk when an API call is invoked. + /// @var external + /// @brief An ID specified by tools to associate external events. + /// See include/rocprofiler-sdk/external_correlation.h +} rocprofiler_async_correlation_id_t; + /** * @brief The NULL value of an internal correlation ID. */ #define ROCPROFILER_CORRELATION_ID_INTERNAL_NONE ROCPROFILER_UINT64_C(0) +/** + * @brief The NULL value of an ancestor correlation ID. + */ +#define ROCPROFILER_CORRELATION_ID_ANCESTOR_NONE ROCPROFILER_UINT64_C(0) + /** * @struct rocprofiler_buffer_id_t * @brief Buffer ID. diff --git a/source/include/rocprofiler-sdk/pc_sampling.h b/source/include/rocprofiler-sdk/pc_sampling.h index 6f8e64a9ba..b366d8b801 100644 --- a/source/include/rocprofiler-sdk/pc_sampling.h +++ b/source/include/rocprofiler-sdk/pc_sampling.h @@ -254,10 +254,13 @@ typedef struct rocprofiler_pc_sampling_record_host_trap_v0_t uint64_t exec_mask; ///< active SIMD lanes when sampled uint64_t timestamp; ///< timestamp when sample is generated uint64_t dispatch_id; ///< originating kernel dispatch ID - rocprofiler_correlation_id_t correlation_id; ///< API launch call id that matches dispatch ID - rocprofiler_dim3_t workgroup_id; ///< wave coordinates within the workgroup - uint32_t wave_in_group : 8; ///< wave position within the workgroup (0-31) - uint32_t reserved0 : 24; ///< wave position within the workgroup (0-31) + rocprofiler_async_correlation_id_t correlation_id; + rocprofiler_dim3_t workgroup_id; ///< wave coordinates within the workgroup + uint32_t wave_in_group : 8; ///< wave position within the workgroup (0-31) + uint32_t reserved0 : 24; ///< wave position within the workgroup (0-31) + + /// @var correlation_id + /// @brief API launch call id that matches dispatch ID } rocprofiler_pc_sampling_record_host_trap_v0_t; /** @} */ diff --git a/source/lib/output/generatePerfetto.cpp b/source/lib/output/generatePerfetto.cpp index c8c20c9b7c..aeb356796d 100644 --- a/source/lib/output/generatePerfetto.cpp +++ b/source/lib/output/generatePerfetto.cpp @@ -379,7 +379,10 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); + TRACE_EVENT_END( sdk::perfetto_category::name, track, itr.end_timestamp); tracing_session->FlushBlocking(); @@ -409,7 +412,10 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); + TRACE_EVENT_END( sdk::perfetto_category::name, track, itr.end_timestamp); tracing_session->FlushBlocking(); @@ -442,7 +448,9 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp); @@ -473,7 +481,9 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp); @@ -504,7 +514,9 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp); @@ -535,7 +547,9 @@ write_perfetto( "operation", itr.operation, "corr_id", - itr.correlation_id.internal); + itr.correlation_id.internal, + "ancestor_id", + itr.correlation_id.ancestor); TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp); diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index fa1b5fd351..b97560e53d 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -1234,7 +1234,7 @@ att_shader_data_callback(rocprofiler_agent_id_t agent, rocprofiler_att_control_flags_t att_dispatch_callback(rocprofiler_agent_id_t /* agent_id */, rocprofiler_queue_id_t /* queue_id */, - rocprofiler_correlation_id_t /* correlation_id */, + rocprofiler_async_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t dispatch_id, void* /*userdata_config*/, diff --git a/source/lib/rocprofiler-sdk/context/correlation_id.cpp b/source/lib/rocprofiler-sdk/context/correlation_id.cpp index 1b73a30fd2..56dfe1863f 100644 --- a/source/lib/rocprofiler-sdk/context/correlation_id.cpp +++ b/source/lib/rocprofiler-sdk/context/correlation_id.cpp @@ -23,6 +23,7 @@ #include "lib/rocprofiler-sdk/context/correlation_id.hpp" #include "lib/common/logging.hpp" #include "lib/common/static_object.hpp" +#include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" @@ -130,7 +131,11 @@ correlation_tracing_service::construct(uint32_t _init_ref_count) auto* corr_id_map = get_correlation_id_map(); if(!corr_id_map) return nullptr; auto& ret = corr_id_map->wlock([](auto& data) -> auto& { return data.emplace_back(); }); - ret = std::make_unique(_init_ref_count, common::get_tid(), _internal_id); + + ret = std::make_unique(_init_ref_count, common::get_tid(), _internal_id); + + if(auto* prev_api_corr_id = get_latest_correlation_id()) + ret->ancestor = prev_api_corr_id->internal; get_latest_correlation_id_impl().emplace_back(ret.get()); diff --git a/source/lib/rocprofiler-sdk/context/correlation_id.hpp b/source/lib/rocprofiler-sdk/context/correlation_id.hpp index a4a7d09a50..d4e1267bd2 100644 --- a/source/lib/rocprofiler-sdk/context/correlation_id.hpp +++ b/source/lib/rocprofiler-sdk/context/correlation_id.hpp @@ -62,6 +62,7 @@ struct correlation_id rocprofiler_thread_id_t thread_idx = 0; uint64_t internal = 0; + uint64_t ancestor = 0; uint32_t get_ref_count() const { return m_ref_count.load(); } uint32_t add_ref_count(); @@ -76,9 +77,6 @@ private: std::atomic m_ref_count = {0}; }; -correlation_id* -get_correlation_id(rocprofiler_thread_id_t tid, uint64_t internal_id); - // latest correlation id for thread correlation_id* get_latest_correlation_id(); diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index cfb3a85fe8..17e09d0f8d 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -80,7 +80,7 @@ queue_cb(const context::context* ctx, if(!is_enabled || !info->user_cb) return {no_instrumentation(), true}; auto _corr_id_v = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; + rocprofiler_async_correlation_id_t{.internal = 0, .external = context::null_user_data}; if(const auto* _corr_id = correlation_id) { _corr_id_v.internal = _corr_id->internal; diff --git a/source/lib/rocprofiler-sdk/counters/sample_processing.cpp b/source/lib/rocprofiler-sdk/counters/sample_processing.cpp index e54bac437e..0f942c745b 100644 --- a/source/lib/rocprofiler-sdk/counters/sample_processing.cpp +++ b/source/lib/rocprofiler-sdk/counters/sample_processing.cpp @@ -77,7 +77,8 @@ proccess_completed_cb(completed_cb_params_t&& params) } auto _corr_id_v = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; + rocprofiler_async_correlation_id_t{.internal = 0, .external = context::null_user_data}; + if(const auto* _corr_id = session.correlation_id) { _corr_id_v.internal = _corr_id->internal; diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index 40b13d589a..7a289ca6c0 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -304,15 +304,15 @@ namespace struct expected_dispatch { // To pass back - rocprofiler_profile_config_id_t id = {.handle = 0}; - rocprofiler_queue_id_t queue_id = {.handle = 0}; - rocprofiler_agent_id_t agent_id = {.handle = 0}; - uint64_t kernel_id = 0; - uint64_t dispatch_id = 0; - rocprofiler_correlation_id_t correlation_id = {.internal = 0, .external = {.value = 0}}; - rocprofiler_dim3_t workgroup_size = {0, 0, 0}; - rocprofiler_dim3_t grid_size = {0, 0, 0}; - rocprofiler_profile_config_id_t* config = nullptr; + rocprofiler_profile_config_id_t id = {.handle = 0}; + rocprofiler_queue_id_t queue_id = {.handle = 0}; + rocprofiler_agent_id_t agent_id = {.handle = 0}; + uint64_t kernel_id = 0; + uint64_t dispatch_id = 0; + rocprofiler_async_correlation_id_t correlation_id = {.internal = 0, .external = {.value = 0}}; + rocprofiler_dim3_t workgroup_size = {0, 0, 0}; + rocprofiler_dim3_t grid_size = {0, 0, 0}; + rocprofiler_profile_config_id_t* config = nullptr; }; void diff --git a/source/lib/rocprofiler-sdk/hip/hip.cpp b/source/lib/rocprofiler-sdk/hip/hip.cpp index da17cc86b6..e38c4bbe55 100644 --- a/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -224,6 +224,7 @@ hip_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -241,6 +242,7 @@ hip_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -284,6 +286,7 @@ hip_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/hip/stream.cpp b/source/lib/rocprofiler-sdk/hip/stream.cpp index 7e65dc9b3c..900bba0e9e 100644 --- a/source/lib/rocprofiler-sdk/hip/stream.cpp +++ b/source/lib/rocprofiler-sdk/hip/stream.cpp @@ -176,6 +176,7 @@ FuncT create_write_functor(RetT (*func)(Args...)) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto internal_corr_id = 0; + auto ancestor_corr_id = 0; constexpr auto stream_idx = common::mpl::index_of::value; auto stream = std::get(std::make_tuple(std::forward(args)...)); @@ -194,6 +195,7 @@ FuncT create_write_functor(RetT (*func)(Args...)) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, ROCPROFILER_HIP_STREAM_CREATE, tracer_data); @@ -237,6 +239,7 @@ FuncT create_destroy_functor(RetT (*func)(Args...)) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto internal_corr_id = 0; + auto ancestor_corr_id = 0; auto stream = std::get(std::make_tuple(std::forward(args)...)); @@ -252,6 +255,7 @@ FuncT create_destroy_functor(RetT (*func)(Args...)) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, ROCPROFILER_HIP_STREAM_DESTROY, tracer_data); @@ -296,6 +300,7 @@ FuncT create_read_functor(RetT (*func)(Args...)) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto internal_corr_id = 0; + auto ancestor_corr_id = 0; auto stream = std::get(std::make_tuple(std::forward(args)...)); @@ -306,6 +311,7 @@ FuncT create_read_functor(RetT (*func)(Args...)) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, ROCPROFILER_HIP_STREAM_SET, tracer_data); diff --git a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp index 3eeec71a34..f9205cddec 100644 --- a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp +++ b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -194,7 +194,7 @@ async_copy_data::get_buffered_record(const context_t* _ctx, auto _external_corr_id = (_ctx) ? tracing_data.external_correlation_ids.at(_ctx) : context::null_user_data; - auto _corr_id = rocprofiler_correlation_id_t{correlation_id->internal, _external_corr_id}; + auto _corr_id = rocprofiler_async_correlation_id_t{correlation_id->internal, _external_corr_id}; return common::init_public_api_struct(buffered_data_t{}, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, @@ -400,6 +400,7 @@ async_copy_handler(hsa_signal_value_t signal_value, void* arg) _data->tid, _data->correlation_id->internal, _data->tracing_data.external_correlation_ids, + _data->correlation_id->ancestor, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, _data->direction, record); @@ -712,6 +713,7 @@ async_copy_impl(Args... args) thr_id, _data->correlation_id->internal, tracing_data.external_correlation_ids, + _data->correlation_id->ancestor, ROCPROFILER_CALLBACK_TRACING_MEMORY_COPY, _direction, _tracer_data); diff --git a/source/lib/rocprofiler-sdk/hsa/hsa.cpp b/source/lib/rocprofiler-sdk/hsa/hsa.cpp index f32bfbfc5a..1b603ee23a 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa.cpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa.cpp @@ -340,7 +340,7 @@ hsa_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_hsa_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; - + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, external_corr_id_domain_idx, @@ -356,6 +356,7 @@ hsa_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -399,6 +400,7 @@ hsa_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/hsa/memory_allocation.cpp b/source/lib/rocprofiler-sdk/hsa/memory_allocation.cpp index 5624c8e981..16f0320899 100644 --- a/source/lib/rocprofiler-sdk/hsa/memory_allocation.cpp +++ b/source/lib/rocprofiler-sdk/hsa/memory_allocation.cpp @@ -362,7 +362,8 @@ memory_allocation_data::get_buffered_record(const context_t* _ctx, { auto _external_corr_id = (_ctx) ? tracing_data.external_correlation_ids.at(_ctx) : context::null_user_data; - auto _corr_id = rocprofiler_correlation_id_t{correlation_id->internal, _external_corr_id}; + auto _corr_id = rocprofiler_correlation_id_t{ + correlation_id->internal, _external_corr_id, correlation_id->ancestor}; return common::init_public_api_struct(buffered_data_t{}, ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION, @@ -520,6 +521,7 @@ memory_allocation_impl(Args... args) thr_id, _data.correlation_id->internal, tracing_data.external_correlation_ids, + _data.correlation_id->ancestor, ROCPROFILER_CALLBACK_TRACING_MEMORY_ALLOCATION, rocprofiler_enum, _tracer_data); @@ -561,6 +563,7 @@ memory_allocation_impl(Args... args) _data.tid, _data.correlation_id->internal, _data.tracing_data.external_correlation_ids, + _data.correlation_id->ancestor, ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION, rocprofiler_enum, record); @@ -644,6 +647,7 @@ memory_free_impl(Args... args) thr_id, _data.correlation_id->internal, tracing_data.external_correlation_ids, + _data.correlation_id->ancestor, ROCPROFILER_CALLBACK_TRACING_MEMORY_ALLOCATION, rocprofiler_enum, _tracer_data); @@ -679,6 +683,7 @@ memory_free_impl(Args... args) _data.tid, _data.correlation_id->internal, _data.tracing_data.external_correlation_ids, + _data.correlation_id->ancestor, ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION, rocprofiler_enum, record); diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index b196c58f19..6483df24dd 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -287,6 +287,7 @@ WriteInterceptor(const void* packets, auto thr_id = (corr_id) ? corr_id->thread_idx : common::get_tid(); auto user_data = rocprofiler_user_data_t{.value = 0}; auto internal_corr_id = (corr_id) ? corr_id->internal : 0; + auto ancestor_corr_id = (corr_id) ? corr_id->ancestor : 0; // if we constructed a correlation id, this decrements the reference count after the // underlying function returns @@ -352,6 +353,7 @@ WriteInterceptor(const void* packets, thr_id, internal_corr_id, tracing_data_v.external_correlation_ids, + ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, ROCPROFILER_KERNEL_DISPATCH_ENQUEUE, tracer_data); diff --git a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp index d362180771..441270aa53 100644 --- a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp +++ b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp @@ -455,6 +455,7 @@ impl(Args... args) const auto thr_id = common::get_tid(); auto* corr_id = context::get_latest_correlation_id(); auto internal_corr_id = (corr_id) ? corr_id->internal : 0; + auto ancestor_corr_id = (corr_id) ? corr_id->ancestor : 0; [[maybe_unused]] const auto get_agent_id = [](const hsa_queue_t* hsa_queue) -> rocprofiler_agent_id_t { @@ -500,6 +501,7 @@ impl(Args... args) thr_id, internal_corr_id, tls.external_corr_ids, + ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, OpIdx, tls.callback_data); @@ -553,6 +555,7 @@ impl(Args... args) thr_id, internal_corr_id, tls.external_corr_ids, + ancestor_corr_id, ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, OpIdx, std::move(_buffered_data)); diff --git a/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp b/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp index e99b26ed61..e05a51e754 100644 --- a/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp +++ b/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp @@ -75,6 +75,7 @@ dispatch_complete(queue_info_session_t& session, profiling_time dispatch_time) const auto& _extern_corr_ids = session.tracing_data.external_correlation_ids; auto _tid = session.tid; auto _internal_corr_id = (_corr_id) ? _corr_id->internal : 0; + auto _ancestor_corr_id = (_corr_id) ? _corr_id->ancestor : 0; if(dispatch_time.status == HSA_STATUS_SUCCESS) { @@ -88,6 +89,7 @@ dispatch_complete(queue_info_session_t& session, profiling_time dispatch_time) _tid, _internal_corr_id, _extern_corr_ids, + _ancestor_corr_id, ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, ROCPROFILER_KERNEL_DISPATCH_COMPLETE, tracer_data); @@ -98,7 +100,7 @@ dispatch_complete(queue_info_session_t& session, profiling_time dispatch_time) auto record = kernel_dispatch_record_t{sizeof(kernel_dispatch_record_t), ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, ROCPROFILER_KERNEL_DISPATCH_COMPLETE, - rocprofiler_correlation_id_t{}, + rocprofiler_async_correlation_id_t{}, _tid, callback_record.start_timestamp, callback_record.end_timestamp, @@ -108,6 +110,7 @@ dispatch_complete(queue_info_session_t& session, profiling_time dispatch_time) _tid, _internal_corr_id, _extern_corr_ids, + _ancestor_corr_id, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, ROCPROFILER_KERNEL_DISPATCH_COMPLETE, record); diff --git a/source/lib/rocprofiler-sdk/marker/marker.cpp b/source/lib/rocprofiler-sdk/marker/marker.cpp index 06fa9c44b3..00f9135156 100644 --- a/source/lib/rocprofiler-sdk/marker/marker.cpp +++ b/source/lib/rocprofiler-sdk/marker/marker.cpp @@ -170,6 +170,7 @@ roctx_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -186,6 +187,7 @@ roctx_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -229,6 +231,7 @@ roctx_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/ompt/ompt.cpp b/source/lib/rocprofiler-sdk/ompt/ompt.cpp index 126de13b57..f41ef20997 100644 --- a/source/lib/rocprofiler-sdk/ompt/ompt.cpp +++ b/source/lib/rocprofiler-sdk/ompt/ompt.cpp @@ -739,6 +739,7 @@ ompt_impl::begin(ompt_data_t* data, Args... args) auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -756,6 +757,7 @@ ompt_impl::begin(ompt_data_t* data, Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -813,6 +815,7 @@ ompt_impl::end(ompt_data_t* data, Args... args) auto* corr_id = state->corr_id; auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; ROCP_FATAL_IF(common::get_tid() != state->thr_id) << "MIsmatch of OMPT begin/end thread id: " @@ -878,6 +881,7 @@ ompt_impl::end(ompt_data_t* data, Args... args) state->thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); @@ -919,6 +923,7 @@ ompt_impl::event_common(Args... args) auto tracer_data = common::init_public_api_struct(callback_ompt_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); uint64_t internal_corr_id = corr_id->internal; + uint64_t ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -935,6 +940,7 @@ ompt_impl::event_common(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -951,6 +957,7 @@ ompt_impl::event_common(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp b/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp index 7d15403b3b..c866589832 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp @@ -49,8 +49,8 @@ namespace Parser { struct dispatch_correlation_ids_t { - rocprofiler_dispatch_id_t dispatch_id; - rocprofiler_correlation_id_t correlation_id; + rocprofiler_dispatch_id_t dispatch_id; + rocprofiler_async_correlation_id_t correlation_id; }; /** @@ -107,6 +107,14 @@ struct std::hash namespace Parser { +// 64B for performance reasons +constexpr auto pcs_parser_sample_record_size = 64; +static_assert(sizeof(generic_sample_t) == pcs_parser_sample_record_size); +static_assert(sizeof(generic_sample_t) == sizeof(perf_sample_snapshot_v1)); +static_assert(sizeof(generic_sample_t) == sizeof(perf_sample_host_trap_v1)); +static_assert(sizeof(generic_sample_t) == sizeof(upcoming_samples_t)); +static_assert(sizeof(generic_sample_t) == sizeof(dispatch_pkt_id_t)); + /** * Coordinates DispatchMap and DoorBellMap to reconstruct the original correlation_id * from the correlation_id seen by the trap handler. diff --git a/source/lib/rocprofiler-sdk/pc_sampling/parser/rocr.h b/source/lib/rocprofiler-sdk/pc_sampling/parser/rocr.h index e711a6d00a..541e52ecaa 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/parser/rocr.h +++ b/source/lib/rocprofiler-sdk/pc_sampling/parser/rocr.h @@ -72,8 +72,8 @@ typedef struct uint64_t write_index; uint64_t read_index; /// both internal and external correlation ID. - rocprofiler_correlation_id_t correlation_id; - rocprofiler_dispatch_id_t dispatch_id; + rocprofiler_async_correlation_id_t correlation_id; + rocprofiler_dispatch_id_t dispatch_id; } dispatch_pkt_id_t; typedef struct diff --git a/source/lib/rocprofiler-sdk/pc_sampling/parser/stochastic_records.h b/source/lib/rocprofiler-sdk/pc_sampling/parser/stochastic_records.h index a1bd2c8b35..285c3feb4a 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/parser/stochastic_records.h +++ b/source/lib/rocprofiler-sdk/pc_sampling/parser/stochastic_records.h @@ -169,10 +169,10 @@ typedef struct rocprofiler_pc_sampling_record_stochastic_v0_t rocprofiler_pc_t pc; ///< information about sampled program counter uint64_t exec_mask; ///< active SIMD lanes at the moment sampling rocprofiler_dim3_t workgroup_id; ///< wave coordinates within the workgroup - uint32_t wave_count; /// active waves on the CU at the moment of sampling - uint64_t timestamp; ///< timestamp when sample is generated - uint64_t dispatch_id; ///< originating kernel dispatch ID - rocprofiler_correlation_id_t correlation_id; + uint32_t wave_count; /// active waves on the CU at the moment of sampling + uint64_t timestamp; ///< timestamp when sample is generated + uint64_t dispatch_id; ///< originating kernel dispatch ID + rocprofiler_async_correlation_id_t correlation_id; rocprofiler_pc_sampling_snapshot_v0_t snapshot; ///< @see ::rocprofiler_pc_sampling_snapshot_v0_t rocprofiler_pc_sampling_memory_counters_t diff --git a/source/lib/rocprofiler-sdk/rccl/rccl.cpp b/source/lib/rocprofiler-sdk/rccl/rccl.cpp index eb55507fbd..322ccfe63e 100644 --- a/source/lib/rocprofiler-sdk/rccl/rccl.cpp +++ b/source/lib/rocprofiler-sdk/rccl/rccl.cpp @@ -189,6 +189,7 @@ rccl_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->internal; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -205,6 +206,7 @@ rccl_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -248,6 +250,7 @@ rccl_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp b/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp index c06b488e9e..6dc74b4886 100644 --- a/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp +++ b/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp @@ -188,6 +188,7 @@ rocdecode_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -204,6 +205,7 @@ rocdecode_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -247,6 +249,7 @@ rocdecode_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp b/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp index 3e6e1361c9..d21aec4976 100644 --- a/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp +++ b/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp @@ -188,6 +188,7 @@ rocjpeg_api_impl::functor(Args... args) auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; + auto ancestor_corr_id = corr_id->ancestor; tracing::populate_external_correlation_ids(external_corr_ids, thr_id, @@ -204,6 +205,7 @@ rocjpeg_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::callback_domain_idx, info_type::operation_idx, tracer_data); @@ -247,6 +249,7 @@ rocjpeg_api_impl::functor(Args... args) thr_id, internal_corr_id, external_corr_ids, + ancestor_corr_id, info_type::buffered_domain_idx, info_type::operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/runtime_initialization.cpp b/source/lib/rocprofiler-sdk/runtime_initialization.cpp index 50fd587651..391316baf0 100644 --- a/source/lib/rocprofiler-sdk/runtime_initialization.cpp +++ b/source/lib/rocprofiler-sdk/runtime_initialization.cpp @@ -119,7 +119,8 @@ initialize(rocprofiler_runtime_initialization_operation_t operation_idx, constexpr auto callback_domain_idx = ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION; constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_RUNTIME_INITIALIZATION; - constexpr auto corr_id = rocprofiler_correlation_id_t{0, rocprofiler_user_data_t{.value = 0}}; + constexpr auto corr_id = + rocprofiler_correlation_id_t{0, rocprofiler_user_data_t{.value = 0}, 0}; ROCP_INFO << pretty_name_by_id(operation_idx) << " has been initialized"; @@ -138,6 +139,7 @@ initialize(rocprofiler_runtime_initialization_operation_t operation_idx, thr_id, corr_id.internal, data.external_correlation_ids, + corr_id.ancestor, callback_domain_idx, operation_idx, tracer_data); @@ -151,6 +153,7 @@ initialize(rocprofiler_runtime_initialization_operation_t operation_idx, thr_id, corr_id.internal, data.external_correlation_ids, + corr_id.ancestor, buffered_domain_idx, operation_idx, buffer_record); diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index 8171a07afe..fef6adfd8b 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -319,10 +319,13 @@ DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue, rocprofiler_user_data_t* user_data, const context::correlation_id* corr_id) { - rocprofiler_correlation_id_t rocprof_corr_id = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; + rocprofiler_async_correlation_id_t rocprof_corr_id = + rocprofiler_async_correlation_id_t{.internal = 0, .external = context::null_user_data}; - if(corr_id) rocprof_corr_id.internal = corr_id->internal; + if(corr_id) + { + rocprof_corr_id.internal = corr_id->internal; + } // TODO: Get external std::shared_lock lk(agents_map_mut); diff --git a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index 1120f9a89b..7f0243ddce 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -152,7 +152,7 @@ TEST(thread_trace, configure_test) params.size(), [](rocprofiler_agent_id_t, rocprofiler_queue_id_t, - rocprofiler_correlation_id_t, + rocprofiler_async_correlation_id_t, rocprofiler_kernel_id_t, rocprofiler_dispatch_id_t, void*, @@ -209,7 +209,7 @@ TEST(thread_trace, perfcounters_configure_test) params.size(), [](rocprofiler_agent_id_t, rocprofiler_queue_id_t, - rocprofiler_correlation_id_t, + rocprofiler_async_correlation_id_t, rocprofiler_kernel_id_t, rocprofiler_dispatch_id_t, void*, diff --git a/source/lib/rocprofiler-sdk/tracing/tracing.hpp b/source/lib/rocprofiler-sdk/tracing/tracing.hpp index 055243dc56..8fb6e55084 100644 --- a/source/lib/rocprofiler-sdk/tracing/tracing.hpp +++ b/source/lib/rocprofiler-sdk/tracing/tracing.hpp @@ -25,6 +25,7 @@ #include "lib/common/mpl.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/context/correlation_id.hpp" #include "lib/rocprofiler-sdk/tracing/fwd.hpp" #include @@ -202,6 +203,7 @@ execute_phase_none_callbacks(callback_context_data_vec_t& callback_conte rocprofiler_thread_id_t thr_id, uint64_t internal_corr_id, const external_correlation_id_map_t& external_corr_ids, + uint64_t ancestor_corr_id, rocprofiler_callback_tracing_kind_t domain, rocprofiler_tracing_operation_t operation, TracerDataT& tracer_data) @@ -215,7 +217,8 @@ execute_phase_none_callbacks(callback_context_data_vec_t& callback_conte auto& user_data = itr.user_data; const auto& extern_corr_id_v = external_corr_ids.at(ctx); - auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; + auto corr_id_v = + rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v, ancestor_corr_id}; record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, thr_id, corr_id_v, @@ -235,6 +238,7 @@ execute_phase_enter_callbacks(callback_context_data_vec_t& callback_cont rocprofiler_thread_id_t thr_id, uint64_t internal_corr_id, const external_correlation_id_map_t& external_corr_ids, + uint64_t ancestor_corr_id, rocprofiler_callback_tracing_kind_t domain, rocprofiler_tracing_operation_t operation, TracerDataT& tracer_data) @@ -248,7 +252,8 @@ execute_phase_enter_callbacks(callback_context_data_vec_t& callback_cont auto& user_data = itr.user_data; const auto& extern_corr_id_v = external_corr_ids.at(ctx); - auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; + auto corr_id_v = + rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v, ancestor_corr_id}; record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, thr_id, corr_id_v, @@ -279,8 +284,8 @@ execute_phase_exit_callbacks(callback_context_data_vec_t& callback_conte auto& user_data = itr.user_data; const auto& extern_corr_id_v = external_corr_ids.at(ctx); - auto corr_id_v = - rocprofiler_correlation_id_t{record.correlation_id.internal, extern_corr_id_v}; + auto corr_id_v = rocprofiler_correlation_id_t{ + record.correlation_id.internal, extern_corr_id_v, record.correlation_id.ancestor}; record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, record.thread_id, corr_id_v, @@ -300,15 +305,33 @@ execute_buffer_record_emplace(const buffered_context_data_vec_t& buffered_cont rocprofiler_thread_id_t thr_id, uint64_t internal_corr_id, const external_correlation_id_map_t& external_corr_ids, + uint64_t ancestor_corr_id, rocprofiler_buffer_tracing_kind_t domain, OperationT operation, BufferRecordT&& base_record) { + using record_corr_id_t = decltype(base_record.correlation_id); + base_record.thread_id = thr_id; base_record.kind = domain; base_record.operation = operation; + // external correlation will be updated right before record is placed in buffer - base_record.correlation_id = rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; + if constexpr(std::is_same_v) + { + base_record.correlation_id = + rocprofiler_correlation_id_t{internal_corr_id, empty_user_data, ancestor_corr_id}; + } + else if constexpr(std::is_same_v) + { + base_record.correlation_id = + rocprofiler_async_correlation_id_t{internal_corr_id, empty_user_data}; + } + else + { + static_assert(common::mpl::assert_false::value, + "Invalid correlation ID type"); + } for(const auto& itr : buffered_contexts) { diff --git a/tests/async-copy-tracing/validate.py b/tests/async-copy-tracing/validate.py index 17daef6c6e..14d7218f4c 100644 --- a/tests/async-copy-tracing/validate.py +++ b/tests/async-copy-tracing/validate.py @@ -34,6 +34,19 @@ def node_exists(name, data, min_len=1): assert len(data[name]) >= min_len, f"{name}:\n{data}" +def get_operation(record, kind_name, op_name=None): + for idx, itr in enumerate(record["names"]): + if kind_name == itr["kind"]: + if op_name is None: + return idx, itr["operations"] + else: + for oidx, oname in enumerate(itr["operations"]): + if op_name == oname: + return oidx + + return None + + def test_data_structure(input_data): """verify minimum amount of expected data is present""" data = input_data @@ -334,6 +347,125 @@ def test_async_copy_direction(input_data): ) == 0, f"TID={tid}:\n\t{async_dir_cnt}" +def test_ancestor_ids(input_data): + from rocprofiler_sdk.pytest_utils.dotdict import dotdict + + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + buffer_records = sdk_data.buffer_records + + hip_memcopy_id = get_operation(buffer_records, "HIP_RUNTIME_API", "hipMemcpyAsync") + has_async_memcopy_id = get_operation( + buffer_records, "HSA_AMD_EXT_API", "hsa_amd_memory_async_copy_on_engine" + ) + + def by_thread_id(trace_item, op_id, corr_id_name): + """ + Returns a dict like so, where corr_id_name == "ancestor" or "internal" + { + tid: { + corr_id-1: [ list of objects with value of corr_id_name == corr_id-1 ] + corr_id-2: [ list of objects with value of corr_id_name == corr_id-2 ] + ... + } + } + """ + ret = {} + + for x in trace_item: + if x.operation == op_id: + corr_id = x.correlation_id[corr_id_name] + + if x.thread_id not in ret.keys(): + ret[x.thread_id] = {} + + if corr_id in ret[x.thread_id].keys(): + ret[x.thread_id][corr_id].append(x) + else: + ret[x.thread_id][corr_id] = [x] + return dotdict(ret) + + # dict with tid -> internal id -> record + hip_memcopies = by_thread_id( + buffer_records.hip_api_traces, hip_memcopy_id, "internal" + ) + # dict with tid -> ancestor id -> record + hsa_memcopies = by_thread_id( + buffer_records.hsa_api_traces, has_async_memcopy_id, "ancestor" + ) + + # We expect the same keys in hip.thread_id and hsa.thread_id, because they are parent <-> child corr ids + # both must have same thread ids + # For each tid, hip.ancestor and hsa.internal keys must be same + assert ( + hip_memcopies.keys() == hsa_memcopies.keys() + ), "TIDs are not same for hsa and hip calls" + for tid in hip_memcopies.keys(): + assert ( + hip_memcopies[tid].keys() == hsa_memcopies[tid].keys() + ), "Parent-child keys are not equal" + + memcopies = buffer_records.memory_copies + + num_hsa_memcopies = sum( + [ + len(val) + for tid in hsa_memcopies.keys() + for _, val in hsa_memcopies[tid].items() + ] + ) + assert ( + len(memcopies) == num_hsa_memcopies + ), "Expected number of memcopies to be same as number of async HSA (hsa_amd_memory_async_copy_on_engine) calls" + + print(" > hip_memcopy_id : ", hip_memcopy_id) + print(" > has_async_memcopy_id : ", has_async_memcopy_id) + print( + " > len(hip_memcopies) : ", num_hsa_memcopies, len(hip_memcopies), len(memcopies) + ) + + for tid in hip_memcopies: + # We expect only 1 record with this internal id, per thread + for corr_id, records in hip_memcopies[tid].items(): + assert ( + len(records) == 1 + ), "Expected only one record with this (thread_id, corr_id)" + + for _, tid_records in hsa_memcopies.items(): + for _, records in tid_records.items(): + for record in records: + child_memcpy_record = [ + x + for x in buffer_records.memory_copies + if x.correlation_id.internal == record.correlation_id.internal + ] + assert ( + len(child_memcpy_record) == 1 + ), "Expected only one memcpy record with this internal ID" + child_memcpy_record = child_memcpy_record[0] + + assert record.thread_id in hsa_memcopies.keys() + assert record.thread_id in hip_memcopies.keys() + assert ( + record.correlation_id.ancestor + in hip_memcopies[record.thread_id].keys() + ) + # This is basically the same as the Expected only one record with this (thread_id, corr_id) test + assert ( + len(hip_memcopies[record.thread_id][record.correlation_id.ancestor]) + == 1 + ), "Expected only one record with this (thread_id, corr_id)" + + ancestor_hip_record = hip_memcopies[record.thread_id][ + record.correlation_id.ancestor + ][0] + assert ( + ancestor_hip_record.correlation_id.internal + == record.correlation_id.ancestor + ) + assert ancestor_hip_record.thread_id == record.thread_id + + def test_retired_correlation_ids(input_data): data = input_data sdk_data = data["rocprofiler-sdk-json-tool"] diff --git a/tests/thread-trace/multi_dispatch.cpp b/tests/thread-trace/multi_dispatch.cpp index fcdc99a860..09c27b7ddf 100644 --- a/tests/thread-trace/multi_dispatch.cpp +++ b/tests/thread-trace/multi_dispatch.cpp @@ -39,7 +39,7 @@ rocprofiler_client_id_t* client_id = nullptr; rocprofiler_att_control_flags_t dispatch_callback(rocprofiler_agent_id_t /* agent */, rocprofiler_queue_id_t /* queue_id */, - rocprofiler_correlation_id_t /* correlation_id */, + rocprofiler_async_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t /* kernel_id */, rocprofiler_dispatch_id_t /* dispatch_id */, void* userdata, diff --git a/tests/thread-trace/single_dispatch.cpp b/tests/thread-trace/single_dispatch.cpp index c5432bc325..8ad0ec0298 100644 --- a/tests/thread-trace/single_dispatch.cpp +++ b/tests/thread-trace/single_dispatch.cpp @@ -38,7 +38,7 @@ rocprofiler_client_id_t* client_id = nullptr; rocprofiler_att_control_flags_t dispatch_callback(rocprofiler_agent_id_t /* agent */, rocprofiler_queue_id_t /* queue_id */, - rocprofiler_correlation_id_t /* correlation_id */, + rocprofiler_async_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t /* dispatch_id */, void* userdata,