[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
[ROCm/rocprofiler-sdk commit: f27f76716e]
Este cometimento está contido em:
cometido por
GitHub
ascendente
8d3f0c6881
cometimento
fbcb5025e7
+7
-7
@@ -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.
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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);
|
||||
|
||||
+7
-7
@@ -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;
|
||||
|
||||
/**
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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;
|
||||
|
||||
/** @} */
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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*/,
|
||||
|
||||
@@ -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());
|
||||
|
||||
|
||||
@@ -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();
|
||||
|
||||
+1
-1
@@ -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;
|
||||
|
||||
+2
-1
@@ -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;
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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));
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
+10
-2
@@ -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.
|
||||
|
||||
@@ -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
|
||||
|
||||
+4
-4
@@ -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
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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);
|
||||
|
||||
+2
-2
@@ -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*,
|
||||
|
||||
@@ -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)
|
||||
{
|
||||
|
||||
@@ -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"]
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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,
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador