AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

[SDK] Add Stack IDs (#269)

Bladeren bron 
* 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
This commit is contained in:
Kuricheti, Mythreya
2025-03-20 22:52:48 -07:00
gecommit door GitHub
bovenliggende ae0db8cee5
commit f27f76716e
36 gewijzigde bestanden met toevoegingen van 341 en 67 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h
+7 -7
Bestand weergeven
@@ -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.
source/include/rocprofiler-sdk/buffer_tracing.h
+2 -2
Bestand weergeven
@@ -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
source/include/rocprofiler-sdk/cxx/serialization.hpp
+9
Bestand weergeven
@@ -131,6 +131,15 @@ save(ArchiveT& ar, rocprofiler_counter_id_t data)
template <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_async_correlation_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(internal);
ROCP_SDK_SAVE_DATA_VALUE("external", external.value);
source/include/rocprofiler-sdk/dispatch_counting_service.h
+7 -7
Bestand weergeven
@@ -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;
/**
source/include/rocprofiler-sdk/fwd.h
+30
Bestand weergeven
@@ -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.
source/include/rocprofiler-sdk/pc_sampling.h
+7 -4
Bestand weergeven
@@ -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;
/** @} */
source/lib/output/generatePerfetto.cpp
+20 -6
Bestand weergeven
@@ -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<sdk::category::hsa_api>::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<sdk::category::hip_api>::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<sdk::category::marker_api>::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<sdk::category::rccl_api>::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<sdk::category::rocdecode_api>::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<sdk::category::rocjpeg_api>::name,
track,
itr.end_timestamp);
source/lib/rocprofiler-sdk-tool/tool.cpp
+1 -1
Bestand weergeven
@@ -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*/,
source/lib/rocprofiler-sdk/context/correlation_id.cpp
+6 -1
Bestand weergeven
@@ -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<correlation_id>(_init_ref_count, common::get_tid(), _internal_id);
ret = std::make_unique<correlation_id>(_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());
source/lib/rocprofiler-sdk/context/correlation_id.hpp
+1 -3
Bestand weergeven
@@ -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<uint32_t> 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();
source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp
+1 -1
Bestand weergeven
@@ -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;
source/lib/rocprofiler-sdk/counters/sample_processing.cpp
+2 -1
Bestand weergeven
@@ -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;
source/lib/rocprofiler-sdk/counters/tests/core.cpp
+9 -9
Bestand weergeven
@@ -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
source/lib/rocprofiler-sdk/hip/hip.cpp
+3
Bestand weergeven
@@ -224,6 +224,7 @@ hip_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/hip/stream.cpp
+6
Bestand weergeven
@@ -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<hipStream_t*, function_args_type>::value;
auto stream = std::get<stream_idx>(std::make_tuple(std::forward<Args>(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<stream_idx>(std::make_tuple(std::forward<Args>(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<stream_idx>(std::make_tuple(std::forward<Args>(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);
source/lib/rocprofiler-sdk/hsa/async_copy.cpp
+3 -1
Bestand weergeven
@@ -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);
source/lib/rocprofiler-sdk/hsa/hsa.cpp
+3 -1
Bestand weergeven
@@ -340,7 +340,7 @@ hsa_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/hsa/memory_allocation.cpp
+6 -1
Bestand weergeven
@@ -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);
source/lib/rocprofiler-sdk/hsa/queue.cpp
+2
Bestand weergeven
@@ -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);
source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp
+3
Bestand weergeven
@@ -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));
source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp
+4 -1
Bestand weergeven
@@ -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);
source/lib/rocprofiler-sdk/marker/marker.cpp
+3
Bestand weergeven
@@ -170,6 +170,7 @@ roctx_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/ompt/ompt.cpp
+7
Bestand weergeven
@@ -739,6 +739,7 @@ ompt_impl<OpIdx>::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<OpIdx>::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<OpIdx>::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<OpIdx>::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<OpIdx>::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<OpIdx>::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<OpIdx>::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);
source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp
+10 -2
Bestand weergeven
@@ -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<Parser::DispatchPkt>
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.
source/lib/rocprofiler-sdk/pc_sampling/parser/rocr.h
+2 -2
Bestand weergeven
@@ -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
source/lib/rocprofiler-sdk/pc_sampling/parser/stochastic_records.h
+4 -4
Bestand weergeven
@@ -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
source/lib/rocprofiler-sdk/rccl/rccl.cpp
+3
Bestand weergeven
@@ -189,6 +189,7 @@ rccl_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp
+3
Bestand weergeven
@@ -188,6 +188,7 @@ rocdecode_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp
+3
Bestand weergeven
@@ -188,6 +188,7 @@ rocjpeg_api_impl<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::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);
source/lib/rocprofiler-sdk/runtime_initialization.cpp
+4 -1
Bestand weergeven
@@ -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);
source/lib/rocprofiler-sdk/thread_trace/att_core.cpp
+6 -3
Bestand weergeven
@@ -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<std::shared_mutex> lk(agents_map_mut);
source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp
+2 -2
Bestand weergeven
@@ -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*,
source/lib/rocprofiler-sdk/tracing/tracing.hpp
+28 -5
Bestand weergeven
@@ -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 <rocprofiler-sdk/fwd.h>
@@ -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<rocprofiler_correlation_id_t, record_corr_id_t>)
{
base_record.correlation_id =
rocprofiler_correlation_id_t{internal_corr_id, empty_user_data, ancestor_corr_id};
}
else if constexpr(std::is_same_v<rocprofiler_async_correlation_id_t, record_corr_id_t>)
{
base_record.correlation_id =
rocprofiler_async_correlation_id_t{internal_corr_id, empty_user_data};
}
else
{
static_assert(common::mpl::assert_false<record_corr_id_t>::value,
"Invalid correlation ID type");
}
for(const auto& itr : buffered_contexts)
{