[rocprofv3] Unconditionally collect stream and kernel rename data in rocprofv3 for rocpd (#171)
* Remove config checks for stream and kernel rename data collection * Updated csv generation to check if kernel rename is on before calling get_kernel_name * Update metadata to use kernel_rename bool argument * Formatting + unconditionally store kernel name in rocpd * Readded kernel rename parameter after rebase * Fixed rebase conflicts * Updated comment in line with github comments * Added check in rocpd csv.cpp to output kernel name if region name is empty * Add test for kernel rename --------- Co-authored-by: Ian Trowbridge <Ian.Trowbridge@amd.com>
Этот коммит содержится в:
systems-assistant[bot]
коммит произвёл
GitHub
GitHub
родитель
156de36a92
Коммит
b60c0ceddd
@@ -288,11 +288,12 @@ generate_csv(const output_config&
|
||||
{
|
||||
for(auto record : data.get(ditr))
|
||||
{
|
||||
auto row_ss = std::stringstream{};
|
||||
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
|
||||
record.correlation_id.external.value);
|
||||
auto row_ss = std::stringstream{};
|
||||
const auto* kernel_info =
|
||||
tool_metadata.get_kernel_symbol(record.dispatch_info.kernel_id);
|
||||
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
|
||||
cfg.kernel_rename,
|
||||
record.correlation_id.external.value);
|
||||
auto lds_block_size_v =
|
||||
(kernel_info->group_segment_size + (lds_block_size - 1)) & ~(lds_block_size - 1);
|
||||
|
||||
@@ -637,7 +638,8 @@ generate_csv(const output_config& cfg,
|
||||
record.thread_id,
|
||||
magnitude(record.dispatch_data.dispatch_info.grid_size),
|
||||
record.dispatch_data.dispatch_info.kernel_id,
|
||||
tool_metadata.get_kernel_name(kernel_id, correlation_id.external.value),
|
||||
tool_metadata.get_kernel_name(
|
||||
kernel_id, cfg.kernel_rename, correlation_id.external.value),
|
||||
magnitude(record.dispatch_data.dispatch_info.workgroup_size),
|
||||
lds_block_size_v,
|
||||
record.dispatch_data.dispatch_info.private_segment_size,
|
||||
|
||||
@@ -695,8 +695,8 @@ write_otf2(const output_config& cfg,
|
||||
const auto* sym = _get_kernel_sym_data(info);
|
||||
CHECK(sym != nullptr);
|
||||
|
||||
auto name =
|
||||
tool_metadata.get_kernel_name(info.kernel_id, itr.correlation_id.external.value);
|
||||
auto name = tool_metadata.get_kernel_name(
|
||||
info.kernel_id, cfg.kernel_rename, itr.correlation_id.external.value);
|
||||
_hash_data.emplace(
|
||||
get_hash_id(name),
|
||||
region_info{std::string{name}, OTF2_REGION_ROLE_FUNCTION, OTF2_PARADIGM_HIP});
|
||||
|
||||
@@ -647,6 +647,11 @@ write_rocpd(
|
||||
itr.demangled_kernel_name,
|
||||
itr.truncated_kernel_name);
|
||||
|
||||
for(const auto& itr : tool_metadata.kernel_rename_map.get())
|
||||
{
|
||||
add_string_entry(_metadata, itr.first);
|
||||
}
|
||||
|
||||
for(const auto& itr : tool_metadata.get_code_objects())
|
||||
if(itr.uri != nullptr) add_string_entry(_metadata, itr.uri);
|
||||
|
||||
@@ -1020,10 +1025,12 @@ write_rocpd(
|
||||
}
|
||||
|
||||
dispatch_evt_ids.at(dispatch_id) = evt_id;
|
||||
|
||||
// Unconditionally collect kernel rename data if it is available. rocpd needs to be able
|
||||
// to use kernel rename option after data has already been collected, so the kernel
|
||||
// rename data needs to be stored in generated db.
|
||||
auto region_name =
|
||||
(corr_id.external.value > 0 && (enable_duplicate_check || kernel_id > 0))
|
||||
? tool_metadata.get_kernel_name(kernel_id, corr_id.external.value)
|
||||
? tool_metadata.get_kernel_name(kernel_id, true, corr_id.external.value)
|
||||
: std::string_view{};
|
||||
|
||||
auto agent_node_id = tool_metadata.get_agent(info.agent_id)->node_id;
|
||||
|
||||
@@ -62,7 +62,7 @@ get_stats(const stats_map_t& data_v)
|
||||
} // namespace
|
||||
|
||||
stats_entry_t
|
||||
generate_stats(const output_config& /*cfg*/,
|
||||
generate_stats(const output_config& cfg,
|
||||
const metadata& tool_metadata,
|
||||
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& data)
|
||||
{
|
||||
@@ -72,6 +72,7 @@ generate_stats(const output_config& /*cfg*/,
|
||||
for(auto record : data.get(ditr))
|
||||
{
|
||||
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
|
||||
cfg.kernel_rename,
|
||||
record.correlation_id.external.value);
|
||||
|
||||
kernel_stats[kernel_name] += (record.end_timestamp - record.start_timestamp);
|
||||
|
||||
@@ -696,19 +696,22 @@ metadata::get_marker_message(uint64_t corr_id) const
|
||||
}
|
||||
|
||||
std::string_view
|
||||
metadata::get_kernel_name(uint64_t kernel_id, uint64_t rename_id) const
|
||||
metadata::get_kernel_name(uint64_t kernel_id, bool rename_kernel, uint64_t rename_id) const
|
||||
{
|
||||
auto string_entry = kernel_rename_map.rlock(
|
||||
[](auto& _data, uint64_t _val) {
|
||||
for(const auto& itr : _data)
|
||||
if(itr.second == _val) return itr.first;
|
||||
return std::string_view{};
|
||||
},
|
||||
rename_id);
|
||||
if(!string_entry.empty())
|
||||
if(rename_kernel)
|
||||
{
|
||||
if(const auto* _name = common::get_string_entry(string_entry))
|
||||
return std::string_view{*_name};
|
||||
auto string_entry = kernel_rename_map.rlock(
|
||||
[](auto& _data, uint64_t _val) {
|
||||
for(const auto& itr : _data)
|
||||
if(itr.second == _val) return itr.first;
|
||||
return std::string_view{};
|
||||
},
|
||||
rename_id);
|
||||
if(!string_entry.empty())
|
||||
{
|
||||
if(const auto* _name = common::get_string_entry(string_entry))
|
||||
return std::string_view{*_name};
|
||||
}
|
||||
}
|
||||
|
||||
const auto* _kernel_data = get_kernel_symbol(kernel_id);
|
||||
|
||||
@@ -231,7 +231,9 @@ struct metadata
|
||||
const std::vector<std::string>& _command_line = {});
|
||||
|
||||
std::string_view get_marker_message(uint64_t corr_id) const;
|
||||
std::string_view get_kernel_name(uint64_t kernel_id, uint64_t rename_id) const;
|
||||
std::string_view get_kernel_name(uint64_t kernel_id,
|
||||
bool rename_kernel,
|
||||
uint64_t rename_id) const;
|
||||
std::string_view get_kind_name(rocprofiler_callback_tracing_kind_t kind) const;
|
||||
std::string_view get_kind_name(rocprofiler_buffer_tracing_kind_t kind) const;
|
||||
std::string_view get_operation_name(rocprofiler_callback_tracing_kind_t kind,
|
||||
|
||||
@@ -234,8 +234,8 @@ write_kernel_csv(
|
||||
CsvType::KERNEL_DISPATCH,
|
||||
kernel_dispatch_gen,
|
||||
[](CsvManager& cm, CsvType type, const rocpd::types::kernel_dispatch& kernel) {
|
||||
std::string kernel_identifier = cm.config.kernel_rename ? kernel.region : kernel.name;
|
||||
|
||||
std::string kernel_identifier =
|
||||
(cm.config.kernel_rename && !kernel.region.empty()) ? kernel.region : kernel.name;
|
||||
std::string agent_identifier = create_agent_index(cm.config.agent_index_value,
|
||||
kernel.agent_abs_index,
|
||||
kernel.agent_log_index,
|
||||
|
||||
@@ -488,11 +488,9 @@ set_kernel_rename_and_stream_correlation_id(rocprofiler_thread_id_t thr_id,
|
||||
// Check whether services are enabled
|
||||
const bool kernel_rename_service_enabled =
|
||||
kind == ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH &&
|
||||
tool::get_config().kernel_rename && thread_dispatch_rename != nullptr &&
|
||||
!thread_dispatch_rename->empty();
|
||||
thread_dispatch_rename != nullptr && !thread_dispatch_rename->empty();
|
||||
|
||||
const bool hip_stream_enabled =
|
||||
!tool::get_config().group_by_queue && rocprofiler::tool::stream::stream_stack_not_null();
|
||||
const bool hip_stream_enabled = rocprofiler::tool::stream::stream_stack_not_null();
|
||||
|
||||
if(!kernel_rename_service_enabled && !hip_stream_enabled) return 1;
|
||||
|
||||
@@ -566,7 +564,7 @@ kernel_rename_callback(rocprofiler_callback_tracing_record_t record,
|
||||
rocprofiler_user_data_t* user_data,
|
||||
void* data)
|
||||
{
|
||||
if(!tool::get_config().kernel_rename || thread_dispatch_rename == nullptr) return;
|
||||
if(thread_dispatch_rename == nullptr) return;
|
||||
|
||||
if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_RANGE_API)
|
||||
{
|
||||
@@ -617,8 +615,7 @@ hip_stream_display_callback(rocprofiler_callback_tracing_record_t record,
|
||||
rocprofiler_user_data_t* user_data,
|
||||
void* data)
|
||||
{
|
||||
if(tool::get_config().group_by_queue || record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_STREAM)
|
||||
return;
|
||||
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);
|
||||
@@ -2078,67 +2075,60 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
|
||||
start_context(counter_collection_ctx, "counter collection");
|
||||
}
|
||||
|
||||
if(tool::get_config().kernel_rename)
|
||||
{
|
||||
auto rename_ctx = rocprofiler_context_id_t{0};
|
||||
auto marker_core_api_kinds = std::array<rocprofiler_tracing_operation_t, 2>{
|
||||
ROCPROFILER_MARKER_CORE_RANGE_API_ID_roctxMarkA,
|
||||
ROCPROFILER_MARKER_CORE_RANGE_API_ID_roctxThreadRangeA,
|
||||
};
|
||||
auto rename_ctx = rocprofiler_context_id_t{0};
|
||||
auto marker_core_api_kinds = std::array<rocprofiler_tracing_operation_t, 2>{
|
||||
ROCPROFILER_MARKER_CORE_RANGE_API_ID_roctxMarkA,
|
||||
ROCPROFILER_MARKER_CORE_RANGE_API_ID_roctxThreadRangeA,
|
||||
};
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&rename_ctx), "failed to create context");
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&rename_ctx), "failed to create context");
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
|
||||
rename_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_RANGE_API,
|
||||
marker_core_api_kinds.data(),
|
||||
marker_core_api_kinds.size(),
|
||||
callbacks.kernel_rename,
|
||||
nullptr),
|
||||
"callback tracing service failed to configure");
|
||||
ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
|
||||
rename_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_RANGE_API,
|
||||
marker_core_api_kinds.data(),
|
||||
marker_core_api_kinds.size(),
|
||||
callbacks.kernel_rename,
|
||||
nullptr),
|
||||
"callback tracing service failed to configure");
|
||||
|
||||
start_context(rename_ctx, "kernel rename");
|
||||
}
|
||||
start_context(rename_ctx, "kernel rename");
|
||||
|
||||
if(!tool::get_config().group_by_queue)
|
||||
{
|
||||
// Track stream ID information via callback service
|
||||
auto hip_stream_display_ctx = rocprofiler_context_id_t{0};
|
||||
// Track stream ID information via callback service
|
||||
auto hip_stream_display_ctx = rocprofiler_context_id_t{0};
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&hip_stream_display_ctx),
|
||||
"failed to create hip stream context");
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&hip_stream_display_ctx),
|
||||
"failed to create hip stream context");
|
||||
|
||||
ROCPROFILER_CALL(
|
||||
rocprofiler_configure_callback_tracing_service(hip_stream_display_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_HIP_STREAM,
|
||||
nullptr,
|
||||
0,
|
||||
callbacks.hip_stream,
|
||||
nullptr),
|
||||
"hip stream tracing configure failed");
|
||||
ROCPROFILER_CALL(
|
||||
rocprofiler_configure_callback_tracing_service(hip_stream_display_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_HIP_STREAM,
|
||||
nullptr,
|
||||
0,
|
||||
callbacks.hip_stream,
|
||||
nullptr),
|
||||
"hip stream tracing configure failed");
|
||||
|
||||
start_context(hip_stream_display_ctx, "hip stream");
|
||||
start_context(hip_stream_display_ctx, "hip stream");
|
||||
|
||||
// Track if HIP runtime has been initialized via runtime_intialization service
|
||||
auto runtime_initialization_ctx = rocprofiler_context_id_t{0};
|
||||
// Track if HIP runtime has been initialized via runtime_intialization service
|
||||
auto runtime_initialization_ctx = rocprofiler_context_id_t{0};
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&runtime_initialization_ctx),
|
||||
"failed to create runtime initialization context");
|
||||
ROCPROFILER_CALL(rocprofiler_create_context(&runtime_initialization_ctx),
|
||||
"failed to create runtime initialization context");
|
||||
|
||||
ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
|
||||
runtime_initialization_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION,
|
||||
nullptr,
|
||||
0,
|
||||
runtime_initialization_callback,
|
||||
nullptr),
|
||||
"runtime initialization tracing configure failed");
|
||||
ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
|
||||
runtime_initialization_ctx,
|
||||
ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION,
|
||||
nullptr,
|
||||
0,
|
||||
runtime_initialization_callback,
|
||||
nullptr),
|
||||
"runtime initialization tracing configure failed");
|
||||
|
||||
start_context(runtime_initialization_ctx, "runtime initialization");
|
||||
}
|
||||
start_context(runtime_initialization_ctx, "runtime initialization");
|
||||
|
||||
if((tool::get_config().kernel_rename || !tool::get_config().group_by_queue) &&
|
||||
tool::get_config().benchmark_mode != tool::config::benchmark::execution_profile)
|
||||
if(tool::get_config().benchmark_mode != tool::config::benchmark::execution_profile)
|
||||
{
|
||||
auto external_corr_id_request_kinds =
|
||||
std::array<rocprofiler_external_correlation_id_request_kind_t, 4>{
|
||||
|
||||
@@ -1498,6 +1498,7 @@ get_ids(int kind, std::index_sequence<Kind, Kinds...>)
|
||||
return get_ids(kind, std::index_sequence<Kinds...>{});
|
||||
|
||||
ROCP_CI_LOG(WARNING) << fmt::format("KFD events get_ids: Unknown Kind {}", kind);
|
||||
return {};
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
Ссылка в новой задаче
Block a user