From 633a80ec1b32ce47ea388bd454c07d6db5dfcea6 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Fri, 15 Dec 2023 12:44:50 -0600 Subject: [PATCH] Fix rocprof and rocprofiler-sdk-tool (#304) * Fix rocprof and rocprofiler-sdk-tool - removed redundant options for rocprofv3 - fixed extension of CSV files - fixed memory error (segfault) when tracing kernels - fixed constructor of output_file - using string after move - output_file writes csv header - implemented tool_fini - flush buffer and stop context * Minor updates to rocprofv3 * Update source/bin/CMakeLists.txt - fix rocprof configure output location [ROCm/rocprofiler-sdk commit: 7d117870d89c973ef101167951648fdd39165fff] --- .../rocprofiler-sdk/source/bin/CMakeLists.txt | 2 +- projects/rocprofiler-sdk/source/bin/rocprofv3 | 13 +-- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 103 ++++++++++++------ 3 files changed, 73 insertions(+), 45 deletions(-) diff --git a/projects/rocprofiler-sdk/source/bin/CMakeLists.txt b/projects/rocprofiler-sdk/source/bin/CMakeLists.txt index 5c4928e54b..266ea9e5ac 100644 --- a/projects/rocprofiler-sdk/source/bin/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/bin/CMakeLists.txt @@ -4,7 +4,7 @@ rocprofiler_activate_clang_tidy() -configure_file(rocprofv3 ${PROJECT_BINARY_DIR} COPYONLY) +configure_file(rocprofv3 ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/rocprofv3 COPYONLY) install( FILES rocprofv3 DESTINATION ${CMAKE_INSTALL_BINDIR} diff --git a/projects/rocprofiler-sdk/source/bin/rocprofv3 b/projects/rocprofiler-sdk/source/bin/rocprofv3 index ea9a50b860..3cb31d17f8 100755 --- a/projects/rocprofiler-sdk/source/bin/rocprofv3 +++ b/projects/rocprofiler-sdk/source/bin/rocprofv3 @@ -14,10 +14,9 @@ GREY='\033[0;90m' RESET='\033[0m' usage() { - echo -e "${RESET}ROCProfilerV2 Run Script Usage:" + echo -e "${RESET}ROCProfilerV3 Run Script Usage:" echo -e "${GREEN}-h | --help ${RESET} For showing this message" - echo -e "${GREEN}--hsa-api ${RESET} For Collecting HSA API Traces" - echo -e "${GREEN}--hsa-activity | --hsa-trace ${RESET} For Collecting HSA API Activities Traces" + echo -e "${GREEN}--hsa-trace ${RESET} For Collecting HSA API Traces" echo -e "${GREEN}--kernel-trace ${RESET} For Collecting Kernel Dispatch Traces" echo -e "${GREEN}-o | --output-file ${RESET} For the output file name" echo -e "\t#${GREY} usage e.g:(with current dir): rocprofv3 --hip-trace -o " @@ -56,13 +55,9 @@ while [ 1 ]; do fi shift shift - elif [ "$1" == "--hsa-api" ]; then + elif [ "$1" == "--hsa-trace" ]; then export ROCPROFILER_HSA_API_TRACE=1 shift - elif [[ "$1" == "--hsa-activity" || "$1" == "--hsa-trace" ]]; then - export ROCPROFILER_HSA_API_TRACE=1 - export ROCPROFILER_HSA_ACTIVITY_TRACE=1 - shift elif [ "$1" == "--kernel-trace" ]; then export ROCPROFILER_KERNEL_TRACE=1 shift @@ -80,4 +75,4 @@ done export ROCPROFILER_OUTPUT_PATH -ROCP_TOOL_LIBRARIES=${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so $* +ROCP_TOOL_LIBRARIES=${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so "${@}" diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 66aad11de0..9d6fb2ec76 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -21,6 +21,7 @@ // SOFTWARE. #include "helper.hpp" +#include "rocprofiler-sdk/context.h" #include "trace_buffer.hpp" #include "lib/common/environment.hpp" @@ -46,13 +47,14 @@ TRACE_BUFFER_INSTANTIATE(); namespace { auto tool_buffer = rocprofiler_buffer_id_t{}; +auto context_id = rocprofiler_context_id_t{}; auto output_path = fs::path{common::get_env("ROCPROFILER_OUTPUT_PATH", fs::current_path().string())}; auto output_file_name = common::get_env("ROCPROFILER_OUTPUT_FILE_NAME", std::to_string(getpid()) + "-"); std::pair -get_output_stream(const std::string& fname, const std::string& ext = ".ext") +get_output_stream(const std::string& fname, const std::string& ext = ".csv") { if(output_path.string().empty()) return {&std::clog, [](auto*&) {}}; @@ -65,8 +67,7 @@ get_output_stream(const std::string& fname, const std::string& ext = ".ext") auto output_file = output_path / (output_file_name + fname + ext); auto* _ofs = new std::ofstream{output_file}; if(!_ofs && !*_ofs) - throw std::runtime_error{ - fmt::format("Failed to open {} for output", (output_path / output_file_name).string())}; + throw std::runtime_error{fmt::format("Failed to open {} for output", output_file.string())}; std::cout << "Results File: " << output_file << std::endl; return {_ofs, [](std::ostream*& v) { if(v) dynamic_cast(v)->close(); @@ -88,10 +89,18 @@ as_hex(Tp _v, size_t _width = 16) struct output_file { - output_file(std::string name) - : m_name(std::move(name)) + output_file(std::string name, std::vector&& header) + : m_name{std::move(name)} { - std::tie(m_stream, m_dtor) = get_output_stream(name); + std::tie(m_stream, m_dtor) = get_output_stream(m_name); + auto ss = std::stringstream{}; + for(auto&& itr : header) + { + ss << "," << itr; + } + + // write the csv header + if(!ss.str().empty()) *m_stream << ss.str().substr(1) << '\n'; } ~output_file() { m_dtor(m_stream); } @@ -109,6 +118,8 @@ struct output_file std::ostream& operator<<(std::ostream& (*func)(std::ostream&) ) { return (*m_stream) << func; } + operator bool() const { return m_stream != nullptr; } + private: using stream_dtor_t = void (*)(std::ostream*&); @@ -120,14 +131,33 @@ private: auto& get_hsa_api_file() { - static auto _v = output_file{"hsa_api_trace"}; + static auto _v = + output_file{"hsa_api_trace", {"KERNEL_NAME", "BEGIN_TS", "END_TS", "CORRELATION_ID"}}; return _v; } auto& get_kernel_trace_file() { - static auto _v = output_file{"kernel_trace"}; + static auto _v = output_file{"kernel_trace", + {"AGENT_ID", + "QUEUE_ID", + "KERNEL_ID", + "KERNEL_NAME", + "CONTEXT_ID", + "BUFFER_ID", + "CORRELATION_ID", + "KIND", + "START_TS", + "END_TS", + "PRIVATE_SEGMENT_SIZE", + "GROUP_SEGMENT_SIZE", + "WORKGROUP_SIZE_X", + "WORKGROUP_SIZE_Y", + "WORKGROUP_SIZE_Z", + "GRID_SIZE_X", + "GRID_SIZE_Y", + "GRID_SIZE_Z"}}; return _v; } @@ -166,21 +196,15 @@ TraceBuffer hsa_api_buffer("HSA API", rocprofiler_tool_callback_name_info_t name_info; -void -tool_fini(void* tool_data) -{ - (void) (tool_data); -} - void TracerFlushRecord(void* data, rocprofiler_callback_tracing_kind_t kind) { if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) { auto* entry = reinterpret_cast(data); - get_hsa_api_file() << entry->api_name << " " << entry->begin_timestamp << ":" - << entry->end_timestamp << " " << entry->record.correlation_id.internal - << '\n'; + get_hsa_api_file() << "\"" << entry->api_name << "\"" + << "," << entry->begin_timestamp << ":" << entry->end_timestamp << " " + << entry->record.correlation_id.internal << '\n'; } } void @@ -203,7 +227,6 @@ rocprofiler_tracing_callback(rocprofiler_callback_tracing_record_t record, { user_data->value = timestamp; } - else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) { const auto* info_name_str = name_info.operation_names[record.kind][record.operation]; @@ -212,7 +235,6 @@ rocprofiler_tracing_callback(rocprofiler_callback_tracing_record_t record, entry.valid.store(TRACE_ENTRY_COMPLETE, std::memory_order_release); } } - else if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) { // To be implemented @@ -284,6 +306,7 @@ kernel_tracing_callback(rocprofiler_context_id_t context, throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the " "array of headers. this should never happen"}; + auto kernel_trace_ss = std::stringstream{}; for(size_t i = 0; i < num_headers; ++i) { auto* header = headers[i]; @@ -309,21 +332,25 @@ kernel_tracing_callback(rocprofiler_context_id_t context, kernel_name = kernel_data.at(record->kernel_id); } - get_kernel_trace_file() - << "agent_id=" << record->agent_id.handle - << ", queue_id=" << record->queue_id.handle << ", kernel_id=" << record->kernel_id - << ", kernel=" << kernel_name << ", context=" << context.handle - << ", buffer_id=" << buffer_id.handle << ", cid=" << record->correlation_id.internal - << ", extern_cid=" << record->correlation_id.external.value - << ", kind=" << record->kind << ", start=" << record->start_timestamp - << ", stop=" << record->end_timestamp - << ", private_segment_size=" << record->private_segment_size - << ", group_segment_size=" << record->group_segment_size << ", workgroup_size=(" - << record->workgroup_size.x << "," << record->workgroup_size.y << "," - << record->workgroup_size.z << "), grid_size=(" << record->grid_size.x << "," - << record->grid_size.y << "," << record->grid_size.z << ")" << '\n'; + kernel_trace_ss << record->agent_id.handle << "," << record->queue_id.handle << "," + << record->kernel_id << ",\"" << kernel_name << "\"," << context.handle + << "," << buffer_id.handle << "," << record->correlation_id.internal + << "," << record->kind << "," << record->start_timestamp << "," + << record->end_timestamp << "," << record->private_segment_size << "," + << record->group_segment_size << "," << record->workgroup_size.x << "," + << record->workgroup_size.y << "," << record->workgroup_size.z << "," + << record->grid_size.x << "," << record->grid_size.y << "," + << record->grid_size.z << '\n'; } } + + static auto _sync = std::mutex{}; + auto _lk = std::unique_lock{_sync}; + if(get_kernel_trace_file()) + get_kernel_trace_file() << kernel_trace_ss.str(); + else + std::cerr << "kernel trace file already closed: " << kernel_trace_ss.str(); + (void) (user_data); } @@ -389,9 +416,6 @@ get_callback_id_names() int tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data) { - // Add the rocporfiler_call macro - rocprofiler_context_id_t context_id; - name_info = get_callback_id_names(); ROCPROFILER_CALL(rocprofiler_create_context(&context_id), "create context failed"); @@ -415,6 +439,7 @@ tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data) tool_data, &tool_buffer), "buffer creation"); + ROCPROFILER_CALL( rocprofiler_configure_buffer_tracing_service( context_id, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, tool_buffer), @@ -439,6 +464,14 @@ tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data) return 0; } +void +tool_fini(void* tool_data) +{ + rocprofiler_flush_buffer(tool_buffer); + rocprofiler_stop_context(context_id); + (void) (tool_data); +} + extern "C" rocprofiler_tool_configure_result_t* rocprofiler_configure(uint32_t /*version*/, const char* /*runtime_version*/,