AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Adding JSON support (#860)

Browse Source 
* 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 <bits/utility.h> 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<Tp>

* 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 <source-tree>/rocprofv3.sh -> <binary-tree>/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 <Sriraksha.Nagaraj@amd.com>

[ROCm/rocprofiler-sdk commit: 92b7326910]
This commit is contained in:
Jonathan R. Madsen
2024-05-22 00:53:42 -05:00
committed by GitHub
parent 01cdd61d70
commit 0ba9f26a6a
60 changed files with 3900 additions and 2191 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake
+1
View File
@@ -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
projects/rocprofiler-sdk/external/CMakeLists.txt
Vendored
+14 -16
View File
@@ -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 $<BUILD_INTERFACE:CEREAL_THREAD_SAFE=1>)
target_include_directories(
rocprofiler-cereal
INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/cereal/include>)
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 $<BUILD_INTERFACE:CEREAL_THREAD_SAFE=1>)
target_include_directories(
rocprofiler-cereal
INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/cereal/include>)
# doxygen-awesome
if(ROCPROFILER_BUILD_DOCS)
rocprofiler_checkout_git_submodule(
projects/rocprofiler-sdk/source/bin/CMakeLists.txt
+14 -2
View File
@@ -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)
projects/rocprofiler-sdk/source/bin/rocprofv3 → projects/rocprofiler-sdk/source/bin/rocprofv3.sh
+9
View File
@@ -55,6 +55,7 @@ usage() {
echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d <out_dir> -o <file_name> <executable>${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 <out_dir> <executable>${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
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp
+734
View File
@@ -23,16 +23,745 @@
#pragma once
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/external_correlation.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/internal_threading.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <cereal/archives/binary.hpp>
#include <cereal/archives/json.hpp>
#include <cereal/archives/portable_binary.hpp>
#include <cereal/cereal.hpp>
#include <cereal/types/array.hpp>
#include <cereal/types/atomic.hpp>
#include <cereal/types/bitset.hpp>
#include <cereal/types/chrono.hpp>
#include <cereal/types/common.hpp>
#include <cereal/types/complex.hpp>
#include <cereal/types/deque.hpp>
#include <cereal/types/functional.hpp>
#include <cereal/types/list.hpp>
#include <cereal/types/map.hpp>
#include <cereal/types/memory.hpp>
#include <cereal/types/optional.hpp>
#include <cereal/types/polymorphic.hpp>
#include <cereal/types/queue.hpp>
#include <cereal/types/set.hpp>
#include <cereal/types/stack.hpp>
#include <cereal/types/string.hpp>
#include <cereal/types/tuple.hpp>
#include <cereal/types/unordered_map.hpp>
#include <cereal/types/unordered_set.hpp>
#include <cereal/types/utility.hpp>
#include <cereal/types/variant.hpp>
#include <cereal/types/vector.hpp>
#include <string>
#include <string_view>
#include <vector>
#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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_context_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_agent_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, hsa_agent_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_queue_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_counter_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_correlation_id_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(internal);
ROCP_SDK_SAVE_DATA_VALUE("external", external.value);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, const hsa_queue_t& data)
{
ar(make_nvp("queue_id", data.id));
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data)
{
ar(make_nvp("queue_id", *data.queue));
}
template <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data)
{
ar(make_nvp("queue_id", *data.queue));
}
template <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_marker_api_retval_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(int64_t_retval);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_hip_api_retval_t data)
{
ROCP_SDK_SAVE_DATA_FIELD(hipError_t_retval);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT, typename Tp>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, HSA_MEMORYPROPERTY data)
{
ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable);
ROCP_SDK_SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, HSA_ENGINE_VERSION data)
{
ROCP_SDK_SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA);
ROCP_SDK_SAVE_DATA_BITFIELD("uCodeRes", uCodeRes);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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<std::remove_pointer_t<decltype(value)>>;
auto vec = std::vector<value_type>{};
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT, typename EnumT, typename ValueT>
void
save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info<EnumT, ValueT>& data)
@@ -56,3 +785,8 @@ save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info_impl<EnumT, ValueT
ar(cereal::make_nvp("operations", _ops));
}
} // namespace cereal
#undef ROCP_SDK_SAVE_DATA_FIELD
#undef ROCP_SDK_SAVE_DATA_VALUE
#undef ROCP_SDK_SAVE_DATA_CSTR
#undef ROCP_SDK_SAVE_DATA_BITFIELD
projects/rocprofiler-sdk/source/lib/common/container/ring_buffer.hpp
+2 -1
View File
@@ -281,7 +281,8 @@ ring_buffer::retrieve() const
template <typename Tp>
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();
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt
+25 -5
View File
@@ -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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp
+115
View File
@@ -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 <fmt/format.h>
namespace rocprofiler
{
namespace tool
{
using float_type = double;
using stats_data_t = statistics<uint64_t, float_type>;
template <typename Tp, domain_type DomainT>
struct buffered_output
{
using ring_buffer_type = rocprofiler::common::container::ring_buffer<Tp>;
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<Tp> element_data = {};
stats_data_t stats = {};
private:
bool enabled = false;
};
template <typename Tp, domain_type DomainT>
buffered_output<Tp, DomainT>::buffered_output(bool _enabled)
: enabled{_enabled}
{}
template <typename Tp, domain_type DomainT>
void
buffered_output<Tp, DomainT>::flush()
{
if(!enabled) return;
flush_tmp_buffer<ring_buffer_type>(buffer_type_v);
}
template <typename Tp, domain_type DomainT>
void
buffered_output<Tp, DomainT>::read()
{
if(!enabled) return;
flush();
element_data = get_buffer_elements(read_tmp_file<ring_buffer_type>(buffer_type_v));
}
template <typename Tp, domain_type DomainT>
void
buffered_output<Tp, DomainT>::clear()
{
if(!enabled) return;
element_data.clear();
}
template <typename Tp, domain_type DomainT>
void
buffered_output<Tp, DomainT>::destroy()
{
if(!enabled) return;
clear();
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<ring_buffer_type>(buffer_type_v);
_tmp_buf->destroy();
delete _tmp_buf;
delete _tmp_file;
}
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp
+63 -10
View File
@@ -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<std::string>{};
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<std::string_view>{"CSV", "JSON"};
for(const auto& itr : entries)
{
LOG_IF(FATAL, supported_formats.count(itr) == 0)
<< "Unsupported output format type: " << itr;
}
}
std::vector<output_key>
output_keys(std::string _tag)
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp
+2 -1
View File
@@ -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<std::string> kernel_names = {};
std::set<std::string> counters = {};
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp
+29 -19
View File
@@ -38,27 +38,36 @@ namespace tool
{
namespace csv
{
template <typename TupleT, size_t... Idx>
struct numerical_formatter
{
template <typename Tp>
std::ostream& operator()(std::ostream& ofs, const Tp& _val) const
{
using value_type = common::mpl::unqualified_type_t<Tp>;
if constexpr(std::is_floating_point<value_type>::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 <typename FmtT = numerical_formatter, typename TupleT, size_t... Idx>
std::ostream&
write_csv_entry(std::ostream& ofs, TupleT&& _data, std::index_sequence<Idx...>)
{
auto _write = [&ofs](size_t idx, auto&& _val) {
using value_type = common::mpl::unqualified_type_t<decltype(_val)>;
if(idx > 0) ofs << ",";
if constexpr(std::is_floating_point<value_type>::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_type>::value) ofs << "\"";
ofs << _val;
if constexpr(common::mpl::is_string_type<value_type>::value) ofs << "\"";
}
if constexpr(common::mpl::is_string_type<value_type>::value) ofs << "\"";
FmtT{}(ofs, _val) << _val;
if constexpr(common::mpl::is_string_type<value_type>::value) ofs << "\"";
};
(_write(Idx, std::get<Idx>(_data)), ...);
@@ -70,21 +79,22 @@ struct csv_encoder
{
static constexpr auto columns = NumCols;
template <typename... Args,
template <typename FmtT = numerical_formatter,
typename... Args,
typename Tp = void,
std::enable_if_t<sizeof...(Args) == columns, int> = 0>
static auto write_row(std::ostream& ofs, Args&&... args)
{
write_csv_entry(
write_csv_entry<FmtT>(
ofs, std::make_tuple(std::forward<Args>(args)...), std::make_index_sequence<columns>{});
return csv_encoder<columns>{};
}
template <typename Tp, size_t N>
template <typename FmtT = numerical_formatter, typename Tp, size_t N>
static auto write_row(std::ostream& ofs, const std::array<Tp, N>& arr)
{
static_assert(N == columns, "Error! too many/few args passed");
write_csv_entry(ofs, arr, std::make_index_sequence<columns>{});
write_csv_entry<FmtT>(ofs, arr, std::make_index_sequence<columns>{});
return csv_encoder<columns>{};
}
};
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp
+87
View File
@@ -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 <utility>
namespace
{
template <domain_type DomainT>
struct domain_type_name;
#define DEFINE_BUFFER_TYPE_NAME(ENUM_VALUE, COLUMN_NAME, FILENAME) \
template <> \
struct domain_type_name<domain_type::ENUM_VALUE> \
{ \
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 <size_t Idx, size_t... TailIdx>
std::string_view
get_domain_file_name(domain_type _buffer_type, std::index_sequence<Idx, TailIdx...>)
{
if(static_cast<size_t>(_buffer_type) == Idx)
return domain_type_name<static_cast<domain_type>(Idx)>::filename;
if constexpr(sizeof...(TailIdx) > 0)
return get_domain_file_name(_buffer_type, std::index_sequence<TailIdx...>{});
return std::string_view{};
}
template <size_t Idx, size_t... IdxTail>
std::string_view
get_domain_column_name(domain_type buffer_type, std::index_sequence<Idx, IdxTail...>)
{
if(static_cast<size_t>(buffer_type) == Idx)
return domain_type_name<static_cast<domain_type>(Idx)>::column_name;
if constexpr(sizeof...(IdxTail) > 0)
return get_domain_column_name(buffer_type, std::index_sequence<IdxTail...>{});
return std::string_view{};
}
} // namespace
std::string_view
get_domain_file_name(domain_type _buffer_type)
{
constexpr auto buffer_type_last_v = static_cast<size_t>(domain_type::LAST);
return get_domain_file_name(_buffer_type, std::make_index_sequence<buffer_type_last_v>{});
}
std::string_view
get_domain_column_name(domain_type buffer_type)
{
constexpr auto last_v = static_cast<size_t>(domain_type::LAST);
return get_domain_column_name(buffer_type, std::make_index_sequence<last_v>{});
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp
+43
View File
@@ -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 <string_view>
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);
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp
+497 -257
View File
@@ -23,11 +23,14 @@
#include "generateCSV.hpp"
#include "csv.hpp"
#include "helper.hpp"
#include "lib/rocprofiler-sdk-tool/config.hpp"
#include "statistics.hpp"
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <iomanip>
#include <string_view>
#include <utility>
namespace rocprofiler
@@ -36,44 +39,171 @@ namespace tool
{
namespace
{
using float_type = double;
using stats_data_t = statistics<uint64_t, float_type>;
using stats_map_t = std::map<std::string_view, stats_data_t>;
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<float_type>(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 <typename Tp>
std::ostream& operator()(std::ostream& ofs, const Tp& _val) const
{
using value_type = common::mpl::unqualified_type_t<Tp>;
if constexpr(std::is_floating_point<value_type>::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<Tp, percentage>::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<std::pair<std::string_view, stats_data_t>>{};
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<stats_formatter>(_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<rocprofiler_agent_v0_t>& data)
generate_csv(tool_table* /*tool_functions*/, std::vector<rocprofiler_agent_v0_t>& 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<rocprofiler_agent_v0_t>& 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<rocprofiler_agent_v0_t>& 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<kernel_dispatch_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& 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<hip_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& 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<hsa_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& 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<memory_copy_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& 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<marker_api_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& 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<rocprofiler_tool_counter_collection_record_t>& 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<std::string, uint64_t>{};
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<std::string, uint64_t>{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<counter_collection_ring_buffer_t>& 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<std::string, uint64_t>{};
for(const auto& count : record->profiler_record)
{
auto rec = static_cast<rocprofiler_record_counter_t>(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<std::string, uint64_t>{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<scratch_memory_ring_buffer_t>& data)
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& 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<domain_type, stats_data_t>& 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<stats_formatter>(_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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp
+32 -18
View File
@@ -23,6 +23,7 @@
#pragma once
#include "helper.hpp"
#include "statistics.hpp"
#include <rocprofiler-sdk/agent.h>
@@ -30,28 +31,41 @@ namespace rocprofiler
{
namespace tool
{
using float_type = double;
using stats_data_t = statistics<uint64_t, float_type>;
void
generate_csv(tool_table* tool_functions, std::vector<rocprofiler_agent_v0_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<kernel_dispatch_ring_buffer_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_tool_counter_collection_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<hip_ring_buffer_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<hsa_ring_buffer_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<memory_copy_ring_buffer_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<marker_api_ring_buffer_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<counter_collection_ring_buffer_t>& data);
void
generate_csv(tool_table* tool_functions, std::vector<scratch_memory_ring_buffer_t>& data);
generate_csv(tool_table* tool_functions, std::unordered_map<domain_type, stats_data_t>& data);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp
+139
View File
@@ -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 <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <utility>
namespace rocprofiler
{
namespace tool
{
void
write_json(tool_table* tool_functions,
uint64_t pid,
std::vector<rocprofiler_agent_v0_t> agent_data,
std::vector<rocprofiler_tool_counter_info_t> counter_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* hip_api_deque,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_deque,
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_deque,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_deque,
std::deque<rocprofiler_tool_counter_collection_record_t>* counter_collection_deque,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_deque,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp
+45
View File
@@ -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<rocprofiler_agent_v0_t> agent_data,
std::vector<rocprofiler_tool_counter_info_t> counter_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* hip_api_deque,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_deque,
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_deque,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_deque,
std::deque<rocprofiler_tool_counter_collection_record_t>* counter_collection_deque,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_deque,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_api_deque);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp
+5 -123
View File
@@ -24,6 +24,7 @@
#include "config.hpp"
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <atomic>
#include <iostream>
@@ -33,133 +34,14 @@
#include <unordered_set>
#include <utility>
rocprofiler_tool_buffer_name_info_t
::rocprofiler::sdk::buffer_name_info_t<std::string_view>
get_buffer_id_names()
{
static auto supported = std::unordered_set<rocprofiler_buffer_tracing_kind_t>{
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<rocprofiler_tool_buffer_name_info_t*>(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<rocprofiler_tool_buffer_name_info_t*>(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<void*>(data)),
"query buffer failed");
}
return 0;
};
ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kinds(tracing_kind_cb,
static_cast<void*>(&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<std::string_view>
get_callback_id_names()
{
static auto supported = std::unordered_set<rocprofiler_callback_tracing_kind_t>{
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<rocprofiler_tool_callback_name_info_t*>(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<rocprofiler_tool_callback_name_info_t*>(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<void*>(data)),
"query callback failed");
}
return 0;
};
ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kinds(tracing_kind_cb,
static_cast<void*>(&cb_name_info)),
"iterate_callback failed");
return cb_name_info;
return ::rocprofiler::sdk::get_callback_tracing_names();
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp
+264 -79
View File
@@ -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 <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/cxx/serialization.hpp>
#include <amd_comgr/amd_comgr.h>
#include <hsa/amd_hsa_kernel_code.h>
@@ -39,6 +46,8 @@
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <hsa/hsa_ven_amd_loader.h>
#include <glog/logging.h>
#include <cxxabi.h>
#include <sys/syscall.h>
#include <sys/types.h>
@@ -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<rocprofiler_buffer_tracing_kind_t, const char*>;
std::unordered_map<rocprofiler_buffer_tracing_kind_t, std::string>;
using rocprofiler_tool_buffer_kind_operation_names_t =
std::unordered_map<rocprofiler_buffer_tracing_kind_t,
std::unordered_map<uint32_t, const char*>>;
std::unordered_map<uint32_t, std::string>>;
using marker_message_map_t = std::unordered_map<uint64_t, std::string>;
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<rocprofiler_callback_tracing_kind_t, const char*>;
using rocprofiler_tool_callback_kind_operation_names_t =
std::unordered_map<rocprofiler_callback_tracing_kind_t,
std::unordered_map<uint32_t, const char*>>;
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<rocprofiler_kernel_id_t, kernel_symbol_data>;
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<rocprofiler_counter_dimension_id_t>;
using dimension_info_vec_t = std::vector<rocprofiler_record_dimension_info_t>;
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<rocprofiler_counter_dimension_id_t> dimension_ids = {};
std::vector<rocprofiler_record_dimension_info_t> dimension_info = {};
};
rocprofiler_tool_buffer_name_info_t
rocprofiler::sdk::buffer_name_info_t<std::string_view>
get_buffer_id_names();
rocprofiler_tool_callback_name_info_t
::rocprofiler::sdk::callback_name_info_t<std::string_view>
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<uint64_t, std::string>
get_callback_roctx_msg();
rocprofiler_timestamp_t start_timestamp;
rocprofiler_timestamp_t end_timestamp;
std::vector<kernel_symbol_data>
get_kernel_symbol_data();
std::vector<rocprofiler_callback_tracing_code_object_load_data_t>
get_code_object_data();
std::vector<rocprofiler_tool_counter_info_t>
get_tool_counter_info();
std::vector<rocprofiler_record_dimension_info_t>
get_tool_counter_dimension_info();
enum tracing_marker_kind
{
MARKER_API_CORE = 0,
MARKER_API_CONTROL,
MARKER_API_NAME,
MARKER_API_LAST,
};
template <size_t CommonV>
struct marker_tracing_kind_conversion;
#define MAP_TRACING_KIND_CONVERSION(COMMON, CALLBACK, BUFFERED) \
template <> \
struct marker_tracing_kind_conversion<COMMON> \
{ \
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 <typename TracingKindT, size_t Idx, size_t... Tail>
auto
convert_marker_tracing_kind(TracingKindT val, std::index_sequence<Idx, Tail...>)
{
if(marker_tracing_kind_conversion<Idx>{} == val)
{
return marker_tracing_kind_conversion<Idx>{}.convert(val);
}
if constexpr(sizeof...(Tail) > 0)
return convert_marker_tracing_kind(val, std::index_sequence<Tail...>{});
return marker_tracing_kind_conversion<MARKER_API_LAST>{}.convert(val);
}
template <typename TracingKindT>
auto
convert_marker_tracing_kind(TracingKindT val)
{
return convert_marker_tracing_kind(val, std::make_index_sequence<MARKER_API_LAST>{});
}
struct rocprofiler_tool_dimension_pos_t
{
uint64_t dimension_id;
size_t instance;
template <typename ArchiveT>
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 <typename ArchiveT>
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<rocprofiler_record_counter_t, 128> 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<rocprofiler_tool_record_counter_t> 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 <typename ArchiveT>
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<rocprofiler_buffer_tracing_hip_api_record_t>;
using hsa_ring_buffer_t =
rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_hsa_api_record_t>;
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<rocprofiler_buffer_tracing_memory_copy_record_t>;
using counter_collection_buffer_t =
rocprofiler::common::container::ring_buffer<rocprofiler_record_counter_t>;
using marker_api_ring_buffer_t =
rocprofiler::common::container::ring_buffer<rocprofiler_tool_marker_record_t>;
using counter_collection_ring_buffer_t =
rocprofiler::common::container::ring_buffer<rocprofiler_tool_counter_collection_record_t>;
using scratch_memory_ring_buffer_t =
rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_scratch_memory_record_t>;
namespace rocprofiler
{
namespace tool
{
template <typename Tp, domain_type DomainT>
struct buffered_output;
}
} // namespace rocprofiler
using hip_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_hip_api_record_t,
domain_type::HIP>;
using hsa_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_hsa_api_record_t,
domain_type::HSA>;
using kernel_dispatch_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_kernel_dispatch_record_t,
domain_type::KERNEL_DISPATCH>;
using memory_copy_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_memory_copy_record_t,
domain_type::MEMORY_COPY>;
using marker_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_marker_api_record_t,
domain_type::MARKER>;
using counter_collection_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_tool_counter_collection_record_t,
domain_type::COUNTER_COLLECTION>;
using scratch_memory_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler_buffer_tracing_scratch_memory_record_t,
domain_type::SCRATCH_MEMORY>;
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 <typename Tp, template <typename, typename...> class ContainerT, typename... ParamsT>
ContainerT<Tp>
get_buffer_elements(ContainerT<rocprofiler::common::container::ring_buffer<Tp>, ParamsT...>&& data)
{
auto ret = ContainerT<Tp>{};
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 <typename ArchiveT>
void
save(ArchiveT& ar, const kernel_symbol_data& data)
{
cereal::save(ar, static_cast<const rocprofiler_kernel_symbol_data_t&>(data));
SAVE_DATA_FIELD(formatted_kernel_name);
SAVE_DATA_FIELD(demangled_kernel_name);
SAVE_DATA_FIELD(truncated_kernel_name);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, const rocprofiler_tool_counter_info_t& data)
{
SAVE_DATA_FIELD(agent_id);
cereal::save(ar, static_cast<const rocprofiler_counter_info_v0_t&>(data));
SAVE_DATA_FIELD(dimension_ids);
}
#undef SAVE_DATA_FIELD
} // namespace cereal
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/main.c
+143
View File
@@ -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 <dlfcn.h>
#include <stdbool.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <unistd.h>
//
// 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);
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp
+16 -8
View File
@@ -24,6 +24,7 @@
#include "config.hpp"
#include "lib/common/logging.hpp"
#include <fmt/core.h>
#include <fmt/format.h>
namespace rocprofiler
@@ -32,8 +33,8 @@ namespace tool
{
namespace fs = common::filesystem;
std::pair<std::ostream*, void (*)(std::ostream*&)>
get_output_stream(const std::string& fname, const std::string& ext)
std::pair<std::ostream*, output_stream_dtor_t>
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) {
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp
+9 -9
View File
@@ -41,8 +41,10 @@ namespace rocprofiler
{
namespace tool
{
std::pair<std::ostream*, void (*)(std::ostream*&)>
get_output_stream(const std::string& fname, const std::string& ext = ".csv");
using output_stream_dtor_t = void (*)(std::ostream*&);
std::pair<std::ostream*, output_stream_dtor_t>
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 <size_t N>
@@ -80,7 +80,7 @@ output_file::output_file(std::string name,
std::array<std::string_view, N>&& 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)
{
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.cpp
+36
View File
@@ -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 <fmt/format.h>
#include <utility>
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)));
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp
+113
View File
@@ -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 <fmt/format.h>
#include <deque>
#include <mutex>
#include <string>
#include <tuple>
#include <utility>
template <typename Tp>
using ring_buffer_t = rocprofiler::common::container::ring_buffer<Tp>;
std::string
compose_tmp_file_name(domain_type buffer_type);
template <typename Tp>
std::tuple<Tp*, tmp_file*>
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 <typename Tp>
void
offload_buffer(domain_type type)
{
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<Tp>(type);
auto _lk = std::lock_guard<std::mutex>(_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 <typename Tp>
void
write_ring_buffer(Tp _v, domain_type type)
{
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<ring_buffer_t<Tp>>(type);
auto* ptr = _tmp_buf->request(false);
if(ptr == nullptr)
{
offload_buffer<ring_buffer_t<Tp>>(type);
ptr = _tmp_buf->request(false);
CHECK(ptr != nullptr);
}
*ptr = std::move(_v);
}
template <typename Tp>
void
flush_tmp_buffer(domain_type type)
{
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<Tp>(type);
if(!_tmp_buf->is_empty()) offload_buffer<Tp>(type);
}
template <typename Tp>
std::deque<Tp>
read_tmp_file(domain_type type)
{
auto _data = std::deque<Tp>{};
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<Tp>(type);
auto _lk = std::lock_guard<std::mutex>{_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;
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
+582 -768
View File
File diff suppressed because it is too large Load Diff
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp
+8 -1
View File
@@ -191,7 +191,14 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args... args)
constexpr auto external_corr_id_domain_idx =
hip_domain_info<TableIdx>::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>(args)...);
if constexpr(!std::is_void<RetT>::value)
return _ret;
else
return;
}
constexpr auto ref_count = 2;
auto thr_id = common::get_tid();
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp
-2
View File
@@ -306,8 +306,6 @@ hsa_api_impl<TableIdx, OpIdx>::functor(Args... args)
constexpr auto external_corr_id_domain_idx =
hsa_domain_info<TableIdx>::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>(args)...);
projects/rocprofiler-sdk/tests/common/CMakeLists.txt
+5 -4
View File
@@ -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 $<BUILD_INTERFACE:CEREAL_THREAD_SAFE=1>)
@@ -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})
projects/rocprofiler-sdk/tests/common/serialization.hpp
+2 -723
View File
@@ -23,726 +23,5 @@
#pragma once
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/external_correlation.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/internal_threading.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <cereal/archives/binary.hpp>
#include <cereal/archives/json.hpp>
#include <cereal/archives/portable_binary.hpp>
#include <cereal/cereal.hpp>
#include <cereal/types/array.hpp>
#include <cereal/types/atomic.hpp>
#include <cereal/types/bitset.hpp>
#include <cereal/types/chrono.hpp>
#include <cereal/types/common.hpp>
#include <cereal/types/complex.hpp>
#include <cereal/types/deque.hpp>
#include <cereal/types/functional.hpp>
#include <cereal/types/list.hpp>
#include <cereal/types/map.hpp>
#include <cereal/types/memory.hpp>
#include <cereal/types/optional.hpp>
#include <cereal/types/polymorphic.hpp>
#include <cereal/types/queue.hpp>
#include <cereal/types/set.hpp>
#include <cereal/types/stack.hpp>
#include <cereal/types/string.hpp>
#include <cereal/types/tuple.hpp>
#include <cereal/types/unordered_map.hpp>
#include <cereal/types/unordered_set.hpp>
#include <cereal/types/utility.hpp>
#include <cereal/types/variant.hpp>
#include <cereal/types/vector.hpp>
#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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_context_id_t data)
{
SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_agent_id_t data)
{
SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, hsa_agent_t data)
{
SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_queue_id_t data)
{
SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_counter_id_t data)
{
SAVE_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_correlation_id_t data)
{
SAVE_DATA_FIELD(internal);
SAVE_DATA_VALUE("external", external.value);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_dim3_t data)
{
SAVE_DATA_FIELD(x);
SAVE_DATA_FIELD(y);
SAVE_DATA_FIELD(z);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data)
{
SAVE_DATA_FIELD(uint64_t_retval);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, const hsa_queue_t& data)
{
ar(make_nvp("queue_id", data.id));
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data)
{
ar(make_nvp("queue_id", *data.queue));
}
template <typename ArchiveT>
void
save(ArchiveT& ar, hsa_amd_event_scratch_free_end_t data)
{
ar(make_nvp("queue_id", *data.queue));
SAVE_DATA_FIELD(flags);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data)
{
ar(make_nvp("queue_id", *data.queue));
}
template <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_marker_api_retval_t data)
{
SAVE_DATA_FIELD(int64_t_retval);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_hip_api_retval_t data)
{
SAVE_DATA_FIELD(hipError_t_retval);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT, typename Tp>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_record_counter_t data)
{
SAVE_DATA_FIELD(id);
SAVE_DATA_FIELD(counter_value);
SAVE_DATA_FIELD(dispatch_id);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data)
{
save_buffer_tracing_api_record(ar, data);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, HSA_MEMORYPROPERTY data)
{
SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable);
SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, HSA_ENGINE_VERSION data)
{
SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA);
SAVE_DATA_BITFIELD("uCodeRes", uCodeRes);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
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<std::remove_pointer_t<decltype(value)>>;
auto vec = std::vector<value_type>{};
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 <typename ArchiveT>
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 <rocprofiler-sdk/cxx/serialization.hpp>
projects/rocprofiler-sdk/tests/pytest-packages/pytest_utils/__init__.py
+20
View File
@@ -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)
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt
-3
View File
@@ -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)
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt
+10 -9
View File
@@ -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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
${CMAKE_CURRENT_BINARY_DIR}/input.txt -T -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1
-o pmc1 $<TARGET_FILE:vector-ops>)
-o pmc1 --output-format csv,json $<TARGET_FILE:vector-ops>)
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}")
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/conftest.py
+31
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini → projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/pytest.ini
+1 -3
View File
@@ -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
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py
+54
View File
@@ -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)
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt
+5 -7
View File
@@ -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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-cc -o
out $<TARGET_FILE:simple-transpose>)
${CMAKE_CURRENT_BINARY_DIR}/input.txt --output-format CSV,JSON -d
${CMAKE_CURRENT_BINARY_DIR}/%argt%-cc -o out $<TARGET_FILE:simple-transpose>)
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
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py → projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/conftest.py
+15 -12
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py
+2 -2
View File
@@ -1,5 +1,5 @@
import pandas as pd
import os
#!/usr/bin/env python3
import sys
import pytest
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt
+1 -4
View File
@@ -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
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/conftest.py
-1
View File
@@ -2,7 +2,6 @@
import csv
import pytest
import pandas as pd
def pytest_addoption(parser):
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/list_metrics/validate.py
+2 -2
View File
@@ -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"] != ""
projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt
+9 -8
View File
@@ -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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace --kernel-trace -d
${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $<TARGET_FILE:multiqueue_testapp>)
${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format json:csv
$<TARGET_FILE:multiqueue_testapp>)
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
projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/conftest.py
+16
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py
+89
View File
@@ -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)
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt
+38 -16
View File
@@ -16,6 +16,15 @@ add_test(
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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
$<TARGET_FILE:hip-in-libraries>)
add_test(
NAME rocprofv3-test-trace-hip-in-libraries-json-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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 $<TARGET_FILE:hip-in-libraries>)
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}")
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/conftest.py
+17
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py
+143
View File
@@ -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*, float const*, float const*, int, int)",
"void (anonymous namespace)::divide_kernel<float>(float*, float const*, float const*, int, int)",
"void (anonymous namespace)::multiply_kernel<float>(float*, float const*, float const*, int, int)",
"void (anonymous namespace)::subtract_kernel<float>(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,
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt
+25 -20
View File
@@ -20,31 +20,36 @@ add_test(
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace -i
${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace
-o pmc3 $<TARGET_FILE:simple-transpose>)
-o pmc3 --output-format JSON:CSV $<TARGET_FILE:simple-transpose>)
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()
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py
+176 -5
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt
+6 -4
View File
@@ -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
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
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py
+12 -33
View File
@@ -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__":
projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt
+14 -12
View File
@@ -15,7 +15,7 @@ add_test(
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> -M --hsa-trace --kernel-trace
--memory-copy-trace --marker-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o
out $<TARGET_FILE:simple-transpose>)
out --output-format csv,json $<TARGET_FILE:simple-transpose>)
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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --sys-trace -d
${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out
${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out --output-format csv,json
$<TARGET_FILE:simple-transpose>)
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
projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py
+16
View File
@@ -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)))
projects/rocprofiler-sdk/tests/rocprofv3/tracing/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py
+136 -2
View File
@@ -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)
projects/rocprofiler-sdk/tests/tools/CMakeLists.txt
+3 -2
View File
@@ -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"