Adding JSON support (#860)
* Adding json support minor bugs Fixing tests Fixing formatting issues Fixing test test fix Misc testing fixes Use rocprofiler/cxx/name_info in rocprofiler-sdk-tool fixes to reduce the Json file size Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/helper.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/tool.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Update source/lib/rocprofiler-sdk-tool/generateJSON.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> misc fixes Removing int cast for JSON tests formatting removing a condition test on Navi3 adding debug info Misc fix * CSV updates - fix stats - numerical formatter support for customizing write_csv_entry - misc formatting - get_marker_stats_file * Misc tests/rocprofv3/counter-collection/input2 fixes - rocprofiler_configure_pytest_files in rocprofv3/counter-collection/input2 - removed state code from merge in rocprofv3/counter-collection/input2 * Tool: "Agent-id" -> "Agent_Id" - consistency * Tool update - remove rocprofiler_tool_marker_record_t - add marker_tracing_kind_conversion - fix memory leak in write_json - minor update to get_output_stream - rework handling of marker records * Update tests/pytest-packages/pytest_utils/__init__.py - add collapse_dict_list function for converting a dictionary value that is a list of length one into a directly mapped value * Update tests/rocprofv3/**/conftest.py - use collapse_dict_list when reading in JSONs * Update tests/rocprofv3/counter-collection/input1/validate.py - relax testing requirements gfx1102 (AQLProfile bugs) - in addition to relaxed testing requirements for gfx1101 * Update tests/rocprofv3/tracing/validate.py - fix removal of PID in every marker record * Update tests/rocprofv3/tracing-plus-cc - remove test design that relies on iterating subdirectories * Wrapper around __libc_start_main - Ensures finalization happens before main returns - Update tests/rocprofv3/tracing/validate.py - wrapper around __libc_start_main changed roctx calls * Combine include/rocprofiler-sdk/cxx/serialization.hpp and include/rocprofiler-sdk/external/serialization.hpp - tests/common/serialization.hpp simply includes include/rocprofiler-sdk/cxx/serialization.hpp now * Update lib/rocprofiler-sdk/hip/hip.cpp - tracing function immediately returns when fini_status is non-zero * Update lib/rocprofiler-sdk/hsa/hsa.cpp - remove logging of tracing function when fini_status is non-zero * Update lib/rocprofiler-sdk-tool/CMakeLists.txt - remove rocprofv3_trigger_list_metrics.cpp from TOOL_SOURCES * Update tests/rocprofv3/tracing-plus-cc/CMakeLists.txt - fix depends * Domain statistics * Update tests/rocprofv3/tracing-plus-cc/CMakeLists.txt - do not set ROCP_LOG_LEVEL in env * Remove erroneous <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>
This commit is contained in:
Jonathan R. Madsen
committed by
GitHub
GitHub
parent
83e2d7d8af
commit
92b7326910
@@ -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)
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
|
||||
@@ -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();
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
@@ -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)
|
||||
|
||||
@@ -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 = {};
|
||||
|
||||
@@ -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>{};
|
||||
}
|
||||
};
|
||||
|
||||
@@ -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>{});
|
||||
}
|
||||
@@ -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);
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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();
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
@@ -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) {
|
||||
|
||||
@@ -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)
|
||||
{
|
||||
|
||||
@@ -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)));
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -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();
|
||||
|
||||
@@ -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)...);
|
||||
|
||||
Reference in New Issue
Block a user