From bef14ad1b24332b6d8a8f4e4a8e94e99cc234962 Mon Sep 17 00:00:00 2001 From: SrirakshaNag <104580803+SrirakshaNag@users.noreply.github.com> Date: Tue, 9 Apr 2024 05:25:28 -0500 Subject: [PATCH] rocprofiler-sdk-tool library intermediate binary output (#734) * Support for binary temporary files * clang formatting * formating ring buffer.hpp * Update source/lib/common/container/ring_buffer.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * fixing bugs * fix loop range * Fix for v3 test failures * bug fix * fix bug * fix memory leaks * destructing agent_info * Update CMakeLists.txt * clang-tidy fixes * Fix data race on destructor of rocprofiler_agent_t map in rocprofiler-sdk-tool library * Create lib/rocproifler-sdk-tool/tmp_file.* - move tmp_file class into separate header/implementation * Agent Info CSV in rocprofiler-sdk-tool - update tests to use agent_info.csv instead of rocminfo * Update lib/rocprofiler-sdk-tool/tool.cpp - use logical_node_id instead of node_id * Adding stats file * Adding tests for stats * Update scratch memory support - convert scratch memory support to use binary output * Tool Update: scratch memory stats + extended statistics - replace generate_*_csv with generate_csv overloads - added generate_csv for scratch memory - enable stats for scratch memory - replace ROCPROF_*_STATS env variables with ROCPROF_STATS env variable * rocprofv3 update - simple --stats option - add scratch memory trace to --sys-trace * Update tests/rocprofv3/tracing-hip-in-libraries - extend validate.py to test stats data - fix conftest.py for memory_copy_stats_data * Code coverage fixes - invoke __gcov_dump to ensure that code coverage is flushed after finalization --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Ammar ELWazir Co-authored-by: Jonathan R. Madsen --- cmake/rocprofiler_build_settings.cmake | 2 +- source/bin/rocprofv3 | 7 +- source/lib/common/container/ring_buffer.hpp | 12 +- .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 16 +- source/lib/rocprofiler-sdk-tool/config.cpp | 6 +- source/lib/rocprofiler-sdk-tool/config.hpp | 15 +- source/lib/rocprofiler-sdk-tool/csv.hpp | 23 +- .../lib/rocprofiler-sdk-tool/generateCSV.cpp | 446 +++++++++ .../lib/rocprofiler-sdk-tool/generateCSV.hpp | 57 ++ source/lib/rocprofiler-sdk-tool/helper.cpp | 1 + source/lib/rocprofiler-sdk-tool/helper.hpp | 115 ++- .../lib/rocprofiler-sdk-tool/statistics.hpp | 231 +++++ source/lib/rocprofiler-sdk-tool/tmp_file.cpp | 131 +++ source/lib/rocprofiler-sdk-tool/tmp_file.hpp | 51 + source/lib/rocprofiler-sdk-tool/tool.cpp | 911 +++++++++++++----- source/lib/rocprofiler-sdk/hip/hip.cpp | 2 +- .../counter-collection/input2/validate.py | 13 +- .../tracing-hip-in-libraries/CMakeLists.txt | 19 +- .../tracing-hip-in-libraries/conftest.py | 89 ++ .../tracing-hip-in-libraries/validate.py | 93 +- tests/rocprofv3/tracing-plus-cc/validate.py | 9 +- tests/rocprofv3/tracing/CMakeLists.txt | 14 +- tests/rocprofv3/tracing/conftest.py | 17 + tests/rocprofv3/tracing/validate.py | 81 +- 24 files changed, 2038 insertions(+), 323 deletions(-) create mode 100644 source/lib/rocprofiler-sdk-tool/generateCSV.cpp create mode 100644 source/lib/rocprofiler-sdk-tool/generateCSV.hpp create mode 100644 source/lib/rocprofiler-sdk-tool/statistics.hpp create mode 100644 source/lib/rocprofiler-sdk-tool/tmp_file.cpp create mode 100644 source/lib/rocprofiler-sdk-tool/tmp_file.hpp diff --git a/cmake/rocprofiler_build_settings.cmake b/cmake/rocprofiler_build_settings.cmake index 585ade4f4f..d10644e041 100644 --- a/cmake/rocprofiler_build_settings.cmake +++ b/cmake/rocprofiler_build_settings.cmake @@ -76,7 +76,7 @@ target_compile_options( INTERFACE $<$:$<$:-rdynamic>> $<$:$<$:-rdynamic>>) -if(NOT APPLE) +if(NOT APPLE AND NOT ROCPROFILER_ENABLE_CLANG_TIDY) target_link_options(rocprofiler-debug-flags INTERFACE $<$:-rdynamic>) endif() diff --git a/source/bin/rocprofv3 b/source/bin/rocprofv3 index a4ed407108..a471848ad7 100755 --- a/source/bin/rocprofv3 +++ b/source/bin/rocprofv3 @@ -42,7 +42,8 @@ usage() { echo -e "${GREEN}--hsa-image-trace ${RESET} For Collecting HSA API Traces (Image-extenson API)" echo -e "${GREEN}--hsa-finalizer-trace ${RESET} For Collecting HSA API Traces (Finalizer-extension API)" echo -e "" - echo -e "${GREEN}--sys-trace ${RESET} For Collecting HIP,HSA, Memory Copy, (marker)ROCTx and Kernel dispatch traces\n" + echo -e "${GREEN}--sys-trace ${RESET} For Collecting HIP, HSA, Marker (ROCTx), Memory copy, Scratch memory, and Kernel dispatch traces\n" + echo -e "${GREEN}--stats ${RESET} For Collecting statistics of enabled tracing types\n" echo -e "" echo -e "${GREEN}-o | --output-file ${RESET} For the output file name" echo -e "\t#${GREY} usage (with current dir): rocprofv3 --hsa-trace -o " @@ -177,6 +178,10 @@ while true; do export ROCPROF_MARKER_API_TRACE=1 export ROCPROF_KERNEL_TRACE=1 export ROCPROF_MEMORY_COPY_TRACE=1 + export ROCPROF_SCRATCH_MEMORY_TRACE=1 + shift + elif [ "$1" == "--stats" ]; then + export ROCPROF_STATS=1 shift elif [ "$1" == "--" ]; then shift diff --git a/source/lib/common/container/ring_buffer.hpp b/source/lib/common/container/ring_buffer.hpp index 14e5b64f4c..8959aac4f0 100644 --- a/source/lib/common/container/ring_buffer.hpp +++ b/source/lib/common/container/ring_buffer.hpp @@ -92,7 +92,7 @@ struct ring_buffer /// memory is uninitialized. Typically used by allocators. If Tp is a class type, /// be sure to use a placement new instead of a memcpy. template - Tp* request(); + Tp* request(bool wrap = true); /// Read class-type data from buffer (uses placement new). template @@ -210,11 +210,11 @@ ring_buffer::write(Tp* in, std::enable_if_t::value, int>) // template Tp* -ring_buffer::request() +ring_buffer::request(bool wrap) { if(m_ptr == nullptr) return nullptr; - return request(sizeof(Tp)); + return reinterpret_cast(request(sizeof(Tp), wrap)); } // template @@ -266,7 +266,7 @@ ring_buffer::retrieve() const { if(m_ptr == nullptr) return nullptr; - return retrieve(sizeof(Tp)); + return reinterpret_cast(retrieve(sizeof(Tp))); } // } // namespace base @@ -320,7 +320,7 @@ struct ring_buffer : private base::ring_buffer Tp* read(Tp* _dest) const { return base_type::read(_dest).second; } /// Get an uninitialized address at tail of buffer. - Tp* request() { return base_type::request(); } + Tp* request(bool wrap = true) { return base_type::request(wrap); } /// Read data from head of buffer. Tp* retrieve() { return base_type::retrieve(); } @@ -337,6 +337,8 @@ struct ring_buffer : private base::ring_buffer /// Returns if the buffer is full. bool is_full() const { return (base_type::free() < sizeof(Tp)); } + bool clear() { return base_type::clear(); } + template auto emplace(Args&&... args) { diff --git a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index 2b7d5b2229..315a4aa2da 100644 --- a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -2,10 +2,15 @@ # Tool library used by rocprofiler # +rocprofiler_activate_clang_tidy() + +set(TOOL_HEADERS config.hpp csv.hpp generateCSV.hpp helper.hpp output_file.hpp + statistics.hpp tmp_file.hpp) +set(TOOL_SOURCES config.cpp generateCSV.cpp helper.cpp output_file.cpp tmp_file.cpp + tool.cpp) + add_library(rocprofiler-sdk-tool SHARED) -target_sources( - rocprofiler-sdk-tool PRIVATE config.hpp config.cpp csv.hpp helper.hpp helper.cpp - output_file.hpp output_file.cpp tool.cpp) +target_sources(rocprofiler-sdk-tool PRIVATE ${TOOL_SOURCES} ${TOOL_HEADERS}) add_subdirectory(plugins) @@ -15,6 +20,11 @@ target_link_libraries( rocprofiler::rocprofiler-build-flags rocprofiler::rocprofiler-memcheck rocprofiler::rocprofiler-common-library) +if(ROCPROFILER_BUILD_CODECOV) + target_compile_definitions(rocprofiler-sdk-tool PRIVATE ROCPROFILER_CODECOV=1) + target_link_libraries(rocprofiler-sdk-tool PRIVATE gcov) +endif() + set_target_properties( rocprofiler-sdk-tool PROPERTIES LIBRARY_OUTPUT_DIRECTORY diff --git a/source/lib/rocprofiler-sdk-tool/config.cpp b/source/lib/rocprofiler-sdk-tool/config.cpp index 2d95319a83..4a61211049 100644 --- a/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/source/lib/rocprofiler-sdk-tool/config.cpp @@ -48,13 +48,13 @@ namespace tool { namespace { -auto launch_time = std::make_unique(std::time_t{std::time(nullptr)}); +auto launch_time = new std::time_t{std::time(nullptr)}; std::string get_local_datetime(const char* dt_format, std::time_t* dt_curr) { char mbstr[512]; - if(!dt_curr) dt_curr = launch_time.get(); + if(!dt_curr) dt_curr = launch_time; if(std::strftime(mbstr, sizeof(mbstr), dt_format, std::localtime(dt_curr)) != 0) return std::string{mbstr}; @@ -267,7 +267,7 @@ output_keys(std::string _tag) } } - auto* _launch_time = launch_time.get(); + auto* _launch_time = launch_time; auto _time_format = get_env("ROCP_TIME_FORMAT", "%F_%H.%M"); auto _mpi_size = get_mpi_size(); diff --git a/source/lib/rocprofiler-sdk-tool/config.hpp b/source/lib/rocprofiler-sdk-tool/config.hpp index 6e1d446e2d..da799d1f22 100644 --- a/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/source/lib/rocprofiler-sdk-tool/config.hpp @@ -72,10 +72,13 @@ struct config bool hip_compiler_api_trace = get_env("ROCPROF_HIP_COMPILER_API_TRACE", false); bool list_metrics = get_env("ROCPROF_LIST_METRICS", false); bool list_metrics_output_file = get_env("ROCPROF_OUTPUT_LIST_METRICS_FILE", false); + bool stats = get_env("ROCPROF_STATS", false); int mpi_size = get_mpi_size(); int mpi_rank = get_mpi_rank(); - std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); - std::string output_file = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid())); + std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); + std::string output_file = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid())); + std::string output_format = get_env("ROCPROF_OUTPUT_FORMAT", "CSV"); + std::string tmp_directory = get_env("ROCPROF_TMPDIR", fs::current_path().string()); std::vector kernel_names = {}; std::set counters = {}; }; @@ -86,14 +89,14 @@ get_config() { if constexpr(ContextT == config_context::global) { - static auto _v = config{}; - return _v; + static auto* _v = new config{}; + return *_v; } else { // context specific config copied from global config - static auto _v = get_config(); - return _v; + static auto* _v = new config{get_config()}; + return *_v; } } diff --git a/source/lib/rocprofiler-sdk-tool/csv.hpp b/source/lib/rocprofiler-sdk-tool/csv.hpp index f9da7ff12f..986d7e9b75 100644 --- a/source/lib/rocprofiler-sdk-tool/csv.hpp +++ b/source/lib/rocprofiler-sdk-tool/csv.hpp @@ -26,8 +26,11 @@ #include #include +#include +#include #include #include +#include namespace rocprofiler { @@ -40,10 +43,22 @@ std::ostream& write_csv_entry(std::ostream& ofs, TupleT&& _data, std::index_sequence) { auto _write = [&ofs](size_t idx, auto&& _val) { + using value_type = common::mpl::unqualified_type_t; if(idx > 0) ofs << ","; - if constexpr(rocprofiler::common::mpl::is_string_type::value) ofs << "\""; - ofs << _val; - if constexpr(rocprofiler::common::mpl::is_string_type::value) ofs << "\""; + if constexpr(std::is_floating_point::value) + { + constexpr value_type one = 1; + if(_val >= one) + ofs << std::setprecision(6) << std::fixed << _val; + else + ofs << std::setprecision(6) << std::scientific << _val; + } + else + { + if constexpr(common::mpl::is_string_type::value) ofs << "\""; + ofs << _val; + if constexpr(common::mpl::is_string_type::value) ofs << "\""; + } }; (_write(Idx, std::get(_data)), ...); @@ -75,6 +90,7 @@ struct csv_encoder }; using api_csv_encoder = csv_encoder<7>; +using agent_info_csv_encoder = csv_encoder<53>; using kernel_trace_csv_encoder = csv_encoder<16>; using counter_collection_csv_encoder = csv_encoder<15>; using memory_copy_csv_encoder = csv_encoder<7>; @@ -82,6 +98,7 @@ using marker_csv_encoder = csv_encoder<7>; using list_basic_metrics_csv_encoder = csv_encoder<5>; using list_derived_metrics_csv_encoder = csv_encoder<5>; using scratch_memory_encoder = csv_encoder<8>; +using stats_csv_encoder = csv_encoder<8>; } // namespace csv } // namespace tool } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-tool/generateCSV.cpp b/source/lib/rocprofiler-sdk-tool/generateCSV.cpp new file mode 100644 index 0000000000..dd4266bfca --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/generateCSV.cpp @@ -0,0 +1,446 @@ +// MIT License +// +// Copyright (c) 2023 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. + +#include "generateCSV.hpp" +#include "csv.hpp" +#include "helper.hpp" +#include "statistics.hpp" + +#include +#include + +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace +{ +using float_type = double; +using stats_data_t = statistics; +using stats_map_t = std::map; + +void +write_stats(output_file& ofs, timestamps_t* app_timestamps, const stats_map_t& data) +{ + auto _ss = std::stringstream{}; + float_type app_duration = (app_timestamps->app_end_time - app_timestamps->app_start_time); + for(const auto& [id, value] : data) + { + auto duration_ns = value.get_sum(); + auto calls = value.get_count(); + const auto& name = id; + float_type avg_ns = value.get_mean(); + float_type percentage = (duration_ns / app_duration) * static_cast(100); + + rocprofiler::tool::csv::stats_csv_encoder::write_row(_ss, + name, + calls, + duration_ns, + avg_ns, + percentage, + value.get_min(), + value.get_max(), + value.get_stddev()); + } + ofs << _ss.str(); +} +} // namespace + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + std::sort(data.begin(), data.end(), [](rocprofiler_agent_v0_t lhs, rocprofiler_agent_v0_t rhs) { + return lhs.node_id < rhs.node_id; + }); + + for(auto& itr : data) + { + auto _type = std::string_view{}; + if(itr.type == ROCPROFILER_AGENT_TYPE_CPU) + _type = "CPU"; + else if(itr.type == ROCPROFILER_AGENT_TYPE_GPU) + _type = "GPU"; + else + _type = "UNK"; + + auto agent_info_ss = std::stringstream{}; + rocprofiler::tool::csv::agent_info_csv_encoder::write_row(agent_info_ss, + itr.node_id, + itr.logical_node_id, + _type, + itr.cpu_cores_count, + itr.simd_count, + itr.cpu_core_id_base, + itr.simd_id_base, + itr.max_waves_per_simd, + itr.lds_size_in_kb, + itr.gds_size_in_kb, + itr.num_gws, + itr.wave_front_size, + itr.num_xcc, + itr.cu_count, + itr.array_count, + itr.num_shader_banks, + itr.simd_arrays_per_engine, + itr.cu_per_simd_array, + itr.simd_per_cu, + itr.max_slots_scratch_cu, + itr.gfx_target_version, + itr.vendor_id, + itr.device_id, + itr.location_id, + itr.domain, + itr.drm_render_minor, + itr.num_sdma_engines, + itr.num_sdma_xgmi_engines, + itr.num_sdma_queues_per_engine, + itr.num_cp_queues, + itr.max_engine_clk_ccompute, + itr.max_engine_clk_fcompute, + itr.sdma_fw_version.Value, + itr.fw_version.Value, + itr.capability.Value, + itr.cu_per_engine, + itr.max_waves_per_cu, + itr.family_id, + itr.workgroup_max_size, + itr.grid_max_size, + itr.local_mem_size, + itr.hive_id, + itr.gpu_id, + itr.workgroup_max_dim.x, + itr.workgroup_max_dim.y, + itr.workgroup_max_dim.z, + itr.grid_max_dim.x, + itr.grid_max_dim.y, + itr.grid_max_dim.z, + itr.name, + itr.vendor_name, + itr.product_name, + itr.model_name); + tool_functions->tool_get_agent_info_file_fn() << agent_info_ss.str(); + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + auto kernel_stats = stats_map_t{}; + for(auto& buf : data) + { + while(true) + { + auto kernel_trace_ss = std::stringstream{}; + rocprofiler_buffer_tracing_kernel_dispatch_record_t* record = buf.retrieve(); + if(record == nullptr) break; + auto kernel_name = tool_functions->tool_get_kernel_name_fn(record->kernel_id); + rocprofiler::tool::csv::kernel_trace_csv_encoder::write_row( + kernel_trace_ss, + tool_functions->tool_get_domain_name_fn(record->kind), + tool_functions->tool_get_agent_node_id_fn(record->agent_id), + record->queue_id.handle, + record->kernel_id, + kernel_name, + record->correlation_id.internal, + 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); + + if(tool::get_config().stats) + kernel_stats[kernel_name] += (record->end_timestamp - record->start_timestamp); + + tool_functions->tool_get_kernel_trace_file_fn() << kernel_trace_ss.str(); + } + } + + if(tool::get_config().stats) + write_stats(tool_functions->tool_get_kernel_stats_file_fn(), + tool_functions->tool_get_app_timestamps_fn(), + kernel_stats); +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + auto hip_stats = stats_map_t{}; + for(auto& buf : data) + { + while(true) + { + auto hip_trace_ss = std::stringstream{}; + rocprofiler_buffer_tracing_hip_api_record_t* record = buf.retrieve(); + if(record == nullptr) break; + auto api_name = + tool_functions->tool_get_operation_name_fn(record->kind, record->operation); + rocprofiler::tool::csv::api_csv_encoder::write_row( + hip_trace_ss, + tool_functions->tool_get_domain_name_fn(record->kind), + api_name, + getpid(), + record->thread_id, + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp); + + if(tool::get_config().stats) + hip_stats[api_name] += (record->end_timestamp - record->start_timestamp); + + tool_functions->tool_get_hip_api_trace_file_fn() << hip_trace_ss.str(); + } + } + + if(tool::get_config().stats) + { + write_stats(tool_functions->tool_get_hip_stats_file_fn(), + tool_functions->tool_get_app_timestamps_fn(), + hip_stats); + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + auto hsa_stats = stats_map_t{}; + for(auto& buf : data) + { + while(true) + { + auto hsa_trace_ss = std::stringstream{}; + rocprofiler_buffer_tracing_hsa_api_record_t* record = buf.retrieve(); + if(record == nullptr) break; + auto api_name = + tool_functions->tool_get_operation_name_fn(record->kind, record->operation); + rocprofiler::tool::csv::api_csv_encoder::write_row( + hsa_trace_ss, + tool_functions->tool_get_domain_name_fn(record->kind), + api_name, + getpid(), + record->thread_id, + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp); + + if(tool::get_config().stats) + hsa_stats[api_name] += (record->end_timestamp - record->start_timestamp); + + tool_functions->tool_get_hsa_api_trace_file_fn() << hsa_trace_ss.str(); + } + } + + if(tool::get_config().stats) + { + write_stats(tool_functions->tool_get_hsa_stats_file_fn(), + tool_functions->tool_get_app_timestamps_fn(), + hsa_stats); + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + auto memory_copy_stats = stats_map_t{}; + for(auto& buf : data) + { + while(true) + { + auto memory_copy_trace_ss = std::stringstream{}; + rocprofiler_buffer_tracing_memory_copy_record_t* record = buf.retrieve(); + if(record == nullptr) break; + auto api_name = + tool_functions->tool_get_operation_name_fn(record->kind, record->operation); + rocprofiler::tool::csv::memory_copy_csv_encoder::write_row( + memory_copy_trace_ss, + tool_functions->tool_get_domain_name_fn(record->kind), + api_name, + tool_functions->tool_get_agent_node_id_fn(record->src_agent_id), + tool_functions->tool_get_agent_node_id_fn(record->dst_agent_id), + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp); + + if(tool::get_config().stats) + memory_copy_stats[api_name] += (record->end_timestamp - record->start_timestamp); + + tool_functions->tool_get_memory_copy_trace_file_fn() << memory_copy_trace_ss.str(); + } + } + + if(tool::get_config().stats) + { + write_stats(tool_functions->tool_get_memory_copy_stats_file_fn(), + tool_functions->tool_get_app_timestamps_fn(), + memory_copy_stats); + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + for(auto& buf : data) + { + while(true) + { + auto marker_api_trace_ss = std::stringstream{}; + rocprofiler_tool_marker_record_t* record = buf.retrieve(); + if(record == nullptr) break; + if(record->kind == ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API && + ((record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA && + record->phase == ROCPROFILER_CALLBACK_PHASE_EXIT) || + (record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePop && + record->phase == ROCPROFILER_CALLBACK_PHASE_ENTER) || + (record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStop && + record->phase == ROCPROFILER_CALLBACK_PHASE_ENTER))) + { + tool::csv::marker_csv_encoder::write_row( + marker_api_trace_ss, + tool_functions->tool_get_callback_kind_fn(record->kind), + tool_functions->tool_get_roctx_msg_fn(record->cid), + record->pid, + record->tid, + record->cid, + record->start_timestamp, + record->end_timestamp); + } + else + { + tool::csv::marker_csv_encoder::write_row( + marker_api_trace_ss, + tool_functions->tool_get_callback_kind_fn(record->kind), + tool_functions->tool_get_callback_op_name_fn(record->kind, record->op), + record->pid, + record->tid, + record->cid, + record->start_timestamp, + record->end_timestamp); + } + tool_functions->tool_get_marker_api_trace_file_fn() << marker_api_trace_ss.str(); + } + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + for(auto& buf : data) + { + while(true) + { + rocprofiler_tool_counter_collection_record_t* record = buf.retrieve(); + if(record == nullptr) break; + auto kernel_id = record->dispatch_data.kernel_id; + auto counter_name_value = std::map{}; + for(const auto& count : record->profiler_record) + { + auto rec = static_cast(count); + std::string counter_name = tool_functions->tool_get_counter_info_name_fn(rec.id); + auto search = counter_name_value.find(counter_name); + if(search == counter_name_value.end()) + counter_name_value.emplace( + std::pair{counter_name, rec.counter_value}); + else + search->second = search->second + rec.counter_value; + } + + const auto& correlation_id = record->dispatch_data.correlation_id; + + auto magnitude = [](rocprofiler_dim3_t dims) { return (dims.x * dims.y * dims.z); }; + auto counter_collection_ss = std::stringstream{}; + for(auto& itr : counter_name_value) + { + tool::csv::counter_collection_csv_encoder::write_row( + counter_collection_ss, + correlation_id.internal, + record->dispatch_index, + tool_functions->tool_get_agent_node_id_fn(record->dispatch_data.agent_id), + record->dispatch_data.queue_id.handle, + record->pid, + record->thread_id, + magnitude(record->dispatch_data.grid_size), + tool_functions->tool_get_kernel_name_fn(kernel_id), + magnitude(record->dispatch_data.workgroup_size), + record->lds_block_size_v, + record->private_segment_size, + record->arch_vgpr_count, + record->sgpr_count, + itr.first, + itr.second); + } + tool_functions->tool_get_counter_collection_file_fn() << counter_collection_ss.str(); + } + } +} + +void +generate_csv(tool_table* tool_functions, std::vector& data) +{ + auto* ofs = tool_functions->tool_get_scratch_memory_file_fn(); + if(!ofs) throw std::runtime_error{"error creating scratch memory output file"}; + + auto scratch_memory_stats = stats_map_t{}; + for(auto& buf : data) + { + while(true) + { + rocprofiler_buffer_tracing_scratch_memory_record_t* record = buf.retrieve(); + if(record == nullptr) break; + + auto scratch_memory_trace = std::stringstream{}; + auto kind_name = tool_functions->tool_get_domain_name_fn(record->kind); + auto op_name = + tool_functions->tool_get_operation_name_fn(record->kind, record->operation); + + tool::csv::scratch_memory_encoder::write_row( + scratch_memory_trace, + kind_name, + op_name, + tool_functions->tool_get_agent_node_id_fn(record->agent_id), + record->queue_id.handle, + record->thread_id, + record->flags, + record->start_timestamp, + record->end_timestamp); + + if(tool::get_config().stats) + scratch_memory_stats[op_name] += (record->end_timestamp - record->start_timestamp); + + (*ofs) << scratch_memory_trace.str(); + } + } + + if(tool::get_config().stats) + { + auto* stats_ofs = tool_functions->tool_get_scratch_memory_stats_file_fn(); + if(!stats_ofs) throw std::runtime_error{"error creating scratch memory stats output file"}; + write_stats(*stats_ofs, tool_functions->tool_get_app_timestamps_fn(), scratch_memory_stats); + } +} +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-tool/generateCSV.hpp b/source/lib/rocprofiler-sdk-tool/generateCSV.hpp new file mode 100644 index 0000000000..9cf235fd42 --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/generateCSV.hpp @@ -0,0 +1,57 @@ +// MIT License +// +// Copyright (c) 2023 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. + +#pragma once + +#include "helper.hpp" + +#include + +namespace rocprofiler +{ +namespace tool +{ +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); + +void +generate_csv(tool_table* tool_functions, std::vector& data); +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-tool/helper.cpp b/source/lib/rocprofiler-sdk-tool/helper.cpp index c69bf1ad92..837a12ad37 100644 --- a/source/lib/rocprofiler-sdk-tool/helper.cpp +++ b/source/lib/rocprofiler-sdk-tool/helper.cpp @@ -30,6 +30,7 @@ #include #include #include +#include #include #include #include diff --git a/source/lib/rocprofiler-sdk-tool/helper.hpp b/source/lib/rocprofiler-sdk-tool/helper.hpp index 035406302d..c793416184 100644 --- a/source/lib/rocprofiler-sdk-tool/helper.hpp +++ b/source/lib/rocprofiler-sdk-tool/helper.hpp @@ -22,8 +22,11 @@ #pragma once +#include "lib/common/container/ring_buffer.hpp" +#include "lib/common/container/small_vector.hpp" #include "lib/common/defines.hpp" #include "lib/common/filesystem.hpp" +#include "output_file.hpp" #include #include @@ -60,7 +63,8 @@ rocprofiler_get_status_string(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__)); \ std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ << " failed with error code " << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \ - << ": " << status_msg << std::endl; \ + << ": " << status_msg << "\n" \ + << std::flush; \ std::stringstream errmsg{}; \ errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ << status_msg << ")"; \ @@ -77,6 +81,9 @@ using rocprofiler_tool_buffer_kind_operation_names_t = std::unordered_map>; +namespace common = ::rocprofiler::common; +namespace tool = ::rocprofiler::tool; + struct rocprofiler_tool_buffer_name_info_t { rocprofiler_tool_buffer_kind_names_t kind_names = {}; @@ -100,3 +107,109 @@ get_buffer_id_names(); rocprofiler_tool_callback_name_info_t get_callback_id_names(); + +struct rocprofiler_tool_marker_record_t +{ + rocprofiler_callback_tracing_kind_t kind; + uint32_t op = 0; + uint32_t phase = 0; + uint64_t pid = 0; + uint64_t tid = 0; + uint64_t cid = 0; + + rocprofiler_timestamp_t start_timestamp; + rocprofiler_timestamp_t end_timestamp; +}; + +struct rocprofiler_tool_counter_collection_record_t +{ + rocprofiler_profile_counting_dispatch_data_t dispatch_data; + common::container::small_vector profiler_record; + uint64_t pid = 0; + uint64_t id = 0; + uint64_t thread_id = 0; + uint64_t dispatch_index = 0; + uint64_t private_segment_size = 0; + uint64_t arch_vgpr_count = 0; + uint64_t sgpr_count = 0; + uint64_t lds_block_size_v = 0; +}; + +struct timestamps_t +{ + rocprofiler_timestamp_t app_start_time; + rocprofiler_timestamp_t app_end_time; +}; + +using hip_ring_buffer_t = + rocprofiler::common::container::ring_buffer; +using hsa_ring_buffer_t = + rocprofiler::common::container::ring_buffer; +using kernel_dispatch_ring_buffer_t = rocprofiler::common::container::ring_buffer< + rocprofiler_buffer_tracing_kernel_dispatch_record_t>; +using memory_copy_ring_buffer_t = + rocprofiler::common::container::ring_buffer; +using counter_collection_buffer_t = + rocprofiler::common::container::ring_buffer; +using marker_api_ring_buffer_t = + rocprofiler::common::container::ring_buffer; +using counter_collection_ring_buffer_t = + rocprofiler::common::container::ring_buffer; +using scratch_memory_ring_buffer_t = + rocprofiler::common::container::ring_buffer; + +using tool_get_agent_node_id_fn_t = uint64_t (*)(rocprofiler_agent_id_t); +using tool_get_app_timestamps_fn_t = timestamps_t* (*) (); +using tool_get_kernel_name_fn_t = std::string_view (*)(uint64_t); +using tool_get_domain_name_fn_t = std::string_view (*)(rocprofiler_buffer_tracing_kind_t); +using tool_get_operation_name_fn_t = std::string_view (*)(rocprofiler_buffer_tracing_kind_t, + rocprofiler_tracing_operation_t); +using tool_get_callback_kind_name_fn_t = std::string_view (*)(rocprofiler_callback_tracing_kind_t); +using tool_get_callback_op_name_fn_t = std::string_view (*)(rocprofiler_callback_tracing_kind_t, + uint32_t); +using tool_get_roctx_msg_fn_t = std::string_view (*)(uint64_t); +using tool_get_counter_info_name_fn_t = std::string (*)(uint64_t); +using tool_get_output_file_ref_fn_t = rocprofiler::tool::output_file& (*) (); +using tool_get_output_file_ptr_fn_t = rocprofiler::tool::output_file*& (*) (); + +enum class buffer_type_t +{ + ROCPROFILER_TOOL_BUFFER_HSA = 0, + ROCPROFILER_TOOL_BUFFER_HIP, + ROCPROFILER_TOOL_BUFFER_MEMORY_COPY, + ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION, + ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH, + ROCPROFILER_TOOL_BUFFER_MARKER_API, + ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY, +}; + +struct tool_table +{ + // node id + tool_get_agent_node_id_fn_t tool_get_agent_node_id_fn = nullptr; + // timestamps + tool_get_app_timestamps_fn_t tool_get_app_timestamps_fn = nullptr; + // names and messages + tool_get_kernel_name_fn_t tool_get_kernel_name_fn = nullptr; + tool_get_domain_name_fn_t tool_get_domain_name_fn = nullptr; + tool_get_operation_name_fn_t tool_get_operation_name_fn = nullptr; + tool_get_counter_info_name_fn_t tool_get_counter_info_name_fn = nullptr; + tool_get_callback_kind_name_fn_t tool_get_callback_kind_fn = nullptr; + tool_get_callback_op_name_fn_t tool_get_callback_op_name_fn = nullptr; + tool_get_roctx_msg_fn_t tool_get_roctx_msg_fn = nullptr; + // trace files + tool_get_output_file_ref_fn_t tool_get_agent_info_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_kernel_trace_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_hsa_api_trace_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_hip_api_trace_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_marker_api_trace_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_counter_collection_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_memory_copy_trace_file_fn = nullptr; + tool_get_output_file_ptr_fn_t tool_get_scratch_memory_file_fn = nullptr; + // stats files + tool_get_output_file_ref_fn_t tool_get_kernel_stats_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_hip_stats_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_hsa_stats_file_fn = nullptr; + tool_get_output_file_ref_fn_t tool_get_memory_copy_stats_file_fn = nullptr; + tool_get_output_file_ptr_fn_t tool_get_scratch_memory_stats_file_fn = nullptr; +}; diff --git a/source/lib/rocprofiler-sdk-tool/statistics.hpp b/source/lib/rocprofiler-sdk-tool/statistics.hpp new file mode 100644 index 0000000000..655b8ba56b --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/statistics.hpp @@ -0,0 +1,231 @@ +// MIT License +// +// Copyright (c) 2023 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. + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +/// \struct statistics +/// \tparam Tp data type for statistical accumulation +/// \tparam Fp floating point data type to use for division +/// \brief A generic class for statistical accumulation. +/// +template +struct statistics +{ +public: + using value_type = Tp; + using float_type = Fp; + using this_type = statistics; + static_assert(std::is_arithmetic::value, "only supports arithmetic types"); + +public: + inline statistics() = default; + inline ~statistics() = default; + inline statistics(const statistics&) = default; + inline statistics(statistics&&) noexcept = default; + inline statistics& operator=(const statistics&) = default; + inline statistics& operator=(statistics&&) noexcept = default; + + explicit statistics(value_type val) + : m_cnt(1) + , m_sum(val) + , m_sqr(val * val) + , m_min(val) + , m_max(val) + {} + + statistics& operator=(value_type val) + { + m_cnt = 1; + m_sum = val; + m_min = val; + m_max = val; + m_sqr = (val * val); + return *this; + } + +public: + // Accumulated values + inline int64_t get_count() const { return m_cnt; } + inline value_type get_min() const { return m_min; } + inline value_type get_max() const { return m_max; } + inline value_type get_sum() const { return m_sum; } + inline value_type get_sqr() const { return m_sqr; } + inline float_type get_mean() const { return static_cast(m_sum) / m_cnt; } + inline float_type get_variance() const + { + if(m_cnt < 2) return (m_sum - m_sum); + + auto _sum_of_squared_samples = m_sqr; + auto _sum_squared_mean = (m_sum * m_sum) / static_cast(m_cnt); + return (_sum_of_squared_samples - _sum_squared_mean) / static_cast(m_cnt - 1); + } + + inline float_type get_stddev() const { return ::std::sqrt(::std::abs(get_variance())); } + + // Modifications + inline void reset() + { + m_cnt = 0; + m_sum = value_type{}; + m_sqr = value_type{}; + m_min = value_type{}; + m_max = value_type{}; + } + +public: + // Operators (value_type) + inline statistics& operator+=(value_type val) + { + if(m_cnt == 0) + { + m_sum = val; + m_sqr = (val * val); + m_min = val; + m_max = val; + } + else + { + m_sum += val; + m_sqr += (val * val); + m_min = ::std::min(m_min, val); + m_max = ::std::max(m_max, val); + } + ++m_cnt; + + return *this; + } + + inline statistics& operator-=(value_type val) + { + if(m_cnt > 1) --m_cnt; + m_sum -= val; + m_sqr -= (val * val); + m_min -= val; + m_max -= val; + return *this; + } + + inline statistics& operator*=(value_type val) + { + m_sum *= val; + m_sqr *= (val * val); + m_min *= val; + m_max *= val; + return *this; + } + + inline statistics& operator/=(value_type val) + { + m_sum /= val; + m_sqr /= (val * val); + m_min /= val; + m_max /= val; + return *this; + } + +public: + // Operators (this_type) + inline statistics& operator+=(const statistics& rhs) + { + if(m_cnt == 0) + { + m_sum = rhs.m_sum; + m_sqr = rhs.m_sqr; + m_min = rhs.m_min; + m_max = rhs.m_max; + } + else + { + m_sum += rhs.m_sum; + m_sqr += rhs.m_sqr; + m_min = ::std::min(m_min, rhs.m_min); + m_max = ::std::max(m_max, rhs.m_max); + } + m_cnt += rhs.m_cnt; + return *this; + } + + // Operators (this_type) + inline statistics& operator-=(const statistics& rhs) + { + if(m_cnt > 0) + { + m_sum -= rhs.m_sum; + m_sqr -= rhs.m_sqr; + m_min = ::std::min(m_min, rhs.m_min); + m_max = ::std::max(m_max, rhs.m_max); + } + return *this; + } + +private: + // summation of each history^1 + int64_t m_cnt = 0; + value_type m_sum = value_type{}; + value_type m_sqr = value_type{}; + value_type m_min = value_type{}; + value_type m_max = value_type{}; + +public: + // friend operator for addition + friend statistics operator+(const statistics& lhs, const statistics& rhs) + { + return statistics(lhs) += rhs; + } + + friend statistics operator-(const statistics& lhs, const statistics& rhs) + { + return statistics(lhs) -= rhs; + } +}; +} // namespace tool +} // namespace rocprofiler + +namespace std +{ +template +::rocprofiler::tool::statistics +max(::rocprofiler::tool::statistics lhs, const Tp& rhs) +{ + return lhs.get_max(rhs); +} + +template +::rocprofiler::tool::statistics +min(::rocprofiler::tool::statistics lhs, const Tp& rhs) +{ + return lhs.get_min(rhs); +} +} // namespace std diff --git a/source/lib/rocprofiler-sdk-tool/tmp_file.cpp b/source/lib/rocprofiler-sdk-tool/tmp_file.cpp new file mode 100644 index 0000000000..095da445c6 --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/tmp_file.cpp @@ -0,0 +1,131 @@ +// MIT License +// +// Copyright (c) 2023 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. + +#include "tmp_file.hpp" +#include "config.hpp" + +#include "lib/common/filesystem.hpp" + +namespace fs = ::rocprofiler::common::filesystem; + +bool +tmp_file::fopen(const char* _mode) +{ + if(!fs::exists(filename)) + { + // if the filepath does not exist, open in out mode to create it + std::ofstream _ofs{filename}; + } + file = std::fopen(filename.c_str(), _mode); + if(file) fd = ::fileno(file); + + return (file != nullptr && fd > 0); +} + +tmp_file::tmp_file(std::string _filename) +: filename(std::move(_filename)) +{} + +tmp_file::~tmp_file() +{ + close(); + remove(); +} + +bool +tmp_file::flush() +{ + if(stream.is_open()) + { + stream.flush(); + } + else if(file != nullptr) + { + int _ret = fflush(file); + int _cnt = 0; + while(_ret == EAGAIN || _ret == EINTR) + { + // std::this_thread::sleep_for(std::chrono::milliseconds{ 100 }); + _ret = fflush(file); + if(++_cnt > 10) break; + } + return (_ret == 0); + } + + return true; +} + +bool +tmp_file::close() +{ + flush(); + + if(stream.is_open()) + { + stream.close(); + return !stream.is_open(); + } + else if(file != nullptr) + { + auto _ret = fclose(file); + if(_ret == 0) + { + file = nullptr; + fd = -1; + } + return (_ret == 0); + } + + return true; +} + +bool +tmp_file::open(std::ios::openmode _mode) +{ + if(!fs::exists(filename)) + { + // if the filepath does not exist, open in out mode to create it + std::ofstream _ofs{}; + _ofs.open(filename, std::ofstream::binary | std::ofstream::out); + } + + stream.open(filename, _mode); + return (stream.is_open() && stream.good()); +} + +bool +tmp_file::remove() +{ + close(); + if(fs::exists(filename)) + { + auto _ret = ::remove(filename.c_str()); + return (_ret == 0); + } + + return true; +} + +tmp_file::operator bool() const +{ + return (stream.is_open() && stream.good()) || (file != nullptr && fd > 0); +} diff --git a/source/lib/rocprofiler-sdk-tool/tmp_file.hpp b/source/lib/rocprofiler-sdk-tool/tmp_file.hpp new file mode 100644 index 0000000000..04348fde7b --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/tmp_file.hpp @@ -0,0 +1,51 @@ +// MIT License +// +// Copyright (c) 2023 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. + +#pragma once + +#include +#include +#include +#include +#include + +struct tmp_file +{ + tmp_file(std::string _filename); + ~tmp_file(); + + bool open(std::ios::openmode = std::ios::binary | std::ios::in | std::ios::out); + bool fopen(const char* = "r+"); + bool flush(); + bool close(); + bool remove(); + + explicit operator bool() const; + + std::string filename = {}; + std::string subdirectory = {}; + std::fstream stream = {}; + FILE* file = nullptr; + int fd = -1; + std::set file_pos = {}; + std::mutex file_mutex = {}; +}; diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index 146ead262c..5812ca3836 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -22,8 +22,10 @@ #include "config.hpp" #include "csv.hpp" +#include "generateCSV.hpp" #include "helper.hpp" #include "output_file.hpp" +#include "tmp_file.hpp" #include "lib/common/demangle.hpp" #include "lib/common/environment.hpp" @@ -54,6 +56,13 @@ #include #include +#if defined(ROCPROFILER_CODECOV) && ROCPROFILER_CODECOV > 0 +extern "C" { +extern void +__gcov_dump(void); +} +#endif + namespace common = ::rocprofiler::common; namespace tool = ::rocprofiler::tool; @@ -110,6 +119,31 @@ add_destructor(Tp*& ptr) std::call_once(_once, []() { add_destructor(PTR); }); \ } +tool::output_file*& +hsa_stats_file() +{ + static auto* _v = new tool::output_file{"hsa_stats", + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_hsa_stats_file() +{ + return get_dereference(hsa_stats_file()); +} + tool::output_file*& get_hsa_api_file() { @@ -126,6 +160,37 @@ get_hsa_api_file() return _v; } +tool::output_file& +get_hsa_api_trace_file() +{ + return get_dereference(get_hsa_api_file()); +} + +tool::output_file*& +hip_stats_file() +{ + static auto* _v = new tool::output_file{"hip_stats", + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_hip_stats_file() +{ + return get_dereference(hip_stats_file()); +} + tool::output_file*& get_hip_api_file() { @@ -142,8 +207,107 @@ get_hip_api_file() return _v; } +tool::output_file& +get_hip_api_trace_file() +{ + return get_dereference(get_hip_api_file()); +} + tool::output_file*& -get_kernel_trace_file() +get_agent_info_file_impl() +{ + static auto* _v = new tool::output_file{"agent_info", + tool::csv::agent_info_csv_encoder{}, + {"Node_Id", + "Logical_Node_Id", + "Agent_Type", + "Cpu_Cores_Count", + "Simd_Count", + "Cpu_Core_Id_Base", + "Simd_Id_Base", + "Max_Waves_Per_Simd", + "Lds_Size_In_Kb", + "Gds_Size_In_Kb", + "Num_Gws", + "Wave_Front_Size", + "Num_Xcc", + "Cu_Count", + "Array_Count", + "Num_Shader_Banks", + "Simd_Arrays_Per_Engine", + "Cu_Per_Simd_Array", + "Simd_Per_Cu", + "Max_Slots_Scratch_Cu", + "Gfx_Target_Version", + "Vendor_Id", + "Device_Id", + "Location_Id", + "Domain", + "Drm_Render_Minor", + "Num_Sdma_Engines", + "Num_Sdma_Xgmi_Engines", + "Num_Sdma_Queues_Per_Engine", + "Num_Cp_Queues", + "Max_Engine_Clk_Ccompute", + "Max_Engine_Clk_Fcompute", + "Sdma_Fw_Version", + "Fw_Version", + "Capability", + "Cu_Per_Engine", + "Max_Waves_Per_Cu", + "Family_Id", + "Workgroup_Max_Size", + "Grid_Max_Size", + "Local_Mem_Size", + "Hive_Id", + "Gpu_Id", + "Workgroup_Max_Dim_X", + "Workgroup_Max_Dim_Y", + "Workgroup_Max_Dim_Z", + "Grid_Max_Dim_X", + "Grid_Max_Dim_Y", + "Grid_Max_Dim_Z", + "Name", + "Vendor_Name", + "Product_Name", + "Model_Name"}}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_agent_info_file() +{ + return get_dereference(get_agent_info_file_impl()); +} + +tool::output_file*& +kernel_stats_file() +{ + static auto* _v = new tool::output_file{"kernel_stats", + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_kernel_stats_file() +{ + return get_dereference(kernel_stats_file()); +} + +tool::output_file*& +get_kernel_file() { static auto* _v = new tool::output_file{"kernel_trace", tool::csv::kernel_trace_csv_encoder{}, @@ -167,6 +331,12 @@ get_kernel_trace_file() return _v; } +tool::output_file& +get_kernel_trace_file() +{ + return get_dereference(get_kernel_file()); +} + tool::output_file*& get_counter_collection_file() { @@ -192,6 +362,12 @@ get_counter_collection_file() return _v; } +tool::output_file& +get_counter_file() +{ + return get_dereference(get_counter_collection_file()); +} + tool::output_file*& get_memory_copy_trace_file() { @@ -208,7 +384,60 @@ get_memory_copy_trace_file() return _v; } -auto& +tool::output_file& +get_memory_copy_file() +{ + return get_dereference(get_memory_copy_trace_file()); +} + +tool::output_file*& +memory_copy_stats_file() +{ + static auto* _v = new tool::output_file{"memory_copy_stats", + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_memory_copy_stats_file() +{ + return get_dereference(memory_copy_stats_file()); +} + +tool::output_file*& +get_marker_api_file() +{ + static auto* _v = new tool::output_file{"marker_api_trace", + tool::csv::marker_csv_encoder{}, + {"Domain", + "Function", + "Process_Id", + "Thread_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + ADD_DESTRUCTOR(_v); + return _v; +} + +tool::output_file& +get_marker_api_trace_file() +{ + return get_dereference(get_marker_api_file()); +} + +tool::output_file*& get_scratch_memory_trace_file() { static auto* _v = new tool::output_file{"scratch_memory_trace", @@ -228,17 +457,20 @@ get_scratch_memory_trace_file() } tool::output_file*& -get_marker_api_file() +get_scratch_memory_stats_file() { - static auto* _v = new tool::output_file{"marker_api_trace", - tool::csv::marker_csv_encoder{}, - {"Domain", - "Function", - "Process_Id", - "Thread_Id", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp"}}; + static auto* _v = new tool::output_file{"scratch_memory_stats", + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; ADD_DESTRUCTOR(_v); return _v; } @@ -335,16 +567,22 @@ using kernel_symbol_data_map_t = std::unordered_map; using counter_dimension_info_map_t = std::unordered_map>; -using agent_info_map_t = std::unordered_map; +using agent_info_map_t = std::unordered_map; +using marker_msg_t = std::unordered_map; auto code_obj_data = common::Synchronized{}; -auto kernel_data = common::Synchronized{}; +auto kernel_data_mutex = std::mutex{}; +auto* kernel_data = as_pointer(kernel_symbol_data_map_t{}); +auto marker_msg_mutex = std::mutex{}; +auto* marker_msg_cid = as_pointer(marker_msg_t{}); auto counter_dimension_data = common::Synchronized{}; auto target_kernels = common::Synchronized{}; auto dispatch_index = std::atomic{0}; auto* buffered_name_info = as_pointer(get_buffer_id_names()); auto* callback_name_info = as_pointer(get_callback_id_names()); auto* agent_info = as_pointer(agent_info_map_t{}); +auto* tool_functions = as_pointer(tool_table{}); +auto* stats_timestamp = as_pointer(timestamps_t{}); bool add_kernel_target(uint64_t _kern_id) @@ -388,6 +626,134 @@ flush() ROCP_INFO << "Buffers flushed"; } +std::string +get_file_name(buffer_type_t buffer_type) +{ + std::string filename; + switch(buffer_type) + { + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA: filename = "hsa_trace"; break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP: filename = "hip_trace"; break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API: filename = "marker_trace"; break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_MEMORY_COPY: filename = "memory_copy"; break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION: + filename = "counter_collection"; + break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH: + filename = "kernel_dispatch"; + break; + case buffer_type_t::ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY: + filename = "scratch_memory"; + break; + default: LOG_IF(FATAL, "buffer tpe invalid"); + } + return filename; +} + +std::string +compose_tmp_file_name(buffer_type_t buffer_type) +{ + static std::mutex _mutex{}; + std::unique_lock _lk{_mutex}; + std::string ext = ".dat"; + std::string explicit_path = rocprofiler::tool::get_config().tmp_directory; + std::string subdirectory = rocprofiler::tool::format(explicit_path + '/' + "%ppid%_"); + std::string fname = get_file_name(buffer_type); + std::string filename = subdirectory + fname + ext; + return filename; +} + +template +std::tuple +get_tmp_file_buffer(buffer_type_t type) +{ + static Tp* _buffer = new Tp(rocprofiler::common::units::get_page_size()); + static tmp_file* _tmp_file = new tmp_file(compose_tmp_file_name(type)); + return std::tuple(_buffer, _tmp_file); +} + +template +void +offload_buffer(buffer_type_t type) +{ + Tp* _tmp_buf = nullptr; + tmp_file* _tmp_file = nullptr; + std::tie(_tmp_buf, _tmp_file) = get_tmp_file_buffer(type); + auto _lk = std::lock_guard(_tmp_file->file_mutex); + auto& _fs = _tmp_file->stream; + [[maybe_unused]] static auto _success = _tmp_file->open(); + _tmp_file->file_pos.emplace(_fs.tellg()); + _tmp_buf->save(_fs); + _tmp_buf->clear(); + assert(_tmp_buf->is_empty() == true); +} + +template +void +write_ring_buffer(Tb* _v, buffer_type_t type) +{ + Tp* _tmp_buf = nullptr; + tmp_file* _tmp_file = nullptr; + std::tie(_tmp_buf, _tmp_file) = get_tmp_file_buffer(type); + auto* ptr = _tmp_buf->request(false); + if(ptr == nullptr) + { + offload_buffer(type); + ptr = _tmp_buf->request(false); + assert(ptr != nullptr); + } + *ptr = std::move(*reinterpret_cast(_v)); +} + +template +void +flush_tmp_buffer(buffer_type_t type) +{ + Tp* _tmp_buf = nullptr; + tmp_file* _tmp_file = nullptr; + std::tie(_tmp_buf, _tmp_file) = get_tmp_file_buffer(type); + if(!_tmp_buf->is_empty()) offload_buffer(type); +} + +template +void +read_tmp_file(buffer_type_t type, std::vector& _data) +{ + Tp* _tmp_buf = nullptr; + tmp_file* _tmp_file; + std::tie(_tmp_buf, _tmp_file) = get_tmp_file_buffer(type); + auto _lk = std::lock_guard{_tmp_file->file_mutex}; + auto& _fs = _tmp_file->stream; + if(_fs.is_open()) _fs.close(); + _tmp_file->open(std::ios::binary | std::ios::in); + for(auto itr : _tmp_file->file_pos) + { + _fs.seekg(itr); // set to the absolute position + if(_fs.eof()) break; + Tp _buffer; + _buffer.load(_fs); + _data.emplace_back(std::move(_buffer)); + } +} + +std::string_view +get_callback_kind(rocprofiler_callback_tracing_kind_t kind) +{ + return CHECK_NOTNULL(callback_name_info)->kind_names.at(kind); +} + +std::string_view +get_callback_op_name(rocprofiler_callback_tracing_kind_t kind, uint32_t op) +{ + return CHECK_NOTNULL(callback_name_info)->operation_names.at(kind).at(op); +} + +std::string_view +get_roctx_msg(uint64_t cid) +{ + return CHECK_NOTNULL(marker_msg_cid)->at(cid); +} + void cntrl_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* user_data, @@ -411,26 +777,24 @@ cntrl_tracing_callback(rocprofiler_callback_tracing_record_t record, auto ts = rocprofiler_timestamp_t{}; rocprofiler_get_timestamp(&ts); - const auto* kind_name = CHECK_NOTNULL(callback_name_info)->kind_names.at(record.kind); if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) { user_data->value = ts; } else { - const auto* op_name = CHECK_NOTNULL(callback_name_info) - ->operation_names.at(record.kind) - .at(record.operation); - auto ss = std::stringstream{}; - tool::csv::marker_csv_encoder::write_row(ss, - kind_name, - op_name, - getpid(), - rocprofiler::common::get_tid(), - record.correlation_id.internal, - user_data->value, - ts); - get_dereference(get_marker_api_file()) << ss.str(); + rocprofiler_tool_marker_record_t marker_record; + marker_record.kind = record.kind; + marker_record.phase = record.phase; + marker_record.op = record.operation; + marker_record.pid = getpid(); + marker_record.tid = rocprofiler::common::get_tid(); + marker_record.cid = record.correlation_id.internal; + marker_record.start_timestamp = user_data->value; + marker_record.end_timestamp = ts; + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API; + write_ring_buffer( + &marker_record, buffer_type); } } } @@ -452,21 +816,27 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, auto ts = rocprofiler_timestamp_t{}; rocprofiler_get_timestamp(&ts); - const auto* kind_name = CHECK_NOTNULL(callback_name_info)->kind_names.at(record.kind); if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA) { if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) { - auto ss = std::stringstream{}; - tool::csv::marker_csv_encoder::write_row(ss, - kind_name, - marker_data->args.roctxMarkA.message, - getpid(), - rocprofiler::common::get_tid(), - record.correlation_id.internal, - ts, - ts); - get_dereference(get_marker_api_file()) << ss.str(); + { + std::lock_guard lk(marker_msg_mutex); + std::string msg = marker_data->args.roctxMarkA.message; + CHECK_NOTNULL(marker_msg_cid)->emplace(record.correlation_id.internal, msg); + } + rocprofiler_tool_marker_record_t marker_record; + marker_record.kind = record.kind; + marker_record.op = record.operation; + marker_record.phase = record.phase; + marker_record.pid = getpid(); + marker_record.tid = rocprofiler::common::get_tid(); + marker_record.cid = record.correlation_id.internal; + marker_record.start_timestamp = ts; + marker_record.end_timestamp = ts; + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API; + write_ring_buffer( + &marker_record, buffer_type); } } else if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePushA) @@ -492,11 +862,23 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, auto val = stacked_range.back(); stacked_range.pop_back(); - - auto ss = std::stringstream{}; - tool::csv::marker_csv_encoder::write_row( - ss, kind_name, val.message, val.pid, val.tid, val.cid, val.data.value, ts); - get_dereference(get_marker_api_file()) << ss.str(); + { + std::lock_guard lk(marker_msg_mutex); + std::string msg = val.message; + CHECK_NOTNULL(marker_msg_cid)->emplace(val.cid, msg); + } + rocprofiler_tool_marker_record_t marker_record; + marker_record.kind = record.kind; + marker_record.op = record.operation; + marker_record.phase = record.phase; + marker_record.pid = val.pid; + marker_record.tid = val.tid; + marker_record.cid = val.cid; + marker_record.start_timestamp = val.data.value; + marker_record.end_timestamp = ts; + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API; + write_ring_buffer( + &marker_record, buffer_type); } } else if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStartA) @@ -509,7 +891,11 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, _entry.cid = record.correlation_id.internal; _entry.data.value = ts; _entry.message = marker_data->args.roctxRangeStartA.message; - + { + std::lock_guard lk(marker_msg_mutex); + std::string msg = _entry.message; + CHECK_NOTNULL(marker_msg_cid)->emplace(_entry.cid, msg); + } global_range.wlock( [_id, &_entry](auto& map) { map.emplace(_id, std::move(_entry)); }); } @@ -521,18 +907,18 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, auto _id = marker_data->args.roctxRangeStop.id; auto&& _entry = global_range.rlock( [](const auto& map, auto _key) { return map.at(_key); }, _id); - - auto ss = std::stringstream{}; - tool::csv::marker_csv_encoder::write_row(ss, - kind_name, - _entry.message, - _entry.pid, - 0, - _entry.cid, - _entry.data.value, - ts); - get_dereference(get_marker_api_file()) << ss.str(); - + rocprofiler_tool_marker_record_t marker_record; + marker_record.kind = record.kind; + marker_record.op = record.operation; + marker_record.phase = record.phase; + marker_record.pid = _entry.pid; + marker_record.tid = 0; + marker_record.cid = _entry.cid; + marker_record.start_timestamp = _entry.data.value; + marker_record.end_timestamp = ts; + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API; + write_ring_buffer( + &marker_record, buffer_type); global_range.wlock([](auto& map, auto _key) { return map.erase(_key); }, _id); } } @@ -544,19 +930,18 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, } else { - const auto* op_name = CHECK_NOTNULL(callback_name_info) - ->operation_names.at(record.kind) - .at(record.operation); - auto ss = std::stringstream{}; - tool::csv::marker_csv_encoder::write_row(ss, - kind_name, - op_name, - getpid(), - rocprofiler::common::get_tid(), - record.correlation_id.internal, - user_data->value, - ts); - get_dereference(get_marker_api_file()) << ss.str(); + rocprofiler_tool_marker_record_t marker_record; + marker_record.kind = record.kind; + marker_record.op = record.operation; + marker_record.phase = record.phase; + marker_record.pid = getpid(); + marker_record.tid = rocprofiler::common::get_tid(); + marker_record.cid = record.correlation_id.internal; + marker_record.start_timestamp = user_data->value; + marker_record.end_timestamp = ts; + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API; + write_ring_buffer( + &marker_record, buffer_type); } } } @@ -598,12 +983,13 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, auto* sym_data = static_cast(record.payload); if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD) { - auto itr = kernel_data.wlock( - [](kernel_symbol_data_map_t& kdata, rocprofiler_kernel_symbol_data_t* sym_data_v) { - return kdata.emplace(sym_data_v->kernel_id, kernel_symbol_data{*sym_data_v}); - }, - CHECK_NOTNULL(sym_data)); - + std::pair itr; + { + std::lock_guard lk(kernel_data_mutex); + itr = CHECK_NOTNULL(kernel_data) + ->emplace(sym_data->kernel_id, + kernel_symbol_data{get_dereference(sym_data)}); + } LOG_IF(WARNING, !itr.second) << "duplicate kernel symbol data for kernel_id=" << sym_data->kernel_id; @@ -649,6 +1035,30 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, (void) data; } +std::string_view +get_kernel_name(uint64_t kernel_id) +{ + return CHECK_NOTNULL(kernel_data)->at(kernel_id).formatted_kernel_name; +} + +std::string_view +get_domain_name(rocprofiler_buffer_tracing_kind_t record_kind) +{ + return CHECK_NOTNULL(buffered_name_info)->kind_names.at(record_kind); +} + +uint64_t +get_agent_node_id(rocprofiler_agent_id_t agent_id) +{ + return agent_info->at(agent_id).logical_node_id; +} + +std::string_view +get_operation_name(rocprofiler_buffer_tracing_kind_t kind, rocprofiler_tracing_operation_t op) +{ + return CHECK_NOTNULL(buffered_name_info)->operation_names.at(kind).at(op); +} + void buffered_tracing_callback(rocprofiler_context_id_t /*context*/, rocprofiler_buffer_id_t /*buffer_id*/, @@ -675,38 +1085,14 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, { auto* record = static_cast( header->payload); - std::string kernel_name = kernel_data.rlock( - [](const kernel_symbol_data_map_t& kdata, - rocprofiler_buffer_tracing_kernel_dispatch_record_t* record_v) { - auto _name = kdata.at(record_v->kernel_id).formatted_kernel_name; - return _name; - }, - record); - auto kernel_trace_ss = std::stringstream{}; - // TODO(aelwazir): To be changed back to use node id once ROCR fixes - // the hsa_agents to use the real node id - tool::csv::kernel_trace_csv_encoder::write_row( - kernel_trace_ss, - CHECK_NOTNULL(buffered_name_info)->kind_names.at(record->kind), - agent_info->at(record->agent_id)->logical_node_id, - record->queue_id.handle, - record->kernel_id, - std::move(kernel_name), - record->correlation_id.internal, - 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); + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH; - get_dereference(get_kernel_trace_file()) << kernel_trace_ss.str(); + write_ring_buffer(record, + buffer_type); } + else if(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API || header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API || header->kind == ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API || @@ -715,84 +1101,38 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, auto* record = static_cast(header->payload); - auto hsa_trace_ss = std::stringstream{}; - tool::csv::api_csv_encoder::write_row( - hsa_trace_ss, - CHECK_NOTNULL(buffered_name_info)->kind_names.at(record->kind), - CHECK_NOTNULL(buffered_name_info) - ->operation_names.at(record->kind) - .at(record->operation), - getpid(), - record->thread_id, - record->correlation_id.internal, - record->start_timestamp, - record->end_timestamp); - - get_dereference(get_hsa_api_file()) << hsa_trace_ss.str(); + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA; + write_ring_buffer( + record, buffer_type); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) { auto* record = static_cast(header->payload); - auto memory_copy_trace_ss = std::stringstream{}; - // TODO(aelwazir): To be changed back to use node id once ROCR fixes - // the hsa_agents to use the real node id - tool::csv::memory_copy_csv_encoder::write_row( - memory_copy_trace_ss, - CHECK_NOTNULL(buffered_name_info)->kind_names.at(record->kind), - CHECK_NOTNULL(buffered_name_info) - ->operation_names.at(record->kind) - .at(record->operation), - agent_info->at(record->src_agent_id)->logical_node_id, - agent_info->at(record->dst_agent_id)->logical_node_id, - record->correlation_id.internal, - record->start_timestamp, - record->end_timestamp); - - get_dereference(get_memory_copy_trace_file()) << memory_copy_trace_ss.str(); + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MEMORY_COPY; + write_ring_buffer(record, + buffer_type); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY) { auto* record = static_cast( header->payload); - auto scratch_memory_trace = std::stringstream{}; - tool::csv::scratch_memory_encoder::write_row( - scratch_memory_trace, - CHECK_NOTNULL(buffered_name_info)->kind_names.at(record->kind), - CHECK_NOTNULL(buffered_name_info) - ->operation_names.at(record->kind) - .at(record->operation), - agent_info->at(record->agent_id)->node_id, - record->queue_id.handle, - record->thread_id, - record->flags, - record->start_timestamp, - record->end_timestamp); - - get_dereference(get_scratch_memory_trace_file()) << scratch_memory_trace.str(); + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY; + write_ring_buffer(record, + buffer_type); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API || header->kind == ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API) { auto* record = static_cast(header->payload); - - auto hip_trace_ss = std::stringstream{}; - tool::csv::api_csv_encoder::write_row( - hip_trace_ss, - CHECK_NOTNULL(buffered_name_info)->kind_names.at(record->kind), - CHECK_NOTNULL(buffered_name_info) - ->operation_names.at(record->kind) - .at(record->operation), - getpid(), - record->thread_id, - record->correlation_id.internal, - record->start_timestamp, - record->end_timestamp); - - get_dereference(get_hip_api_file()) << hip_trace_ss.str(); + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP; + write_ring_buffer( + record, buffer_type); } else { @@ -925,6 +1265,21 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, } } +std::string +get_counter_info_name(uint64_t record_id) +{ + auto info = rocprofiler_counter_info_v0_t{}; + auto counter_id = rocprofiler_counter_id_t{}; + ROCPROFILER_CALL(rocprofiler_query_record_counter_id(record_id, &counter_id), + "query record counter id"); + ROCPROFILER_CALL(rocprofiler_query_counter_info(rocprofiler_counter_id_t{counter_id}, + ROCPROFILER_COUNTER_INFO_VERSION_0, + static_cast(&info)), + "query counter info"); + std::string counter_name = info.name; + return counter_name; +} + void counter_record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, rocprofiler_record_counter_t* record_data, @@ -934,72 +1289,39 @@ counter_record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_da { auto kernel_id = dispatch_data.kernel_id; const auto* cnt_dispatch_data_v = static_cast(user_data.ptr); - const auto* kernel_info = kernel_data.rlock( - [](const kernel_symbol_data_map_t& kdata, uint64_t kid) -> const auto* { - return &kdata.at(kid); - }, - kernel_id); + + rocprofiler_tool_counter_collection_record_t counter_record; + counter_record.dispatch_data = dispatch_data; + counter_record.dispatch_index = cnt_dispatch_data_v->dispatch_index; + counter_record.thread_id = cnt_dispatch_data_v->thread_id; + counter_record.pid = getpid(); + const kernel_symbol_data* kernel_info = nullptr; + + { + std::lock_guard lk(kernel_data_mutex); + kernel_info = &(CHECK_NOTNULL(kernel_data)->at(kernel_id)); + } + + auto lds_block_size_v = + (kernel_info->group_segment_size + (lds_block_size - 1)) & ~(lds_block_size - 1); + + counter_record.private_segment_size = kernel_info->private_segment_size; + counter_record.arch_vgpr_count = kernel_info->arch_vgpr_count; + counter_record.sgpr_count = kernel_info->sgpr_count; + counter_record.lds_block_size_v = lds_block_size_v; LOG_IF(FATAL, !kernel_info) << "missing kernel information for kernel_id=" << kernel_id; LOG_IF(ERROR, record_count == 0) << "zero record count for kernel_id=" << kernel_id << " (name=" << kernel_info->kernel_name << ")"; - auto counter_name_value = std::map{}; for(size_t count = 0; count < record_count; count++) - { - auto profiler_record = static_cast(record_data[count]); - auto counter_id = rocprofiler_counter_id_t{}; - auto info = rocprofiler_counter_info_v0_t{}; + counter_record.profiler_record.push_back( + static_cast(record_data[count])); - ROCPROFILER_CALL(rocprofiler_query_record_counter_id(profiler_record.id, &counter_id), - "query record counter id"); - ROCPROFILER_CALL( - rocprofiler_query_counter_info( - counter_id, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), - "query counter info"); - - auto search = counter_name_value.find(info.name); - if(search == counter_name_value.end()) - counter_name_value.emplace( - std::pair{info.name, profiler_record.counter_value}); - else - search->second = search->second + profiler_record.counter_value; - } - - auto lds_block_size_v = - (kernel_info->group_segment_size + (lds_block_size - 1)) & ~(lds_block_size - 1); - - const auto& correlation_id = dispatch_data.correlation_id; - - auto magnitude = [](rocprofiler_dim3_t dims) { return (dims.x * dims.y * dims.z); }; - - for(auto& itr : counter_name_value) - { - using csv_encoder = tool::csv::counter_collection_csv_encoder; - - auto counter_collection_ss = std::stringstream{}; - // TODO(aelwazir): To be changed back to use node id once ROCR fixes - // the hsa_agents to use the real node id - csv_encoder::write_row(counter_collection_ss, - correlation_id.internal, - cnt_dispatch_data_v->dispatch_index, - agent_info->at(dispatch_data.agent_id)->logical_node_id, - dispatch_data.queue_id.handle, - getpid(), - cnt_dispatch_data_v->thread_id, - magnitude(dispatch_data.grid_size), - kernel_info->formatted_kernel_name, - magnitude(dispatch_data.workgroup_size), - lds_block_size_v, - kernel_info->private_segment_size, - kernel_info->arch_vgpr_count, - kernel_info->sgpr_count, - itr.first, - itr.second); - - get_dereference(get_counter_collection_file()) << counter_collection_ss.str(); - } + buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION; + write_ring_buffer(&counter_record, buffer_type); delete cnt_dispatch_data_v; } @@ -1118,6 +1440,52 @@ list_metrics_iterate_agents(rocprofiler_agent_version_t, rocprofiler_client_finalize_t client_finalizer = nullptr; rocprofiler_client_id_t* client_identifier = nullptr; +timestamps_t* +get_app_timestamps() +{ + return stats_timestamp; +} + +void +init_tool_table() +{ + // agent and timestamp functions + tool_functions->tool_get_agent_node_id_fn = get_agent_node_id; + tool_functions->tool_get_app_timestamps_fn = get_app_timestamps; + + // name functions + tool_functions->tool_get_domain_name_fn = get_domain_name; + tool_functions->tool_get_kernel_name_fn = get_kernel_name; + tool_functions->tool_get_operation_name_fn = get_operation_name; + tool_functions->tool_get_counter_info_name_fn = get_counter_info_name; + tool_functions->tool_get_callback_kind_fn = get_callback_kind; + tool_functions->tool_get_callback_op_name_fn = get_callback_op_name; + tool_functions->tool_get_roctx_msg_fn = get_roctx_msg; + + // trace files + tool_functions->tool_get_agent_info_file_fn = get_agent_info_file; + tool_functions->tool_get_kernel_trace_file_fn = get_kernel_trace_file; + tool_functions->tool_get_hsa_api_trace_file_fn = get_hsa_api_trace_file; + tool_functions->tool_get_hip_api_trace_file_fn = get_hip_api_trace_file; + tool_functions->tool_get_memory_copy_trace_file_fn = get_memory_copy_file; + tool_functions->tool_get_counter_collection_file_fn = get_counter_file; + tool_functions->tool_get_marker_api_trace_file_fn = get_marker_api_trace_file; + tool_functions->tool_get_scratch_memory_file_fn = get_scratch_memory_trace_file; + + // stats files + tool_functions->tool_get_kernel_stats_file_fn = get_kernel_stats_file; + tool_functions->tool_get_hip_stats_file_fn = get_hip_stats_file; + tool_functions->tool_get_hsa_stats_file_fn = get_hsa_stats_file; + tool_functions->tool_get_memory_copy_stats_file_fn = get_memory_copy_stats_file; + tool_functions->tool_get_scratch_memory_stats_file_fn = get_scratch_memory_stats_file; +} + +void +fini_tool_table() +{ + *tool_functions = tool_table{}; +} + int tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) { @@ -1126,6 +1494,10 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) constexpr uint64_t buffer_size = 4096; constexpr uint64_t buffer_watermark = 4096; + rocprofiler_get_timestamp(&(stats_timestamp->app_start_time)); + + init_tool_table(); + ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "create context failed"); auto code_obj_ctx = rocprofiler_context_id_t{}; @@ -1338,15 +1710,90 @@ api_registration_callback(rocprofiler_intercept_table_t, "Iterate rocporfiler agents") } +namespace +{ +template +void +generate_output(buffer_type_t buffer_type) +{ + auto _data = std::vector{}; + flush_tmp_buffer(buffer_type); + read_tmp_file(buffer_type, _data); + + if(tool::get_config().output_format == "CSV") + rocprofiler::tool::generate_csv(tool_functions, _data); + + auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer(buffer_type); + _tmp_buf->destroy(); + delete _tmp_buf; + delete _tmp_file; +} +} // namespace + void tool_fini(void* tool_data) { client_identifier = nullptr; client_finalizer = nullptr; - flush(); - rocprofiler_stop_context(get_client_ctx()); + rocprofiler_get_timestamp(&(stats_timestamp->app_end_time)); + rocprofiler_stop_context(get_client_ctx()); + flush(); + + std::string_view output_format = tool::get_config().output_format; + + if(output_format == "CSV") + { + auto _agents = std::vector{}; + _agents.reserve(agent_info->size()); + for(auto& itr : *agent_info) + _agents.emplace_back(itr.second); + rocprofiler::tool::generate_csv(tool_functions, _agents); + } + + if(tool::get_config().kernel_trace) + { + generate_output( + buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH); + } + + if(tool::get_config().hsa_core_api_trace || tool::get_config().hsa_amd_ext_api_trace || + tool::get_config().hsa_image_ext_api_trace || tool::get_config().hsa_finalizer_ext_api_trace) + { + generate_output(buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA); + } + + if(tool::get_config().hip_runtime_api_trace || tool::get_config().hip_compiler_api_trace) + { + generate_output(buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP); + } + + if(tool::get_config().memory_copy_trace) + { + generate_output( + buffer_type_t::ROCPROFILER_TOOL_BUFFER_MEMORY_COPY); + } + + if(tool::get_config().marker_api_trace) + { + generate_output( + buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API); + } + + if(tool::get_config().counter_collection) + { + generate_output( + buffer_type_t::ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION); + } + + if(tool::get_config().scratch_memory) + { + generate_output( + buffer_type_t::ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY); + } + + fini_tool_table(); if(destructors) { for(const auto& itr : *destructors) @@ -1356,6 +1803,9 @@ tool_fini(void* tool_data) } (void) (tool_data); +#if defined(ROCPROFILER_CODECOV) && ROCPROFILER_CODECOV > 0 + __gcov_dump(); +#endif } } // namespace @@ -1384,9 +1834,18 @@ rocprofiler_configure(uint32_t version, uint32_t minor = (version % 10000) / 100; uint32_t patch = version % 100; + ::atexit([]() { + if(client_finalizer && client_identifier) client_finalizer(*client_identifier); + }); + // ensure these pointers are not leaked add_destructor(buffered_name_info); add_destructor(callback_name_info); + add_destructor(marker_msg_cid); + add_destructor(kernel_data); + add_destructor(tool_functions); + add_destructor(agent_info); + add_destructor(stats_timestamp); if(tool::get_config().list_metrics) { @@ -1403,7 +1862,7 @@ rocprofiler_configure(uint32_t version, for(size_t i = 0; i < num_agents; ++i) { auto* agent = static_cast(agents[i]); - agent_info->emplace(agent->id, agent); + agent_info->emplace(agent->id, *agent); } return ROCPROFILER_STATUS_SUCCESS; }, diff --git a/source/lib/rocprofiler-sdk/hip/hip.cpp b/source/lib/rocprofiler-sdk/hip/hip.cpp index 6f370ff25d..64fef33c6c 100644 --- a/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -253,7 +253,7 @@ hip_api_impl::functor(Args... args) if(callback_contexts.empty() && buffered_contexts.empty()) { - auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); if constexpr(!std::is_void::value) return _ret; else diff --git a/tests/rocprofv3/counter-collection/input2/validate.py b/tests/rocprofv3/counter-collection/input2/validate.py index 534412ff25..5d7390b313 100644 --- a/tests/rocprofv3/counter-collection/input2/validate.py +++ b/tests/rocprofv3/counter-collection/input2/validate.py @@ -31,12 +31,13 @@ def test_validate_counter_collection_pmc2(input_dir: pd.DataFrame): file_path = os.path.join(subdirectory_path, file_name) assert os.path.isfile(file_path), f"{file_path} is not a file." - with open(file_path, "r") as file: - df = pd.read_csv(file) - # check if kernel-name is present - assert len(df["Kernel_Name"]) > 0 - # check if counter value is positive - assert len(df["Counter_Value"]) > 0 + if "agent_info.csv" not in file_path: + with open(file_path, "r") as file: + df = pd.read_csv(file) + # check if kernel-name is present + assert len(df["Kernel_Name"]) > 0 + # check if counter value is positive + assert len(df["Counter_Value"]) > 0 if __name__ == "__main__": diff --git a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt index 0d850dd442..b6ea22ad4a 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt +++ b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt @@ -15,7 +15,7 @@ add_test( COMMAND $ --hip-runtime-trace --hip-compiler-trace --hsa-core-trace --hsa-amd-trace --hsa-image-trace - --hsa-finalizer-trace --kernel-trace --memory-copy-trace -d + --hsa-finalizer-trace --kernel-trace --memory-copy-trace --stats -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV @@ -51,13 +51,26 @@ add_test( --kernel-input ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv --memory-copy-input - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv + --agent-input + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv + --kernel-stats + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv + --hip-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv + --hsa-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv + --memory-copy-stats + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv) set(VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_api_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv) set_tests_properties( rocprofv3-test-trace-hip-in-libraries-validate diff --git a/tests/rocprofv3/tracing-hip-in-libraries/conftest.py b/tests/rocprofv3/tracing-hip-in-libraries/conftest.py index 6ca73a0aaf..4d44b2e484 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/conftest.py +++ b/tests/rocprofv3/tracing-hip-in-libraries/conftest.py @@ -6,6 +6,11 @@ import pytest def pytest_addoption(parser): + parser.addoption( + "--agent-input", + action="store", + help="Path to agent info CSV file.", + ) parser.addoption( "--hsa-input", action="store", @@ -31,6 +36,38 @@ def pytest_addoption(parser): action="store", help="Path to HIP runtime and compiler API tracing CSV file.", ) + parser.addoption( + "--hip-stats", + action="store", + help="Path to HIP stats CSV file.", + ) + parser.addoption( + "--hsa-stats", + action="store", + help="Path to HSA stats CSV file.", + ) + parser.addoption( + "--kernel-stats", + action="store", + help="Path to kernel stats CSV file.", + ) + parser.addoption( + "--memory-copy-stats", + action="store", + help="Path to memory copy stats CSV file.", + ) + + +@pytest.fixture +def agent_info_input_data(request): + filename = request.config.getoption("--agent-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data @pytest.fixture @@ -92,3 +129,55 @@ def hip_input_data(request): data.append(row) return data + + +@pytest.fixture +def hip_stats_data(request): + filename = request.config.getoption("--hip-stats") + data = [] + if os.path.exists(filename): + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def hsa_stats_data(request): + filename = request.config.getoption("--hsa-stats") + data = [] + if os.path.exists(filename): + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def kernel_stats_data(request): + filename = request.config.getoption("--kernel-stats") + data = [] + if os.path.exists(filename): + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def memory_copy_stats_data(request): + filename = request.config.getoption("--memory-copy-stats") + data = [] + if os.path.exists(filename): + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data diff --git a/tests/rocprofv3/tracing-hip-in-libraries/validate.py b/tests/rocprofv3/tracing-hip-in-libraries/validate.py index 4a66d42d3f..11aada4f03 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/validate.py +++ b/tests/rocprofv3/tracing-hip-in-libraries/validate.py @@ -3,7 +3,6 @@ import re import sys import pytest -import subprocess class dim3(object): @@ -16,7 +15,29 @@ class dim3(object): return (self.x, self.y, self.z) -def test_api_trace(hsa_input_data, hip_input_data): +def validate_average(row): + avg_ns = float(row["AverageNs"]) + tot_ns = int(row["TotalDurationNs"]) + num_cnt = float(row["Calls"]) + assert abs(avg_ns - (tot_ns / num_cnt)) < 1.0e-3, f"{row}" + + +def validate_stats(row): + min_v = int(row["MinNs"]) + max_v = int(row["MaxNs"]) + avg_v = float(row["AverageNs"]) + cnt_v = int(row["Calls"]) + stddev_v = float(row["StdDev"]) + + assert min_v > 0, f"{row}" + assert max_v > 0, f"{row}" + assert min_v < max_v if cnt_v > 1 else min_v == max_v, f"{row}" + assert min_v < avg_v if cnt_v > 1 else min_v == int(avg_v), f"{row}" + assert max_v > avg_v if cnt_v > 1 else max_v == int(avg_v), f"{row}" + assert stddev_v > 0.0 if cnt_v > 1 else int(stddev_v) == 0, f"{row}" + + +def test_api_trace(hsa_input_data, hip_input_data, hip_stats_data): functions = [] correlation_ids = [] for row in hsa_input_data: @@ -77,8 +98,15 @@ def test_api_trace(hsa_input_data, hip_input_data): ): assert itr in functions + for row in hip_stats_data: + assert int(row["TotalDurationNs"]) > 0 + assert int(row["Calls"]) > 0 + validate_average(row) + assert float(row["Percentage"]) > 0.0 + validate_stats(row) -def test_kernel_trace(kernel_input_data): + +def test_kernel_trace(kernel_input_data, kernel_stats_data): valid_kernel_names = sorted( [ "(anonymous namespace)::transpose(int const*, int*, int, int)", @@ -125,27 +153,42 @@ def test_kernel_trace(kernel_input_data): kernels = sorted(list(set(kernels))) assert kernels == valid_kernel_names + for row in kernel_stats_data: + assert int(row["TotalDurationNs"]) > 0 + assert int(row["Calls"]) > 0 + validate_average(row) + assert float(row["Percentage"]) > 0.0 + validate_stats(row) + + +def test_memory_copy_trace( + agent_info_input_data, + memory_copy_input_data, + hsa_input_data, + hsa_stats_data, + memory_copy_stats_data, +): + def get_agent(node_id): + for row in agent_info_input_data: + if row["Logical_Node_Id"] == node_id: + return row + return None -def test_memory_copy_trace(memory_copy_input_data, hsa_input_data): for row in memory_copy_input_data: assert row["Kind"] == "MEMORY_COPY" assert row["Direction"] in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + + src_agent = get_agent(row["Source_Agent_Id"]) + dst_agent = get_agent(row["Destination_Agent_Id"]) + assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}" + if row["Direction"] == "HOST_TO_DEVICE": - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Source_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("CPU")) >= 0 + assert src_agent["Agent_Type"] == "CPU" + assert dst_agent["Agent_Type"] == "GPU" elif row["Direction"] == "DEVICE_TO_HOST": - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Destination_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("CPU")) >= 0 + assert src_agent["Agent_Type"] == "GPU" + assert dst_agent["Agent_Type"] == "CPU" + assert int(row["Correlation_Id"]) > 0 assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) @@ -155,6 +198,20 @@ def test_memory_copy_trace(memory_copy_input_data, hsa_input_data): valid_length += 1 assert len(memory_copy_input_data) == valid_length + for row in hsa_stats_data: + assert int(row["TotalDurationNs"]) > 0 + assert int(row["Calls"]) > 0 + validate_average(row) + assert float(row["Percentage"]) > 0.0 + validate_stats(row) + + for row in memory_copy_stats_data: + assert int(row["TotalDurationNs"]) > 0 + assert int(row["Calls"]) > 0 + validate_average(row) + assert float(row["Percentage"]) > 0.0 + validate_stats(row) + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) diff --git a/tests/rocprofv3/tracing-plus-cc/validate.py b/tests/rocprofv3/tracing-plus-cc/validate.py index e35a8c7cd8..6ed9a73a8d 100644 --- a/tests/rocprofv3/tracing-plus-cc/validate.py +++ b/tests/rocprofv3/tracing-plus-cc/validate.py @@ -31,10 +31,11 @@ def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame): file_path = os.path.join(subdirectory_path, file_name) assert os.path.isfile(file_path), f"{file_path} is not a file." - with open(file_path, "r") as file: - df = pd.read_csv(file) - # check if either kernel-name/FUNCTION is present - assert "Kernel_Name" in df.columns or "Function" in df.columns + if "agent_info.csv" not in file_path: + with open(file_path, "r") as file: + df = pd.read_csv(file) + # check if either kernel-name/FUNCTION is present + assert "Kernel_Name" in df.columns or "Function" in df.columns if __name__ == "__main__": diff --git a/tests/rocprofv3/tracing/CMakeLists.txt b/tests/rocprofv3/tracing/CMakeLists.txt index 3e6d58e15a..08d72760fe 100644 --- a/tests/rocprofv3/tracing/CMakeLists.txt +++ b/tests/rocprofv3/tracing/CMakeLists.txt @@ -56,13 +56,16 @@ add_test( --memory-copy-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv --marker-input - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv + --agent-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv) set(VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_kernel_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv) set_tests_properties( rocprofv3-test-trace-validate @@ -111,13 +114,16 @@ add_test( --memory-copy-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_memory_copy_trace.csv --marker-input - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv + --agent-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv) set(SYS_VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_kernel_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv) set_tests_properties( rocprofv3-test-systrace-validate diff --git a/tests/rocprofv3/tracing/conftest.py b/tests/rocprofv3/tracing/conftest.py index e3f88e1884..64c9d2b0a3 100644 --- a/tests/rocprofv3/tracing/conftest.py +++ b/tests/rocprofv3/tracing/conftest.py @@ -5,6 +5,11 @@ import pytest def pytest_addoption(parser): + parser.addoption( + "--agent-input", + action="store", + help="Path to agent info CSV file.", + ) parser.addoption( "--hsa-input", action="store", @@ -32,6 +37,18 @@ def pytest_addoption(parser): ) +@pytest.fixture +def agent_info_input_data(request): + filename = request.config.getoption("--agent-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + @pytest.fixture def hsa_input_data(request): filename = request.config.getoption("--hsa-input") diff --git a/tests/rocprofv3/tracing/validate.py b/tests/rocprofv3/tracing/validate.py index cac07e4465..9451099c08 100644 --- a/tests/rocprofv3/tracing/validate.py +++ b/tests/rocprofv3/tracing/validate.py @@ -1,10 +1,27 @@ #!/usr/bin/env python3 import sys -import subprocess import pytest +def test_agent_info(agent_info_input_data): + logical_node_id = max([int(itr["Logical_Node_Id"]) for itr in agent_info_input_data]) + + assert logical_node_id + 1 == len(agent_info_input_data) + + for row in agent_info_input_data: + agent_type = row["Agent_Type"] + assert agent_type in ("CPU", "GPU") + if agent_type == "CPU": + assert int(row["Cpu_Cores_Count"]) > 0 + assert int(row["Simd_Count"]) == 0 + assert int(row["Max_Waves_Per_Simd"]) == 0 + else: + assert int(row["Cpu_Cores_Count"]) == 0 + assert int(row["Simd_Count"]) > 0 + assert int(row["Max_Waves_Per_Simd"]) > 0 + + def test_hsa_api_trace(hsa_input_data): functions = [] correlation_ids = [] @@ -60,49 +77,37 @@ def test_kernel_trace(kernel_input_data): assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) -def test_memory_copy_trace(memory_copy_input_data): +def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data): + + def get_agent(node_id): + for row in agent_info_input_data: + if row["Logical_Node_Id"] == node_id: + return row + return None + for row in memory_copy_input_data: assert row["Kind"] == "MEMORY_COPY" assert len(memory_copy_input_data) == 2 - row = memory_copy_input_data[0] - assert row["Direction"] == "HOST_TO_DEVICE" - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Source_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("CPU")) >= 0 - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Destination_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("GPU")) >= 0 - assert int(row["Correlation_Id"]) > 0 - assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + def test_row(idx, direction): + assert direction in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + row = memory_copy_input_data[idx] + assert row["Direction"] == direction + src_agent = get_agent(row["Source_Agent_Id"]) + dst_agent = get_agent(row["Destination_Agent_Id"]) + assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}" + if direction == "HOST_TO_DEVICE": + assert src_agent["Agent_Type"] == "CPU" + assert dst_agent["Agent_Type"] == "GPU" + else: + assert src_agent["Agent_Type"] == "GPU" + assert dst_agent["Agent_Type"] == "CPU" + assert int(row["Correlation_Id"]) > 0 + assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) - row = memory_copy_input_data[1] - assert row["Direction"] == "DEVICE_TO_HOST" - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Source_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("GPU")) >= 0 - output = subprocess.check_output( - 'rocminfo | grep "Node: *' - + row["Destination_Agent_Id"] - + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'', - shell=True, - ) - assert int(str(output).find("CPU")) >= 0 - assert int(row["Correlation_Id"]) > 0 - assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + test_row(0, "HOST_TO_DEVICE") + test_row(1, "DEVICE_TO_HOST") def test_marker_api_trace(marker_input_data):