From 0ba9f26a6a47eb37b33c4981e2c27153af4fce07 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Wed, 22 May 2024 00:53:42 -0500 Subject: [PATCH] Adding JSON support (#860) * Adding json support minor bugs Fixing tests Fixing formatting issues Fixing test test fix Misc testing fixes Use rocprofiler/cxx/name_info in rocprofiler-sdk-tool fixes to reduce the Json file size Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> misc fixes Removing int cast for JSON tests formatting removing a condition test on Navi3 adding debug info Misc fix * CSV updates - fix stats - numerical formatter support for customizing write_csv_entry - misc formatting - get_marker_stats_file * Misc tests/rocprofv3/counter-collection/input2 fixes - rocprofiler_configure_pytest_files in rocprofv3/counter-collection/input2 - removed state code from merge in rocprofv3/counter-collection/input2 * Tool: "Agent-id" -> "Agent_Id" - consistency * Tool update - remove rocprofiler_tool_marker_record_t - add marker_tracing_kind_conversion - fix memory leak in write_json - minor update to get_output_stream - rework handling of marker records * Update tests/pytest-packages/pytest_utils/__init__.py - add collapse_dict_list function for converting a dictionary value that is a list of length one into a directly mapped value * Update tests/rocprofv3/**/conftest.py - use collapse_dict_list when reading in JSONs * Update tests/rocprofv3/counter-collection/input1/validate.py - relax testing requirements gfx1102 (AQLProfile bugs) - in addition to relaxed testing requirements for gfx1101 * Update tests/rocprofv3/tracing/validate.py - fix removal of PID in every marker record * Update tests/rocprofv3/tracing-plus-cc - remove test design that relies on iterating subdirectories * Wrapper around __libc_start_main - Ensures finalization happens before main returns - Update tests/rocprofv3/tracing/validate.py - wrapper around __libc_start_main changed roctx calls * Combine include/rocprofiler-sdk/cxx/serialization.hpp and include/rocprofiler-sdk/external/serialization.hpp - tests/common/serialization.hpp simply includes include/rocprofiler-sdk/cxx/serialization.hpp now * Update lib/rocprofiler-sdk/hip/hip.cpp - tracing function immediately returns when fini_status is non-zero * Update lib/rocprofiler-sdk/hsa/hsa.cpp - remove logging of tracing function when fini_status is non-zero * Update lib/rocprofiler-sdk-tool/CMakeLists.txt - remove rocprofv3_trigger_list_metrics.cpp from TOOL_SOURCES * Update tests/rocprofv3/tracing-plus-cc/CMakeLists.txt - fix depends * Domain statistics * Update tests/rocprofv3/tracing-plus-cc/CMakeLists.txt - do not set ROCP_LOG_LEVEL in env * Remove erroneous include * Restructure tool source + reduce tool table + support multiple formats - buffered_output struct for handling output - support multiple output formats, e.g. --output-format csv,json - rename buffer_type_t -> domain_type - simplified generation of CSV output files - removed rocprofiler_tool_marker_record_t * Update lib/common/container/ring_buffer.hpp - value_type alias in ring_buffer * Remove all but one json-execute tests - generate CSV and JSON in same run * Fix include for domain_type.cpp * Update tests/rocprofv3/tracing-plus-cc/input.txt - only specify counters which can be found on gfx8, gfx9, gfx10, gfx11, etc. - use :device= syntax * Update lib/rocprofiler-sdk-tool/config.cpp - support :device=N syntax for counters file - improve stripping comments in PMC files - only read after pmc: * Rework tool library counter collection - fatal error if all requested counters for device are not found - support :device= syntax * Update tests/rocprofv3/tracing-plus-cc/input.txt - removed L2CacheHit (not supported on mi300) * Disable JSON tests in tests/rocprofv3 * Update include/rocprofiler-sdk/cxx/serialization.hpp - support rocprofiler_record_dimension_info_t * Update tool JSON schema - remove domain_type::CODE_OBJECT - rocprofiler_tool_agent_v0_t - rocprofiler_agent_v0_t + counters - rocprofiler_tool_counter_info_t - get_code_object_data() * Update JSON schema for tool * Update lib/rocprofiler-sdk-tool/tool.cpp - fix ROCP_WARNING_IF * rocprofv3 -> rocprofv3.sh - install rocprofv3.sh into sbin - configure_file /rocprofv3.sh -> /bin/rocprofv3 * Update tool counter collection - rocprofiler_tool_record_counter_t - rocprofiler_tool_counter_collection_record_t * Update tests/rocprofv3/counter-collection/input1/CMakeLists.txt - use rocprofiler_configure_pytest_files for validate.py, conftest.py, and input.txt * Update tests/rocprofv3/counter-collection/input1/validate.py - re-enable test_validate_counter_collection_pmc1_json * Update tests/rocprofv3/counter-collection/input2/validate.py - remove unused code * Update tests/rocprofv3/counter-collection/input2/validate.py - remove unused code * Update tests/rocprofv3/hsa-queue-dependency/validate.py - re-enable JSON tests * Misc tests/rocprofv3 CMake updates * Update tests/rocprofv3/tracing/validate.py - re-enable JSON tests * Update tests/rocprofv3/tracing-hip-in-libraries/validate.py - re-enable JSON tests * Update tests/rocprofv3/tracing/validate.py - remove unused node_exists function * Update tests/rocprofv3/tracing/validate.py - fix test_marker_api_trace_json --------- Co-authored-by: Sriraksha Nagaraj [ROCm/rocprofiler-sdk commit: 92b7326910d6614c887b190eeddcce1a1fb5a408] --- .../cmake/rocprofiler_interfaces.cmake | 1 + .../rocprofiler-sdk/external/CMakeLists.txt | 30 +- .../rocprofiler-sdk/source/bin/CMakeLists.txt | 16 +- .../source/bin/{rocprofv3 => rocprofv3.sh} | 9 + .../rocprofiler-sdk/cxx/serialization.hpp | 734 +++++++++ .../lib/common/container/ring_buffer.hpp | 3 +- .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 30 +- .../rocprofiler-sdk-tool/buffered_output.hpp | 115 ++ .../lib/rocprofiler-sdk-tool/config.cpp | 73 +- .../lib/rocprofiler-sdk-tool/config.hpp | 3 +- .../source/lib/rocprofiler-sdk-tool/csv.hpp | 48 +- .../lib/rocprofiler-sdk-tool/domain_type.cpp | 87 ++ .../lib/rocprofiler-sdk-tool/domain_type.hpp | 43 + .../lib/rocprofiler-sdk-tool/generateCSV.cpp | 754 +++++---- .../lib/rocprofiler-sdk-tool/generateCSV.hpp | 50 +- .../lib/rocprofiler-sdk-tool/generateJSON.cpp | 139 ++ .../lib/rocprofiler-sdk-tool/generateJSON.hpp | 45 + .../lib/rocprofiler-sdk-tool/helper.cpp | 128 +- .../lib/rocprofiler-sdk-tool/helper.hpp | 343 ++++- .../source/lib/rocprofiler-sdk-tool/main.c | 143 ++ .../lib/rocprofiler-sdk-tool/output_file.cpp | 24 +- .../lib/rocprofiler-sdk-tool/output_file.hpp | 18 +- .../rocprofiler-sdk-tool/tmp_file_buffer.cpp | 36 + .../rocprofiler-sdk-tool/tmp_file_buffer.hpp | 113 ++ .../source/lib/rocprofiler-sdk-tool/tool.cpp | 1350 +++++++---------- .../source/lib/rocprofiler-sdk/hip/hip.cpp | 9 +- .../source/lib/rocprofiler-sdk/hsa/hsa.cpp | 2 - .../tests/common/CMakeLists.txt | 9 +- .../tests/common/serialization.hpp | 725 +-------- .../pytest-packages/pytest_utils/__init__.py | 20 + .../counter-collection/CMakeLists.txt | 3 - .../counter-collection/input1/CMakeLists.txt | 19 +- .../counter-collection/input1/conftest.py | 31 + .../{ => input1}/pytest.ini | 4 +- .../counter-collection/input1/validate.py | 54 + .../counter-collection/input2/CMakeLists.txt | 12 +- .../{ => input2}/conftest.py | 27 +- .../counter-collection/input2/pytest.ini | 5 + .../counter-collection/input2/validate.py | 4 +- .../list_metrics/CMakeLists.txt | 5 +- .../list_metrics/conftest.py | 1 - .../list_metrics/pytest.ini | 5 + .../list_metrics/validate.py | 4 +- .../hsa-queue-dependency/CMakeLists.txt | 17 +- .../hsa-queue-dependency/conftest.py | 16 + .../rocprofv3/hsa-queue-dependency/pytest.ini | 5 + .../hsa-queue-dependency/validate.py | 89 ++ .../tracing-hip-in-libraries/CMakeLists.txt | 54 +- .../tracing-hip-in-libraries/conftest.py | 17 + .../tracing-hip-in-libraries/pytest.ini | 5 + .../tracing-hip-in-libraries/validate.py | 143 ++ .../rocprofv3/tracing-plus-cc/CMakeLists.txt | 45 +- .../rocprofv3/tracing-plus-cc/conftest.py | 181 ++- .../tests/rocprofv3/tracing-plus-cc/input.txt | 10 +- .../rocprofv3/tracing-plus-cc/validate.py | 45 +- .../tests/rocprofv3/tracing/CMakeLists.txt | 26 +- .../tests/rocprofv3/tracing/conftest.py | 16 + .../tests/rocprofv3/tracing/pytest.ini | 5 + .../tests/rocprofv3/tracing/validate.py | 138 +- .../tests/tools/CMakeLists.txt | 5 +- 60 files changed, 3900 insertions(+), 2191 deletions(-) rename projects/rocprofiler-sdk/source/bin/{rocprofv3 => rocprofv3.sh} (97%) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/main.c create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/conftest.py rename projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/{ => input1}/pytest.ini (52%) rename projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/{ => input2}/conftest.py (69%) create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/tracing/pytest.ini diff --git a/projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake b/projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake index 3f25c76a4b..3efe0eb895 100644 --- a/projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake +++ b/projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake @@ -18,6 +18,7 @@ rocprofiler_add_interface_library(rocprofiler-threading "Enables multithreading INTERNAL) rocprofiler_add_interface_library(rocprofiler-perfetto "Enables Perfetto support" INTERNAL) +rocprofiler_add_interface_library(rocprofiler-cereal "Enables Cereal support" INTERNAL) rocprofiler_add_interface_library(rocprofiler-compile-definitions "Compile definitions" INTERNAL) rocprofiler_add_interface_library(rocprofiler-static-libgcc diff --git a/projects/rocprofiler-sdk/external/CMakeLists.txt b/projects/rocprofiler-sdk/external/CMakeLists.txt index a7944339ed..c62a4d3a0e 100644 --- a/projects/rocprofiler-sdk/external/CMakeLists.txt +++ b/projects/rocprofiler-sdk/external/CMakeLists.txt @@ -62,22 +62,6 @@ if(ROCPROFILER_BUILD_TESTS) find_package(GTest REQUIRED) target_link_libraries(rocprofiler-gtest INTERFACE GTest::gtest) endif() - - # checkout submodule if not already checked out or clone repo if no .gitmodules file - rocprofiler_checkout_git_submodule( - RECURSIVE - RELATIVE_PATH external/cereal - WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} - REPO_URL https://github.com/jrmadsen/cereal.git - REPO_BRANCH "rocprofiler") - - add_library(rocprofiler-cereal INTERFACE) - add_library(rocprofiler::cereal ALIAS rocprofiler-cereal) - target_compile_definitions(rocprofiler-cereal - INTERFACE $) - target_include_directories( - rocprofiler-cereal - INTERFACE $) endif() if(ROCPROFILER_BUILD_GLOG) @@ -142,6 +126,20 @@ if(NOT TARGET PTL::ptl-static) add_subdirectory(ptl EXCLUDE_FROM_ALL) endif() +# checkout submodule if not already checked out or clone repo if no .gitmodules file +rocprofiler_checkout_git_submodule( + RECURSIVE + RELATIVE_PATH external/cereal + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + REPO_URL https://github.com/jrmadsen/cereal.git + REPO_BRANCH "rocprofiler") + +target_compile_definitions(rocprofiler-cereal + INTERFACE $) +target_include_directories( + rocprofiler-cereal + INTERFACE $) + # doxygen-awesome if(ROCPROFILER_BUILD_DOCS) rocprofiler_checkout_git_submodule( diff --git a/projects/rocprofiler-sdk/source/bin/CMakeLists.txt b/projects/rocprofiler-sdk/source/bin/CMakeLists.txt index 16d3988cd1..d88cb1a209 100644 --- a/projects/rocprofiler-sdk/source/bin/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/bin/CMakeLists.txt @@ -4,10 +4,22 @@ rocprofiler_activate_clang_tidy() -configure_file(rocprofv3 ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/rocprofv3 COPYONLY) +configure_file(rocprofv3.sh ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_SBINDIR}/rocprofv3.sh + COPYONLY) + +configure_file(rocprofv3.sh ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/rocprofv3 + COPYONLY) + install( - FILES rocprofv3 + FILES ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/rocprofv3 DESTINATION ${CMAKE_INSTALL_BINDIR} PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE COMPONENT tools) + +install( + FILES ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_SBINDIR}/rocprofv3.sh + DESTINATION ${CMAKE_INSTALL_SBINDIR} + PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ + WORLD_EXECUTE + COMPONENT tools) diff --git a/projects/rocprofiler-sdk/source/bin/rocprofv3 b/projects/rocprofiler-sdk/source/bin/rocprofv3.sh similarity index 97% rename from projects/rocprofiler-sdk/source/bin/rocprofv3 rename to projects/rocprofiler-sdk/source/bin/rocprofv3.sh index ebb0c935d7..692dcda182 100755 --- a/projects/rocprofiler-sdk/source/bin/rocprofv3 +++ b/projects/rocprofiler-sdk/source/bin/rocprofv3.sh @@ -55,6 +55,7 @@ usage() { echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d -o ${RESET}\n" echo -e "" echo -e "${GREEN}-d | --output-directory ${RESET} For adding output path where the output files will be saved" + echo -e "${GREEN} | --output-format ${RESET} For adding output format" echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d ${RESET}" echo -e "" echo -e "${GREEN}-M | --mangled-kernels ${RESET} Do not demangle the kernel names" @@ -142,6 +143,14 @@ while true; do fi shift shift + elif [[ "$1" == "--output-format" ]]; then + if [ "$2" ]; then + export ROCPROF_OUTPUT_FORMAT=$2 + else + usage 1 + fi + shift + shift elif [ "$1" == "--hsa-trace" ]; then export ROCPROF_HSA_CORE_API_TRACE=1 export ROCPROF_HSA_AMD_EXT_API_TRACE=1 diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp index 217938abb5..91985887b6 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp @@ -23,16 +23,745 @@ #pragma once +#include +#include +#include +#include +#include +#include #include +#include +#include +#include #include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include #include #include #include +#define ROCP_SDK_SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD)) +#define ROCP_SDK_SAVE_DATA_VALUE(NAME, VALUE) ar(make_nvp(NAME, data.VALUE)) +#define ROCP_SDK_SAVE_DATA_CSTR(FIELD) \ + ar(make_nvp(#FIELD, std::string{data.FIELD ? data.FIELD : ""})) +#define ROCP_SDK_SAVE_DATA_BITFIELD(NAME, VALUE) \ + { \ + auto _val = data.VALUE; \ + ar(make_nvp(NAME, _val)); \ + } + namespace cereal { +template +void +save(ArchiveT& ar, rocprofiler_context_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, hsa_agent_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_queue_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_correlation_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(internal); + ROCP_SDK_SAVE_DATA_VALUE("external", external.value); +} + +template +void +save(ArchiveT& ar, rocprofiler_dim3_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(x); + ROCP_SDK_SAVE_DATA_FIELD(y); + ROCP_SDK_SAVE_DATA_FIELD(z); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_load_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_FIELD(rocp_agent); + ROCP_SDK_SAVE_DATA_FIELD(hsa_agent); + ROCP_SDK_SAVE_DATA_CSTR(uri); + ROCP_SDK_SAVE_DATA_FIELD(load_base); + ROCP_SDK_SAVE_DATA_FIELD(load_size); + ROCP_SDK_SAVE_DATA_FIELD(load_delta); + ROCP_SDK_SAVE_DATA_FIELD(storage_type); + if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) + { + ROCP_SDK_SAVE_DATA_FIELD(storage_file); + } + else if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY) + { + ROCP_SDK_SAVE_DATA_FIELD(memory_base); + ROCP_SDK_SAVE_DATA_FIELD(memory_size); + } +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kernel_id); + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_CSTR(kernel_name); + ROCP_SDK_SAVE_DATA_FIELD(kernel_object); + ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_alignment); + ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); +} + +template +void +save(ArchiveT& ar, const hsa_queue_t& data) +{ + ar(make_nvp("queue_id", data.id)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_alloc_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_alloc_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(num_slots); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_free_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_marker_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(int64_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_marker_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_hip_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(hipError_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_hip_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(flags); + ROCP_SDK_SAVE_DATA_FIELD(args_kind); +} + +template +void +save(ArchiveT& ar, rocprofiler_kernel_dispatch_info_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(kernel_id); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_size); + ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(grid_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_memory_copy_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(num_records); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(context_id); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(phase); +} + +template +void +save_buffer_tracing_api_record(ArchiveT& ar, Tp data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_record_counter_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(counter_value); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_copy_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(bytes); +} + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(pid); + + switch(data.operation) + { + case ROCPROFILER_PAGE_MIGRATION_PAGE_FAULT: + { + ar(make_nvp("page_fault", data.page_fault)); + break; + } + case ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE: + { + ar(make_nvp("page_migrate", data.page_migrate)); + break; + } + case ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND: + { + ar(make_nvp("queue_suspend", data.queue_suspend)); + break; + } + case ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU: + { + ar(make_nvp("unmap_from_gpu", data.unmap_from_gpu)); + break; + } + case ROCPROFILER_PAGE_MIGRATION_NONE: + case ROCPROFILER_PAGE_MIGRATION_LAST: + { + throw std::runtime_error{"unsupported page migration operation type"}; + break; + } + } +} + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_page_fault_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(node_id); + ROCP_SDK_SAVE_DATA_FIELD(address); + ROCP_SDK_SAVE_DATA_FIELD(read_fault); + ROCP_SDK_SAVE_DATA_FIELD(migrated); +} + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_page_migrate_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(start_addr); + ROCP_SDK_SAVE_DATA_FIELD(end_addr); + ROCP_SDK_SAVE_DATA_FIELD(from_node); + ROCP_SDK_SAVE_DATA_FIELD(to_node); + ROCP_SDK_SAVE_DATA_FIELD(prefetch_node); + ROCP_SDK_SAVE_DATA_FIELD(preferred_node); + ROCP_SDK_SAVE_DATA_FIELD(trigger); +} + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_queue_suspend_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(node_id); + ROCP_SDK_SAVE_DATA_FIELD(trigger); + ROCP_SDK_SAVE_DATA_FIELD(rescheduled); +} + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_unmap_from_gpu_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(node_id); + ROCP_SDK_SAVE_DATA_FIELD(start_addr); + ROCP_SDK_SAVE_DATA_FIELD(end_addr); + ROCP_SDK_SAVE_DATA_FIELD(trigger); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_scratch_memory_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_correlation_id_retirement_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(internal_correlation_id); +} + +template +void +save(ArchiveT& ar, HsaCacheType data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("Data", ui32.Data); + ROCP_SDK_SAVE_DATA_BITFIELD("Instruction", ui32.Instruction); + ROCP_SDK_SAVE_DATA_BITFIELD("CPU", ui32.CPU); + ROCP_SDK_SAVE_DATA_BITFIELD("HSACU", ui32.HSACU); +} + +template +void +save(ArchiveT& ar, HSA_LINKPROPERTY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("Override", ui32.Override); + ROCP_SDK_SAVE_DATA_BITFIELD("NonCoherent", ui32.NonCoherent); + ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics32bit", ui32.NoAtomics32bit); + ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics64bit", ui32.NoAtomics64bit); + ROCP_SDK_SAVE_DATA_BITFIELD("NoPeerToPeerDMA", ui32.NoPeerToPeerDMA); +} + +template +void +save(ArchiveT& ar, HSA_CAPABILITY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); + ROCP_SDK_SAVE_DATA_BITFIELD("HSAMMUPresent", ui32.HSAMMUPresent); + ROCP_SDK_SAVE_DATA_BITFIELD("SharedWithGraphics", ui32.SharedWithGraphics); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueSizePowerOfTwo", ui32.QueueSizePowerOfTwo); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueSize32bit", ui32.QueueSize32bit); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueIdleEvent", ui32.QueueIdleEvent); + ROCP_SDK_SAVE_DATA_BITFIELD("VALimit", ui32.VALimit); + ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsSupported", ui32.WatchPointsSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsTotalBits", ui32.WatchPointsTotalBits); + ROCP_SDK_SAVE_DATA_BITFIELD("DoorbellType", ui32.DoorbellType); + ROCP_SDK_SAVE_DATA_BITFIELD("AQLQueueDoubleMap", ui32.AQLQueueDoubleMap); + ROCP_SDK_SAVE_DATA_BITFIELD("DebugTrapSupported", ui32.DebugTrapSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchTrapOverrideSupported", + ui32.WaveLaunchTrapOverrideSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchModeSupported", ui32.WaveLaunchModeSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("PreciseMemoryOperationsSupported", + ui32.PreciseMemoryOperationsSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("DEPRECATED_SRAM_EDCSupport", ui32.DEPRECATED_SRAM_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("Mem_EDCSupport", ui32.Mem_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("RASEventNotify", ui32.RASEventNotify); + ROCP_SDK_SAVE_DATA_BITFIELD("ASICRevision", ui32.ASICRevision); + ROCP_SDK_SAVE_DATA_BITFIELD("SRAM_EDCSupport", ui32.SRAM_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("SVMAPISupported", ui32.SVMAPISupported); + ROCP_SDK_SAVE_DATA_BITFIELD("CoherentHostAccess", ui32.CoherentHostAccess); + ROCP_SDK_SAVE_DATA_BITFIELD("DebugSupportedFirmware", ui32.DebugSupportedFirmware); + ROCP_SDK_SAVE_DATA_BITFIELD("Reserved", ui32.Reserved); +} + +template +void +save(ArchiveT& ar, HSA_MEMORYPROPERTY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); + ROCP_SDK_SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile); +} + +template +void +save(ArchiveT& ar, HSA_ENGINE_VERSION data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA); + ROCP_SDK_SAVE_DATA_BITFIELD("uCodeRes", uCodeRes); +} + +template +void +save(ArchiveT& ar, HSA_ENGINE_ID data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("uCode", ui32.uCode); + ROCP_SDK_SAVE_DATA_BITFIELD("Major", ui32.Major); + ROCP_SDK_SAVE_DATA_BITFIELD("Minor", ui32.Minor); + ROCP_SDK_SAVE_DATA_BITFIELD("Stepping", ui32.Stepping); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_cache_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(processor_id_low); + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(level); + ROCP_SDK_SAVE_DATA_FIELD(cache_line_size); + ROCP_SDK_SAVE_DATA_FIELD(cache_lines_per_tag); + ROCP_SDK_SAVE_DATA_FIELD(association); + ROCP_SDK_SAVE_DATA_FIELD(latency); + ROCP_SDK_SAVE_DATA_FIELD(type); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_io_link_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(type); + ROCP_SDK_SAVE_DATA_FIELD(version_major); + ROCP_SDK_SAVE_DATA_FIELD(version_minor); + ROCP_SDK_SAVE_DATA_FIELD(node_from); + ROCP_SDK_SAVE_DATA_FIELD(node_to); + ROCP_SDK_SAVE_DATA_FIELD(weight); + ROCP_SDK_SAVE_DATA_FIELD(min_latency); + ROCP_SDK_SAVE_DATA_FIELD(max_latency); + ROCP_SDK_SAVE_DATA_FIELD(min_bandwidth); + ROCP_SDK_SAVE_DATA_FIELD(max_bandwidth); + ROCP_SDK_SAVE_DATA_FIELD(recommended_transfer_size); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_mem_bank_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(heap_type); + ROCP_SDK_SAVE_DATA_FIELD(flags); + ROCP_SDK_SAVE_DATA_FIELD(width); + ROCP_SDK_SAVE_DATA_FIELD(mem_clk_max); + ROCP_SDK_SAVE_DATA_FIELD(size_in_bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_configuration_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(method); + ROCP_SDK_SAVE_DATA_FIELD(unit); + ROCP_SDK_SAVE_DATA_FIELD(min_interval); + ROCP_SDK_SAVE_DATA_FIELD(max_interval); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, const rocprofiler_agent_v0_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(type); + ROCP_SDK_SAVE_DATA_FIELD(cpu_cores_count); + ROCP_SDK_SAVE_DATA_FIELD(simd_count); + ROCP_SDK_SAVE_DATA_FIELD(mem_banks_count); + ROCP_SDK_SAVE_DATA_FIELD(caches_count); + ROCP_SDK_SAVE_DATA_FIELD(io_links_count); + ROCP_SDK_SAVE_DATA_FIELD(cpu_core_id_base); + ROCP_SDK_SAVE_DATA_FIELD(simd_id_base); + ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_simd); + ROCP_SDK_SAVE_DATA_FIELD(lds_size_in_kb); + ROCP_SDK_SAVE_DATA_FIELD(gds_size_in_kb); + ROCP_SDK_SAVE_DATA_FIELD(num_gws); + ROCP_SDK_SAVE_DATA_FIELD(wave_front_size); + ROCP_SDK_SAVE_DATA_FIELD(num_xcc); + ROCP_SDK_SAVE_DATA_FIELD(cu_count); + ROCP_SDK_SAVE_DATA_FIELD(array_count); + ROCP_SDK_SAVE_DATA_FIELD(num_shader_banks); + ROCP_SDK_SAVE_DATA_FIELD(simd_arrays_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(cu_per_simd_array); + ROCP_SDK_SAVE_DATA_FIELD(simd_per_cu); + ROCP_SDK_SAVE_DATA_FIELD(max_slots_scratch_cu); + ROCP_SDK_SAVE_DATA_FIELD(gfx_target_version); + ROCP_SDK_SAVE_DATA_FIELD(vendor_id); + ROCP_SDK_SAVE_DATA_FIELD(device_id); + ROCP_SDK_SAVE_DATA_FIELD(location_id); + ROCP_SDK_SAVE_DATA_FIELD(domain); + ROCP_SDK_SAVE_DATA_FIELD(drm_render_minor); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_engines); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_xgmi_engines); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_queues_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(num_cp_queues); + ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_ccompute); + ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_fcompute); + ROCP_SDK_SAVE_DATA_FIELD(sdma_fw_version); + ROCP_SDK_SAVE_DATA_FIELD(fw_version); + ROCP_SDK_SAVE_DATA_FIELD(capability); + ROCP_SDK_SAVE_DATA_FIELD(cu_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_cu); + ROCP_SDK_SAVE_DATA_FIELD(family_id); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_size); + ROCP_SDK_SAVE_DATA_FIELD(grid_max_size); + ROCP_SDK_SAVE_DATA_FIELD(local_mem_size); + ROCP_SDK_SAVE_DATA_FIELD(hive_id); + ROCP_SDK_SAVE_DATA_FIELD(gpu_id); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_dim); + ROCP_SDK_SAVE_DATA_FIELD(grid_max_dim); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(vendor_name); + ROCP_SDK_SAVE_DATA_CSTR(product_name); + ROCP_SDK_SAVE_DATA_CSTR(model_name); + ROCP_SDK_SAVE_DATA_FIELD(num_pc_sampling_configs); + ROCP_SDK_SAVE_DATA_FIELD(node_id); + ROCP_SDK_SAVE_DATA_FIELD(logical_node_id); + + auto generate = [&](auto name, const auto* value, uint64_t size) { + using value_type = std::remove_const_t>; + auto vec = std::vector{}; + vec.reserve(size); + for(uint64_t i = 0; i < size; ++i) + vec.emplace_back(value[i]); + ar(make_nvp(name, vec)); + }; + + generate("mem_banks", data.mem_banks, data.mem_banks_count); + generate("caches", data.caches, data.caches_count); + generate("io_links", data.io_links, data.io_links_count); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_info_v0_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); + ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(description); + ROCP_SDK_SAVE_DATA_CSTR(block); + ROCP_SDK_SAVE_DATA_CSTR(expression); +} + +template +void +save(ArchiveT& ar, rocprofiler_record_dimension_info_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(instance_size); + ROCP_SDK_SAVE_DATA_CSTR(name); +} + template void save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info& data) @@ -56,3 +785,8 @@ save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info_impl struct ring_buffer : private base::ring_buffer { - using base_type = base::ring_buffer; + using base_type = base::ring_buffer; + using value_type = Tp; static size_t get_items_per_page(); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index 0f75238a29..ed2c13621b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -4,10 +4,30 @@ 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) +set(TOOL_HEADERS + buffered_output.hpp + config.hpp + csv.hpp + domain_type.hpp + generateCSV.hpp + generateJSON.hpp + helper.hpp + output_file.hpp + statistics.hpp + tmp_file_buffer.hpp + tmp_file.hpp) + +set(TOOL_SOURCES + config.cpp + domain_type.cpp + generateCSV.cpp + generateJSON.cpp + helper.cpp + main.c + output_file.cpp + tmp_file_buffer.cpp + tmp_file.cpp + tool.cpp) add_library(rocprofiler-sdk-tool SHARED) target_sources(rocprofiler-sdk-tool PRIVATE ${TOOL_SOURCES} ${TOOL_HEADERS}) @@ -18,7 +38,7 @@ target_link_libraries( rocprofiler-sdk-tool PRIVATE rocprofiler::rocprofiler-shared-library rocprofiler::rocprofiler-headers rocprofiler::rocprofiler-build-flags rocprofiler::rocprofiler-memcheck - rocprofiler::rocprofiler-common-library) + rocprofiler::rocprofiler-common-library rocprofiler::rocprofiler-cereal) set_target_properties( rocprofiler-sdk-tool diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp new file mode 100644 index 0000000000..adbac73975 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp @@ -0,0 +1,115 @@ +// 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 "statistics.hpp" +#include "tmp_file_buffer.hpp" + +#include "lib/common/container/ring_buffer.hpp" +#include "lib/common/logging.hpp" + +#include + +namespace rocprofiler +{ +namespace tool +{ +using float_type = double; +using stats_data_t = statistics; + +template +struct buffered_output +{ + using ring_buffer_type = rocprofiler::common::container::ring_buffer; + static constexpr auto buffer_type_v = DomainT; + + explicit buffered_output(bool _enabled); + ~buffered_output() = default; + buffered_output(const buffered_output&) = delete; + buffered_output(buffered_output&&) noexcept = delete; + buffered_output& operator=(const buffered_output&) = default; + buffered_output& operator=(buffered_output&&) noexcept = default; + + void flush(); + void read(); + void clear(); + void destroy(); + + operator bool() const { return enabled; } + + std::deque element_data = {}; + stats_data_t stats = {}; + +private: + bool enabled = false; +}; + +template +buffered_output::buffered_output(bool _enabled) +: enabled{_enabled} +{} + +template +void +buffered_output::flush() +{ + if(!enabled) return; + + flush_tmp_buffer(buffer_type_v); +} + +template +void +buffered_output::read() +{ + if(!enabled) return; + + flush(); + + element_data = get_buffer_elements(read_tmp_file(buffer_type_v)); +} + +template +void +buffered_output::clear() +{ + if(!enabled) return; + + element_data.clear(); +} + +template +void +buffered_output::destroy() +{ + if(!enabled) return; + + clear(); + auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer(buffer_type_v); + _tmp_buf->destroy(); + delete _tmp_buf; + delete _tmp_file; +} +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp index cc10733a33..367e783e6d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp @@ -128,7 +128,7 @@ handle_special_chars(std::string& str) { // Iterate over the string and replace any special characters with a space. auto pos = std::string::npos; - while((pos = str.find_first_of("!@#$%&(),*+-./;<=>?@{}^`~|:")) != std::string::npos) + while((pos = str.find_first_of("!@#$%&(),*+-./;<>?@{}^`~|")) != std::string::npos) str.at(pos) = ' '; } @@ -186,21 +186,36 @@ parse_counters(std::string line) if(line.empty()) return counters; - // trim line for any white spaces + // strip the comment + if(auto pos = std::string::npos; (pos = line.find('#')) != std::string::npos) + line = line.substr(0, pos); + + // trim line for any white spaces after comment strip trim(line); - if(!(line[0] == '#' || line.find("pmc") == std::string::npos)) + // check to see if comment stripping + trim resulted in empty line + if(line.empty()) return counters; + + constexpr auto pmc_qualifier = std::string_view{"pmc:"}; + auto pos = std::string::npos; + + // should we handle an "pmc:" not being present? Seems like it should be a fatal error + if((pos = line.find(pmc_qualifier)) != std::string::npos) { + // strip out pmc qualifier + line = line.substr(pos + pmc_qualifier.length()); + handle_special_chars(line); - std::stringstream input_line(line); - std::string counter; - while(getline(input_line, counter, ' ')) + auto input_ss = std::stringstream{line}; + while(true) { - if(counter.substr(0, 3) != "pmc" && has_counter_format(counter)) - { + auto counter = std::string{}; + input_ss >> counter; + if(counter.empty()) + break; + else if(counter != pmc_qualifier && has_counter_format(counter)) counters.emplace(counter); - } } } @@ -227,7 +242,45 @@ get_mpi_rank() config::config() : kernel_names{parse_kernel_names(get_env("ROCPROF_KERNEL_NAMES", std::string{}))} , counters{parse_counters(get_env("ROCPROF_COUNTERS", std::string{}))} -{} +{ + auto output_format = get_env("ROCPROF_OUTPUT_FORMAT", "CSV"); + + for(auto& itr : output_format) + itr = toupper(itr); + + for(auto itr : {',', ';', ':'}) + { + auto pos = std::string::npos; + do + { + pos = output_format.find(itr); + if(pos != std::string::npos) output_format.replace(pos, 1, " "); + } while(pos != std::string::npos); + } + + auto entries = std::set{}; + auto parser = std::stringstream{output_format}; + + while(true) + { + auto _val = std::string{}; + parser >> _val; + if(!_val.empty()) + entries.emplace(_val); + else + break; + } + + csv_output = entries.count("CSV") > 0 || entries.empty(); + json_output = entries.count("JSON") > 0; + + const auto supported_formats = std::set{"CSV", "JSON"}; + for(const auto& itr : entries) + { + LOG_IF(FATAL, supported_formats.count(itr) == 0) + << "Unsupported output format type: " << itr; + } +} std::vector output_keys(std::string _tag) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp index 95678ca730..dc5cca7d1d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp @@ -73,11 +73,12 @@ struct config 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); + bool csv_output = false; + bool json_output = 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_format = get_env("ROCPROF_OUTPUT_FORMAT", "CSV"); std::string tmp_directory = get_env("ROCPROF_TMPDIR", output_path); std::vector kernel_names = {}; std::set counters = {}; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp index 986d7e9b75..b4472e22af 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp @@ -38,27 +38,36 @@ namespace tool { namespace csv { -template +struct numerical_formatter +{ + template + std::ostream& operator()(std::ostream& ofs, const Tp& _val) const + { + using value_type = common::mpl::unqualified_type_t; + + if constexpr(std::is_floating_point::value) + { + constexpr value_type one = 1; + if(_val >= one) + ofs << std::setprecision(6) << std::fixed; + else + ofs << std::setprecision(8) << std::scientific; + } + + return ofs; + } +}; + +template 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(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 << "\""; - } + if constexpr(common::mpl::is_string_type::value) ofs << "\""; + FmtT{}(ofs, _val) << _val; + if constexpr(common::mpl::is_string_type::value) ofs << "\""; }; (_write(Idx, std::get(_data)), ...); @@ -70,21 +79,22 @@ struct csv_encoder { static constexpr auto columns = NumCols; - template = 0> static auto write_row(std::ostream& ofs, Args&&... args) { - write_csv_entry( + write_csv_entry( ofs, std::make_tuple(std::forward(args)...), std::make_index_sequence{}); return csv_encoder{}; } - template + template static auto write_row(std::ostream& ofs, const std::array& arr) { static_assert(N == columns, "Error! too many/few args passed"); - write_csv_entry(ofs, arr, std::make_index_sequence{}); + write_csv_entry(ofs, arr, std::make_index_sequence{}); return csv_encoder{}; } }; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp new file mode 100644 index 0000000000..94750bb26f --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp @@ -0,0 +1,87 @@ +// 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 "domain_type.hpp" + +#include + +namespace +{ +template +struct domain_type_name; + +#define DEFINE_BUFFER_TYPE_NAME(ENUM_VALUE, COLUMN_NAME, FILENAME) \ + template <> \ + struct domain_type_name \ + { \ + static constexpr auto column_name = COLUMN_NAME; \ + static constexpr auto filename = FILENAME; \ + }; + +DEFINE_BUFFER_TYPE_NAME(HSA, "HSA_API", "hsa_trace") +DEFINE_BUFFER_TYPE_NAME(HIP, "HIP_API", "hip_trace") +DEFINE_BUFFER_TYPE_NAME(MEMORY_COPY, "MEMORY_COPY", "memory_copy") +DEFINE_BUFFER_TYPE_NAME(COUNTER_COLLECTION, "COUNTER_COLLECTION", "counter_collection") +DEFINE_BUFFER_TYPE_NAME(KERNEL_DISPATCH, "KERNEL_DISPATCH", "kernel_dispatch") +DEFINE_BUFFER_TYPE_NAME(MARKER, "MARKER", "marker_trace") +DEFINE_BUFFER_TYPE_NAME(SCRATCH_MEMORY, "SCRATCH_MEMORY", "scratch_memory") + +#undef DEFINE_BUFFER_TYPE_NAME + +template +std::string_view +get_domain_file_name(domain_type _buffer_type, std::index_sequence) +{ + if(static_cast(_buffer_type) == Idx) + return domain_type_name(Idx)>::filename; + if constexpr(sizeof...(TailIdx) > 0) + return get_domain_file_name(_buffer_type, std::index_sequence{}); + return std::string_view{}; +} + +template +std::string_view +get_domain_column_name(domain_type buffer_type, std::index_sequence) +{ + if(static_cast(buffer_type) == Idx) + return domain_type_name(Idx)>::column_name; + if constexpr(sizeof...(IdxTail) > 0) + return get_domain_column_name(buffer_type, std::index_sequence{}); + + return std::string_view{}; +} +} // namespace + +std::string_view +get_domain_file_name(domain_type _buffer_type) +{ + constexpr auto buffer_type_last_v = static_cast(domain_type::LAST); + + return get_domain_file_name(_buffer_type, std::make_index_sequence{}); +} + +std::string_view +get_domain_column_name(domain_type buffer_type) +{ + constexpr auto last_v = static_cast(domain_type::LAST); + return get_domain_column_name(buffer_type, std::make_index_sequence{}); +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp new file mode 100644 index 0000000000..8fa92dc403 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp @@ -0,0 +1,43 @@ +// 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 + +enum class domain_type +{ + HSA = 0, + HIP, + MEMORY_COPY, + COUNTER_COLLECTION, + KERNEL_DISPATCH, + MARKER, + SCRATCH_MEMORY, + LAST, +}; + +std::string_view +get_domain_file_name(domain_type val); + +std::string_view +get_domain_column_name(domain_type _buffer_type); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp index d26d1e4ffc..cd8da07a8a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp @@ -23,11 +23,14 @@ #include "generateCSV.hpp" #include "csv.hpp" #include "helper.hpp" +#include "lib/rocprofiler-sdk-tool/config.hpp" #include "statistics.hpp" #include #include +#include +#include #include namespace rocprofiler @@ -36,44 +39,171 @@ 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) +struct percentage { - 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); + float_type value = {}; - 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()); + friend std::ostream& operator<<(std::ostream& os, percentage val) { return (os << val.value); } +}; + +struct stats_formatter +{ + template + std::ostream& operator()(std::ostream& ofs, const Tp& _val) const + { + using value_type = common::mpl::unqualified_type_t; + + if constexpr(std::is_floating_point::value) + { + constexpr value_type one_hundredth = 1.0e-2; + if(_val > one_hundredth) + ofs << std::setprecision(6) << std::fixed; + else + ofs << std::setprecision(8) << std::scientific; + } + else if constexpr(std::is_same::value) + { + constexpr float_type one = 1.0; + constexpr float_type one_hundredth = 1.0e-2; + if(_val.value >= one) + ofs << std::setprecision(2) << std::fixed; + else if(_val.value > one_hundredth) + ofs << std::setprecision(4) << std::fixed; + else + ofs << std::setprecision(3) << std::scientific; + } + + return ofs; } - ofs << _ss.str(); +}; + +tool::output_file +get_stats_output_file(std::string name) +{ + return tool::output_file{std::move(name), + tool::csv::stats_csv_encoder{}, + { + "Name", + "Calls", + "TotalDurationNs", + "AverageNs", + "Percentage", + "MinNs", + "MaxNs", + "StdDev", + }}; +} + +stats_data_t +write_stats(output_file&& ofs, const stats_map_t& data_v) +{ + auto data = std::vector>{}; + auto _duration = stats_data_t{}; + for(const auto& [id, value] : data_v) + { + data.emplace_back(id, value); + _duration += value; + } + + std::sort(data.begin(), data.end(), [](const auto& lhs, const auto& rhs) { + return (lhs.second.get_sum() > rhs.second.get_sum()); + }); + + constexpr float_type one_hundred = 100.0; + + const float_type _total_duration = _duration.get_sum(); + for(const auto& [name, value] : data) + { + auto duration_ns = value.get_sum(); + auto calls = value.get_count(); + float_type avg_ns = value.get_mean(); + float_type percent_v = (duration_ns / _total_duration) * one_hundred; + + auto _row = std::stringstream{}; + rocprofiler::tool::csv::stats_csv_encoder::write_row(_row, + name, + calls, + duration_ns, + avg_ns, + percentage{percent_v}, + value.get_min(), + value.get_max(), + value.get_stddev()); + ofs << _row.str() << std::flush; + } + + return _duration; } } // namespace void -generate_csv(tool_table* tool_functions, std::vector& data) +generate_csv(tool_table* /*tool_functions*/, std::vector& data) { + if(data.empty()) return; + std::sort(data.begin(), data.end(), [](rocprofiler_agent_v0_t lhs, rocprofiler_agent_v0_t rhs) { return lhs.node_id < rhs.node_id; }); + auto ofs = 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"}}; + for(auto& itr : data) { auto _type = std::string_view{}; @@ -84,8 +214,8 @@ generate_csv(tool_table* tool_functions, std::vector& da else _type = "UNK"; - auto agent_info_ss = std::stringstream{}; - rocprofiler::tool::csv::agent_info_csv_encoder::write_row(agent_info_ss, + auto row_ss = std::stringstream{}; + rocprofiler::tool::csv::agent_info_csv_encoder::write_row(row_ss, itr.node_id, itr.logical_node_id, _type, @@ -139,309 +269,419 @@ generate_csv(tool_table* tool_functions, std::vector& da itr.vendor_name, itr.product_name, itr.model_name); - tool_functions->tool_get_agent_info_file_fn() << agent_info_ss.str(); + ofs << row_ss.str(); } } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { + if(data.empty()) return stats_data_t{}; + auto kernel_stats = stats_map_t{}; - for(auto& buf : data) + auto ofs = tool::output_file{"kernel_trace", + tool::csv::kernel_trace_csv_encoder{}, + {"Kind", + "Agent_Id", + "Queue_Id", + "Kernel_Id", + "Kernel_Name", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp", + "Private_Segment_Size", + "Group_Segment_Size", + "Workgroup_Size_X", + "Workgroup_Size_Y", + "Workgroup_Size_Z", + "Grid_Size_X", + "Grid_Size_Y", + "Grid_Size_Z"}}; + + for(const auto& record : 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->dispatch_info.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->dispatch_info.agent_id), - record->dispatch_info.queue_id.handle, - record->dispatch_info.kernel_id, - kernel_name, - record->correlation_id.internal, - record->start_timestamp, - record->end_timestamp, - record->dispatch_info.private_segment_size, - record->dispatch_info.group_segment_size, - record->dispatch_info.workgroup_size.x, - record->dispatch_info.workgroup_size.y, - record->dispatch_info.workgroup_size.z, - record->dispatch_info.grid_size.x, - record->dispatch_info.grid_size.y, - record->dispatch_info.grid_size.z); + auto row_ss = std::stringstream{}; + auto kernel_name = tool_functions->tool_get_kernel_name_fn(record.dispatch_info.kernel_id); + rocprofiler::tool::csv::kernel_trace_csv_encoder::write_row( + row_ss, + tool_functions->tool_get_domain_name_fn(record.kind), + tool_functions->tool_get_agent_node_id_fn(record.dispatch_info.agent_id), + record.dispatch_info.queue_id.handle, + record.dispatch_info.kernel_id, + kernel_name, + record.correlation_id.internal, + record.start_timestamp, + record.end_timestamp, + record.dispatch_info.private_segment_size, + record.dispatch_info.group_segment_size, + record.dispatch_info.workgroup_size.x, + record.dispatch_info.workgroup_size.y, + record.dispatch_info.workgroup_size.z, + record.dispatch_info.grid_size.x, + record.dispatch_info.grid_size.y, + record.dispatch_info.grid_size.z); - if(tool::get_config().stats) - kernel_stats[kernel_name] += (record->end_timestamp - record->start_timestamp); + 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(); - } + ofs << row_ss.str(); } + auto _duration = stats_data_t{}; if(tool::get_config().stats) - write_stats(tool_functions->tool_get_kernel_stats_file_fn(), - tool_functions->tool_get_app_timestamps_fn(), - kernel_stats); + _duration = write_stats(get_stats_output_file("kernel_stats"), kernel_stats); + + return _duration; } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { + if(data.empty()) return stats_data_t{}; + auto hip_stats = stats_map_t{}; - for(auto& buf : data) + auto ofs = tool::output_file{"hip_api_trace", + tool::csv::api_csv_encoder{}, + {"Domain", + "Function", + "Process_Id", + "Thread_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + for(const auto& record : 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); + auto row_ss = std::stringstream{}; + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + rocprofiler::tool::csv::api_csv_encoder::write_row( + row_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); + 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(); - } + ofs << row_ss.str(); } + auto _duration = stats_data_t{}; if(tool::get_config().stats) { - write_stats(tool_functions->tool_get_hip_stats_file_fn(), - tool_functions->tool_get_app_timestamps_fn(), - hip_stats); + _duration = write_stats(get_stats_output_file("hip_stats"), hip_stats); } + return _duration; } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { + if(data.empty()) return stats_data_t{}; + auto hsa_stats = stats_map_t{}; - for(auto& buf : data) + auto ofs = tool::output_file{"hsa_api_trace", + tool::csv::api_csv_encoder{}, + {"Domain", + "Function", + "Process_Id", + "Thread_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + + for(const auto& record : 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); + auto row_ss = std::stringstream{}; + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + rocprofiler::tool::csv::api_csv_encoder::write_row( + row_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); + 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(); - } + ofs << row_ss.str(); } + auto _duration = stats_data_t{}; if(tool::get_config().stats) { - write_stats(tool_functions->tool_get_hsa_stats_file_fn(), - tool_functions->tool_get_app_timestamps_fn(), - hsa_stats); + _duration = write_stats(get_stats_output_file("hsa_stats"), hsa_stats); } + return _duration; } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { + if(data.empty()) return stats_data_t{}; + auto memory_copy_stats = stats_map_t{}; - for(auto& buf : data) + auto ofs = tool::output_file{"memory_copy_trace", + tool::csv::memory_copy_csv_encoder{}, + {"Kind", + "Direction", + "Source_Agent_Id", + "Destination_Agent_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + for(const auto& record : 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); + auto row_ss = std::stringstream{}; + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + rocprofiler::tool::csv::memory_copy_csv_encoder::write_row( + row_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); + 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(); - } + ofs << row_ss.str(); } + auto _duration = stats_data_t{}; 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); + _duration = write_stats(get_stats_output_file("memory_copy_stats"), memory_copy_stats); } + return _duration; } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { - for(auto& buf : data) + if(data.empty()) return stats_data_t{}; + + auto marker_stats = stats_map_t{}; + auto ofs = tool::output_file{"marker_api_trace", + tool::csv::marker_csv_encoder{}, + {"Domain", + "Function", + "Process_Id", + "Thread_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + for(const auto& record : data) { - while(true) + auto row_ss = std::stringstream{}; + auto _name = std::string_view{}; + + if(record.kind == ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API && + (record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA || + record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePushA || + record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStartA)) { - 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); - } + _name = tool_functions->tool_get_roctx_msg_fn(record.correlation_id.internal); + } + else + { + _name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + } + + tool::csv::marker_csv_encoder::write_row( + row_ss, + tool_functions->tool_get_domain_name_fn(record.kind), + _name, + getpid(), + record.thread_id, + record.correlation_id.internal, + record.start_timestamp, + record.end_timestamp); + + if(tool::get_config().stats) + marker_stats[_name] += (record.end_timestamp - record.start_timestamp); + + ofs << row_ss.str(); + } + + auto _duration = stats_data_t{}; + if(tool::get_config().stats) + { + _duration = write_stats(get_stats_output_file("marker_stats"), marker_stats); + } + return _duration; +} + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) +{ + if(data.empty()) return stats_data_t{}; + + auto ofs = tool::output_file{"counter_collection", + tool::csv::counter_collection_csv_encoder{}, + {"Correlation_Id", + "Dispatch_Id", + "Agent_Id", + "Queue_Id", + "Process_Id", + "Thread_Id", + "Grid_Size", + "Kernel_Name", + "Workgroup_Size", + "LDS_Block_Size", + "Scratch_Size", + "VGPR_Count", + "SGPR_Count", + "Counter_Name", + "Counter_Value"}}; + for(const auto& record : data) + { + auto kernel_id = record.dispatch_data.dispatch_info.kernel_id; + auto counter_name_value = std::map{}; + for(const auto& count : record.records) + { + auto rec = count.record_counter; + 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 - { - 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(); + search->second = search->second + rec.counter_value; } - } -} -void -generate_csv(tool_table* tool_functions, std::vector& data) -{ - for(auto& buf : data) - { - while(true) + const auto& correlation_id = record.dispatch_data.correlation_id; + + auto magnitude = [](rocprofiler_dim3_t dims) { return (dims.x * dims.y * dims.z); }; + auto row_ss = std::stringstream{}; + for(auto& itr : counter_name_value) { - rocprofiler_tool_counter_collection_record_t* record = buf.retrieve(); - if(record == nullptr) break; - auto kernel_id = record->dispatch_data.dispatch_info.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.dispatch_info.agent_id), - record->dispatch_data.dispatch_info.queue_id.handle, - record->pid, - record->thread_id, - magnitude(record->dispatch_data.dispatch_info.grid_size), - tool_functions->tool_get_kernel_name_fn(kernel_id), - magnitude(record->dispatch_data.dispatch_info.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(); + tool::csv::counter_collection_csv_encoder::write_row( + row_ss, + correlation_id.internal, + record.dispatch_data.dispatch_info.dispatch_id, + tool_functions->tool_get_agent_node_id_fn( + record.dispatch_data.dispatch_info.agent_id), + record.dispatch_data.dispatch_info.queue_id.handle, + getpid(), + record.thread_id, + magnitude(record.dispatch_data.dispatch_info.grid_size), + tool_functions->tool_get_kernel_name_fn(kernel_id), + magnitude(record.dispatch_data.dispatch_info.workgroup_size), + record.lds_block_size_v, + record.dispatch_data.dispatch_info.private_segment_size, + record.arch_vgpr_count, + record.sgpr_count, + itr.first, + itr.second); } + ofs << row_ss.str(); } + return stats_data_t{}; } -void -generate_csv(tool_table* tool_functions, std::vector& data) +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data) { - auto* ofs = tool_functions->tool_get_scratch_memory_file_fn(); - if(!ofs) throw std::runtime_error{"error creating scratch memory output file"}; + if(data.empty()) return stats_data_t{}; + + auto ofs = tool::output_file{"scratch_memory_trace", + tool::csv::scratch_memory_encoder{}, + { + "Kind", + "Operation", + "Agent_Id", + "Queue_Id", + "Thread_Id", + "Alloc_flags", + "Start_Timestamp", + "End_Timestamp", + }}; auto scratch_memory_stats = stats_map_t{}; - for(auto& buf : data) + for(const auto& record : data) { - while(true) - { - rocprofiler_buffer_tracing_scratch_memory_record_t* record = buf.retrieve(); - if(record == nullptr) break; + auto row_ss = 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); - 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( + row_ss, + 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); - 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); - if(tool::get_config().stats) - scratch_memory_stats[op_name] += (record->end_timestamp - record->start_timestamp); - - (*ofs) << scratch_memory_trace.str(); - } + ofs << row_ss.str(); } + auto _duration = stats_data_t{}; 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); + _duration = + write_stats(get_stats_output_file("scratch_memory_stats"), scratch_memory_stats); + } + return _duration; +} + +void +generate_csv(tool_table* /*tool_functions*/, std::unordered_map& data) +{ + using csv_encoder_t = rocprofiler::tool::csv::stats_csv_encoder; + + if(!tool::get_config().stats) return; + + auto _total_stats = stats_data_t{}; + for(const auto& itr : data) + _total_stats += itr.second; + + if(_total_stats.get_count() == 0) return; + + auto ofs = get_stats_output_file("domain_stats"); + + constexpr float_type one_hundred = 100.0; + const float_type _total_duration = _total_stats.get_sum(); + for(const auto& [type, value] : data) + { + auto name = get_domain_column_name(type); + auto duration_ns = value.get_sum(); + auto calls = value.get_count(); + float_type avg_ns = value.get_mean(); + float_type percent_v = (duration_ns / _total_duration) * one_hundred; + + auto _row = std::stringstream{}; + csv_encoder_t::write_row(_row, + name, + calls, + duration_ns, + avg_ns, + percentage{percent_v}, + value.get_min(), + value.get_max(), + value.get_stddev()); + ofs << _row.str() << std::flush; } } } // namespace tool diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp index 9cf235fd42..2b6ef64b39 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp @@ -23,6 +23,7 @@ #pragma once #include "helper.hpp" +#include "statistics.hpp" #include @@ -30,28 +31,41 @@ namespace rocprofiler { namespace tool { +using float_type = double; +using stats_data_t = statistics; + void generate_csv(tool_table* tool_functions, std::vector& data); -void -generate_csv(tool_table* tool_functions, std::vector& data); +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& data); + +stats_data_t +generate_csv(tool_table* tool_functions, + const std::deque& 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); +generate_csv(tool_table* tool_functions, std::unordered_map& data); } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp new file mode 100644 index 0000000000..9c38ee20bf --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp @@ -0,0 +1,139 @@ +// 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 "generateJSON.hpp" +#include "config.hpp" +#include "helper.hpp" +#include "output_file.hpp" + +#include +#include + +#include + +namespace rocprofiler +{ +namespace tool +{ +void +write_json(tool_table* tool_functions, + uint64_t pid, + std::vector agent_data, + std::vector counter_data, + std::deque* hip_api_deque, + std::deque* hsa_api_deque, + std::deque* kernel_dispatch_deque, + std::deque* memory_copy_deque, + std::deque* counter_collection_deque, + std::deque* marker_api_deque, + std::deque* scratch_api_deque) + +{ + using JSONOutputArchive = cereal::MinimalJSONOutputArchive; + + constexpr auto json_prec = 32; + constexpr auto json_indent = JSONOutputArchive::Options::IndentChar::space; + auto json_opts = JSONOutputArchive::Options{json_prec, json_indent, 1}; + auto filename = std::string_view{"results"}; + auto [output_stream, cleanup] = get_output_stream(filename, ".json"); + + { + auto json_ar = JSONOutputArchive{*output_stream, json_opts}; + json_ar.setNextName("rocprofiler-sdk-tool"); + json_ar.startNode(); + + json_ar.makeArray(); + json_ar.startNode(); + + // metadata + { + json_ar.setNextName("metadata"); + json_ar.startNode(); + auto* timestamps = tool_functions->tool_get_app_timestamps_fn(); + json_ar(cereal::make_nvp("pid", pid)); + json_ar(cereal::make_nvp("init_time", timestamps->app_start_time)); + json_ar(cereal::make_nvp("fini_time", timestamps->app_end_time)); + json_ar.finishNode(); + } + + json_ar(cereal::make_nvp("agents", agent_data)); + json_ar(cereal::make_nvp("counters", counter_data)); + + { + auto callback_name_info = get_callback_id_names(); + auto buffer_name_info = get_buffer_id_names(); + auto counter_dims = get_tool_counter_dimension_info(); + auto marker_msg_data = get_callback_roctx_msg(); + + json_ar.setNextName("strings"); + json_ar.startNode(); + json_ar(cereal::make_nvp("callback_records", callback_name_info)); + json_ar(cereal::make_nvp("buffer_records", buffer_name_info)); + json_ar(cereal::make_nvp("marker_api", marker_msg_data)); + + { + json_ar.setNextName("counters"); + json_ar.startNode(); + json_ar(cereal::make_nvp("dimension_ids", counter_dims)); + json_ar.finishNode(); + } + + json_ar.finishNode(); + } + + { + auto kern_sym_data = get_kernel_symbol_data(); + auto code_obj_data = get_code_object_data(); + + json_ar(cereal::make_nvp("code_objects", code_obj_data)); + json_ar(cereal::make_nvp("kernel_symbols", kern_sym_data)); + } + + { + json_ar.setNextName("callback_records"); + json_ar.startNode(); + json_ar(cereal::make_nvp("counter_collection", *counter_collection_deque)); + json_ar.finishNode(); + } + + { + json_ar.setNextName("buffer_records"); + json_ar.startNode(); + json_ar(cereal::make_nvp("kernel_dispatches", *kernel_dispatch_deque)); + json_ar(cereal::make_nvp("hip_api", *hip_api_deque)); + json_ar(cereal::make_nvp("hsa_api", *hsa_api_deque)); + json_ar(cereal::make_nvp("memory_copy", *memory_copy_deque)); + json_ar(cereal::make_nvp("scratch_api", *scratch_api_deque)); + json_ar(cereal::make_nvp("marker_api", *marker_api_deque)); + json_ar.finishNode(); + } + + json_ar.finishNode(); // end array + json_ar.finishNode(); + } + + *output_stream << std::flush; + if(cleanup) cleanup(output_stream); +} + +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp new file mode 100644 index 0000000000..b36ac84450 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp @@ -0,0 +1,45 @@ +// 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" + +namespace rocprofiler +{ +namespace tool +{ +void +write_json(tool_table* tool_functions, + uint64_t pid, + std::vector agent_data, + std::vector counter_data, + std::deque* hip_api_deque, + std::deque* hsa_api_deque, + std::deque* kernel_dispatch_deque, + std::deque* memory_copy_deque, + std::deque* counter_collection_deque, + std::deque* marker_api_deque, + std::deque* scratch_api_deque); + +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp index ec816a0506..7e5b41593c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp @@ -24,6 +24,7 @@ #include "config.hpp" #include +#include #include #include @@ -33,133 +34,14 @@ #include #include -rocprofiler_tool_buffer_name_info_t +::rocprofiler::sdk::buffer_name_info_t get_buffer_id_names() { - static auto supported = std::unordered_set{ - ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, - ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API, - ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API, - ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API, - ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, - ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API, - ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API, - ROCPROFILER_BUFFER_TRACING_MARKER_CONTROL_API, - ROCPROFILER_BUFFER_TRACING_MARKER_NAME_API, - ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, - ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, - }; - - auto cb_name_info = rocprofiler_tool_buffer_name_info_t{}; - // - // callback for each kind operation - // - static auto tracing_kind_operation_cb = - [](rocprofiler_buffer_tracing_kind_t kindv, uint32_t operation, void* data_v) { - auto* name_info_v = static_cast(data_v); - - if(supported.count(kindv) > 0) - { - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_operation_name( - kindv, operation, &name, nullptr), - "query buffer failed"); - if(name) name_info_v->operation_names[kindv][operation] = name; - } - - return 0; - }; - - // - // callback for each kind (i.e. domain) - // - static auto tracing_kind_cb = [](rocprofiler_buffer_tracing_kind_t kind, void* data) { - // store the buffer kind name - auto* name_info_v = static_cast(data); - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_name(kind, &name, nullptr), - "query buffer failed"); - - if(name) name_info_v->kind_names[kind] = name; - - if(supported.count(kind) > 0) - { - ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kind_operations( - kind, tracing_kind_operation_cb, static_cast(data)), - "query buffer failed"); - } - - return 0; - }; - - ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kinds(tracing_kind_cb, - static_cast(&cb_name_info)), - "iterate_buffer failed"); - - return cb_name_info; + return ::rocprofiler::sdk::get_buffer_tracing_names(); } -rocprofiler_tool_callback_name_info_t +::rocprofiler::sdk::callback_name_info_t get_callback_id_names() { - static auto supported = std::unordered_set{ - ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API, - ROCPROFILER_CALLBACK_TRACING_HSA_AMD_EXT_API, - ROCPROFILER_CALLBACK_TRACING_HSA_IMAGE_EXT_API, - ROCPROFILER_CALLBACK_TRACING_HSA_FINALIZE_EXT_API, - ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API, - ROCPROFILER_CALLBACK_TRACING_HIP_COMPILER_API, - ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API, - ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API, - ROCPROFILER_CALLBACK_TRACING_MARKER_NAME_API, - ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, - }; - - auto cb_name_info = rocprofiler_tool_callback_name_info_t{}; - // - // callback for each kind operation - // - static auto tracing_kind_operation_cb = - [](rocprofiler_callback_tracing_kind_t kindv, uint32_t operation, void* data_v) { - auto* name_info_v = static_cast(data_v); - - if(supported.count(kindv) > 0) - { - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name( - kindv, operation, &name, nullptr), - "query callback failed"); - if(name) name_info_v->operation_names[kindv][operation] = name; - } - - return 0; - }; - - // - // callback for each kind (i.e. domain) - // - static auto tracing_kind_cb = [](rocprofiler_callback_tracing_kind_t kind, void* data) { - // store the callback kind name - auto* name_info_v = static_cast(data); - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_name(kind, &name, nullptr), - "query callback failed"); - - if(name) name_info_v->kind_names[kind] = name; - - if(supported.count(kind) > 0) - { - ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operations( - kind, tracing_kind_operation_cb, static_cast(data)), - "query callback failed"); - } - - return 0; - }; - - ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kinds(tracing_kind_cb, - static_cast(&cb_name_info)), - "iterate_callback failed"); - - return cb_name_info; + return ::rocprofiler::sdk::get_callback_tracing_names(); } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp index c793416184..ce6d3135f5 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp @@ -22,14 +22,21 @@ #pragma once +#include "domain_type.hpp" #include "lib/common/container/ring_buffer.hpp" #include "lib/common/container/small_vector.hpp" #include "lib/common/defines.hpp" +#include "lib/common/demangle.hpp" #include "lib/common/filesystem.hpp" #include "output_file.hpp" +#include +#include +#include #include #include +#include +#include #include #include @@ -39,6 +46,8 @@ #include #include +#include + #include #include #include @@ -76,63 +85,208 @@ constexpr size_t BUFFER_SIZE_BYTES = 4096; constexpr size_t WATERMARK = (BUFFER_SIZE_BYTES / 2); using rocprofiler_tool_buffer_kind_names_t = - std::unordered_map; + std::unordered_map; using rocprofiler_tool_buffer_kind_operation_names_t = std::unordered_map>; + std::unordered_map>; + +using marker_message_map_t = std::unordered_map; +using rocprofiler_kernel_symbol_data_t = + rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; namespace common = ::rocprofiler::common; namespace tool = ::rocprofiler::tool; -struct rocprofiler_tool_buffer_name_info_t +struct kernel_symbol_data : rocprofiler_kernel_symbol_data_t { - rocprofiler_tool_buffer_kind_names_t kind_names = {}; - rocprofiler_tool_buffer_kind_operation_names_t operation_names = {}; + using base_type = rocprofiler_kernel_symbol_data_t; + + kernel_symbol_data(const base_type& _base) + : base_type{_base} + , formatted_kernel_name{tool::format_name(CHECK_NOTNULL(_base.kernel_name))} + , demangled_kernel_name{common::cxx_demangle(CHECK_NOTNULL(_base.kernel_name))} + , truncated_kernel_name{common::truncate_name(demangled_kernel_name)} + {} + + kernel_symbol_data(); + ~kernel_symbol_data() = default; + kernel_symbol_data(const kernel_symbol_data&) = default; + kernel_symbol_data(kernel_symbol_data&&) noexcept = default; + kernel_symbol_data& operator=(const kernel_symbol_data&) = default; + kernel_symbol_data& operator=(kernel_symbol_data&&) noexcept = default; + + std::string formatted_kernel_name = {}; + std::string demangled_kernel_name = {}; + std::string truncated_kernel_name = {}; }; -using rocprofiler_tool_callback_kind_names_t = - std::unordered_map; -using rocprofiler_tool_callback_kind_operation_names_t = - std::unordered_map>; +inline kernel_symbol_data::kernel_symbol_data() +: base_type{0, 0, 0, "", 0, 0, 0, 0, 0, 0, 0, 0} +{} -struct rocprofiler_tool_callback_name_info_t +using kernel_symbol_data_map_t = std::unordered_map; + +struct rocprofiler_tool_counter_info_t : rocprofiler_counter_info_v0_t { - rocprofiler_tool_callback_kind_names_t kind_names = {}; - rocprofiler_tool_callback_kind_operation_names_t operation_names = {}; + using parent_type = rocprofiler_counter_info_v0_t; + using dimension_id_vec_t = std::vector; + using dimension_info_vec_t = std::vector; + + rocprofiler_tool_counter_info_t(rocprofiler_agent_id_t _agent_id, + parent_type _info, + dimension_id_vec_t&& _dim_ids, + dimension_info_vec_t&& _dim_info) + : parent_type{_info} + , agent_id{_agent_id} + , dimension_ids{std::move(_dim_ids)} + , dimension_info{std::move(_dim_info)} + {} + + ~rocprofiler_tool_counter_info_t() = default; + rocprofiler_tool_counter_info_t(const rocprofiler_tool_counter_info_t&) = default; + rocprofiler_tool_counter_info_t(rocprofiler_tool_counter_info_t&&) noexcept = default; + rocprofiler_tool_counter_info_t& operator=(const rocprofiler_tool_counter_info_t&) = default; + rocprofiler_tool_counter_info_t& operator=(rocprofiler_tool_counter_info_t&&) noexcept = + default; + + rocprofiler_agent_id_t agent_id = {}; + std::vector dimension_ids = {}; + std::vector dimension_info = {}; }; -rocprofiler_tool_buffer_name_info_t +rocprofiler::sdk::buffer_name_info_t get_buffer_id_names(); -rocprofiler_tool_callback_name_info_t +::rocprofiler::sdk::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; +std::map +get_callback_roctx_msg(); - rocprofiler_timestamp_t start_timestamp; - rocprofiler_timestamp_t end_timestamp; +std::vector +get_kernel_symbol_data(); + +std::vector +get_code_object_data(); + +std::vector +get_tool_counter_info(); + +std::vector +get_tool_counter_dimension_info(); + +enum tracing_marker_kind +{ + MARKER_API_CORE = 0, + MARKER_API_CONTROL, + MARKER_API_NAME, + MARKER_API_LAST, +}; + +template +struct marker_tracing_kind_conversion; + +#define MAP_TRACING_KIND_CONVERSION(COMMON, CALLBACK, BUFFERED) \ + template <> \ + struct marker_tracing_kind_conversion \ + { \ + static constexpr auto callback_value = CALLBACK; \ + static constexpr auto buffered_value = BUFFERED; \ + \ + bool operator==(rocprofiler_callback_tracing_kind_t val) const \ + { \ + return (callback_value == val); \ + } \ + \ + bool operator==(rocprofiler_buffer_tracing_kind_t val) const \ + { \ + return (buffered_value == val); \ + } \ + \ + auto convert(rocprofiler_callback_tracing_kind_t) const { return buffered_value; } \ + auto convert(rocprofiler_buffer_tracing_kind_t) const { return callback_value; } \ + }; + +MAP_TRACING_KIND_CONVERSION(MARKER_API_CORE, + ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API, + ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API) +MAP_TRACING_KIND_CONVERSION(MARKER_API_CONTROL, + ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API, + ROCPROFILER_BUFFER_TRACING_MARKER_CONTROL_API) +MAP_TRACING_KIND_CONVERSION(MARKER_API_NAME, + ROCPROFILER_CALLBACK_TRACING_MARKER_NAME_API, + ROCPROFILER_BUFFER_TRACING_MARKER_NAME_API) +MAP_TRACING_KIND_CONVERSION(MARKER_API_LAST, + ROCPROFILER_CALLBACK_TRACING_LAST, + ROCPROFILER_BUFFER_TRACING_LAST) + +template +auto +convert_marker_tracing_kind(TracingKindT val, std::index_sequence) +{ + if(marker_tracing_kind_conversion{} == val) + { + return marker_tracing_kind_conversion{}.convert(val); + } + + if constexpr(sizeof...(Tail) > 0) + return convert_marker_tracing_kind(val, std::index_sequence{}); + + return marker_tracing_kind_conversion{}.convert(val); +} + +template +auto +convert_marker_tracing_kind(TracingKindT val) +{ + return convert_marker_tracing_kind(val, std::make_index_sequence{}); +} + +struct rocprofiler_tool_dimension_pos_t +{ + uint64_t dimension_id; + size_t instance; + + template + void save(ArchiveT& ar) const + { + ar(cereal::make_nvp("dimension_id", dimension_id)); + ar(cereal::make_nvp("instance", instance)); + } +}; + +struct rocprofiler_tool_record_counter_t +{ + rocprofiler_counter_id_t counter_id = {}; + rocprofiler_record_counter_t record_counter = {}; + + template + void save(ArchiveT& ar) const + { + ar(cereal::make_nvp("counter_id", counter_id)); + ar(cereal::make_nvp("value", record_counter.counter_value)); + } }; 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; + rocprofiler_profile_counting_dispatch_data_t dispatch_data = {}; + std::vector records = {}; + uint64_t thread_id = 0; + uint64_t arch_vgpr_count = 0; + uint64_t sgpr_count = 0; + uint64_t lds_block_size_v = 0; + + template + void save(ArchiveT& ar) const + { + ar(cereal::make_nvp("dispatch_data", dispatch_data)); + ar(cereal::make_nvp("records", records)); + ar(cereal::make_nvp("thread_id", thread_id)); + ar(cereal::make_nvp("arch_vgpr_count", arch_vgpr_count)); + ar(cereal::make_nvp("sgpr_count", sgpr_count)); + ar(cereal::make_nvp("lds_block_size_v", lds_block_size_v)); + } }; struct timestamps_t @@ -141,22 +295,36 @@ struct timestamps_t 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; +namespace rocprofiler +{ +namespace tool +{ +template +struct buffered_output; +} +} // namespace rocprofiler + +using hip_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using hsa_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using kernel_dispatch_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using memory_copy_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using marker_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using counter_collection_buffered_output_t = + ::rocprofiler::tool::buffered_output; +using scratch_memory_buffered_output_t = + ::rocprofiler::tool::buffered_output; using tool_get_agent_node_id_fn_t = uint64_t (*)(rocprofiler_agent_id_t); using tool_get_app_timestamps_fn_t = timestamps_t* (*) (); @@ -169,19 +337,6 @@ using tool_get_callback_op_name_fn_t = std::string_view (*)(rocprofiler_callba 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 { @@ -197,19 +352,49 @@ struct tool_table 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; }; + +/// converts a container of ring buffers of element Tp into a single container of elements +template class ContainerT, typename... ParamsT> +ContainerT +get_buffer_elements(ContainerT, ParamsT...>&& data) +{ + auto ret = ContainerT{}; + for(auto& buf : data) + { + Tp* record = nullptr; + do + { + record = buf.retrieve(); + if(record) ret.emplace_back(*record); + } while(record != nullptr); + } + + return ret; +} + +namespace cereal +{ +#define SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD)) + +template +void +save(ArchiveT& ar, const kernel_symbol_data& data) +{ + cereal::save(ar, static_cast(data)); + SAVE_DATA_FIELD(formatted_kernel_name); + SAVE_DATA_FIELD(demangled_kernel_name); + SAVE_DATA_FIELD(truncated_kernel_name); +} + +template +void +save(ArchiveT& ar, const rocprofiler_tool_counter_info_t& data) +{ + SAVE_DATA_FIELD(agent_id); + cereal::save(ar, static_cast(data)); + SAVE_DATA_FIELD(dimension_ids); +} + +#undef SAVE_DATA_FIELD +} // namespace cereal diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/main.c b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/main.c new file mode 100644 index 0000000000..7e1434b83a --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/main.c @@ -0,0 +1,143 @@ +// 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. + +#define _GNU_SOURCE + +#define ROCPROFV3_PUBLIC_API __attribute__((visibility("default"))); +#define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal"))); + +#include +#include +#include +#include +#include +#include + +// +// local type definitions +// +typedef int (*main_func_t)(int, char**, char**); +typedef int (*start_main_t)(int (*)(int, char**, char**), + int, + char**, + int (*)(int, char**, char**), + void (*)(void), + void (*)(void), + void*); + +// +// local function declarations +// +int +rocprofv3_libc_start_main(int (*)(int, char**, char**), + int, + char**, + int (*)(int, char**, char**), + void (*)(void), + void (*)(void), + void*) ROCPROFV3_INTERNAL_API; + +int +__libc_start_main(int (*)(int, char**, char**), + int, + char**, + int (*)(int, char**, char**), + void (*)(void), + void (*)(void), + void*) ROCPROFV3_PUBLIC_API; + +// +// external function declarations +// +extern void +rocprofv3_set_main(main_func_t main_func) ROCPROFV3_INTERNAL_API; + +extern int +rocprofv3_main(int argc, char** argv, char** envp) ROCPROFV3_INTERNAL_API; + +int +rocprofv3_libc_start_main(int (*_main)(int, char**, char**), + int _argc, + char** _argv, + int (*_init)(int, char**, char**), + void (*_fini)(void), + void (*_rtld_fini)(void), + void* _stack_end) +{ + // prevent re-entry + static int _reentry = 0; + if(_reentry > 0) + { + fprintf(stderr, + "[%i][%s:%i] recursive call into %s\n", + getpid(), + basename(__FILE__), + __LINE__, + __FUNCTION__); + fflush(stderr); + return -1; + } + _reentry = 1; + + // Save the real main function address + rocprofv3_set_main(_main); + + // Find the real __libc_start_main + start_main_t next_main = (start_main_t) dlsym(RTLD_NEXT, "__libc_start_main"); + + if(next_main) + { + // call rocprofv3 main function wrapper + return next_main(rocprofv3_main, _argc, _argv, _init, _fini, _rtld_fini, _stack_end); + } + + // grab address __libc_start_main overload + start_main_t dflt_main = (start_main_t) dlsym(RTLD_DEFAULT, "__libc_start_main"); + + // get the address of this function (approximately) + void* this_func = __builtin_extract_return_addr(__builtin_return_address(0)); + + fprintf(stderr, + "[%s:%i][%s] Error! rocprofv3 could not find __libc_start_main! " + "__builtin_return_address(0)=%p, RTLD_DEFAULT=%p, RTLD_NEXT=%p\n", + basename(__FILE__), + __LINE__, + __FUNCTION__, + this_func, + (void*) dflt_main, + (void*) next_main); + fflush(stderr); + + return -1; +} + +int +__libc_start_main(int (*_main)(int, char**, char**), + int _argc, + char** _argv, + int (*_init)(int, char**, char**), + void (*_fini)(void), + void (*_rtld_fini)(void), + void* _stack_end) +{ + return rocprofv3_libc_start_main(_main, _argc, _argv, _init, _fini, _rtld_fini, _stack_end); +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp index 9e69339654..f54403d574 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp @@ -24,6 +24,7 @@ #include "config.hpp" #include "lib/common/logging.hpp" +#include #include namespace rocprofiler @@ -32,8 +33,8 @@ namespace tool { namespace fs = common::filesystem; -std::pair -get_output_stream(const std::string& fname, const std::string& ext) +std::pair +get_output_stream(std::string_view fname, std::string_view ext) { auto cfg_output_path = tool::format(tool::get_config().output_path); @@ -44,8 +45,14 @@ get_output_stream(const std::string& fname, const std::string& ext) else if(cfg_output_path.empty()) return {&std::clog, [](auto*&) {}}; - auto output_path = fs::path{cfg_output_path}; - auto output_file_name = tool::format(tool::get_config().output_file); + // add a period to provided file extension if necessary + constexpr auto period = std::string_view{"."}; + constexpr auto noperiod = std::string_view{}; + const auto _ext = + fmt::format("{}{}", (!ext.empty() && ext.find('.') != 0) ? period : noperiod, ext); + + auto output_path = fs::path{cfg_output_path}; + auto output_prefix = tool::format(tool::get_config().output_file); if(fs::exists(output_path) && !fs::is_directory(fs::status(output_path))) throw std::runtime_error{ @@ -53,11 +60,12 @@ get_output_stream(const std::string& fname, const std::string& ext) output_path.string())}; if(!fs::exists(output_path)) fs::create_directories(output_path); - auto output_file = tool::format(output_path / (output_file_name + "_" + fname + ext)); - auto* _ofs = new std::ofstream{output_file}; - if(!_ofs && !*_ofs) - throw std::runtime_error{fmt::format("Failed to open {} for output", output_file)}; + auto output_file = + tool::format(output_path / fmt::format("{}_{}{}", output_prefix, fname, _ext)); + auto* _ofs = new std::ofstream{output_file}; + + LOG_IF(FATAL, !_ofs && !*_ofs) << fmt::format("Failed to open {} for output", output_file); ROCP_ERROR << "Opened result file: " << output_file; return {_ofs, [](std::ostream*& v) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp index 04d8ddbe6b..fa34dbcb8e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp @@ -41,8 +41,10 @@ namespace rocprofiler { namespace tool { -std::pair -get_output_stream(const std::string& fname, const std::string& ext = ".csv"); +using output_stream_dtor_t = void (*)(std::ostream*&); + +std::pair +get_output_stream(std::string_view fname, std::string_view ext); struct output_file { @@ -66,12 +68,10 @@ struct output_file operator bool() const { return m_stream != nullptr; } private: - using stream_dtor_t = void (*)(std::ostream*&); - - const std::string m_name = {}; - std::mutex m_mutex = {}; - std::ostream* m_stream = nullptr; - stream_dtor_t m_dtor = [](std::ostream*&) {}; + const std::string m_name = {}; + std::mutex m_mutex = {}; + std::ostream* m_stream = nullptr; + output_stream_dtor_t m_dtor = [](std::ostream*&) {}; }; template @@ -80,7 +80,7 @@ output_file::output_file(std::string name, std::array&& header) : m_name{std::move(name)} { - std::tie(m_stream, m_dtor) = get_output_stream(m_name); + std::tie(m_stream, m_dtor) = get_output_stream(m_name, ".csv"); for(auto& itr : header) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.cpp new file mode 100644 index 0000000000..1fa0a7c1e8 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.cpp @@ -0,0 +1,36 @@ +// 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_buffer.hpp" + +#include + +#include + +std::string +compose_tmp_file_name(domain_type buffer_type) +{ + return rocprofiler::tool::format(fmt::format("{}/.rocprofv3/{}-{}.dat", + rocprofiler::tool::get_config().tmp_directory, + "%ppid%-%pid%", + get_domain_file_name(buffer_type))); +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp new file mode 100644 index 0000000000..d4b71f45d2 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp @@ -0,0 +1,113 @@ +// 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 "tmp_file.hpp" + +#include "lib/common/container/ring_buffer.hpp" +#include "lib/common/logging.hpp" +#include "lib/common/units.hpp" + +#include + +#include +#include +#include +#include +#include + +template +using ring_buffer_t = rocprofiler::common::container::ring_buffer; + +std::string +compose_tmp_file_name(domain_type buffer_type); + +template +std::tuple +get_tmp_file_buffer(domain_type 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(domain_type type) +{ + auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer(type); + auto _lk = std::lock_guard(_tmp_file->file_mutex); + [[maybe_unused]] static auto _success = _tmp_file->open(); + auto& _fs = _tmp_file->stream; + _tmp_file->file_pos.emplace(_fs.tellg()); + _tmp_buf->save(_fs); + _tmp_buf->clear(); + CHECK(_tmp_buf->is_empty() == true); +} + +template +void +write_ring_buffer(Tp _v, domain_type type) +{ + auto [_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); + CHECK(ptr != nullptr); + } + *ptr = std::move(_v); +} + +template +void +flush_tmp_buffer(domain_type type) +{ + auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer(type); + if(!_tmp_buf->is_empty()) offload_buffer(type); +} + +template +std::deque +read_tmp_file(domain_type type) +{ + auto _data = std::deque{}; + + auto [_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)); + } + + return _data; +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 69215e9283..807023c18c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -20,14 +20,16 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#include "buffered_output.hpp" #include "config.hpp" #include "csv.hpp" +#include "domain_type.hpp" #include "generateCSV.hpp" +#include "generateJSON.hpp" #include "helper.hpp" #include "output_file.hpp" #include "tmp_file.hpp" -#include "lib/common/demangle.hpp" #include "lib/common/environment.hpp" #include "lib/common/filesystem.hpp" #include "lib/common/logging.hpp" @@ -41,6 +43,7 @@ #include #include #include +#include #include #include @@ -105,369 +108,13 @@ 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() -{ - static auto* _v = new tool::output_file{"hsa_api_trace", - tool::csv::api_csv_encoder{}, - {"Domain", - "Function", - "Process_Id", - "Thread_Id", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp"}}; - ADD_DESTRUCTOR(_v); - 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() -{ - static auto* _v = new tool::output_file{"hip_api_trace", - tool::csv::api_csv_encoder{}, - {"Domain", - "Function", - "Process_Id", - "Thread_Id", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp"}}; - ADD_DESTRUCTOR(_v); - return _v; -} - -tool::output_file& -get_hip_api_trace_file() -{ - return get_dereference(get_hip_api_file()); -} - -tool::output_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{}, - {"Kind", - "Agent_Id", - "Queue_Id", - "Kernel_Id", - "Kernel_Name", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp", - "Private_Segment_Size", - "Group_Segment_Size", - "Workgroup_Size_X", - "Workgroup_Size_Y", - "Workgroup_Size_Z", - "Grid_Size_X", - "Grid_Size_Y", - "Grid_Size_Z"}}; - ADD_DESTRUCTOR(_v); - return _v; -} - -tool::output_file& -get_kernel_trace_file() -{ - return get_dereference(get_kernel_file()); -} - -tool::output_file*& -get_counter_collection_file() -{ - static auto* _v = new tool::output_file{"counter_collection", - tool::csv::counter_collection_csv_encoder{}, - {"Correlation_Id", - "Dispatch_Id", - "Agent_Id", - "Queue_Id", - "Process_Id", - "Thread_Id", - "Grid_Size", - "Kernel_Name", - "Workgroup_Size", - "LDS_Block_Size", - "Scratch_Size", - "VGPR_Count", - "SGPR_Count", - "Counter_Name", - "Counter_Value"}}; - - ADD_DESTRUCTOR(_v); - return _v; -} - -tool::output_file& -get_counter_file() -{ - return get_dereference(get_counter_collection_file()); -} - -tool::output_file*& -get_memory_copy_trace_file() -{ - static auto* _v = new tool::output_file{"memory_copy_trace", - tool::csv::memory_copy_csv_encoder{}, - {"Kind", - "Direction", - "Source_Agent_Id", - "Destination_Agent_Id", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp"}}; - ADD_DESTRUCTOR(_v); - return _v; -} - -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", - tool::csv::scratch_memory_encoder{}, - { - "Kind", - "Operation", - "Agent_Id", - "Queue_Id", - "Thread_Id", - "Alloc_flags", - "Start_Timestamp", - "End_Timestamp", - }}; - ADD_DESTRUCTOR(_v); - return _v; -} - -tool::output_file*& -get_scratch_memory_stats_file() -{ - 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; -} - tool::output_file*& get_list_basic_metrics_file() { static auto* _v = new tool::output_file{"basic_metrics", tool::csv::list_basic_metrics_csv_encoder{}, - {"Agent-id", "Name", "Description", "Block", "Dimensions"}}; + {"Agent_Id", "Name", "Description", "Block", "Dimensions"}}; ADD_DESTRUCTOR(_v); return _v; } @@ -478,22 +125,13 @@ get_list_derived_metrics_file() static auto* _v = new tool::output_file{"derived_metrics", tool::csv::list_derived_metrics_csv_encoder{}, - {"Agent-id", "Name", "Description", "Expression", "Dimensions"}}; + {"Agent_Id", "Name", "Description", "Expression", "Dimensions"}}; ADD_DESTRUCTOR(_v); return _v; } #undef ADD_DESTRUCTOR -struct marker_entry -{ - uint64_t cid = 0; - pid_t pid = getpid(); - pid_t tid = rocprofiler::common::get_tid(); - rocprofiler_user_data_t data = {}; - std::string message = {}; -}; - struct buffer_ids { rocprofiler_buffer_id_t hsa_api_trace = {}; @@ -522,24 +160,6 @@ get_buffers() } using rocprofiler_code_object_data_t = rocprofiler_callback_tracing_code_object_load_data_t; -using rocprofiler_kernel_symbol_data_t = - rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; - -struct kernel_symbol_data : rocprofiler_kernel_symbol_data_t -{ - using base_type = rocprofiler_kernel_symbol_data_t; - - kernel_symbol_data(const base_type& _base) - : base_type{_base} - , formatted_kernel_name{tool::format_name(CHECK_NOTNULL(_base.kernel_name))} - , demangled_kernel_name{common::cxx_demangle(CHECK_NOTNULL(_base.kernel_name))} - , truncated_kernel_name{common::truncate_name(demangled_kernel_name)} - {} - - std::string formatted_kernel_name = {}; - std::string demangled_kernel_name = {}; - std::string truncated_kernel_name = {}; -}; template Tp* @@ -548,22 +168,24 @@ as_pointer(Tp&& _val) return new Tp{std::forward(_val)}; } -using code_object_data_map_t = std::unordered_map; -using kernel_symbol_data_map_t = std::unordered_map; -using targeted_kernels_set_t = std::unordered_set; +template +Tp* +as_pointer() +{ + return new Tp{}; +} + +using code_object_data_map_t = std::unordered_map; +using targeted_kernels_set_t = std::unordered_set; using counter_dimension_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_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 code_obj_data = as_pointer>(); +auto* kernel_data = as_pointer>(); +auto* marker_msg_data = as_pointer>(); 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{}); @@ -612,127 +234,25 @@ flush() ROCP_INFO << "Buffers flushed"; } -std::string -get_file_name(buffer_type_t buffer_type) -{ - switch(buffer_type) - { - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA: return "hsa_trace"; break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP: return "hip_trace"; break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_MARKER_API: return "marker_trace"; break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_MEMORY_COPY: return "memory_copy"; break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION: - return "counter_collection"; - break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH: - return "kernel_dispatch"; - break; - case buffer_type_t::ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY: return "scratch_memory"; break; - } - - ROCP_FATAL << "buffer type " << static_cast>(buffer_type) - << " not supported"; - return std::string{}; -} - -std::string -compose_tmp_file_name(buffer_type_t buffer_type) -{ - return rocprofiler::tool::format(fmt::format("{}/.rocprofv3/{}-{}.dat", - rocprofiler::tool::get_config().tmp_directory, - "%ppid%-%pid%", - get_file_name(buffer_type))); -} - -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); - [[maybe_unused]] static auto _success = _tmp_file->open(); - auto& _fs = _tmp_file->stream; - _tmp_file->file_pos.emplace(_fs.tellg()); - _tmp_buf->save(_fs); - _tmp_buf->clear(); - CHECK(_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); - CHECK(ptr != nullptr); - } - *ptr = std::move(*_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); + return CHECK_NOTNULL(callback_name_info)->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); + return CHECK_NOTNULL(callback_name_info)->at(kind, op); } std::string_view get_roctx_msg(uint64_t cid) { - return CHECK_NOTNULL(marker_msg_cid)->at(cid); + return CHECK_NOTNULL(marker_msg_data) + ->rlock( + [](const auto& _data, uint64_t _cid_v) -> std::string_view { return _data.at(_cid_v); }, + cid); } void @@ -764,18 +284,15 @@ cntrl_tracing_callback(rocprofiler_callback_tracing_record_t record, } else { - 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; + auto marker_record = rocprofiler_buffer_tracing_marker_api_record_t{}; + marker_record.size = sizeof(rocprofiler_buffer_tracing_marker_api_record_t); + marker_record.kind = convert_marker_tracing_kind(record.kind); + marker_record.operation = record.operation; + marker_record.thread_id = record.thread_id; + marker_record.correlation_id = record.correlation_id; 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); + write_ring_buffer(marker_record, domain_type::MARKER); } } } @@ -785,9 +302,10 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* user_data, void* data) { - static thread_local auto stacked_range = std::vector{}; - static auto global_range = - common::Synchronized>{}; + static thread_local auto stacked_range = + std::vector{}; + static auto global_range = common::Synchronized< + std::unordered_map>{}; if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API) { @@ -801,23 +319,23 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, { if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) { - { - 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; + CHECK_NOTNULL(marker_msg_data) + ->wlock( + [](auto& _data, uint64_t _cid_v, std::string&& _msg) { + _data.emplace(_cid_v, std::move(_msg)); + }, + record.correlation_id.internal, + std::string{marker_data->args.roctxMarkA.message}); + + auto marker_record = rocprofiler_buffer_tracing_marker_api_record_t{}; + marker_record.size = sizeof(rocprofiler_buffer_tracing_marker_api_record_t); + marker_record.kind = convert_marker_tracing_kind(record.kind); + marker_record.operation = record.operation; + marker_record.thread_id = record.thread_id; + marker_record.correlation_id = record.correlation_id; 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); + write_ring_buffer(marker_record, domain_type::MARKER); } } else if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePushA) @@ -826,10 +344,24 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, { if(marker_data->args.roctxRangePushA.message) { - auto& val = stacked_range.emplace_back(); - val.message = marker_data->args.roctxRangePushA.message; - val.data.value = ts; - val.cid = record.correlation_id.internal; + CHECK_NOTNULL(marker_msg_data) + ->wlock( + [](auto& _data, uint64_t _cid_v, std::string&& _msg) { + _data.emplace(_cid_v, std::move(_msg)); + }, + record.correlation_id.internal, + std::string{marker_data->args.roctxRangePushA.message}); + + auto marker_record = rocprofiler_buffer_tracing_marker_api_record_t{}; + marker_record.size = sizeof(rocprofiler_buffer_tracing_marker_api_record_t); + marker_record.kind = convert_marker_tracing_kind(record.kind); + marker_record.operation = record.operation; + marker_record.thread_id = record.thread_id; + marker_record.correlation_id = record.correlation_id; + marker_record.start_timestamp = ts; + marker_record.end_timestamp = 0; + + stacked_range.emplace_back(marker_record); } } } @@ -843,23 +375,9 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, auto val = stacked_range.back(); stacked_range.pop_back(); - { - 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); + + val.end_timestamp = ts; + write_ring_buffer(val, domain_type::MARKER); } } else if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStartA) @@ -867,18 +385,30 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT && marker_data->args.roctxRangeStartA.message) { - auto _id = marker_data->retval.roctx_range_id_t_retval; - auto _entry = marker_entry{}; - _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); - } + CHECK_NOTNULL(marker_msg_data) + ->wlock( + [](auto& _data, uint64_t _cid_v, std::string&& _msg) { + _data.emplace(_cid_v, std::move(_msg)); + }, + record.correlation_id.internal, + std::string{marker_data->args.roctxRangeStartA.message}); + + auto marker_record = rocprofiler_buffer_tracing_marker_api_record_t{}; + marker_record.size = sizeof(rocprofiler_buffer_tracing_marker_api_record_t); + marker_record.kind = convert_marker_tracing_kind(record.kind); + marker_record.operation = record.operation; + marker_record.thread_id = record.thread_id; + marker_record.correlation_id = record.correlation_id; + marker_record.start_timestamp = ts; + marker_record.end_timestamp = 0; + + auto _id = marker_data->retval.roctx_range_id_t_retval; global_range.wlock( - [_id, &_entry](auto& map) { map.emplace(_id, std::move(_entry)); }); + [](auto& map, roctx_range_id_t _range_id, auto&& _record) { + map.emplace(_range_id, std::move(_record)); + }, + _id, + marker_record); } } else if(record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStop) @@ -888,18 +418,9 @@ 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); - 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); + + _entry.end_timestamp = ts; + write_ring_buffer(_entry, domain_type::MARKER); global_range.wlock([](auto& map, auto _key) { return map.erase(_key); }, _id); } } @@ -911,18 +432,15 @@ callback_tracing_callback(rocprofiler_callback_tracing_record_t record, } else { - 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; + auto marker_record = rocprofiler_buffer_tracing_marker_api_record_t{}; + marker_record.size = sizeof(rocprofiler_buffer_tracing_marker_api_record_t); + marker_record.kind = convert_marker_tracing_kind(record.kind); + marker_record.operation = record.operation; + marker_record.thread_id = record.thread_id; + marker_record.correlation_id = record.correlation_id; 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); + write_ring_buffer(marker_record, domain_type::MARKER); } } } @@ -937,20 +455,20 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* user_data, void* data) { + auto ts = rocprofiler_timestamp_t{}; + ROCPROFILER_CALL(rocprofiler_get_timestamp(&ts), "get timestamp"); if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT && record.operation == ROCPROFILER_CODE_OBJECT_LOAD) { if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD) { auto* obj_data = static_cast(record.payload); - if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD) - { - code_obj_data.wlock( - [](code_object_data_map_t& cdata, rocprofiler_code_object_data_t* obj_data_v) { - cdata.emplace(obj_data_v->code_object_id, *obj_data_v); - }, - CHECK_NOTNULL(obj_data)); - } + + code_obj_data->wlock( + [](code_object_data_map_t& cdata, rocprofiler_code_object_data_t* obj_data_v) { + cdata.emplace(obj_data_v->code_object_id, *obj_data_v); + }, + CHECK_NOTNULL(obj_data)); } else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD) { @@ -964,13 +482,11 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, auto* sym_data = static_cast(record.payload); if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD) { - 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)}); - } + auto itr = kernel_data->wlock([sym_data](auto& _data) { + return _data.emplace(sym_data->kernel_id, + kernel_symbol_data{get_dereference(sym_data)}); + }); + ROCP_WARNING_IF(!itr.second) << "duplicate kernel symbol data for kernel_id=" << sym_data->kernel_id; @@ -1019,13 +535,15 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, std::string_view get_kernel_name(uint64_t kernel_id) { - return CHECK_NOTNULL(kernel_data)->at(kernel_id).formatted_kernel_name; + return CHECK_NOTNULL(kernel_data)->rlock([kernel_id](const auto& _data) -> std::string_view { + return _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); + return CHECK_NOTNULL(buffered_name_info)->at(record_kind); } uint64_t @@ -1037,7 +555,7 @@ get_agent_node_id(rocprofiler_agent_id_t agent_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); + return CHECK_NOTNULL(buffered_name_info)->at(kind, op); } void @@ -1050,10 +568,6 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, { ROCP_INFO << "Executing buffered tracing callback for " << num_headers << " headers"; - ROCP_ERROR_IF(headers == nullptr) - << "rocprofiler invoked a buffer callback with a null pointer to the array of headers. " - "this should never happen"; - if(!headers) return; for(size_t i = 0; i < num_headers; ++i) @@ -1067,11 +581,7 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, auto* record = static_cast( header->payload); - buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH; - - write_ring_buffer(record, - buffer_type); + write_ring_buffer(*record, domain_type::KERNEL_DISPATCH); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API || @@ -1082,38 +592,29 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, auto* record = static_cast(header->payload); - buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA; - write_ring_buffer( - record, buffer_type); + write_ring_buffer(*record, domain_type::HSA); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) { auto* record = static_cast(header->payload); - buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_MEMORY_COPY; - write_ring_buffer(record, - buffer_type); + write_ring_buffer(*record, domain_type::MEMORY_COPY); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY) { auto* record = static_cast( header->payload); - buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY; - write_ring_buffer(record, - buffer_type); + write_ring_buffer(*record, domain_type::SCRATCH_MEMORY); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API || header->kind == ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API) { auto* record = static_cast(header->payload); - buffer_type_t buffer_type = buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP; - write_ring_buffer( - record, buffer_type); + + write_ring_buffer(*record, domain_type::HIP); } else { @@ -1138,27 +639,154 @@ dimensions_info_callback(rocprofiler_counter_id_t id, { auto* dimensions_info = static_cast*>(user_data); + dimensions_info->reserve(num_dims); for(size_t j = 0; j < num_dims; j++) - dimensions_info->push_back(dim_info[j]); + dimensions_info->emplace_back(dim_info[j]); } else { counter_dimension_data.wlock( [&id, &dim_info, &num_dims](counter_dimension_info_map_t& counter_dimension_data_v) { - std::vector dimensions; - for(size_t dim = 0; dim < num_dims; dim++) - dimensions.emplace_back(dim_info[dim]); - counter_dimension_data_v.emplace(std::make_pair(id.handle, dimensions)); + if(counter_dimension_data_v.find(id.handle) == counter_dimension_data_v.end()) + { + auto dimensions = std::vector{}; + dimensions.reserve(num_dims); + for(size_t dim = 0; dim < num_dims; ++dim) + dimensions.emplace_back(dim_info[dim]); + counter_dimension_data_v.emplace(id.handle, std::move(dimensions)); + } }); } return ROCPROFILER_STATUS_SUCCESS; } +struct tool_agent +{ + int64_t device_id = 0; + const rocprofiler_agent_v0_t* agent = nullptr; +}; + +using tool_agent_vec_t = std::vector; + +auto +get_gpu_agents() +{ + auto _gpu_agents = tool_agent_vec_t{}; + + ROCPROFILER_CALL( + rocprofiler_query_available_agents( + ROCPROFILER_AGENT_INFO_VERSION_0, + [](rocprofiler_agent_version_t, const void** agents, size_t num_agents, void* _data) { + auto* _gpu_agents_v = static_cast(_data); + for(size_t i = 0; i < num_agents; ++i) + { + auto* agent = static_cast(agents[i]); + if(agent->type == ROCPROFILER_AGENT_TYPE_GPU) + _gpu_agents_v->emplace_back(tool_agent{0, agent}); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + sizeof(rocprofiler_agent_t), + &_gpu_agents), + "Iterate rocporfiler agents") + + // make sure they are sorted by node id + std::sort(_gpu_agents.begin(), _gpu_agents.end(), [](const auto& lhs, const auto& rhs) { + return CHECK_NOTNULL(lhs.agent)->node_id < CHECK_NOTNULL(rhs.agent)->node_id; + }); + + int64_t _dev_id = 0; + for(auto& itr : _gpu_agents) + itr.device_id = _dev_id++; + + return _gpu_agents; +} + +auto +get_agent_counter_info(const tool_agent_vec_t& _agents) +{ + using value_type = + std::unordered_map>; + + auto _data = value_type{}; + + for(auto itr : _agents) + { + ROCPROFILER_CALL( + rocprofiler_iterate_agent_supported_counters( + itr.agent->id, + [](rocprofiler_agent_id_t id, + rocprofiler_counter_id_t* counters, + size_t num_counters, + void* user_data) { + auto* data_v = static_cast(user_data); + for(size_t i = 0; i < num_counters; ++i) + { + // populate global map + ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions( + counters[i], dimensions_info_callback, nullptr), + "iterate_dimension_info"); + + auto _info = rocprofiler_counter_info_v0_t{}; + auto _dim_ids = std::vector{}; + auto _dim_info = std::vector{}; + + ROCPROFILER_CALL( + rocprofiler_query_counter_info( + counters[i], ROCPROFILER_COUNTER_INFO_VERSION_0, &_info), + "Could not query counter_id"); + + // populate local vector + ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions( + counters[i], dimensions_info_callback, &_dim_info), + "iterate_dimension_info"); + + _dim_ids.reserve(_dim_info.size()); + for(auto ditr : _dim_info) + _dim_ids.emplace_back(ditr.id); + + (*data_v)[id].emplace_back( + id, _info, std::move(_dim_ids), std::move(_dim_info)); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + &_data), + "iterate agent supported counters"); + + std::sort(_data.at(itr.agent->id).begin(), + _data.at(itr.agent->id).end(), + [](const auto& lhs, const auto& rhs) { return (lhs.id.handle < rhs.id.handle); }); + + for(auto& citr : _data.at(itr.agent->id)) + { + std::sort(citr.dimension_ids.begin(), citr.dimension_ids.end()); + std::sort(citr.dimension_info.begin(), + citr.dimension_info.end(), + [](const auto& lhs, const auto& rhs) { return (lhs.id < rhs.id); }); + } + } + + return _data; +} + +const tool_agent* +get_tool_agent(rocprofiler_agent_id_t id, const tool_agent_vec_t& data) +{ + for(const auto& itr : data) + { + if(id == itr.agent->id) return &itr; + } + + return nullptr; +} + // this function creates a rocprofiler profile config on the first entry auto get_agent_profile(rocprofiler_agent_id_t agent_id) { - static auto data = common::Synchronized{}; + static auto data = common::Synchronized{}; + static const auto gpu_agents = get_gpu_agents(); + static const auto gpu_agents_counter_info = get_agent_counter_info(gpu_agents); auto profile = std::optional{}; data.ulock( @@ -1172,36 +800,61 @@ get_agent_profile(rocprofiler_agent_id_t agent_id) return false; }, [agent_id, &profile](agent_counter_map_t& data_v) { - auto counters_v = counter_vec_t{}; - ROCPROFILER_CALL( - rocprofiler_iterate_agent_supported_counters( - agent_id, - [](rocprofiler_agent_id_t, - rocprofiler_counter_id_t* counters, - size_t num_counters, - void* user_data) { - auto* vec = static_cast(user_data); - for(size_t i = 0; i < num_counters; i++) - { - ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions( - counters[i], dimensions_info_callback, nullptr), - "iterate_dimension_info"); + auto counters_v = counter_vec_t{}; + auto found_v = std::vector{}; + const auto* tool_agent_v = get_tool_agent(agent_id, gpu_agents); + auto expected_v = tool::get_config().counters.size(); - rocprofiler_counter_info_v0_t info; + constexpr auto device_qualifier = std::string_view{":device="}; + for(const auto& itr : tool::get_config().counters) + { + auto name_v = itr; + if(auto pos = std::string::npos; + (pos = itr.find(device_qualifier)) != std::string::npos) + { + name_v = itr.substr(0, pos); + auto dev_id_s = itr.substr(pos + device_qualifier.length()); - ROCPROFILER_CALL( - rocprofiler_query_counter_info(counters[i], - ROCPROFILER_COUNTER_INFO_VERSION_0, - static_cast(&info)), - "Could not query counter_id"); + LOG_IF(FATAL, + dev_id_s.empty() || + dev_id_s.find_first_not_of("0123456789") != std::string::npos) + << "invalid device qualifier format (':device=N) where N is the GPU id: " + << itr; - if(tool::get_config().counters.count(info.name) > 0) - vec->emplace_back(counters[i]); - } - return ROCPROFILER_STATUS_SUCCESS; - }, - static_cast(&counters_v)), - "iterate agent supported counters"); + auto dev_id_v = std::stol(dev_id_s); + // skip this counter if the counter is for a specific device id (which doesn't + // this agent's device id) + if(dev_id_v != tool_agent_v->device_id) + { + --expected_v; // is not expected + continue; + } + } + + // search the gpu agent counter info for a counter with a matching name + for(const auto& citr : gpu_agents_counter_info.at(agent_id)) + { + if(name_v == std::string_view{citr.name}) + { + counters_v.emplace_back(citr.id); + found_v.emplace_back(itr); + } + } + } + + if(expected_v != counters_v.size()) + { + auto requested_counters = fmt::format("{}", + fmt::join(tool::get_config().counters.begin(), + tool::get_config().counters.end(), + ", ")); + auto found_counters = + fmt::format("{}", fmt::join(found_v.begin(), found_v.end(), ", ")); + LOG(FATAL) << "Unable to find all counters for agent " + << tool_agent_v->agent->node_id << " (gpu-" << tool_agent_v->device_id + << ", " << tool_agent_v->agent->name << ") in [" << requested_counters + << "]. Found: [" << found_counters << "]"; + } if(!counters_v.empty()) { @@ -1219,12 +872,6 @@ get_agent_profile(rocprofiler_agent_id_t agent_id) return profile; } -struct counter_dispatch_data -{ - uint64_t thread_id = 0; - uint64_t dispatch_index = 0; -}; - void dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, rocprofiler_profile_config_id_t* config, @@ -1240,9 +887,8 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, } else if(auto profile = get_agent_profile(agent_id)) { - *config = *profile; - user_data->ptr = new counter_dispatch_data{.thread_id = common::get_tid(), - .dispatch_index = ++dispatch_index}; + *config = *profile; + user_data->value = common::get_tid(); } } @@ -1268,28 +914,24 @@ counter_record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_da rocprofiler_user_data_t user_data, void* /*callback_data_args*/) { - auto kernel_id = dispatch_data.dispatch_info.kernel_id; - const auto* cnt_dispatch_data_v = static_cast(user_data.ptr); + static const auto gpu_agents = get_gpu_agents(); + static const auto gpu_agents_counter_info = get_agent_counter_info(gpu_agents); - 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; + auto counter_record = rocprofiler_tool_counter_collection_record_t{}; + auto kernel_id = dispatch_data.dispatch_info.kernel_id; - { - std::lock_guard lk(kernel_data_mutex); - kernel_info = &(CHECK_NOTNULL(kernel_data)->at(kernel_id)); - } + counter_record.dispatch_data = dispatch_data; + counter_record.thread_id = user_data.value; + + const kernel_symbol_data* kernel_info = + kernel_data->rlock([kernel_id](const auto& _data) { return &_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; + 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; ROCP_FATAL_IF(!kernel_info) << "missing kernel information for kernel_id=" << kernel_id; @@ -1297,14 +939,15 @@ counter_record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_da << " (name=" << kernel_info->kernel_name << ")"; for(size_t count = 0; count < record_count; count++) - counter_record.profiler_record.push_back( - static_cast(record_data[count])); + { + auto _counter_id = rocprofiler_counter_id_t{}; + ROCPROFILER_CALL(rocprofiler_query_record_counter_id(record_data[count].id, &_counter_id), + "query record counter id"); + counter_record.records.emplace_back( + rocprofiler_tool_record_counter_t{_counter_id, record_data[count]}); + } - 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; + write_ring_buffer(counter_record, domain_type::COUNTER_COLLECTION); } rocprofiler_status_t @@ -1421,6 +1064,32 @@ list_metrics_iterate_agents(rocprofiler_agent_version_t, rocprofiler_client_finalize_t client_finalizer = nullptr; rocprofiler_client_id_t* client_identifier = nullptr; +void +initialize_rocprofv3() +{ + if(int status = 0; + rocprofiler_is_initialized(&status) == ROCPROFILER_STATUS_SUCCESS && status == 0) + { + ROCPROFILER_CALL(rocprofiler_force_configure(&rocprofiler_configure), + "force configuration"); + } + + LOG_IF(FATAL, !client_identifier) << "nullptr to client identifier!"; + LOG_IF(FATAL, !client_finalizer && !tool::get_config().list_metrics) + << "nullptr to client finalizer!"; // exception for listing metrics +} + +void +finalize_rocprofv3() +{ + if(client_finalizer && client_identifier) + { + client_finalizer(*client_identifier); + client_finalizer = nullptr; + client_identifier = nullptr; + } +} + timestamps_t* get_app_timestamps() { @@ -1442,23 +1111,6 @@ init_tool_table() 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 @@ -1691,25 +1343,23 @@ api_registration_callback(rocprofiler_intercept_table_t, "Iterate rocporfiler agents") } -namespace -{ -template +using stats_data_t = ::rocprofiler::tool::stats_data_t; + +template void -generate_output(buffer_type_t buffer_type) +generate_output(rocprofiler::tool::buffered_output& output_v, + std::unordered_map& contributions_v) { - auto _data = std::vector{}; - flush_tmp_buffer(buffer_type); - read_tmp_file(buffer_type, _data); + if(!output_v) return; - if(tool::get_config().output_format == "CSV") - rocprofiler::tool::generate_csv(tool_functions, _data); + output_v.read(); - auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer(buffer_type); - _tmp_buf->destroy(); - delete _tmp_buf; - delete _tmp_file; + if(tool::get_config().csv_output) + { + output_v.stats = rocprofiler::tool::generate_csv(tool_functions, output_v.element_data); + contributions_v.emplace(output_v.buffer_type_v, output_v.stats); + } } -} // namespace void tool_fini(void* /*tool_data*/) @@ -1722,57 +1372,76 @@ tool_fini(void* /*tool_data*/) rocprofiler_stop_context(get_client_ctx()); flush(); - std::string_view output_format = tool::get_config().output_format; + auto kernel_dispatch_output = + kernel_dispatch_buffered_output_t{tool::get_config().kernel_trace}; + auto hsa_output = hsa_buffered_output_t{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}; + auto hip_output = hip_buffered_output_t{tool::get_config().hip_runtime_api_trace || + tool::get_config().hip_compiler_api_trace}; + auto memory_copy_output = memory_copy_buffered_output_t{tool::get_config().memory_copy_trace}; + auto marker_output = marker_buffered_output_t{tool::get_config().marker_api_trace}; + auto counters_output = + counter_collection_buffered_output_t{tool::get_config().counter_collection}; + auto scratch_memory_output = + scratch_memory_buffered_output_t{tool::get_config().scratch_memory}; - if(output_format == "CSV") + auto node_id_sort = [](const auto& lhs, const auto& rhs) { return lhs.node_id < rhs.node_id; }; + + auto _agents = std::vector{}; + _agents.reserve(agent_info->size()); + for(auto& itr : *agent_info) + _agents.emplace_back(itr.second); + + auto _counters = get_tool_counter_info(); + + std::sort(_agents.begin(), _agents.end(), node_id_sort); + + if(tool::get_config().csv_output) { - 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) + auto contributions = std::unordered_map{}; + + generate_output(kernel_dispatch_output, contributions); + generate_output(hsa_output, contributions); + generate_output(hip_output, contributions); + generate_output(memory_copy_output, contributions); + generate_output(marker_output, contributions); + generate_output(counters_output, contributions); + generate_output(scratch_memory_output, contributions); + + if(tool::get_config().stats && tool::get_config().csv_output) { - generate_output( - buffer_type_t::ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH); + rocprofiler::tool::generate_csv(tool_functions, contributions); } - 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) + if(tool::get_config().json_output) { - generate_output(buffer_type_t::ROCPROFILER_TOOL_BUFFER_HSA); + rocprofiler::tool::write_json(tool_functions, + getpid(), + _agents, + _counters, + &hip_output.element_data, + &hsa_output.element_data, + &kernel_dispatch_output.element_data, + &memory_copy_output.element_data, + &counters_output.element_data, + &marker_output.element_data, + &scratch_memory_output.element_data); } - if(tool::get_config().hip_runtime_api_trace || tool::get_config().hip_compiler_api_trace) - { - generate_output(buffer_type_t::ROCPROFILER_TOOL_BUFFER_HIP); - } + auto destroy_output = [](auto& _buffered_output_v) { _buffered_output_v.destroy(); }; - 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); - } + destroy_output(kernel_dispatch_output); + destroy_output(hsa_output); + destroy_output(hip_output); + destroy_output(memory_copy_output); + destroy_output(marker_output); + destroy_output(counters_output); + destroy_output(scratch_memory_output); fini_tool_table(); if(destructors) @@ -1789,7 +1458,126 @@ tool_fini(void* /*tool_data*/) } } // namespace -extern "C" rocprofiler_tool_configure_result_t* +std::map +get_callback_roctx_msg() +{ + auto _data = marker_msg_data->rlock([](const auto& _data_v) { return _data_v; }); + auto _ret = std::map{}; + for(const auto& itr : _data) + _ret.emplace(itr.first, itr.second); + return _ret; +} + +std::vector +get_kernel_symbol_data() +{ + auto _data = kernel_data->rlock([](const auto& _data_v) { + auto _info = std::vector{}; + _info.reserve(_data_v.size()); + for(const auto& itr : _data_v) + _info.emplace_back(itr.second); + return _info; + }); + + uint64_t kernel_data_size = 0; + for(const auto& itr : _data) + kernel_data_size = std::max(kernel_data_size, itr.kernel_id); + + auto _symbol_data = std::vector{}; + _symbol_data.resize(kernel_data_size + 1, kernel_symbol_data{}); + // index by the kernel id + for(auto& itr : _data) + _symbol_data.at(itr.kernel_id) = std::move(itr); + + return _symbol_data; +} + +std::vector +get_code_object_data() +{ + auto _data = code_obj_data->rlock([](const auto& _data_v) { + auto _info = std::vector{}; + _info.reserve(_data_v.size()); + for(const auto& itr : _data_v) + _info.emplace_back(itr.second); + return _info; + }); + + uint64_t _sz = 0; + for(const auto& itr : _data) + _sz = std::max(_sz, itr.code_object_id); + + auto _code_obj_data = std::vector{}; + _code_obj_data.resize(_sz + 1, rocprofiler_code_object_data_t{}); + // index by the code object id + for(auto& itr : _data) + _code_obj_data.at(itr.code_object_id) = itr; + + return _code_obj_data; +} + +std::vector +get_tool_counter_info() +{ + auto _data = get_agent_counter_info(get_gpu_agents()); + auto _ret = std::vector{}; + for(const auto& itr : _data) + { + for(const auto& iitr : itr.second) + _ret.emplace_back(iitr); + } + return _ret; +} + +std::vector +get_tool_counter_dimension_info() +{ + auto _data = get_agent_counter_info(get_gpu_agents()); + auto _ret = std::vector{}; + for(const auto& itr : _data) + { + for(const auto& iitr : itr.second) + for(const auto& ditr : iitr.dimension_info) + _ret.emplace_back(ditr); + } + + auto _sorter = [](const rocprofiler_record_dimension_info_t& lhs, + const rocprofiler_record_dimension_info_t& rhs) { + return std::tie(lhs.id, lhs.instance_size) < std::tie(rhs.id, rhs.instance_size); + }; + auto _equiv = [](const rocprofiler_record_dimension_info_t& lhs, + const rocprofiler_record_dimension_info_t& rhs) { + return std::tie(lhs.id, lhs.instance_size) == std::tie(rhs.id, rhs.instance_size); + }; + + std::sort(_ret.begin(), _ret.end(), _sorter); + _ret.erase(std::unique(_ret.begin(), _ret.end(), _equiv), _ret.end()); + + return _ret; +} + +namespace +{ +using main_func_t = int (*)(int, char**, char**); + +main_func_t& +get_main_function() +{ + static main_func_t user_main = nullptr; + return user_main; +} +} // namespace + +#define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal"))); + +extern "C" { +void +rocprofv3_set_main(main_func_t main_func) ROCPROFV3_INTERNAL_API; + +int +rocprofv3_main(int argc, char** argv, char** envp) ROCPROFV3_INTERNAL_API; + +rocprofiler_tool_configure_result_t* rocprofiler_configure(uint32_t version, const char* runtime_version, uint32_t priority, @@ -1814,19 +1602,19 @@ 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(marker_msg_data); + add_destructor(code_obj_data); add_destructor(kernel_data); add_destructor(tool_functions); add_destructor(agent_info); add_destructor(stats_timestamp); + // in case main wrapper is not used + ::atexit(finalize_rocprofv3); + if(tool::get_config().list_metrics) { ROCPROFILER_CALL(rocprofiler_at_intercept_table_registration( @@ -1861,3 +1649,29 @@ rocprofiler_configure(uint32_t version, return &cfg; // data passed around all the callbacks } + +void +rocprofv3_set_main(main_func_t main_func) +{ + get_main_function() = main_func; +} + +int +rocprofv3_main(int argc, char** argv, char** envp) +{ + auto logging_cfg = rocprofiler::common::logging_config{.install_failure_handler = true}; + common::init_logging("ROCPROF_LOG_LEVEL", logging_cfg); + FLAGS_colorlogtostderr = true; + + LOG(INFO) << "initializing rocprofv3..."; + initialize_rocprofv3(); + + auto ret = CHECK_NOTNULL(get_main_function())(argc, argv, envp); + + LOG(INFO) << "finalizing rocprofv3..."; + finalize_rocprofv3(); + + LOG(INFO) << "rocprofv3 finished. exit code: " << ret; + return ret; +} +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp index 829b32b332..50a53e06b8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -191,7 +191,14 @@ hip_api_impl::functor(Args... args) constexpr auto external_corr_id_domain_idx = hip_domain_info::external_correlation_id_domain_idx; - ROCP_INFO_IF(registration::get_fini_status() != 0) << "Executing " << info_type::name; + if(registration::get_fini_status() != 0) + { + [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + if constexpr(!std::is_void::value) + return _ret; + else + return; + } constexpr auto ref_count = 2; auto thr_id = common::get_tid(); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp index a43b4f2319..9785271870 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp @@ -306,8 +306,6 @@ hsa_api_impl::functor(Args... args) constexpr auto external_corr_id_domain_idx = hsa_domain_info::external_correlation_id_domain_idx; - ROCP_INFO_IF(registration::get_fini_status() != 0) << "Executing " << info_type::name; - if(registration::get_fini_status() != 0) { [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); diff --git a/projects/rocprofiler-sdk/tests/common/CMakeLists.txt b/projects/rocprofiler-sdk/tests/common/CMakeLists.txt index 92dd453a8a..890bc3eb1e 100644 --- a/projects/rocprofiler-sdk/tests/common/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/common/CMakeLists.txt @@ -41,11 +41,11 @@ if(ROCPROFILER_BUILD_CI OR ROCPROFILER_BUILD_WERROR) endif() # serialization library -if(NOT TARGET rocprofiler::cereal) +if(NOT TARGET rocprofiler::rocprofiler-cereal) get_filename_component(ROCPROFILER_SOURCE_DIR "${PROJECT_SOURCE_DIR}/.." REALPATH) add_library(rocprofiler-cereal INTERFACE) - add_library(rocprofiler::cereal ALIAS rocprofiler-cereal) + add_library(rocprofiler::rocprofiler-cereal ALIAS rocprofiler-cereal) target_compile_definitions(rocprofiler-cereal INTERFACE $) @@ -115,8 +115,9 @@ cmake_path(GET CMAKE_CURRENT_SOURCE_DIR PARENT_PATH COMMON_LIBRARY_INCLUDE_DIR) add_library(rocprofiler-tests-common-library INTERFACE) add_library(rocprofiler::tests-common-library ALIAS rocprofiler-tests-common-library) -target_link_libraries(rocprofiler-tests-common-library - INTERFACE rocprofiler::tests-build-flags rocprofiler::cereal) +target_link_libraries( + rocprofiler-tests-common-library INTERFACE rocprofiler::tests-build-flags + rocprofiler::rocprofiler-cereal) target_compile_features(rocprofiler-tests-common-library INTERFACE cxx_std_17) target_include_directories(rocprofiler-tests-common-library INTERFACE ${COMMON_LIBRARY_INCLUDE_DIR}) diff --git a/projects/rocprofiler-sdk/tests/common/serialization.hpp b/projects/rocprofiler-sdk/tests/common/serialization.hpp index a21ca6bfdc..00022c018d 100644 --- a/projects/rocprofiler-sdk/tests/common/serialization.hpp +++ b/projects/rocprofiler-sdk/tests/common/serialization.hpp @@ -23,726 +23,5 @@ #pragma once -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#define SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD)) -#define SAVE_DATA_VALUE(NAME, VALUE) ar(make_nvp(NAME, data.VALUE)) -#define SAVE_DATA_CSTR(FIELD) ar(make_nvp(#FIELD, std::string{data.FIELD})) -#define SAVE_DATA_BITFIELD(NAME, VALUE) \ - { \ - auto _val = data.VALUE; \ - ar(make_nvp(NAME, _val)); \ - } - -namespace cereal -{ -template -void -save(ArchiveT& ar, rocprofiler_context_id_t data) -{ - SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_id_t data) -{ - SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, hsa_agent_t data) -{ - SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_queue_id_t data) -{ - SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_id_t data) -{ - SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_correlation_id_t data) -{ - SAVE_DATA_FIELD(internal); - SAVE_DATA_VALUE("external", external.value); -} - -template -void -save(ArchiveT& ar, rocprofiler_dim3_t data) -{ - SAVE_DATA_FIELD(x); - SAVE_DATA_FIELD(y); - SAVE_DATA_FIELD(z); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_load_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(code_object_id); - SAVE_DATA_FIELD(rocp_agent); - SAVE_DATA_FIELD(hsa_agent); - SAVE_DATA_CSTR(uri); - SAVE_DATA_FIELD(load_base); - SAVE_DATA_FIELD(load_size); - SAVE_DATA_FIELD(load_delta); - SAVE_DATA_FIELD(storage_type); - if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) - { - SAVE_DATA_FIELD(storage_file); - } - else if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY) - { - SAVE_DATA_FIELD(memory_base); - SAVE_DATA_FIELD(memory_size); - } -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kernel_id); - SAVE_DATA_FIELD(code_object_id); - SAVE_DATA_CSTR(kernel_name); - SAVE_DATA_FIELD(kernel_object); - SAVE_DATA_FIELD(kernarg_segment_size); - SAVE_DATA_FIELD(kernarg_segment_alignment); - SAVE_DATA_FIELD(group_segment_size); - SAVE_DATA_FIELD(private_segment_size); -} - -template -void -save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data) -{ - SAVE_DATA_FIELD(uint64_t_retval); -} - -template -void -save(ArchiveT& ar, const hsa_queue_t& data) -{ - ar(make_nvp("queue_id", data.id)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_alloc_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - SAVE_DATA_FIELD(dispatch_id); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_alloc_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - SAVE_DATA_FIELD(dispatch_id); - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(num_slots); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_free_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_marker_api_retval_t data) -{ - SAVE_DATA_FIELD(int64_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) -{ - SAVE_DATA_FIELD(size); - // SAVE_DATA_FIELD(args); - SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_marker_api_data_t data) -{ - SAVE_DATA_FIELD(size); - // SAVE_DATA_FIELD(args); - SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_hip_api_retval_t data) -{ - SAVE_DATA_FIELD(hipError_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_hip_api_data_t data) -{ - SAVE_DATA_FIELD(size); - // SAVE_DATA_FIELD(args); - SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(agent_id); - SAVE_DATA_FIELD(queue_id); - SAVE_DATA_FIELD(flags); - SAVE_DATA_FIELD(args_kind); -} - -template -void -save(ArchiveT& ar, rocprofiler_kernel_dispatch_info_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(agent_id); - SAVE_DATA_FIELD(queue_id); - SAVE_DATA_FIELD(kernel_id); - SAVE_DATA_FIELD(dispatch_id); - SAVE_DATA_FIELD(private_segment_size); - SAVE_DATA_FIELD(group_segment_size); - SAVE_DATA_FIELD(workgroup_size); - SAVE_DATA_FIELD(group_segment_size); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_memory_copy_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(dst_agent_id); - SAVE_DATA_FIELD(src_agent_id); - SAVE_DATA_FIELD(bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_data_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_record_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(num_records); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) -{ - SAVE_DATA_FIELD(context_id); - SAVE_DATA_FIELD(thread_id); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(phase); -} - -template -void -save_buffer_tracing_api_record(ArchiveT& ar, Tp data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(thread_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_record_counter_t data) -{ - SAVE_DATA_FIELD(id); - SAVE_DATA_FIELD(counter_value); - SAVE_DATA_FIELD(dispatch_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(thread_id); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_copy_record_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(thread_id); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(dst_agent_id); - SAVE_DATA_FIELD(src_agent_id); - SAVE_DATA_FIELD(bytes); -} - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_record_t& data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(pid); - - switch(data.operation) - { - case ROCPROFILER_PAGE_MIGRATION_PAGE_FAULT: - { - ar(make_nvp("page_fault", data.page_fault)); - break; - } - case ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE: - { - ar(make_nvp("page_migrate", data.page_migrate)); - break; - } - case ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND: - { - ar(make_nvp("queue_suspend", data.queue_suspend)); - break; - } - case ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU: - { - ar(make_nvp("unmap_from_gpu", data.unmap_from_gpu)); - break; - } - case ROCPROFILER_PAGE_MIGRATION_NONE: - case ROCPROFILER_PAGE_MIGRATION_LAST: - { - throw std::runtime_error{"unsupported page migration operation type"}; - break; - } - } -} - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_page_fault_record_t& data) -{ - SAVE_DATA_FIELD(node_id); - SAVE_DATA_FIELD(address); - SAVE_DATA_FIELD(read_fault); - SAVE_DATA_FIELD(migrated); -} - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_page_migrate_record_t& data) -{ - SAVE_DATA_FIELD(start_addr); - SAVE_DATA_FIELD(end_addr); - SAVE_DATA_FIELD(from_node); - SAVE_DATA_FIELD(to_node); - SAVE_DATA_FIELD(prefetch_node); - SAVE_DATA_FIELD(preferred_node); - SAVE_DATA_FIELD(trigger); -} - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_queue_suspend_record_t& data) -{ - SAVE_DATA_FIELD(node_id); - SAVE_DATA_FIELD(trigger); - SAVE_DATA_FIELD(rescheduled); -} - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_unmap_from_gpu_record_t& data) -{ - SAVE_DATA_FIELD(node_id); - SAVE_DATA_FIELD(start_addr); - SAVE_DATA_FIELD(end_addr); - SAVE_DATA_FIELD(trigger); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_scratch_memory_record_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(operation); - SAVE_DATA_FIELD(agent_id); - SAVE_DATA_FIELD(queue_id); - SAVE_DATA_FIELD(thread_id); - SAVE_DATA_FIELD(start_timestamp); - SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(correlation_id); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_correlation_id_retirement_record_t data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(kind); - SAVE_DATA_FIELD(timestamp); - SAVE_DATA_FIELD(internal_correlation_id); -} - -template -void -save(ArchiveT& ar, HsaCacheType data) -{ - SAVE_DATA_BITFIELD("Data", ui32.Data); - SAVE_DATA_BITFIELD("Instruction", ui32.Instruction); - SAVE_DATA_BITFIELD("CPU", ui32.CPU); - SAVE_DATA_BITFIELD("HSACU", ui32.HSACU); -} - -template -void -save(ArchiveT& ar, HSA_LINKPROPERTY data) -{ - SAVE_DATA_BITFIELD("Override", ui32.Override); - SAVE_DATA_BITFIELD("NonCoherent", ui32.NonCoherent); - SAVE_DATA_BITFIELD("NoAtomics32bit", ui32.NoAtomics32bit); - SAVE_DATA_BITFIELD("NoAtomics64bit", ui32.NoAtomics64bit); - SAVE_DATA_BITFIELD("NoPeerToPeerDMA", ui32.NoPeerToPeerDMA); -} - -template -void -save(ArchiveT& ar, HSA_CAPABILITY data) -{ - SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); - SAVE_DATA_BITFIELD("HSAMMUPresent", ui32.HSAMMUPresent); - SAVE_DATA_BITFIELD("SharedWithGraphics", ui32.SharedWithGraphics); - SAVE_DATA_BITFIELD("QueueSizePowerOfTwo", ui32.QueueSizePowerOfTwo); - SAVE_DATA_BITFIELD("QueueSize32bit", ui32.QueueSize32bit); - SAVE_DATA_BITFIELD("QueueIdleEvent", ui32.QueueIdleEvent); - SAVE_DATA_BITFIELD("VALimit", ui32.VALimit); - SAVE_DATA_BITFIELD("WatchPointsSupported", ui32.WatchPointsSupported); - SAVE_DATA_BITFIELD("WatchPointsTotalBits", ui32.WatchPointsTotalBits); - SAVE_DATA_BITFIELD("DoorbellType", ui32.DoorbellType); - SAVE_DATA_BITFIELD("AQLQueueDoubleMap", ui32.AQLQueueDoubleMap); - SAVE_DATA_BITFIELD("DebugTrapSupported", ui32.DebugTrapSupported); - SAVE_DATA_BITFIELD("WaveLaunchTrapOverrideSupported", ui32.WaveLaunchTrapOverrideSupported); - SAVE_DATA_BITFIELD("WaveLaunchModeSupported", ui32.WaveLaunchModeSupported); - SAVE_DATA_BITFIELD("PreciseMemoryOperationsSupported", ui32.PreciseMemoryOperationsSupported); - SAVE_DATA_BITFIELD("DEPRECATED_SRAM_EDCSupport", ui32.DEPRECATED_SRAM_EDCSupport); - SAVE_DATA_BITFIELD("Mem_EDCSupport", ui32.Mem_EDCSupport); - SAVE_DATA_BITFIELD("RASEventNotify", ui32.RASEventNotify); - SAVE_DATA_BITFIELD("ASICRevision", ui32.ASICRevision); - SAVE_DATA_BITFIELD("SRAM_EDCSupport", ui32.SRAM_EDCSupport); - SAVE_DATA_BITFIELD("SVMAPISupported", ui32.SVMAPISupported); - SAVE_DATA_BITFIELD("CoherentHostAccess", ui32.CoherentHostAccess); - SAVE_DATA_BITFIELD("DebugSupportedFirmware", ui32.DebugSupportedFirmware); - SAVE_DATA_BITFIELD("Reserved", ui32.Reserved); -} - -template -void -save(ArchiveT& ar, HSA_MEMORYPROPERTY data) -{ - SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); - SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile); -} - -template -void -save(ArchiveT& ar, HSA_ENGINE_VERSION data) -{ - SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA); - SAVE_DATA_BITFIELD("uCodeRes", uCodeRes); -} - -template -void -save(ArchiveT& ar, HSA_ENGINE_ID data) -{ - SAVE_DATA_BITFIELD("uCode", ui32.uCode); - SAVE_DATA_BITFIELD("Major", ui32.Major); - SAVE_DATA_BITFIELD("Minor", ui32.Minor); - SAVE_DATA_BITFIELD("Stepping", ui32.Stepping); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_cache_t data) -{ - SAVE_DATA_FIELD(processor_id_low); - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(level); - SAVE_DATA_FIELD(cache_line_size); - SAVE_DATA_FIELD(cache_lines_per_tag); - SAVE_DATA_FIELD(association); - SAVE_DATA_FIELD(latency); - SAVE_DATA_FIELD(type); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_io_link_t data) -{ - SAVE_DATA_FIELD(type); - SAVE_DATA_FIELD(version_major); - SAVE_DATA_FIELD(version_minor); - SAVE_DATA_FIELD(node_from); - SAVE_DATA_FIELD(node_to); - SAVE_DATA_FIELD(weight); - SAVE_DATA_FIELD(min_latency); - SAVE_DATA_FIELD(max_latency); - SAVE_DATA_FIELD(min_bandwidth); - SAVE_DATA_FIELD(max_bandwidth); - SAVE_DATA_FIELD(recommended_transfer_size); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_mem_bank_t data) -{ - SAVE_DATA_FIELD(heap_type); - SAVE_DATA_FIELD(flags); - SAVE_DATA_FIELD(width); - SAVE_DATA_FIELD(mem_clk_max); - SAVE_DATA_FIELD(size_in_bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_configuration_t data) -{ - SAVE_DATA_FIELD(method); - SAVE_DATA_FIELD(unit); - SAVE_DATA_FIELD(min_interval); - SAVE_DATA_FIELD(max_interval); - SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, const rocprofiler_agent_t& data) -{ - SAVE_DATA_FIELD(size); - SAVE_DATA_FIELD(id); - SAVE_DATA_FIELD(type); - SAVE_DATA_FIELD(cpu_cores_count); - SAVE_DATA_FIELD(simd_count); - SAVE_DATA_FIELD(mem_banks_count); - SAVE_DATA_FIELD(caches_count); - SAVE_DATA_FIELD(io_links_count); - SAVE_DATA_FIELD(cpu_core_id_base); - SAVE_DATA_FIELD(simd_id_base); - SAVE_DATA_FIELD(max_waves_per_simd); - SAVE_DATA_FIELD(lds_size_in_kb); - SAVE_DATA_FIELD(gds_size_in_kb); - SAVE_DATA_FIELD(num_gws); - SAVE_DATA_FIELD(wave_front_size); - SAVE_DATA_FIELD(num_xcc); - SAVE_DATA_FIELD(cu_count); - SAVE_DATA_FIELD(array_count); - SAVE_DATA_FIELD(num_shader_banks); - SAVE_DATA_FIELD(simd_arrays_per_engine); - SAVE_DATA_FIELD(cu_per_simd_array); - SAVE_DATA_FIELD(simd_per_cu); - SAVE_DATA_FIELD(max_slots_scratch_cu); - SAVE_DATA_FIELD(gfx_target_version); - SAVE_DATA_FIELD(vendor_id); - SAVE_DATA_FIELD(device_id); - SAVE_DATA_FIELD(location_id); - SAVE_DATA_FIELD(domain); - SAVE_DATA_FIELD(drm_render_minor); - SAVE_DATA_FIELD(num_sdma_engines); - SAVE_DATA_FIELD(num_sdma_xgmi_engines); - SAVE_DATA_FIELD(num_sdma_queues_per_engine); - SAVE_DATA_FIELD(num_cp_queues); - SAVE_DATA_FIELD(max_engine_clk_ccompute); - SAVE_DATA_FIELD(max_engine_clk_fcompute); - SAVE_DATA_FIELD(sdma_fw_version); - SAVE_DATA_FIELD(fw_version); - SAVE_DATA_FIELD(capability); - SAVE_DATA_FIELD(cu_per_engine); - SAVE_DATA_FIELD(max_waves_per_cu); - SAVE_DATA_FIELD(family_id); - SAVE_DATA_FIELD(workgroup_max_size); - SAVE_DATA_FIELD(grid_max_size); - SAVE_DATA_FIELD(local_mem_size); - SAVE_DATA_FIELD(hive_id); - SAVE_DATA_FIELD(gpu_id); - SAVE_DATA_FIELD(workgroup_max_dim); - SAVE_DATA_FIELD(grid_max_dim); - SAVE_DATA_CSTR(name); - SAVE_DATA_CSTR(vendor_name); - SAVE_DATA_CSTR(product_name); - SAVE_DATA_CSTR(model_name); - SAVE_DATA_FIELD(num_pc_sampling_configs); - SAVE_DATA_FIELD(node_id); - SAVE_DATA_FIELD(logical_node_id); - - auto generate = [&](auto name, const auto* value, uint64_t size) { - using value_type = std::remove_const_t>; - auto vec = std::vector{}; - vec.reserve(size); - for(uint64_t i = 0; i < size; ++i) - vec.emplace_back(value[i]); - ar(make_nvp(name, vec)); - }; - - generate("mem_banks", data.mem_banks, data.mem_banks_count); - generate("caches", data.caches, data.caches_count); - generate("io_links", data.io_links, data.io_links_count); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_info_v0_t data) -{ - SAVE_DATA_FIELD(id); - SAVE_DATA_BITFIELD("is_constant", is_constant); - SAVE_DATA_BITFIELD("is_derived", is_derived); - SAVE_DATA_CSTR(name); - SAVE_DATA_CSTR(description); - SAVE_DATA_CSTR(block); - SAVE_DATA_CSTR(expression); -} -} // namespace cereal - -#undef SAVE_DATA_FIELD +// provided by the library +#include diff --git a/projects/rocprofiler-sdk/tests/pytest-packages/pytest_utils/__init__.py b/projects/rocprofiler-sdk/tests/pytest-packages/pytest_utils/__init__.py index 24d26e36c3..0a88127d07 100644 --- a/projects/rocprofiler-sdk/tests/pytest-packages/pytest_utils/__init__.py +++ b/projects/rocprofiler-sdk/tests/pytest-packages/pytest_utils/__init__.py @@ -21,3 +21,23 @@ # SOFTWARE. from __future__ import absolute_import + + +def collapse_dict_list(data, key="rocprofiler-sdk-tool"): + """Collapse a dictionary entry list into a single mapped value""" + + def check_return(_data): + assert isinstance(_data, dict), "expected dict, type: {}".format( + type(_data).__name__ + ) + return _data + + if ( + key in data.keys() + and len(data.keys()) == 1 + and isinstance(data[key], (list, tuple)) + and len(data[key]) == 1 + ): + return check_return({key: data[key][0]}) + + return check_return(data) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt index 8aa3baed7a..51e51f3e58 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt @@ -2,9 +2,6 @@ # Various counter collection tests # -# copy to binary directory -rocprofiler_configure_pytest_files(COPY conftest.py CONFIG pytest.ini) - add_subdirectory(input1) add_subdirectory(input2) add_subdirectory(list_metrics) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt index fb44ca077f..66479ceea0 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt @@ -10,10 +10,8 @@ project( find_package(rocprofiler-sdk REQUIRED) -foreach(FILENAME validate.py input.txt) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py + input.txt) # pmc1 add_test( @@ -21,7 +19,7 @@ add_test( COMMAND $ -i ${CMAKE_CURRENT_BINARY_DIR}/input.txt -T -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1 - -o pmc1 $) + -o pmc1 --output-format csv,json $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -33,12 +31,15 @@ set_tests_properties( PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-env-pmc1}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") -add_test(NAME rocprofv3-test-counter-collection-pmc1-validate - COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input - ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv) +add_test( + NAME rocprofv3-test-counter-collection-pmc1-validate + COMMAND + ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input + ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv + --json-input ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_results.json) set_tests_properties( rocprofv3-test-counter-collection-pmc1-validate PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS - rocprofv3-test-counter-collection-pmc1-execute FAIL_REGULAR_EXPRESSION + "rocprofv3-test-counter-collection-pmc1-execute" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/conftest.py new file mode 100644 index 0000000000..965bd0958e --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/conftest.py @@ -0,0 +1,31 @@ +#!/usr/bin/env python3 + +import json +import pytest +import pandas as pd + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list + + +def pytest_addoption(parser): + parser.addoption("--input", action="store", help="Path to csv file.") + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) + + +@pytest.fixture +def input_data(request): + filename = request.config.getoption("--input") + with open(filename, "r") as inp: + return pd.read_csv(inp) + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/pytest.ini similarity index 52% rename from projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini rename to projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/pytest.ini index 536c02a5b9..5e1e1c14a0 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/pytest.ini @@ -1,7 +1,5 @@ [pytest] addopts = --durations=20 -rA -s -vv -testpaths = input1/validate.py - input2/validate.py - list_metrics/validate.py +testpaths = validate.py pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py index e880ea9400..7732d1abe2 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py @@ -52,6 +52,60 @@ def test_validate_counter_collection_pmc1(input_data: pd.DataFrame): assert di_expect == di_uniq +def test_validate_counter_collection_pmc1_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + counter_collection_data = data["callback_records"]["counter_collection"] + dispatch_ids = [] + # at present, AQLProfile has bugs when reporting the counters for below architectures + skip_gfx = ("gfx1101", "gfx1102") + + def get_kernel_name(kernel_id): + return data["kernel_symbols"][kernel_id]["formatted_kernel_name"] + + def get_agent(agent_id): + for agent in data["agents"]: + if agent["id"]["handle"] == agent_id["handle"]: + return agent + return None + + def get_counter(counter_id): + for counter in data["counters"]: + if counter["id"]["handle"] == counter_id["handle"]: + return counter + return None + + for counter in counter_collection_data: + dispatch_data = counter["dispatch_data"]["dispatch_info"] + + assert dispatch_data["dispatch_id"] > 0 + assert dispatch_data["agent_id"]["handle"] > 0 + assert dispatch_data["queue_id"]["handle"] > 0 + + agent = get_agent(dispatch_data["agent_id"]) + kernel_name = get_kernel_name(dispatch_data["kernel_id"]) + + assert agent is not None + assert len(kernel_name) > 0 + + dispatch_ids.append(dispatch_data["dispatch_id"]) + if not re.search(r"__amd_rocclr_.*", kernel_name): + for record in counter["records"]: + counter = get_counter(record["counter_id"]) + assert counter is not None, f"record:\n\t{record}" + assert ( + counter["name"] == "SQ_WAVES" + ), f"record:\n\t{record}\ncounter:\n\t{counter}" + if agent["name"] not in skip_gfx: + assert ( + record["value"] > 0 + ), f"record: {record}\ncounter: {counter}\nagent: {agent}" + + di_uniq = list(set(sorted(dispatch_ids))) + # make sure the dispatch ids are unique and ordered + di_expect = [idx + 1 for idx in range(len(dispatch_ids))] + assert di_expect == di_uniq + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt index 5a4b0e2aaa..8e350eeac6 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt @@ -10,18 +10,16 @@ project( find_package(rocprofiler-sdk REQUIRED) -foreach(FILENAME validate.py input.txt) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py + input.txt) # pmc2 add_test( NAME rocprofv3-test-counter-collection-pmc2-execute COMMAND $ -i - ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-cc -o - out $) + ${CMAKE_CURRENT_BINARY_DIR}/input.txt --output-format CSV,JSON -d + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-cc -o out $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -58,7 +56,7 @@ set_tests_properties( LABELS "integration-tests" DEPENDS - rocprofv3-test-counter-collection-pmc2-execute + "rocprofv3-test-counter-collection-pmc2-execute" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}" ATTACHED_FILES_ON_FAIL diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/conftest.py similarity index 69% rename from projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py rename to projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/conftest.py index 17481ec353..08f1d3c298 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/conftest.py @@ -3,11 +3,17 @@ import json import pytest import csv -import pandas as pd + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list def pytest_addoption(parser): - parser.addoption("--input", action="store", help="Path to csv file.") + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) parser.addoption( "--agent-input", action="store", @@ -20,16 +26,6 @@ def pytest_addoption(parser): ) -@pytest.fixture -def input_data(request): - filename = request.config.getoption("--input") - if filename: - with open(filename, "r") as inp: - return pd.read_csv(filename) - else: - return None - - @pytest.fixture def agent_info_input_data(request): filename = request.config.getoption("--agent-input") @@ -52,3 +48,10 @@ def counter_input_data(request): data.append(row) return data + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py index 7d7d09d171..46d62a92f4 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py @@ -1,5 +1,5 @@ -import pandas as pd -import os +#!/usr/bin/env python3 + import sys import pytest diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt index 4fb3e17531..bcbef41f1b 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt @@ -10,10 +10,7 @@ project( find_package(rocprofiler-sdk REQUIRED) -foreach(FILENAME validate.py conftest.py) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py) # basic-metrics add_test(NAME rocprofv3-test-list-metrics-execute diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/conftest.py index 340f9df28f..99db2d57c7 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/conftest.py @@ -2,7 +2,6 @@ import csv import pytest -import pandas as pd def pytest_addoption(parser): diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/validate.py index df3eeb24c4..396842d596 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/validate.py @@ -5,7 +5,7 @@ import pytest def test_validate_list_basic_metrics(basic_metrics_input_data): for row in basic_metrics_input_data: - assert row["Agent-id"].isdigit() == True + assert row["Agent_Id"].isdigit() == True assert row["Name"] != "" assert row["Description"] != "" assert row["Block"] != "" @@ -19,7 +19,7 @@ def test_validate_list_basic_metrics(basic_metrics_input_data): def test_validate_list_derived_metrics(derived_metrics_input_data): for row in derived_metrics_input_data: - assert row["Agent-id"].isdigit() == True + assert row["Agent_Id"].isdigit() == True assert row["Name"] != "" assert row["Description"] != "" assert row["Expression"] != "" diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt index ce916b80da..1de09e4926 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt @@ -13,10 +13,7 @@ string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV set(tracing-env "${PRELOAD_ENV}") -foreach(FILENAME validate.py conftest.py) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py) find_package(rocprofiler-sdk REQUIRED) @@ -25,7 +22,8 @@ add_test( NAME rocprofv3-test-hsa-multiqueue-execute COMMAND $ --hsa-trace --kernel-trace -d - ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $) + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format json:csv + $) set_tests_properties( rocprofv3-test-hsa-multiqueue-execute @@ -38,11 +36,14 @@ add_test( ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --hsa-trace-input ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_hsa_api_trace.csv --kernel-trace-input - ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_kernel_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_kernel_trace.csv + --json-input + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_results.json) set(MULTIQUEUE_VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_hsa_api_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_kernel_api_trace.csv) + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_kernel_api_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_results.json) set_tests_properties( rocprofv3-test-hsa-multiqueue-validate @@ -51,7 +52,7 @@ set_tests_properties( LABELS "integration-tests" DEPENDS - rocprofv3-test-hsa-multiqueue-execute + "rocprofv3-test-hsa-multiqueue-execute" FAIL_REGULAR_EXPRESSION "AssertionError" ATTACHED_FILES_ON_FAIL diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/conftest.py index 6725dbbb32..456e224bc5 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/conftest.py @@ -2,6 +2,10 @@ import csv import pytest +import json + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list def pytest_addoption(parser): @@ -15,6 +19,11 @@ def pytest_addoption(parser): action="store", help="Path to Kernel API tracing CSV file.", ) + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) @pytest.fixture @@ -39,3 +48,10 @@ def kernel_trace_input_data(request): data.append(row) return data + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py index b09250fe04..bb26c341c4 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py @@ -65,6 +65,95 @@ def test_kernel_trace(kernel_trace_input_data): assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) +def test_kernel_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + valid_kernel_names = ["copyA", "copyB", "copyC"] + buffer_records = data["buffer_records"] + buffer_names = data["strings"]["buffer_records"] + kernel_dispatch_data = buffer_records["kernel_dispatches"] + + def get_kernel_name(kernel_id): + return data["kernel_symbols"][kernel_id]["formatted_kernel_name"] + + assert len(kernel_dispatch_data) == 3 + + for dispatch in kernel_dispatch_data: + + assert buffer_names[dispatch["kind"]]["kind"] == "KERNEL_DISPATCH" + dispatch_info = dispatch["dispatch_info"] + assert dispatch_info["agent_id"]["handle"] > 0 + assert dispatch_info["queue_id"]["handle"] > 0 + assert dispatch_info["kernel_id"] > 0 + + kernel_name = get_kernel_name(dispatch_info["kernel_id"]) + assert kernel_name in valid_kernel_names + + assert dispatch["correlation_id"]["internal"] > 0 + assert dispatch_info["workgroup_size"]["x"] == 64 + assert dispatch_info["workgroup_size"]["y"] == 1 + assert dispatch_info["workgroup_size"]["z"] == 1 + assert dispatch_info["grid_size"]["x"] == 64 + assert dispatch_info["grid_size"]["y"] == 1 + assert dispatch_info["grid_size"]["z"] == 1 + assert dispatch["end_timestamp"] >= dispatch["start_timestamp"] + + +def test_hsa_api_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + functions = [] + correlation_ids = [] + + def get_operation_name(kind_id, op_id): + return data["strings"]["buffer_records"][kind_id]["operations"][op_id] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + metadata = data["metadata"] + buffer_records = data["buffer_records"] + + valid_domain_names = ( + "HSA_CORE_API", + "HSA_AMD_EXT_API", + "HSA_IMAGE_EXT_API", + "HSA_FINALIZE_EXT_API", + ) + + assert metadata["pid"] > 0 + hsa_api_data = buffer_records["hsa_api"] + + for itr in hsa_api_data: + kind = get_kind_name(itr["kind"]) + assert kind in valid_domain_names + assert itr["end_timestamp"] >= itr["start_timestamp"] + functions.append(get_operation_name(itr["kind"], itr["operation"])) + correlation_ids.append(itr["correlation_id"]["internal"]) + + correlation_ids = sorted(list(set(correlation_ids))) + + # deterministic call counts + num_queue_create_calls = 2 + num_queue_destroy_calls = 2 + num_hsa_mem_free_calls = 2 + + # signal create/destroy calls + # although the app explicitly only creates 3 signals + # but hsa_init() internally calls hsa_signal_create + num_hsa_signal_create_calls = 4 + num_hsa_signal_destroy_calls = 4 + + # all correlation ids are unique + assert len(correlation_ids) == len(hsa_api_data) + + functions = list(functions) + assert "hsa_shut_down" in functions + assert functions.count("hsa_queue_create") == num_queue_create_calls + assert functions.count("hsa_queue_destroy") == num_queue_destroy_calls + assert functions.count("hsa_memory_free") == num_hsa_mem_free_calls + assert functions.count("hsa_signal_create") == num_hsa_signal_create_calls + assert functions.count("hsa_signal_destroy") == num_hsa_signal_destroy_calls + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt index b6ea22ad4a..cc6771a025 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt @@ -16,6 +16,15 @@ add_test( $ --hip-runtime-trace --hip-compiler-trace --hsa-core-trace --hsa-amd-trace --hsa-image-trace --hsa-finalizer-trace --kernel-trace --memory-copy-trace --stats -d + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format csv + $) + +add_test( + NAME rocprofv3-test-trace-hip-in-libraries-json-execute + COMMAND + $ --hip-runtime-trace + --hip-compiler-trace --hsa-core-trace --hsa-amd-trace --hsa-image-trace + --hsa-finalizer-trace --kernel-trace --memory-copy-trace --output-format JSON -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV @@ -36,10 +45,20 @@ set_tests_properties( "HSA_CORE_API|HSA_AMD_EXT_API|HSA_IMAGE_EXT_API|HSA_FINALIZER_EXT_API|HIP_API|HIP_COMPILER_API|KERNEL_DISPATCH|CODE_OBJECT" ) -foreach(FILENAME validate.py conftest.py) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +set_tests_properties( + rocprofv3-test-trace-hip-in-libraries-json-execute + PROPERTIES + TIMEOUT + 100 + LABELS + "integration-tests" + ENVIRONMENT + "${tracing-env}" + FAIL_REGULAR_EXPRESSION + "HSA_CORE_API|HSA_AMD_EXT_API|HSA_IMAGE_EXT_API|HSA_FINALIZER_EXT_API|HIP_API|HIP_COMPILER_API|KERNEL_DISPATCH|CODE_OBJECT" + ) + +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py) add_test( NAME rocprofv3-test-trace-hip-in-libraries-validate @@ -59,7 +78,8 @@ add_test( --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) + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv + --json-input ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json) set(VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv @@ -70,17 +90,19 @@ set(VALIDATION_FILES ${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) + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json) set_tests_properties( rocprofv3-test-trace-hip-in-libraries-validate - PROPERTIES TIMEOUT - 45 - LABELS - "integration-tests" - DEPENDS - rocprofv3-test-trace-hip-in-libraries-execute - FAIL_REGULAR_EXPRESSION - "AssertionError" - ATTACHED_FILES_ON_FAIL - "${VALIDATION_FILES}") + PROPERTIES + TIMEOUT + 45 + LABELS + "integration-tests" + DEPENDS + "rocprofv3-test-trace-hip-in-libraries-execute;rocprofv3-test-trace-hip-in-libraries-json-execute" + FAIL_REGULAR_EXPRESSION + "AssertionError" + ATTACHED_FILES_ON_FAIL + "${VALIDATION_FILES}") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/conftest.py index 4d44b2e484..d6835b7488 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/conftest.py @@ -3,6 +3,10 @@ import os import csv import pytest +import json + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list def pytest_addoption(parser): @@ -57,6 +61,12 @@ def pytest_addoption(parser): help="Path to memory copy stats CSV file.", ) + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) + @pytest.fixture def agent_info_input_data(request): @@ -181,3 +191,10 @@ def memory_copy_stats_data(request): data.append(row) return data + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py index a0d8e19140..775442ed36 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py @@ -148,6 +148,87 @@ def test_api_trace( validate_stats(row) +def test_api_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + + metadata = data["metadata"] + names = data["strings"]["buffer_records"] + buffer_records = data["buffer_records"] + hsa_data = buffer_records["hsa_api"] + hip_data = buffer_records["hip_api"] + + valid_domain = [ + "HSA_CORE_API", + "HSA_AMD_EXT_API", + "HSA_IMAGE_EXT_API", + "HSA_FINALIZE_EXT_API", + ] + + valid_hip_domain = [ + "HIP_RUNTIME_API", + "HIP_COMPILER_API", + ] + + def get_operation_name(kind_id, op_id): + return names[kind_id]["operations"][op_id] + + def get_kind_name(kind_id): + return names[kind_id]["kind"] + + assert metadata["pid"] > 0 + + functions = [] + correlation_ids = [] + for api in hsa_data: + kind = get_kind_name(api["kind"]) + assert kind in valid_domain + assert api["thread_id"] >= metadata["pid"] + assert api["end_timestamp"] >= api["start_timestamp"] + functions.append(get_operation_name(api["kind"], api["operation"])) + correlation_ids.append(api["correlation_id"]["internal"]) + + for api in hip_data: + kind = get_kind_name(api["kind"]) + assert kind in valid_hip_domain + assert metadata["pid"] > 0 + assert api["thread_id"] == 0 or api["thread_id"] >= metadata["pid"] + assert api["end_timestamp"] >= api["start_timestamp"] + functions.append(get_operation_name(api["kind"], api["operation"])) + correlation_ids.append(api["correlation_id"]["internal"]) + + correlation_ids = sorted(list(set(correlation_ids))) + + # all correlation ids are unique + assert len(correlation_ids) == (len(hsa_data) + len(hip_data)) + # correlation ids are numbered from 1 to N + assert correlation_ids[0] == 1 + assert correlation_ids[-1] == len(correlation_ids) + + functions = list(set(functions)) + for itr in ( + "hsa_amd_memory_async_copy_on_engine", + "hsa_agent_get_info", + "hsa_agent_iterate_isas", + "hsa_signal_create", + "hsa_agent_get_info", + "hsa_executable_symbol_get_info", + ): + assert itr in functions + if hip_data: + for itr in ( + "hipGetLastError", + "hipLaunchKernel", + "hipStreamSynchronize", + "hipMemcpyAsync", + "hipFree", + "hipStreamDestroy", + "hipDeviceSynchronize", + "hipDeviceReset", + "hipSetDevice", + ): + assert itr in functions + + def test_kernel_trace(kernel_input_data, kernel_stats_data): valid_kernel_names = sorted( [ @@ -203,6 +284,68 @@ def test_kernel_trace(kernel_input_data, kernel_stats_data): validate_stats(row) +def test_kernel_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + + buffer_records = data["buffer_records"] + names = data["strings"]["buffer_records"] + + valid_kernel_names = sorted( + [ + "(anonymous namespace)::transpose(int const*, int*, int, int)", + "void (anonymous namespace)::addition_kernel(float*, float const*, float const*, int, int)", + "void (anonymous namespace)::divide_kernel(float*, float const*, float const*, int, int)", + "void (anonymous namespace)::multiply_kernel(float*, float const*, float const*, int, int)", + "void (anonymous namespace)::subtract_kernel(float*, float const*, float const*, int, int)", + ] + ) + + def get_kernel_name(kernel_id): + return data["kernel_symbols"][kernel_id]["formatted_kernel_name"] + + kernels = [] + for row in buffer_records["kernel_dispatches"]: + dispatch_info = row["dispatch_info"] + kernel_name = get_kernel_name(dispatch_info["kernel_id"]) + if re.search(r"__amd_rocclr_.*", kernel_name): + continue + + kernels.append(kernel_name) + + assert names[row["kind"]]["kind"] == "KERNEL_DISPATCH" + assert dispatch_info["agent_id"]["handle"] > 0 + assert dispatch_info["queue_id"]["handle"] > 0 + assert dispatch_info["kernel_id"] > 0 + assert row["correlation_id"]["internal"] > 0 + assert kernel_name in valid_kernel_names, f"row:\n\t{row}" + + workgrp_size = dim3( + dispatch_info["workgroup_size"]["x"], + dispatch_info["workgroup_size"]["y"], + dispatch_info["workgroup_size"]["z"], + ) + grid_size = dim3( + dispatch_info["grid_size"]["x"], + dispatch_info["grid_size"]["y"], + dispatch_info["grid_size"]["z"], + ) + + if kernel_name == "__amd_rocclr_fillBufferAligned": + assert workgrp_size.as_tuple() > (1, 1, 1) + assert grid_size.as_tuple() > (1, 1, 1) + elif "transpose" in kernel_name: + assert workgrp_size.as_tuple() == (32, 32, 1) + assert grid_size.as_tuple() == (9920, 9920, 1) + else: + assert workgrp_size.as_tuple() == (64, 1, 1) + assert grid_size.as_tuple() == (4096, 2048, 1) + + assert int(row["end_timestamp"]) >= int(row["start_timestamp"]) + + kernels = sorted(list(set(kernels))) + assert kernels == valid_kernel_names + + def test_memory_copy_trace( agent_info_input_data, memory_copy_input_data, diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt index 523d8745d3..52cb110d4b 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt @@ -20,31 +20,36 @@ add_test( COMMAND $ --hsa-trace -i ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace - -o pmc3 $) + -o pmc3 --output-format JSON:CSV $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") -if(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") - set(LOG_LEVEL "warning") # info produces memory leak -else() - set(LOG_LEVEL "info") -endif() - -set(cc-tracing-env "${PRELOAD_ENV}" "ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}" - "ROCPROF_LOG_LEVEL=${LOG_LEVEL}") +set(cc-tracing-env "${PRELOAD_ENV}") set_tests_properties( rocprofv3-test-tracing-plus-cc-execute - PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-tracing-env}" - FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") - -add_test(NAME rocprofv3-test-tracing-plus-cc-validate - COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py - --input-dir "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace") - -set_tests_properties( - rocprofv3-test-tracing-plus-cc-validate - PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS - rocprofv3-test-tracing-plus-cc-execute FAIL_REGULAR_EXPRESSION + PROPERTIES TIMEOUT 45 LABELS "integration-tests;application-replay" ENVIRONMENT + "${cc-tracing-env}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +foreach(_DIR "pmc_1" "pmc_2" "pmc_3" "pmc_4") + add_test( + NAME rocprofv3-test-tracing-plus-cc-validate-${_DIR} + COMMAND + ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_results.json" + --hsa-input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_hsa_api_trace.csv" + --agent-input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_agent_info.csv" + --counter-input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_counter_collection.csv" + ) + + set_tests_properties( + rocprofv3-test-tracing-plus-cc-validate-${_DIR} + PROPERTIES TIMEOUT 45 LABELS "integration-tests;application-replay" DEPENDS + "rocprofv3-test-tracing-plus-cc-execute" FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}") +endforeach() diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py index ad013dbe3a..309203403b 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py @@ -1,15 +1,186 @@ #!/usr/bin/env python3 -import json +import os +import csv import pytest +import json import pandas as pd +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list + def pytest_addoption(parser): - parser.addoption("--input-dir", action="store", help="Path to output dir.") + parser.addoption( + "--agent-input", + action="store", + help="Path to agent info CSV file.", + ) + parser.addoption( + "--hsa-input", + action="store", + help="Path to HSA API tracing CSV file.", + ) + parser.addoption( + "--kernel-input", + action="store", + help="Path to kernel tracing CSV file.", + ) + parser.addoption( + "--counter-input", + action="store", + help="Path to counter collection CSV file.", + ) + parser.addoption( + "--memory-copy-input", + action="store", + help="Path to memory-copy tracing CSV file.", + ) + parser.addoption( + "--marker-input", + action="store", + help="Path to marker API tracing CSV file.", + ) + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) @pytest.fixture -def input_dir(request): - dirname = request.config.getoption("--input-dir") - return dirname +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") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def kernel_input_data(request): + filename = request.config.getoption("--kernel-input") + with open(filename, "r") as inp: + return pd.read_csv(inp) + + return None + + +@pytest.fixture +def counter_input_data(request): + filename = request.config.getoption("--counter-input") + with open(filename, "r") as inp: + return pd.read_csv(inp) + + return None + + +@pytest.fixture +def memory_copy_input_data(request): + filename = request.config.getoption("--memory-copy-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def marker_input_data(request): + filename = request.config.getoption("--marker-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def hip_input_data(request): + filename = request.config.getoption("--hip-input") + 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 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 + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt index 2c1d6151b6..3dd6a5cc48 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt @@ -1,6 +1,8 @@ # multi block counters -pmc: GRBM_COUNT GRBM_GUI_ACTIVE SQ_CYCLES SQ_BUSY_CYCLES SQ_WAVES -pmc: TCC_REQ_sum TCC_STREAMING_REQ_sum TCC_HIT_sum TCC_MISS_sum -pmc: TCC_EA_WRREQ_sum TCC_EA_WRREQ_64B_sum TCC_EA_WR_UNCACHED_32B_sum -pmc: TCC_TOO_MANY_EA_WRREQS_STALL_sum TCC_EA_ATOMIC_sum TCC_EA_RDREQ_sum TCC_EA_RDREQ_32B_sum \ No newline at end of file +pmc: GRBM_COUNT +pmc: GRBM_GUI_ACTIVE +pmc: TA_BUSY_avr + +# below line expects that no system has >= 16384 GPUs so they should never be collected +pmc: SQ_WAVES SQ_CYCLES:device=16384 SQ_BUSY_CYCLES:device=65536 diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py index 04edcbeaec..5b6d2f2edb 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py @@ -1,44 +1,23 @@ -import pandas as pd -import os import sys import pytest -def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame): - directory_path = input_dir +def test_validate_counter_collection_plus_tracing( + json_data, counter_input_data, hsa_input_data +): - # Check if the directory is not empty - assert os.path.isdir(directory_path), f"{directory_path} is not a directory." - assert os.listdir(directory_path), f"{directory_path} is empty." - - # Check if there are 4 subdirectories for pmc's - subdirectories = [ - d - for d in os.listdir(directory_path) - if os.path.isdir(os.path.join(directory_path, d)) - ] + # check if either kernel-name/FUNCTION is present assert ( - len(subdirectories) == 4 - ), f"Expected 4 subdirectories, found {len(subdirectories)}." + "Kernel_Name" in counter_input_data.columns + or "Function" in counter_input_data.columns + ) - # Check if each subdirectory has files - for subdirectory in subdirectories: - subdirectory_path = os.path.join(directory_path, subdirectory) - assert os.listdir(subdirectory_path), f"{subdirectory_path} is empty." + data = json_data["rocprofiler-sdk-tool"] + hsa_api = data["buffer_records"]["hsa_api"] + assert len(hsa_input_data) == len(hsa_api) - # Check if each file in the subdirectory has some data - for file_name in os.listdir(subdirectory_path): - file_path = os.path.join(subdirectory_path, file_name) - # ignore hidden folders - if os.path.isdir(file_path) and os.path.basename(file_path).startswith("."): - continue - assert os.path.isfile(file_path), f"{file_path} is not a file." - - 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 + counter_collection_data = data["callback_records"]["counter_collection"] + assert len(counter_collection_data) > 0 if __name__ == "__main__": diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt index 8201742c6f..4df924cd7f 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt @@ -15,7 +15,7 @@ add_test( COMMAND $ -M --hsa-trace --kernel-trace --memory-copy-trace --marker-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o - out $) + out --output-format csv,json $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -42,10 +42,7 @@ set_tests_properties( "HSA_API|HIP_API|HIP_COMPILER_API|MARKER_CORE_API|MARKER_CONTROL_API|MARKER_NAME_API|KERNEL_DISPATCH|CODE_OBJECT" ) -foreach(FILENAME validate.py conftest.py) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} - ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) -endforeach() +rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py) add_test( NAME rocprofv3-test-trace-validate @@ -59,14 +56,16 @@ add_test( --marker-input ${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) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv + --json-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json) 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_agent_info.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json) set_tests_properties( rocprofv3-test-trace-validate @@ -75,7 +74,7 @@ set_tests_properties( LABELS "integration-tests" DEPENDS - rocprofv3-test-trace-execute + "rocprofv3-test-trace-execute" FAIL_REGULAR_EXPRESSION "AssertionError" ATTACHED_FILES_ON_FAIL @@ -88,7 +87,7 @@ add_test( NAME rocprofv3-test-systrace-execute COMMAND $ --sys-trace -d - ${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out --output-format csv,json $) set_tests_properties( @@ -117,14 +116,17 @@ add_test( --marker-input ${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) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv + --json-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json) 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_agent_info.csv) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json) set_tests_properties( rocprofv3-test-systrace-validate @@ -133,7 +135,7 @@ set_tests_properties( LABELS "integration-tests" DEPENDS - rocprofv3-test-systrace-execute + "rocprofv3-test-systrace-execute" FAIL_REGULAR_EXPRESSION "AssertionError" ATTACHED_FILES_ON_FAIL diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py index 64c9d2b0a3..3db9c1af3e 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py @@ -2,6 +2,10 @@ import csv import pytest +import json + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list def pytest_addoption(parser): @@ -35,6 +39,11 @@ def pytest_addoption(parser): action="store", help="Path to HIP runtime and compiler API tracing CSV file.", ) + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) @pytest.fixture @@ -107,3 +116,10 @@ def hip_input_data(request): data.append(row) return data + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py index 9451099c08..f21afe777e 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py @@ -2,6 +2,7 @@ import sys import pytest +import re def test_agent_info(agent_info_input_data): @@ -41,7 +42,7 @@ def test_hsa_api_trace(hsa_input_data): correlation_ids = sorted(list(set(correlation_ids))) hsa_api_calls_offset = 2 # roctxRangePush is first - num_marker_api_calls = 7 # seven marker API calls, only six entries in + num_marker_api_calls = 6 # seven marker API calls, only six entries in # marker csv data because roctxRangePush + roctxRangePop is one entry # all correlation ids are unique @@ -54,6 +55,49 @@ def test_hsa_api_trace(hsa_input_data): assert "hsa_amd_memory_async_copy_on_engine" in functions +def test_hsa_api_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + + def get_operation_name(kind_id, op_id): + return data["strings"]["buffer_records"][kind_id]["operations"][op_id] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + valid_domain_names = ( + "HSA_CORE_API", + "HSA_AMD_EXT_API", + "HSA_IMAGE_EXT_API", + "HSA_FINALIZE_EXT_API", + ) + + hsa_api_data = data["buffer_records"]["hsa_api"] + + functions = [] + correlation_ids = [] + for api in hsa_api_data: + kind = get_kind_name(api["kind"]) + assert kind in valid_domain_names + assert api["end_timestamp"] >= api["start_timestamp"] + functions.append(get_operation_name(api["kind"], api["operation"])) + correlation_ids.append(api["correlation_id"]["internal"]) + + correlation_ids = sorted(list(set(correlation_ids))) + + hsa_api_calls_offset = 2 # roctxRangePush is first + num_marker_api_calls = 6 # seven marker API calls, only six entries in + # marker csv data because roctxRangePush + roctxRangePop is one entry + + # all correlation ids are unique + assert len(correlation_ids) == len(hsa_api_data) + # correlation ids are numbered from 1 to N + assert correlation_ids[0] == hsa_api_calls_offset + assert correlation_ids[-1] == len(correlation_ids) + num_marker_api_calls + + functions = list(set(functions)) + assert "hsa_amd_memory_async_copy_on_engine" in functions + + def test_kernel_trace(kernel_input_data): valid_kernel_names = ( "_Z15matrixTransposePfS_i.kd", @@ -77,8 +121,43 @@ def test_kernel_trace(kernel_input_data): assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) -def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data): +def test_kernel_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + def get_kernel_name(kernel_id): + return data["kernel_symbols"][kernel_id]["formatted_kernel_name"] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + valid_kernel_names = ( + "_Z15matrixTransposePfS_i.kd", + "matrixTranspose(float*, float*, int)", + ) + kernel_dispatch_data = data["buffer_records"]["kernel_dispatches"] + assert len(kernel_dispatch_data) == 1 + for dispatch in kernel_dispatch_data: + dispatch_info = dispatch["dispatch_info"] + kernel_name = get_kernel_name(dispatch_info["kernel_id"]) + + assert get_kind_name(dispatch["kind"]) == "KERNEL_DISPATCH" + assert dispatch["correlation_id"]["internal"] > 0 + assert dispatch_info["agent_id"]["handle"] > 0 + assert dispatch_info["queue_id"]["handle"] > 0 + assert dispatch_info["kernel_id"] > 0 + if not re.search(r"__amd_rocclr_.*", kernel_name): + assert kernel_name in valid_kernel_names + + assert dispatch_info["workgroup_size"]["x"] == 4 + assert dispatch_info["workgroup_size"]["y"] == 4 + assert dispatch_info["workgroup_size"]["z"] == 1 + assert dispatch_info["grid_size"]["x"] == 1024 + assert dispatch_info["grid_size"]["y"] == 1024 + assert dispatch_info["grid_size"]["z"] == 1 + assert dispatch["end_timestamp"] >= dispatch["start_timestamp"] + + +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: @@ -110,6 +189,45 @@ def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data): test_row(1, "DEVICE_TO_HOST") +def test_memory_copy_json_trace(json_data): + data = json_data["rocprofiler-sdk-tool"] + + buffer_records = data["buffer_records"] + agent_data = data["agents"] + memory_copy_data = buffer_records["memory_copy"] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + def get_agent(node_id): + for agent in agent_data: + if agent["id"]["handle"] == node_id["handle"]: + return agent + return None + + assert len(memory_copy_data) == 2 + + def test_row(idx, direction): + assert direction in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + row = memory_copy_data[idx] + src_agent = get_agent(row["src_agent_id"]) + dst_agent = get_agent(row["dst_agent_id"]) + assert get_kind_name(row["kind"]) == "MEMORY_COPY" + assert src_agent is not None, f"{row}" + assert dst_agent is not None, f"{row}" + if direction == "HOST_TO_DEVICE": + assert src_agent["type"] == 1 + assert dst_agent["type"] == 2 + else: + assert src_agent["type"] == 2 + assert dst_agent["type"] == 1 + assert row["correlation_id"]["internal"] > 0 + assert row["end_timestamp"] >= row["start_timestamp"] + + test_row(0, "HOST_TO_DEVICE") + test_row(1, "DEVICE_TO_HOST") + + def test_marker_api_trace(marker_input_data): functions = [] for row in marker_input_data: @@ -129,6 +247,22 @@ def test_marker_api_trace(marker_input_data): assert "main" in functions +def test_marker_api_trace_json(json_data): + data = json_data["rocprofiler-sdk-tool"] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + valid_domain = ("MARKER_CORE_API", "MARKER_CONTROL_API", "MARKER_NAME_API") + + buffer_records = data["buffer_records"] + marker_data = buffer_records["marker_api"] + for marker in marker_data: + assert get_kind_name(marker["kind"]) in valid_domain + assert marker["thread_id"] >= data["metadata"]["pid"] + assert marker["end_timestamp"] >= marker["start_timestamp"] + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt b/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt index a61abd974e..03890cd2f7 100644 --- a/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt @@ -15,8 +15,9 @@ add_library(rocprofiler-sdk-json-tool SHARED) target_sources(rocprofiler-sdk-json-tool PRIVATE json-tool.cpp) target_link_libraries( rocprofiler-sdk-json-tool - PRIVATE rocprofiler::rocprofiler rocprofiler::cereal rocprofiler::tests-build-flags - rocprofiler::tests-common-library rocprofiler::tests-perfetto) + PRIVATE rocprofiler::rocprofiler rocprofiler::rocprofiler-cereal + rocprofiler::tests-build-flags rocprofiler::tests-common-library + rocprofiler::tests-perfetto) set_target_properties( rocprofiler-sdk-json-tool PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib/rocprofiler-sdk"