diff --git a/projects/rocprofiler-systems/CHANGELOG.md b/projects/rocprofiler-systems/CHANGELOG.md index d849c5628e..263c1572c8 100644 --- a/projects/rocprofiler-systems/CHANGELOG.md +++ b/projects/rocprofiler-systems/CHANGELOG.md @@ -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 diff --git a/projects/rocprofiler-systems/docs/how-to/configuring-runtime-options.rst b/projects/rocprofiler-systems/docs/how-to/configuring-runtime-options.rst index 59c1b9099e..9e540243f0 100644 --- a/projects/rocprofiler-systems/docs/how-to/configuring-runtime-options.rst +++ b/projects/rocprofiler-systems/docs/how-to/configuring-runtime-options.rst @@ -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 --------------------- diff --git a/projects/rocprofiler-systems/docs/reference/development-guide.rst b/projects/rocprofiler-systems/docs/reference/development-guide.rst index 6003021484..5bc7385641 100644 --- a/projects/rocprofiler-systems/docs/reference/development-guide.rst +++ b/projects/rocprofiler-systems/docs/reference/development-guide.rst @@ -32,7 +32,7 @@ rocprof-sys-sample: `source/bin/rocprof-sys-sample `` and a modified environment -rocprof-sys-casual: `source/bin/rocprof-sys-causal `_ +rocprof-sys-causal: `source/bin/rocprof-sys-causal `_ --------------------------------------------------------------------------------------------------------------------------------------------------- When there is exactly one causal profiling configuration variant (which enables debugging), diff --git a/projects/rocprofiler-systems/source/lib/core/categories.hpp b/projects/rocprofiler-systems/source/lib/core/categories.hpp index bb9db10b03..57326e2b66 100644 --- a/projects/rocprofiler-systems/source/lib/core/categories.hpp +++ b/projects/rocprofiler-systems/source/lib/core/categories.hpp @@ -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; 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), \ diff --git a/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.cpp b/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.cpp index 917efaf0fc..c889fc2526 100644 --- a/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.cpp +++ b/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.cpp @@ -372,6 +372,22 @@ config_settings(const std::shared_ptr& _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 @@ -550,6 +566,16 @@ get_rocm_events() " ,;\t\n"); } +bool +get_group_by_queue(void) +{ + std::optional _group_by_queue = + config::get_setting_value("ROCPROFSYS_ROCM_GROUP_BY_QUEUE"); + bool _ret = _group_by_queue.value_or(true); + + return _ret; +} + std::vector get_operations(rocprofiler_callback_tracing_kind_t kindv) { diff --git a/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.hpp b/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.hpp index b70c1af633..f4d21a37d1 100644 --- a/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.hpp +++ b/projects/rocprofiler-systems/source/lib/core/rocprofiler-sdk.hpp @@ -70,6 +70,9 @@ get_operations(rocprofiler_buffer_tracing_kind_t kindv); std::vector get_rocm_events(); +bool +get_group_by_queue(); + std::unordered_set get_backtrace_operations(rocprofiler_callback_tracing_kind_t kindv); diff --git a/projects/rocprofiler-systems/source/lib/rocprof-sys-user/rocprofiler-systems/categories.h b/projects/rocprofiler-systems/source/lib/rocprof-sys-user/rocprofiler-systems/categories.h index 2084113032..27fb063ff3 100644 --- a/projects/rocprofiler-systems/source/lib/rocprof-sys-user/rocprofiler-systems/categories.h +++ b/projects/rocprofiler-systems/source/lib/rocprof-sys-user/rocprofiler-systems/categories.h @@ -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, diff --git a/projects/rocprofiler-systems/source/lib/rocprof-sys/library/rocprofiler-sdk.cpp b/projects/rocprofiler-systems/source/lib/rocprof-sys/library/rocprofiler-sdk.cpp index 26af5b6261..324b6941b7 100644 --- a/projects/rocprofiler-systems/source/lib/rocprof-sys/library/rocprofiler-sdk.cpp +++ b/projects/rocprofiler-systems/source/lib/rocprof-sys/library/rocprofiler-sdk.cpp @@ -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 _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 +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( + _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 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{ _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(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_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_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)); } diff --git a/projects/rocprofiler-systems/tests/CMakeLists.txt b/projects/rocprofiler-systems/tests/CMakeLists.txt index 6e3241fb86..b2e46cee4c 100644 --- a/projects/rocprofiler-systems/tests/CMakeLists.txt +++ b/projects/rocprofiler-systems/tests/CMakeLists.txt @@ -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) diff --git a/projects/rocprofiler-systems/tests/rocprof-sys-rocm-hip-stream.cmake b/projects/rocprofiler-systems/tests/rocprof-sys-rocm-hip-stream.cmake new file mode 100644 index 0000000000..84e3a855c7 --- /dev/null +++ b/projects/rocprofiler-systems/tests/rocprof-sys-rocm-hip-stream.cmake @@ -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() diff --git a/projects/rocprofiler-systems/tests/rocprof-sys-testing.cmake b/projects/rocprofiler-systems/tests/rocprof-sys-testing.cmake index 98886fe90e..6e16c96717 100644 --- a/projects/rocprofiler-systems/tests/rocprof-sys-testing.cmake +++ b/projects/rocprofiler-systems/tests/rocprof-sys-testing.cmake @@ -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()