Allow events to be grouped by HIP stream ID (#274)
- Corelate memory_copy and kernel_dispatch events with their HIP stream_id and add stream_id as an annotation in Perfetto.
- By default, group memory_copy and kernel_dispatch events in Perfetto output by their stream_id.
- Add option, with the configuration setting ROCPROFSYS_ROCM_GROUP_BY_QUEUE, to group by HSA queue instead.
---------
Signed-off-by: David Galiffi <David.Galiffi@amd.com>
Co-authored-by: David Galiffi <David.Galiffi@amd.com>
[ROCm/rocprofiler-systems commit: 4b4a846b58]
Tento commit je obsažen v:
ajanicijamd
GitHub
@@ -10,6 +10,8 @@ Full documentation for ROCm Systems Profiler is available at [https://rocm.docs.
|
||||
- How-to document for VCN and JPEG activity sampling and tracing.
|
||||
- Support for tracing Fortran applications.
|
||||
- Support for tracing MPI API in Fortran.
|
||||
- By default, group "kernel dispatch" and "memory copy" events by HIP stream ID in Perfetto traces.
|
||||
- Add the "ROCPROFSYS_ROCM_GROUP_BY_QUEUE" configuration setting to group events by queue, instead.
|
||||
|
||||
### Changed
|
||||
|
||||
|
||||
@@ -217,6 +217,14 @@ The following example:
|
||||
|
||||
ROCPROFSYS_ROCM_EVENTS = GPUBusy SQ_WAVES:device=0 SQ_INSTS_VALU:device=1
|
||||
|
||||
ROCPROFSYS_ROCM_GROUP_BY_QUEUE
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
By default, Perfetto trace will show the HIP streams to which kernel
|
||||
and memory copy operations submitted. With the
|
||||
``ROCPROFSYS_ROCM_GROUP_BY_QUEUE=ON`` setting, the trace will display HSA queues
|
||||
to which these kernel and memory operations were submitted.
|
||||
|
||||
Exploring GPU Metrics
|
||||
---------------------
|
||||
|
||||
|
||||
@@ -32,7 +32,7 @@ rocprof-sys-sample: `source/bin/rocprof-sys-sample <https://github.com/ROCm/rocp
|
||||
* Adds ``librocprof-sys-dl.so`` to ``LD_PRELOAD``
|
||||
* Is launched by using ``execvpe`` with ``<command> <command-args>`` and a modified environment
|
||||
|
||||
rocprof-sys-casual: `source/bin/rocprof-sys-causal <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-causal>`_
|
||||
rocprof-sys-causal: `source/bin/rocprof-sys-causal <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-causal>`_
|
||||
---------------------------------------------------------------------------------------------------------------------------------------------------
|
||||
|
||||
When there is exactly one causal profiling configuration variant (which enables debugging),
|
||||
|
||||
@@ -96,6 +96,7 @@ ROCPROFSYS_DEFINE_CATEGORY(category, rocm_hip_api, ROCPROFSYS_CATEGORY_ROCM_HIP_
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_hsa_api, ROCPROFSYS_CATEGORY_ROCM_HSA_API, "rocm_hsa_api", "ROCm HSA functions")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_kernel_dispatch, ROCPROFSYS_CATEGORY_ROCM_KERNEL_DISPATCH, "rocm_kernel_dispatch", "ROCm Kernel dispatch")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_memory_copy, ROCPROFSYS_CATEGORY_ROCM_MEMORY_COPY, "rocm_memory_copy", "ROCm Async Memory Copy")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_hip_stream, ROCPROFSYS_CATEGORY_ROCM_HIP_STREAM, "rocm_hip_stream", "ROCm HIP Stream")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_scratch_memory, ROCPROFSYS_CATEGORY_ROCM_SCRATCH_MEMORY, "rocm_scratch_memory", "ROCm kernel scratch memory reallocations")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_page_migration, ROCPROFSYS_CATEGORY_ROCM_PAGE_MIGRATION, "rocm_page_migration", "ROCm memory page migration")
|
||||
ROCPROFSYS_DEFINE_CATEGORY(category, rocm_counter_collection, ROCPROFSYS_CATEGORY_ROCM_COUNTER_COLLECTION, "rocm_counter_collection", "ROCm device counter collection")
|
||||
@@ -166,6 +167,7 @@ using name = perfetto_category<Tp...>;
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_hsa_api), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_kernel_dispatch), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_memory_copy), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_hip_stream), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_scratch_memory), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_page_migration), \
|
||||
ROCPROFSYS_PERFETTO_CATEGORY(category::rocm_counter_collection), \
|
||||
|
||||
@@ -372,6 +372,22 @@ config_settings(const std::shared_ptr<settings>& _config)
|
||||
|
||||
for(const auto& itr : buffered_tracing_info)
|
||||
_add_operation_settings(itr.name, itr, buffered_operation_option_names);
|
||||
|
||||
// Add the ROCPROFSYS_ROCM_GROUP_BY_QUEUE setting if the hip_stream domain is present
|
||||
// in supported ROCProfiler-SDK domains.
|
||||
auto _has_hip_stream = std::find(_domain_choices.begin(), _domain_choices.end(),
|
||||
"hip_stream") != _domain_choices.end();
|
||||
|
||||
if(_has_hip_stream)
|
||||
{
|
||||
ROCPROFSYS_CONFIG_SETTING(
|
||||
bool, "ROCPROFSYS_ROCM_GROUP_BY_QUEUE",
|
||||
"By default, Perfetto trace will show the HIP streams to which kernel "
|
||||
"and memory copy operations submitted. With the "
|
||||
"`ROCPROFSYS_ROCM_GROUP_BY_QUEUE` option, the trace will display HSA queues "
|
||||
"to which these kernel and memory operations were submitted.",
|
||||
false, "rocm", "perfetto");
|
||||
}
|
||||
}
|
||||
|
||||
std::unordered_set<rocprofiler_callback_tracing_kind_t>
|
||||
@@ -550,6 +566,16 @@ get_rocm_events()
|
||||
" ,;\t\n");
|
||||
}
|
||||
|
||||
bool
|
||||
get_group_by_queue(void)
|
||||
{
|
||||
std::optional<bool> _group_by_queue =
|
||||
config::get_setting_value<bool>("ROCPROFSYS_ROCM_GROUP_BY_QUEUE");
|
||||
bool _ret = _group_by_queue.value_or(true);
|
||||
|
||||
return _ret;
|
||||
}
|
||||
|
||||
std::vector<int32_t>
|
||||
get_operations(rocprofiler_callback_tracing_kind_t kindv)
|
||||
{
|
||||
|
||||
@@ -70,6 +70,9 @@ get_operations(rocprofiler_buffer_tracing_kind_t kindv);
|
||||
std::vector<std::string>
|
||||
get_rocm_events();
|
||||
|
||||
bool
|
||||
get_group_by_queue();
|
||||
|
||||
std::unordered_set<int32_t>
|
||||
get_backtrace_operations(rocprofiler_callback_tracing_kind_t kindv);
|
||||
|
||||
|
||||
@@ -49,6 +49,7 @@ extern "C"
|
||||
ROCPROFSYS_CATEGORY_ROCM_KERNEL_DISPATCH,
|
||||
ROCPROFSYS_CATEGORY_ROCM_MEMORY_COPY,
|
||||
ROCPROFSYS_CATEGORY_ROCM_SCRATCH_MEMORY,
|
||||
ROCPROFSYS_CATEGORY_ROCM_HIP_STREAM,
|
||||
ROCPROFSYS_CATEGORY_ROCM_PAGE_MIGRATION,
|
||||
ROCPROFSYS_CATEGORY_ROCM_COUNTER_COLLECTION,
|
||||
ROCPROFSYS_CATEGORY_ROCM_MARKER_API,
|
||||
|
||||
+296
-118
@@ -92,6 +92,70 @@ thread_postcreate(rocprofiler_runtime_library_t /*lib*/, void* /*tool_data*/)
|
||||
pop_thread_state();
|
||||
}
|
||||
|
||||
#if(ROCPROFILER_VERSION < 700)
|
||||
/**
|
||||
* @brief Stream ID.
|
||||
*/
|
||||
typedef struct rocprofiler_stream_id_t
|
||||
{
|
||||
uint64_t handle;
|
||||
} rocprofiler_stream_id_t;
|
||||
|
||||
#endif
|
||||
|
||||
auto&
|
||||
get_stream_stack()
|
||||
{
|
||||
static thread_local std::vector<rocprofiler_stream_id_t> _v{ rocprofiler_stream_id_t{
|
||||
0 } };
|
||||
return _v;
|
||||
}
|
||||
|
||||
void
|
||||
stream_id_push(rocprofiler_stream_id_t stream_id)
|
||||
{
|
||||
get_stream_stack().emplace_back(stream_id);
|
||||
}
|
||||
|
||||
rocprofiler_stream_id_t
|
||||
stream_id_top()
|
||||
{
|
||||
auto stream_id = get_stream_stack().back();
|
||||
return stream_id;
|
||||
}
|
||||
|
||||
void
|
||||
stream_id_pop()
|
||||
{
|
||||
get_stream_stack().pop_back();
|
||||
}
|
||||
|
||||
// Stores stream ids and kernel region ids for kernel-rename service and hip stream
|
||||
// display service
|
||||
struct kernel_rename_and_stream_data
|
||||
{
|
||||
uint64_t region_id = 0; // roctx region correlation id
|
||||
rocprofiler_stream_id_t stream_id = { 0 };
|
||||
};
|
||||
|
||||
template <typename Tp>
|
||||
rocprofiler_stream_id_t
|
||||
get_stream_id(Tp* _record)
|
||||
{
|
||||
auto _stream_id = rocprofiler_stream_id_t{ 0 };
|
||||
if(_record->correlation_id.external.ptr != nullptr)
|
||||
{
|
||||
// Extract the stream id
|
||||
auto* _ecid_data = static_cast<kernel_rename_and_stream_data*>(
|
||||
_record->correlation_id.external.ptr);
|
||||
_stream_id = _ecid_data->stream_id;
|
||||
auto _region_id = _ecid_data->region_id;
|
||||
_record->correlation_id.external.value = _region_id;
|
||||
delete _ecid_data;
|
||||
}
|
||||
return _stream_id;
|
||||
}
|
||||
|
||||
// this function creates a rocprofiler profile config on the first entry
|
||||
std::vector<rocprofiler_counter_id_t>
|
||||
create_agent_profile(rocprofiler_agent_id_t agent_id,
|
||||
@@ -369,8 +433,9 @@ tool_tracing_callback_stop(
|
||||
&args);
|
||||
}
|
||||
|
||||
uint64_t _beg_ts = begin_ts;
|
||||
uint64_t _end_ts = ts;
|
||||
uint64_t _beg_ts = begin_ts;
|
||||
uint64_t _end_ts = ts;
|
||||
auto stream_id = stream_id_top();
|
||||
|
||||
tracing::push_perfetto_ts(
|
||||
CategoryT{}, _name.data(), _beg_ts,
|
||||
@@ -381,6 +446,9 @@ tool_tracing_callback_stop(
|
||||
tracing::add_perfetto_annotation(ctx, "begin_ns", _beg_ts);
|
||||
tracing::add_perfetto_annotation(ctx, "corr_id",
|
||||
record.correlation_id.internal);
|
||||
if(stream_id.handle != 0)
|
||||
tracing::add_perfetto_annotation(ctx, "stream_id",
|
||||
stream_id.handle);
|
||||
for(const auto& [key, val] : args)
|
||||
tracing::add_perfetto_annotation(ctx, key, val);
|
||||
|
||||
@@ -717,6 +785,15 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
|
||||
{
|
||||
if(num_headers == 0 || headers == nullptr) return;
|
||||
|
||||
auto _track_desc_stream = [](uint64_t _stream_id) {
|
||||
return JOIN("", "HIP Activity Stream ", _stream_id);
|
||||
};
|
||||
|
||||
bool _group_by_queue = get_group_by_queue();
|
||||
|
||||
static auto _mtx = std::mutex{};
|
||||
auto _lk = std::unique_lock<std::mutex>{ _mtx };
|
||||
|
||||
for(size_t i = 0; i < num_headers; ++i)
|
||||
{
|
||||
auto* header = headers[i];
|
||||
@@ -740,6 +817,12 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
|
||||
auto _queue_id = record->dispatch_info.queue_id;
|
||||
const auto* _agent = tool_data->get_gpu_tool_agent(_agent_id);
|
||||
|
||||
uint64_t _stream_id = get_stream_id(record).handle;
|
||||
ROCPROFSYS_CI_THROW(
|
||||
_stream_id == 0,
|
||||
"Unexpected zero stream_id in kernel dispatch record: %s.",
|
||||
_name.c_str());
|
||||
|
||||
if(get_use_timemory())
|
||||
{
|
||||
const auto& _tinfo = thread_info::get(record->thread_id, SystemTID);
|
||||
@@ -757,59 +840,77 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
|
||||
|
||||
if(get_use_perfetto())
|
||||
{
|
||||
auto _track_desc = [](int32_t _device_id_v, int64_t _queue_id_v) {
|
||||
return JOIN("", "GPU Kernel Dispatch [", _device_id_v, "] Queue ",
|
||||
_queue_id_v);
|
||||
// Lambda to add common perfetto annotations for kernel dispatch
|
||||
auto add_perfetto_annotations = [&](::perfetto::EventContext ctx) {
|
||||
if(config::get_perfetto_annotations())
|
||||
{
|
||||
tracing::add_perfetto_annotation(ctx, "begin_ns", _beg_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "end_ns", _end_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "corr_id", _corr_id);
|
||||
tracing::add_perfetto_annotation(ctx, "stream_id",
|
||||
_stream_id);
|
||||
|
||||
tracing::add_perfetto_annotation(ctx, "queue",
|
||||
_queue_id.handle);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "dispatch_id", record->dispatch_info.dispatch_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "kernel_id", record->dispatch_info.kernel_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "private_segment_size",
|
||||
record->dispatch_info.private_segment_size);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "group_segment_size",
|
||||
record->dispatch_info.group_segment_size);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "workgroup_size",
|
||||
JOIN("", "(",
|
||||
JOIN(',', record->dispatch_info.workgroup_size.x,
|
||||
record->dispatch_info.workgroup_size.y,
|
||||
record->dispatch_info.workgroup_size.z),
|
||||
")"));
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "grid_size",
|
||||
JOIN("", "(",
|
||||
JOIN(',', record->dispatch_info.grid_size.x,
|
||||
record->dispatch_info.grid_size.y,
|
||||
record->dispatch_info.grid_size.z),
|
||||
")"));
|
||||
}
|
||||
};
|
||||
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_kernel_dispatch{}, _track_desc, _agent->device_id,
|
||||
_queue_id.handle);
|
||||
if(_group_by_queue)
|
||||
{
|
||||
auto _track_desc = [](int32_t _device_id_v, int64_t _queue_id_v) {
|
||||
return JOIN("", "GPU Kernel Dispatch [", _device_id_v,
|
||||
"] Queue ", _queue_id_v);
|
||||
};
|
||||
|
||||
tracing::push_perfetto(
|
||||
category::rocm_kernel_dispatch{}, _name.c_str(), _track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
[&](::perfetto::EventContext ctx) {
|
||||
if(config::get_perfetto_annotations())
|
||||
{
|
||||
tracing::add_perfetto_annotation(ctx, "begin_ns",
|
||||
_beg_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "end_ns", _end_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "corr_id",
|
||||
_corr_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "node_id", _agent->agent->logical_node_id);
|
||||
tracing::add_perfetto_annotation(ctx, "queue",
|
||||
_queue_id.handle);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "dispatch_id",
|
||||
record->dispatch_info.dispatch_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "kernel_id", record->dispatch_info.kernel_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "private_segment_size",
|
||||
record->dispatch_info.private_segment_size);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "group_segment_size",
|
||||
record->dispatch_info.group_segment_size);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "workgroup_size",
|
||||
JOIN("", "(",
|
||||
JOIN(',', record->dispatch_info.workgroup_size.x,
|
||||
record->dispatch_info.workgroup_size.y,
|
||||
record->dispatch_info.workgroup_size.z),
|
||||
")"));
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "grid_size",
|
||||
JOIN("", "(",
|
||||
JOIN(',', record->dispatch_info.grid_size.x,
|
||||
record->dispatch_info.grid_size.y,
|
||||
record->dispatch_info.grid_size.z),
|
||||
")"));
|
||||
}
|
||||
});
|
||||
tracing::pop_perfetto(category::rocm_kernel_dispatch{}, _name.c_str(),
|
||||
_track, _end_ns);
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_kernel_dispatch{}, _track_desc,
|
||||
_agent->device_id, _queue_id.handle);
|
||||
|
||||
tracing::push_perfetto(category::rocm_kernel_dispatch{},
|
||||
_name.c_str(), _track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
add_perfetto_annotations);
|
||||
|
||||
tracing::pop_perfetto(category::rocm_kernel_dispatch{},
|
||||
_name.c_str(), _track, _end_ns);
|
||||
}
|
||||
else
|
||||
{
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_hip_stream{}, _track_desc_stream, _stream_id);
|
||||
|
||||
tracing::push_perfetto(category::rocm_hip_stream{}, _name.c_str(),
|
||||
_track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
add_perfetto_annotations);
|
||||
|
||||
tracing::pop_perfetto(category::rocm_hip_stream{}, _name.c_str(),
|
||||
_track, _end_ns);
|
||||
}
|
||||
}
|
||||
}
|
||||
else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
|
||||
@@ -828,6 +929,11 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
|
||||
auto _name =
|
||||
tool_data->buffered_tracing_info.at(record->kind, record->operation);
|
||||
|
||||
uint64_t _stream_id = get_stream_id(record).handle;
|
||||
ROCPROFSYS_CI_THROW(
|
||||
_stream_id == 0,
|
||||
"Unexpected zero stream_id in memory copy record: %s.", _name.data());
|
||||
|
||||
if(get_use_timemory())
|
||||
{
|
||||
const auto& _tinfo = thread_info::get(record->thread_id, SystemTID);
|
||||
@@ -845,36 +951,55 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
|
||||
|
||||
if(get_use_perfetto())
|
||||
{
|
||||
auto _track_desc = [](int32_t _device_id_v,
|
||||
rocprofiler_thread_id_t _tid) {
|
||||
const auto& _tid_v = thread_info::get(_tid, SystemTID);
|
||||
return JOIN("", "GPU Memory Copy to Agent [", _device_id_v,
|
||||
"] Thread ", _tid_v->index_data->sequent_value);
|
||||
auto add_perfetto_annotations = [&](::perfetto::EventContext ctx) {
|
||||
if(config::get_perfetto_annotations())
|
||||
{
|
||||
tracing::add_perfetto_annotation(ctx, "begin_ns", _beg_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "end_ns", _end_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "corr_id", _corr_id);
|
||||
tracing::add_perfetto_annotation(ctx, "stream_id",
|
||||
_stream_id);
|
||||
tracing::add_perfetto_annotation(ctx, "dst_agent",
|
||||
_dst_agent->logical_node_id);
|
||||
tracing::add_perfetto_annotation(ctx, "src_agent",
|
||||
_src_agent->logical_node_id);
|
||||
}
|
||||
};
|
||||
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_memory_copy{}, _track_desc,
|
||||
_dst_agent->logical_node_id, record->thread_id);
|
||||
if(_group_by_queue)
|
||||
{
|
||||
auto _track_desc = [](int32_t _device_id_v,
|
||||
rocprofiler_thread_id_t _tid) {
|
||||
const auto& _tid_v = thread_info::get(_tid, SystemTID);
|
||||
return JOIN("", "GPU Memory Copy to Agent [", _device_id_v,
|
||||
"] Thread ", _tid_v->index_data->sequent_value);
|
||||
};
|
||||
|
||||
tracing::push_perfetto(
|
||||
category::rocm_memory_copy{}, _name.data(), _track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
[&](::perfetto::EventContext ctx) {
|
||||
if(config::get_perfetto_annotations())
|
||||
{
|
||||
tracing::add_perfetto_annotation(ctx, "begin_ns",
|
||||
_beg_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "end_ns", _end_ns);
|
||||
tracing::add_perfetto_annotation(ctx, "corr_id",
|
||||
_corr_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "dst_agent", _dst_agent->logical_node_id);
|
||||
tracing::add_perfetto_annotation(
|
||||
ctx, "src_agent", _src_agent->logical_node_id);
|
||||
}
|
||||
});
|
||||
tracing::pop_perfetto(category::rocm_memory_copy{}, "", _track,
|
||||
_end_ns);
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_memory_copy{}, _track_desc,
|
||||
_dst_agent->logical_node_id, record->thread_id);
|
||||
|
||||
tracing::push_perfetto(category::rocm_memory_copy{}, _name.data(),
|
||||
_track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
add_perfetto_annotations);
|
||||
|
||||
tracing::pop_perfetto(category::rocm_memory_copy{}, "", _track,
|
||||
_end_ns);
|
||||
}
|
||||
else
|
||||
{
|
||||
const auto _track = tracing::get_perfetto_track(
|
||||
category::rocm_hip_stream{}, _track_desc_stream, _stream_id);
|
||||
|
||||
tracing::push_perfetto(category::rocm_hip_stream{}, _name.data(),
|
||||
_track, _beg_ns,
|
||||
::perfetto::Flow::ProcessScoped(_corr_id),
|
||||
add_perfetto_annotations);
|
||||
|
||||
tracing::pop_perfetto(category::rocm_hip_stream{}, "", _track,
|
||||
_end_ns);
|
||||
}
|
||||
}
|
||||
}
|
||||
else
|
||||
@@ -1005,30 +1130,6 @@ dispatch_counting_service_callback(
|
||||
}
|
||||
}
|
||||
|
||||
// int
|
||||
// external_correlation_id_callback(
|
||||
// rocprofiler_thread_id_t /*thr_id*/, rocprofiler_context_id_t /*ctx_id*/,
|
||||
// rocprofiler_external_correlation_id_request_kind_t /*kind*/,
|
||||
// rocprofiler_tracing_operation_t /*op*/, uint64_t /*internal_corr_id*/,
|
||||
// rocprofiler_user_data_t* external_corr_id, void* /*user_data*/)
|
||||
// {
|
||||
// auto* _data = new kernel_dispatch_bundle_t{ "kernel_dispatch" };
|
||||
// _data->push();
|
||||
// external_corr_id->ptr = _data;
|
||||
// return 0;
|
||||
// }
|
||||
|
||||
// void
|
||||
// agent_counter_profile_callback(rocprofiler_context_id_t context_id,
|
||||
// rocprofiler_agent_id_t agent,
|
||||
// rocprofiler_agent_set_profile_callback_t set_config, void*)
|
||||
// {
|
||||
// if(!agent_counter_profiles) return;
|
||||
// if(auto itr = agent_counter_profiles->find(agent);
|
||||
// itr != agent_counter_profiles->end() && itr->second)
|
||||
// set_config(context_id, *itr->second);
|
||||
// }
|
||||
|
||||
bool
|
||||
is_initialized(rocprofiler_context_id_t ctx)
|
||||
{
|
||||
@@ -1069,6 +1170,73 @@ flush()
|
||||
}
|
||||
}
|
||||
|
||||
int
|
||||
set_kernel_rename_and_stream_correlation_id(
|
||||
rocprofiler_thread_id_t /* thr_id */, rocprofiler_context_id_t /* ctx_id */,
|
||||
rocprofiler_external_correlation_id_request_kind_t /* kind */,
|
||||
rocprofiler_tracing_operation_t /* op */, uint64_t /* internal_corr_id */,
|
||||
rocprofiler_user_data_t* external_corr_id, void* /* user_data */)
|
||||
{
|
||||
auto* _info = new kernel_rename_and_stream_data{};
|
||||
|
||||
_info->stream_id = stream_id_top();
|
||||
|
||||
// Set the external correlation id service to point to struct
|
||||
external_corr_id->ptr = _info;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
#if(ROCPROFILER_VERSION >= 700)
|
||||
void
|
||||
tool_hip_stream_callback(rocprofiler_callback_tracing_record_t record,
|
||||
rocprofiler_user_data_t* /* user_data */, void* /* data */)
|
||||
{
|
||||
if(record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_STREAM) return;
|
||||
// Extract stream ID from record
|
||||
auto* stream_handle_data =
|
||||
static_cast<rocprofiler_callback_tracing_hip_stream_data_t*>(record.payload);
|
||||
auto stream_id = stream_handle_data->stream_id;
|
||||
|
||||
// STREAM_HANDLE_CREATE and DESTROY are no-ops
|
||||
if(record.operation == ROCPROFILER_HIP_STREAM_CREATE)
|
||||
{
|
||||
ROCPROFSYS_VERBOSE_F(
|
||||
2, "Entered hip_streams_callback function for ROCPROFILER_HIP_STREAM_CREATE");
|
||||
}
|
||||
else if(record.operation == ROCPROFILER_HIP_STREAM_DESTROY)
|
||||
{
|
||||
ROCPROFSYS_VERBOSE_F(
|
||||
2,
|
||||
"Entered hip_streams_callback function for ROCPROFILER_HIP_STREAM_DESTROY");
|
||||
}
|
||||
else if(record.operation == ROCPROFILER_HIP_STREAM_SET)
|
||||
{
|
||||
// Push the stream ID onto the stream stack before underlying HIP function is
|
||||
// called
|
||||
if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER)
|
||||
{
|
||||
ROCPROFSYS_VERBOSE_F(
|
||||
2, "Entered hip_streams_callback function for ROCPROFILER_HIP_STREAM_SET "
|
||||
"with ROCPROFILER_CALLBACK_PHASE_ENTER");
|
||||
stream_id_push(stream_id);
|
||||
}
|
||||
// Pop stream ID off of stream stack after underlying HIP function is completed
|
||||
else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT)
|
||||
{
|
||||
ROCPROFSYS_VERBOSE_F(
|
||||
2, "Entered hip_stream_callback function for ROCPROFILER_HIP_STREAM_SET "
|
||||
"with ROCPROFILER_CALLBACK_PHASE_EXIT");
|
||||
stream_id_pop();
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
ROCPROFSYS_FAIL_F("Unknown operation for hip_stream_callback!");
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
int
|
||||
tool_init(rocprofiler_client_finalize_t fini_func, void* user_data)
|
||||
{
|
||||
@@ -1128,6 +1296,30 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* user_data)
|
||||
constexpr auto buffer_size = 8192;
|
||||
constexpr auto watermark = 7936;
|
||||
|
||||
// Configure external correlation id request service for kernel dispatch
|
||||
// and memory copy.
|
||||
|
||||
auto external_corr_id_request_kinds =
|
||||
std::array<rocprofiler_external_correlation_id_request_kind_t, 2>{
|
||||
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH,
|
||||
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY
|
||||
};
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_configure_external_correlation_id_request_service(
|
||||
_data->primary_ctx, external_corr_id_request_kinds.data(),
|
||||
external_corr_id_request_kinds.size(),
|
||||
set_kernel_rename_and_stream_correlation_id, _data));
|
||||
|
||||
#if(ROCPROFILER_VERSION >= 700)
|
||||
if((_buffered_domain.count(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) > 0) ||
|
||||
(_buffered_domain.count(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) > 0))
|
||||
{
|
||||
ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
|
||||
_data->primary_ctx, ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, nullptr, 0,
|
||||
tool_hip_stream_callback, nullptr));
|
||||
}
|
||||
#endif
|
||||
|
||||
if(_buffered_domain.count(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) > 0)
|
||||
{
|
||||
ROCPROFILER_CALL(rocprofiler_create_buffer(
|
||||
@@ -1138,16 +1330,6 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* user_data)
|
||||
ROCPROFILER_CALL(rocprofiler_configure_buffer_tracing_service(
|
||||
_data->primary_ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0,
|
||||
_data->kernel_dispatch_buffer));
|
||||
|
||||
// auto external_corr_id_request_kinds =
|
||||
// std::array<rocprofiler_external_correlation_id_request_kind_t, 1>{
|
||||
// ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH
|
||||
// };
|
||||
|
||||
// ROCPROFILER_CALL(rocprofiler_configure_external_correlation_id_request_service(
|
||||
// _data->primary_ctx, external_corr_id_request_kinds.data(),
|
||||
// external_corr_id_request_kinds.size(), external_correlation_id_callback,
|
||||
// _data));
|
||||
}
|
||||
|
||||
if(_buffered_domain.count(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) > 0)
|
||||
@@ -1157,12 +1339,8 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* user_data)
|
||||
ROCPROFILER_BUFFER_POLICY_LOSSLESS, tool_tracing_buffered, tool_data,
|
||||
&_data->memory_copy_buffer));
|
||||
|
||||
auto _ops =
|
||||
rocprofiler_sdk::get_operations(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY);
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_configure_buffer_tracing_service(
|
||||
_data->primary_ctx, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY,
|
||||
(_ops.empty()) ? nullptr : _ops.data(), _ops.size(),
|
||||
_data->primary_ctx, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, nullptr, 0,
|
||||
_data->memory_copy_buffer));
|
||||
}
|
||||
|
||||
|
||||
@@ -48,5 +48,6 @@ include(${CMAKE_CURRENT_LIST_DIR}/rocprof-sys-python-tests.cmake)
|
||||
include(${CMAKE_CURRENT_LIST_DIR}/rocprof-sys-decode-tests.cmake)
|
||||
include(${CMAKE_CURRENT_LIST_DIR}/rocprof-sys-nic-perf.cmake)
|
||||
include(${CMAKE_CURRENT_LIST_DIR}/rocprof-sys-roctx-tests.cmake)
|
||||
include(${CMAKE_CURRENT_LIST_DIR}/rocprof-sys-rocm-hip-stream.cmake)
|
||||
|
||||
add_subdirectory(source)
|
||||
|
||||
@@ -0,0 +1,65 @@
|
||||
# MIT License
|
||||
#
|
||||
# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
#
|
||||
# Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
# of this software and associated documentation files (the "Software"), to deal
|
||||
# in the Software without restriction, including without limitation the rights
|
||||
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
# copies of the Software, and to permit persons to whom the Software is
|
||||
# furnished to do so, subject to the following conditions:
|
||||
#
|
||||
# The above copyright notice and this permission notice shall be included in
|
||||
# all copies or substantial portions of the Software.
|
||||
#
|
||||
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
# -------------------------------------------------------------------------------------- #
|
||||
#
|
||||
# ROCm tests
|
||||
#
|
||||
# -------------------------------------------------------------------------------------- #
|
||||
|
||||
find_package(ROCmVersion)
|
||||
|
||||
if(NOT ROCmVersion_FOUND)
|
||||
message(
|
||||
WARNING
|
||||
"ROCmVersion_FOUND not found, skipping tests in ${CMAKE_CURRENT_LIST_FILE}"
|
||||
)
|
||||
return()
|
||||
endif()
|
||||
|
||||
if(${ROCmVersion_FULL_VERSION} VERSION_GREATER_EQUAL "7.0")
|
||||
message(STATUS "Adding Group-By Tests")
|
||||
|
||||
rocprofiler_systems_add_test(
|
||||
SKIP_REWRITE SKIP_RUNTIME SKIP_BASELINE
|
||||
NAME transpose-group-by-queue
|
||||
TARGET transpose
|
||||
MPI ${TRANSPOSE_USE_MPI}
|
||||
GPU ON
|
||||
NUM_PROCS ${NUM_PROCS}
|
||||
ENVIRONMENT "${_base_environment};ROCPROFSYS_ROCM_GROUP_BY_QUEUE=YES"
|
||||
LABEL "group-by-queue"
|
||||
RUNTIME_TIMEOUT 480
|
||||
)
|
||||
|
||||
rocprofiler_systems_add_test(
|
||||
SKIP_REWRITE SKIP_RUNTIME SKIP_BASELINE
|
||||
NAME transpose-group-by-stream
|
||||
TARGET transpose
|
||||
MPI ${TRANSPOSE_USE_MPI}
|
||||
GPU ON
|
||||
NUM_PROCS ${NUM_PROCS}
|
||||
ENVIRONMENT "${_base_environment};ROCPROFSYS_ROCM_GROUP_BY_QUEUE=NO"
|
||||
LABEL "group-by-queue"
|
||||
RUNTIME_TIMEOUT 480
|
||||
)
|
||||
endif()
|
||||
@@ -532,7 +532,7 @@ function(ROCPROFILER_SYSTEMS_ADD_TEST)
|
||||
cmake_parse_arguments(
|
||||
TEST
|
||||
"SKIP_BASELINE;SKIP_SAMPLING;SKIP_REWRITE;SKIP_RUNTIME"
|
||||
"NAME;TARGET;MPI;GPU;NUM_PROCS;SAMPLING_TIMEOUT;REWRITE_TIMEOUT;RUNTIME_TIMEOUT"
|
||||
"NAME;TARGET;MPI;GPU;NUM_PROCS;SAMPLING_TIMEOUT;REWRITE_TIMEOUT;RUNTIME_TIMEOUT;WILL_FAIL;DISABLED"
|
||||
"${_KWARGS}"
|
||||
${ARGN}
|
||||
)
|
||||
@@ -584,6 +584,14 @@ function(ROCPROFILER_SYSTEMS_ADD_TEST)
|
||||
set(TEST_SAMPLING_TIMEOUT 120)
|
||||
endif()
|
||||
|
||||
if(NOT TEST_DISABLED)
|
||||
set(TEST_DISABLED OFF)
|
||||
endif()
|
||||
|
||||
if(NOT TEST_WILL_FAIL)
|
||||
set(TEST_WILL_FAIL OFF)
|
||||
endif()
|
||||
|
||||
if(NOT DEFINED TEST_ENVIRONMENT OR "${TEST_ENVIRONMENT}" STREQUAL "")
|
||||
set(TEST_ENVIRONMENT "${_test_environment}")
|
||||
endif()
|
||||
@@ -777,6 +785,8 @@ function(ROCPROFILER_SYSTEMS_ADD_TEST)
|
||||
PASS_REGULAR_EXPRESSION "${${_PASS_REGEX}}"
|
||||
FAIL_REGULAR_EXPRESSION "${${_FAIL_REGEX}}"
|
||||
SKIP_REGULAR_EXPRESSION "${${_SKIP_REGEX}}"
|
||||
WILL_FAIL ${TEST_WILL_FAIL}
|
||||
DISABLED ${TEST_DISABLED}
|
||||
${_props}
|
||||
)
|
||||
endif()
|
||||
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele