diff --git a/.gitignore b/.gitignore index 8de3e86695..63c2c605e7 100644 --- a/.gitignore +++ b/.gitignore @@ -35,7 +35,7 @@ /compile_commands.json /build* /.cache -/.vscode +*.vscode # Github Workflows /.github diff --git a/.gitmodules b/.gitmodules index 2f00d4ebc9..2b7f01fec5 100644 --- a/.gitmodules +++ b/.gitmodules @@ -28,3 +28,6 @@ [submodule "external/yaml-cpp"] path = external/yaml-cpp url = https://github.com/jbeder/yaml-cpp.git +[submodule "external/json"] + path = external/json + url = https://github.com/nlohmann/json.git diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index e954613f96..0a804ac764 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -59,6 +59,7 @@ rocprofiler_add_interface_library(rocprofiler-sdk-dw "ElfUtils dw library" INTER rocprofiler_add_interface_library(rocprofiler-sdk-elfio "ELFIO header-only C++ library" INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-yaml-cpp "YAML CPP Parser" INTERNAL) +rocprofiler_add_interface_library(rocprofiler-sdk-json "nlohmann json" INTERNAL) # # interface for libraries (ROCm-specific) diff --git a/external/CMakeLists.txt b/external/CMakeLists.txt index c6a4dc3a9a..bb26a4e677 100644 --- a/external/CMakeLists.txt +++ b/external/CMakeLists.txt @@ -209,3 +209,17 @@ endif() # OTF2 add_subdirectory(otf2) + +rocprofiler_checkout_git_submodule( + RECURSIVE + RELATIVE_PATH external/json + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + REPO_URL https://github.com/nlohmann/json.git + REPO_BRANCH "develop") + +add_subdirectory(json) + +# target_link_libraries(rocprofiler-sdk-json INTERFACE $) +target_include_directories( + rocprofiler-sdk-json SYSTEM + INTERFACE $) diff --git a/external/json b/external/json new file mode 160000 index 0000000000..e41905fcb0 --- /dev/null +++ b/external/json @@ -0,0 +1 @@ +Subproject commit e41905fcb0938a7502d25ea5242c6b41396e21b0 diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index d833ae1bee..2c7bb3edd5 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -30,7 +30,6 @@ add_subdirectory(code_object_tracing) add_subdirectory(counter_collection) add_subdirectory(intercept_table) add_subdirectory(code_object_isa_decode) -add_subdirectory(advanced_thread_trace) add_subdirectory(external_correlation_id_request) add_subdirectory(pc_sampling) if(ROCPROFILER_BUILD_OPENMP_TESTS) diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h index f1db19dfbd..5f46635f4a 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h @@ -48,7 +48,13 @@ ROCPROFILER_EXTERN_C_INIT * @param [in] num_parameters Number of parameters. Zero is allowed. * @param [in] agent_id agent to configure profiling on. * @param [in] shader_callback Callback fn where the collected data will be sent to. - * @param [in] callback_userdata Passed back to user. + * @param [in] callback_userdata Passed back to user in shader_callback. + * @return ::rocprofiler_status_t + * @retval ROCPROFILER_STATUS_SUCCESS on success + * @retval ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED for configuration locked + * @retval ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID for conflicting configurations in the same ctx + * @retval ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND for invalid context id + * @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT for invalid rocprofiler_att_parameter_t */ rocprofiler_status_t rocprofiler_configure_agent_thread_trace_service( @@ -57,7 +63,7 @@ rocprofiler_configure_agent_thread_trace_service( size_t num_parameters, rocprofiler_agent_id_t agent_id, rocprofiler_att_shader_data_callback_t shader_callback, - void* callback_userdata) ROCPROFILER_API; + rocprofiler_user_data_t callback_userdata) ROCPROFILER_API; /** @} */ diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h index f5da72ccc5..7a96cd6abe 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h @@ -68,104 +68,12 @@ typedef struct * @param [in] data_size Number of bytes in "data" * @param [in] userdata Passed back to user from rocprofiler_att_dispatch_callback_t() */ -typedef void (*rocprofiler_att_shader_data_callback_t)(int64_t shader_engine_id, +typedef void (*rocprofiler_att_shader_data_callback_t)(rocprofiler_agent_id_t agent, + int64_t shader_engine_id, void* data, size_t data_size, rocprofiler_user_data_t userdata); -/** - * @brief Callback for rocprofiler to parsed ATT data. - * The caller must copy a desired instruction on isa_instruction and source_reference, - * while obeying the max length passed by the caller. - * If the caller's length is insufficient, then this function writes the minimum sizes to isa_size - * and source_size and returns ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES. - * If call returns ROCPROFILER_STATUS_SUCCESS, isa_size and source_size are written with bytes used. - * @param[out] isa_instruction Where to copy the ISA line to. - * @param[out] isa_memory_size (Auto) The number of bytes to next instruction. 0 for custom ISA. - * @param[inout] isa_size Size of returned ISA string. - * @param[in] marker_id The generated ATT marker for given codeobject ID. - * @param[in] offset The offset from base vaddr for given codeobj ID. - * If marker_id == 0, this parameter is raw virtual address with no codeobj ID information. - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - * @retval ROCPROFILER_STATUS_SUCCESS on success. - * @retval ROCPROFILER_STATUS_ERROR on generic error. - * @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT for invalid offset or invalid marker_id. - * @retval ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES for insufficient isa_size or source_size. - */ -typedef rocprofiler_status_t (*rocprofiler_att_parser_isa_callback_t)(char* isa_instruction, - uint64_t* isa_memory_size, - uint64_t* isa_size, - uint64_t marker_id, - uint64_t offset, - void* userdata); - -/** - * @brief Callback for the ATT parser to retrieve Shader Engine data. - * Returns the amount of data filled. If no more data is available, then callback return 0 - * If the space available in the buffer is less than required for parsing the full data, - * the full data is transfered over multiple calls. - * When all data has been transfered from current shader_engine_id, the caller has the option to - * 1) Return -1 on shader_engine ID and parsing terminates - * 2) Move to the next shader engine. - * @param[out] shader_engine_id The ID of given shader engine. - * @param[out] buffer The buffer to fill up with SE data. - * @param[out] buffer_size The space available in the buffer. - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - * @returns Number of bytes remaining in shader engine. - * @retval 0 if no more SE data is available. Parsing will stop. - * @retval ret Where 0 > ret > buffer_size for partially filled buffer, and caller moves over to - * next SE. - * @retval buffer_size if the buffer does not hold enough data for the current shader engine. - */ -typedef uint64_t (*rocprofiler_att_parser_se_data_callback_t)(int* shader_engine_id, - uint8_t** buffer, - uint64_t* buffer_size, - void* userdata); - -typedef enum -{ - ROCPROFILER_ATT_PARSER_DATA_TYPE_ISA = 0, - ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY, -} rocprofiler_att_parser_data_type_t; - -typedef struct -{ - uint64_t marker_id; - uint64_t offset; - uint64_t hitcount; - uint64_t latency; -} rocprofiler_att_data_type_isa_t; - -typedef struct -{ - uint64_t marker_id; - uint64_t offset; - uint64_t timestamp : 63; - uint64_t enabled : 1; -} rocprofiler_att_data_type_occupancy_t; - -/** - * @brief Callback for rocprofiler to return traces back to rocprofiler. - * @param[in] att_data A datapoint retrieved from thread_trace - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - */ -typedef void (*rocprofiler_att_parser_trace_callback_t)(rocprofiler_att_parser_data_type_t type, - void* att_data, - void* userdata); - -/** - * @brief Iterate over all event coordinates for a given agent_t and event_t. - * @param[in] se_data_callback Callback to return shader engine data from. - * @param[in] trace_callback Callback where the trace data is returned to. - * @param[in] isa_callback Callback to return ISA lines. - * @param[in] userdata Userdata passed back to caller via callback. - */ -rocprofiler_status_t -rocprofiler_att_parse_data(rocprofiler_att_parser_se_data_callback_t se_data_callback, - rocprofiler_att_parser_trace_callback_t trace_callback, - rocprofiler_att_parser_isa_callback_t isa_callback, - void* userdata) ROCPROFILER_API; - /** @} */ ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h index 7e156cdfd0..1056de1655 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -45,15 +45,23 @@ typedef enum /** * @brief Callback to be triggered every kernel dispatch, indicating to start and/or stop ATT + * @param [in] agent_id agent_id. + * @param [in] queue_id queue_id. + * @param [in] correlation_id internal correlation id. + * @param [in] kernel_id kernel_id. + * @param [in] dispatch_id dispatch_id. + * @param [in] userdata_config Userdata passed back from + * rocprofiler_configure_dispatch_thread_trace_service. + * @param [out] userdata_shader Userdata to be passed in shader_callback */ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)( + rocprofiler_agent_id_t agent_id, rocprofiler_queue_id_t queue_id, - const rocprofiler_agent_t* agent, rocprofiler_correlation_id_t correlation_id, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t dispatch_id, - rocprofiler_user_data_t* userdata_shader, - void* userdata_config); + void* userdata_config, + rocprofiler_user_data_t* userdata_shader); /** * @brief Enables the advanced thread trace service for dispatch-based tracing. @@ -64,7 +72,14 @@ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)( * @param [in] num_parameters Number of parameters. Zero is allowed. * @param [in] dispatch_callback Control fn which decides when ATT starts/stop collecting. * @param [in] shader_callback Callback fn where the collected data will be sent to. - * @param [in] callback_userdata Passed back to user. + * @param [in] callback_userdata Passed back to user in dispatch_callback. + * @return ::rocprofiler_status_t + * @retval ROCPROFILER_STATUS_SUCCESS on success + * @retval ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED for configuration locked + * @retval ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID for conflicting configurations in the same ctx + * @retval ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND for invalid context id + * @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT for invalid rocprofiler_att_parameter_t + * @retval ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED if already configured */ rocprofiler_status_t rocprofiler_configure_dispatch_thread_trace_service( diff --git a/source/lib/CMakeLists.txt b/source/lib/CMakeLists.txt index 2237677653..a5019423bb 100644 --- a/source/lib/CMakeLists.txt +++ b/source/lib/CMakeLists.txt @@ -9,6 +9,7 @@ add_subdirectory(rocprofiler-sdk) set(CMAKE_INSTALL_DEFAULT_COMPONENT_NAME "tools") add_subdirectory(rocprofiler-sdk-codeobj) add_subdirectory(rocprofiler-sdk-roctx) +add_subdirectory(rocprofiler-sdk-att) add_subdirectory(rocprofiler-sdk-tool) if(ROCPROFILER_BUILD_TESTS) diff --git a/source/lib/rocprofiler-sdk-att/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/CMakeLists.txt new file mode 100755 index 0000000000..9066827fa0 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/CMakeLists.txt @@ -0,0 +1,47 @@ +# +# ATT decoder wrapper library for use by the rocprofv3 tool +# +set(ATT_TOOL_SOURCE_FILES + waitcnt/analysis.cpp + waitcnt/gfx9.cpp + waitcnt/gfx10.cpp + waitcnt/gfx12.cpp + att_lib_wrapper.cpp + wave.cpp + code.cpp + filenames.cpp + occupancy.cpp + wstates.cpp + profile_interface.cpp + dl.cpp) + +find_package( + amd_comgr + REQUIRED + CONFIG + HINTS + ${rocm_version_DIR} + ${ROCM_PATH} + PATHS + ${rocm_version_DIR} + ${ROCM_PATH} + PATH_SUFFIXES + lib/cmake/amd_comgr) + +add_library(rocprofiler-sdk-att-parser STATIC) +target_sources(rocprofiler-sdk-att-parser PRIVATE ${ATT_TOOL_SOURCE_FILES}) + +target_link_libraries( + rocprofiler-sdk-att-parser + PRIVATE rocprofiler-sdk::rocprofiler-sdk-codeobj + rocprofiler-sdk::rocprofiler-sdk-headers + rocprofiler-sdk::rocprofiler-sdk-json + rocprofiler-sdk::rocprofiler-sdk-common-library + amd_comgr + dw) + +add_subdirectory(waitcnt) + +if(ROCPROFILER_BUILD_TESTS) + add_subdirectory(tests) +endif() diff --git a/source/lib/rocprofiler-sdk-att/att_decoder.h b/source/lib/rocprofiler-sdk-att/att_decoder.h new file mode 100644 index 0000000000..038456fdab --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/att_decoder.h @@ -0,0 +1,160 @@ +// MIT License +// +// Copyright (c) 2024 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 + +#ifdef __cplusplus +# include +# include +extern "C" { +#else +# include +# include +#endif + +typedef enum +{ + ROCPROFILER_ATT_DECODER_STATUS_SUCCESS = 0, + ROCPROFILER_ATT_DECODER_STATUS_ERROR, + ROCPROFILER_ATT_DECODER_STATUS_ERROR_OUT_OF_RESOURCES, + ROCPROFILER_ATT_DECODER_STATUS_ERROR_INVALID_ARGUMENT, + ROCPROFILER_ATT_DECODER_STATUS_ERROR_INVALID_SHADER_DATA, + ROCPROFILER_ATT_DECODER_STATUS_LAST +} rocprofiler_att_decoder_status_t; + +typedef enum +{ + ROCPROFILER_ATT_DECODER_INFO_NONE = 0, + ROCPROFILER_ATT_DECODER_INFO_DATA_LOST, + ROCPROFILER_ATT_DECODER_INFO_STITCH_INCOMPLETE, + ROCPROFILER_ATT_DECODER_INFO_LAST +} rocprofiler_att_decoder_info_t; + +typedef enum +{ + ROCPROFILER_ATT_DECODER_TYPE_GFXIP = 0, + ROCPROFILER_ATT_DECODER_TYPE_OCCUPANCY, + ROCPROFILER_ATT_DECODER_TYPE_PERFEVENT, + ROCPROFILER_ATT_DECODER_TYPE_WAVE, + ROCPROFILER_ATT_DECODER_TYPE_INFO, + ROCPROFILER_ATT_DECODER_TYPE_DEBUG, + ROCPROFILER_ATT_DECODER_TYPE_LAST +} rocprofiler_att_decoder_record_type_t; + +typedef struct +{ + size_t addr; + size_t marker_id; +} pcinfo_t; + +typedef struct +{ + pcinfo_t pc; + uint64_t time; + uint8_t se; + uint8_t cu; + uint8_t simd; + uint8_t slot; + uint32_t start : 1; + uint32_t _rsvd : 31; +} att_occupancy_info_v2_t; + +typedef struct +{ + int32_t type; + int32_t duration; +} att_wave_state_t; + +typedef struct +{ + uint32_t category : 8; + uint32_t stall : 24; + int32_t duration; + int64_t time; + pcinfo_t pc; +} att_wave_instruction_t; + +typedef enum +{ + ATT_WAVE_STATE_EMPTY = 0, + ATT_WAVE_STATE_LAST = 5 +} att_waveslot_state_t; + +typedef enum +{ + ATT_INST_NONE = 0, + ATT_INST_LAST = 11, +} att_wave_inst_category_t; + +typedef struct +{ + uint8_t cu; + uint8_t simd; + uint8_t wave_id; + uint8_t contexts; + + uint32_t _rsvd; + size_t traceID; + + int64_t begin_time; + int64_t end_time; + + size_t timeline_size; + size_t instructions_size; + att_wave_state_t* timeline_array; + att_wave_instruction_t* instructions_array; +} att_wave_data_t; + +typedef rocprofiler_att_decoder_status_t (*rocprofiler_att_decoder_isa_callback_t)( + char* instruction, + uint64_t* memory_size, + uint64_t* size, + pcinfo_t address, + void* userdata); + +typedef rocprofiler_att_decoder_status_t (*rocprofiler_att_decoder_trace_callback_t)( + rocprofiler_att_decoder_record_type_t record_type_id, + int shader_engine_id, + void* trace_events, + uint64_t trace_size, + void* userdata); + +typedef uint64_t (*rocprofiler_att_decoder_se_data_callback_t)(int* shader_engine_id, + uint8_t** buffer, + uint64_t* buffer_size, + void* userdata); + +rocprofiler_att_decoder_status_t +rocprofiler_att_decoder_parse_data(rocprofiler_att_decoder_se_data_callback_t se_data_callback, + rocprofiler_att_decoder_trace_callback_t trace_callback, + rocprofiler_att_decoder_isa_callback_t isa_callback, + void* userdata); + +const char* +rocprofiler_att_decoder_get_info_string(rocprofiler_att_decoder_info_t info); + +const char* +rocprofiler_att_decoder_get_status_string(rocprofiler_att_decoder_status_t status); + +#ifdef __cplusplus +} +#endif diff --git a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp new file mode 100644 index 0000000000..b3209855bf --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp @@ -0,0 +1,197 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" +#include "dl.hpp" +#include "filenames.hpp" +#include "occupancy.hpp" +#include "profile_interface.hpp" +#include "wave.hpp" +#include "wstates.hpp" + +#include "lib/rocprofiler-sdk/registration.hpp" + +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +auto +get_lib_names() +{ + std::vector> lib_names = { + {tool_att_capability_t::ATT_CAPABILITIES_DEBUG, "libatt_decoder_debug.so"}, + {tool_att_capability_t::ATT_CAPABILITIES_TRACE, "libatt_decoder_trace.so"}, + {tool_att_capability_t::ATT_CAPABILITIES_SUMMARY, "libatt_decoder_summary.so"}, + {tool_att_capability_t::ATT_CAPABILITIES_TESTING, "libatt_decoder_testing.so"}, + }; + return lib_names; +} + +ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr
_dl) +: dir(std::move(_dir)) +, dl(std::move(_dl)) +{ + rocprofiler::common::filesystem::create_directories(dir); + table = std::make_shared(); + codefile = std::make_shared(dir, table); + filenames = std::make_shared(dir); + for(size_t i = 0; i < ATT_WAVE_STATE_LAST; i++) + wstates.at(i) = std::make_shared(i, dir); +} + +ATTFileMgr::~ATTFileMgr() { OccupancyFile::OccupancyFile(dir, table, occupancy); } + +void +ATTFileMgr::parseShader(int se_id, const std::vector& data) +{ + WaveConfig config(se_id, filenames, codefile, wstates); + ToolData tooldata(data, config, dl); + + if(config.occupancy.size()) occupancy.emplace(se_id, std::move(config.occupancy)); + + for(auto& [pc, kernel] : config.kernel_names) + codefile->kernel_names.emplace(pc, std::move(kernel)); +} + +int +get_shader_id(const std::string& name) +{ + auto run_pos = name.rfind('_'); + if(run_pos == std::string::npos) throw std::runtime_error("Invalid name"); + + std::string_view stripped = name.substr(0, run_pos); + auto se_number_pos = stripped.rfind('_'); + if(se_number_pos == std::string::npos || se_number_pos + 1 >= stripped.size()) + throw std::runtime_error("Invalid name"); + + return std::stoi(std::string(stripped.substr(se_number_pos + 1))); +} + +std::vector +query_att_decode_capability() +{ + std::vector ret; + + for(auto& [cap, libname] : get_lib_names()) + { + if(auto handle = dlopen(libname, RTLD_NOW | RTLD_LOCAL)) + { + dlclose(handle); + ret.push_back(cap); + } + } + + return ret; +} + +ATTDecoder::ATTDecoder(tool_att_capability_t capability) +{ + for(auto& [cap, libname] : get_lib_names()) + { + if(cap == capability) + { + dl = std::make_shared
(libname); + return; + } + } +}; + +void +ATTDecoder::parse(const Fspath& input_dir, + const Fspath& output_dir, + const std::vector& att_files, + const std::vector& codeobj_files, + const std::string& output_formats) +{ + auto& formats = GlobalDefs::get().output_formats; + formats = output_formats; + std::transform(formats.begin(), formats.end(), formats.begin(), [](unsigned char c) { + return std::tolower(c); + }); + + ATTFileMgr mgr(output_dir, dl); + + for(auto& file : codeobj_files) + { + if(file.name.find("memory://") == 0) + { + ROCP_WARNING << file.name << " was not loaded"; + continue; + } + + try + { + mgr.table->addDecoder((input_dir / file.name).c_str(), file.id, file.addr, file.size); + } catch(std::exception& e) + { + ROCP_ERROR << file.id << ':' << file.name << " - " << e.what() << std::endl; + } + } + + for(auto& shader : att_files) + { + int shader_id = 0; + try + { + shader_id = get_shader_id(shader); + } catch(std::exception& e) + { + ROCP_WARNING << "Could not retrieve shader_id: " << e.what(); + continue; + } + + std::vector shader_data; + + { + std::ifstream file(input_dir / shader, std::ios::binary); + if(!file.is_open()) + { + ROCP_WARNING << "could not open " << shader; + continue; + } + + file.seekg(0, std::ios::end); + shader_data.resize(file.tellg()); + file.seekg(0); + file.read(shader_data.data(), shader_data.size()); + } + + mgr.parseShader(shader_id, shader_data); + } +} + +bool +ATTDecoder::valid() const +{ + return dl && dl->att_parse_data_fn && dl->att_info_fn && dl->att_status_fn; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp new file mode 100644 index 0000000000..cbcd92ac11 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp @@ -0,0 +1,110 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/common/filesystem.hpp" +#include "lib/rocprofiler-sdk-att/util.hpp" + +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +using Fspath = rocprofiler::common::filesystem::path; + +struct CodeobjLoadInfo +{ + std::string name{}; + size_t id{0}; + size_t addr{0}; + size_t size{0}; +}; + +enum tool_att_capability_t +{ + ATT_CAPABILITIES_TESTING, // used for code coverage testing + ATT_CAPABILITIES_SUMMARY, // used for CSV output only + ATT_CAPABILITIES_TRACE, // used for all outputs + ATT_CAPABILITIES_DEBUG, + ATT_CAPABILITIES_LAST = ATT_CAPABILITIES_DEBUG, +}; + +/** + * Query decoder library capability. Returns list of supported capabilities. + */ +std::vector +query_att_decode_capability(); + +class ATTDecoder +{ +public: + ATTDecoder(tool_att_capability_t cap); + + /** + * Parse a list of att files + * @param[in] input directory where att_files and codeobj_files are relative to + * @param[in] output_dir location where ui_ files are generated + * @param[in] att_files list of ATT files, ideally from the same kernel launch + * @param[in] codeobj_files list of code object information loaded at the time of the trace + * @param[in] output_formats List of comma-separated output formats, e.g. "json,csv" + */ + void parse(const Fspath& input_dir, + const Fspath& output_dir, + const std::vector& att_files, + const std::vector& codeobj_files, + const std::string& output_formats); + + bool valid() const; + +protected: + std::shared_ptr dl{nullptr}; +}; + +class ATTFileMgr +{ + using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; + +public: + ATTFileMgr(Fspath _dir, std::shared_ptr
_dl); + ~ATTFileMgr(); + + void parseShader(int se_id, const std::vector& data); + + Fspath dir{}; + + std::shared_ptr dl{nullptr}; + std::shared_ptr codefile{nullptr}; + std::shared_ptr filenames{nullptr}; + std::shared_ptr table{nullptr}; + + std::map> occupancy; + std::array, ATT_WAVE_STATE_LAST> wstates; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/code.cpp b/source/lib/rocprofiler-sdk-att/code.cpp new file mode 100644 index 0000000000..e4f2b2a089 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/code.cpp @@ -0,0 +1,123 @@ +// MIT License +// +// Copyright (c) 2024 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 "code.hpp" +#include +#include "outputfile.hpp" + +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +#define ATT_CSV_NAME "att_output.csv" + +CodeFile::CodeFile(const Fspath& _dir, std::shared_ptr _table) +: dir(_dir) +, filename(_dir / "code.json") +, table(_table) +{} + +CodeFile::~CodeFile() +{ + std::vector>> vec; + vec.reserve(isa_map.size()); + + for(auto& [pc, isa] : isa_map) + if(isa && isa->code_line) vec.emplace_back(pc, std::move(isa)); + + isa_map.clear(); + line_numbers.clear(); + + if(GlobalDefs::get().has_format("csv")) + { + // Write CSV, ordered by id + vaddr + std::sort(vec.begin(), + vec.end(), + [](const std::pair>& a, + const std::pair>& b) { + if(a.first.marker_id == b.first.marker_id) return a.first.addr < b.first.addr; + return a.first.marker_id < b.first.marker_id; + }); + + OutputFile file(dir / ATT_CSV_NAME); + + file << "CodeObj, Vaddr, Instruction, Hitcount, Latency, Source\n"; + for(auto& [pc, line] : vec) + { + if(kernel_names.find(pc) != kernel_names.end()) + { + file << pc.marker_id << ',' << pc.addr << ",\"; " << kernel_names.at(pc).name + << "\",0,0,\"" << kernel_names.at(pc).demangled << "\"\n"; + } + file << pc.marker_id << ',' << pc.addr << ",\"" << line->code_line->inst << "\"," + << line->hitcount << ',' << line->latency << ',' << line->code_line->comment + << '\n'; + } + } + + if(!GlobalDefs::get().has_format("json")) return; + + // Write JSON, ordered by exec line number + std::sort(vec.begin(), + vec.end(), + [](const std::pair>& a, + const std::pair>& b) { + return a.second->line_number < b.second->line_number; + }); + + nlohmann::json jcode; + + for(auto& line : vec) + { + auto& isa = *line.second; + + if(kernel_names.find(line.first) != kernel_names.end()) + { + std::stringstream code; + code << "[\"; " << kernel_names.at(line.first).name << "\", 100, " + << (isa.line_number - 1) << ", \"" << kernel_names.at(line.first).demangled + << "\", " << line.first.marker_id << ", " << line.first.addr << ", 0, 0]"; + + jcode.push_back(nlohmann::json::parse(code.str())); + } + + std::stringstream code; + code << "[\"" << isa.code_line->inst << "\", 0, " << isa.line_number << ", \"" + << isa.code_line->comment << "\", " << line.first.marker_id << ", " << line.first.addr + << ", " << isa.hitcount << ", " << isa.latency << "]"; + + jcode.push_back(nlohmann::json::parse(code.str())); + } + + nlohmann::json json; + json["code"] = jcode; + json["version"] = TOOL_VERSION; + + OutputFile(filename) << json; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/code.hpp b/source/lib/rocprofiler-sdk-att/code.hpp new file mode 100644 index 0000000000..0d1a6ac666 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/code.hpp @@ -0,0 +1,68 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" + +#include +#include +#include +#include +#include "att_decoder.h" +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +struct CodeLine +{ + using Instruction = rocprofiler::sdk::codeobj::disassembly::Instruction; + + int line_number = 0; + int type = 0; + + std::atomic hitcount{0}; + std::atomic latency{0}; + std::shared_ptr code_line{nullptr}; +}; + +class CodeFile +{ + using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; + +public: + CodeFile(const Fspath& dir, std::shared_ptr table); + ~CodeFile(); + + Fspath dir{}; + Fspath filename{}; + std::unordered_map line_numbers{}; + std::map> isa_map{}; + std::map kernel_names{}; + + std::shared_ptr table; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/counters.cpp b/source/lib/rocprofiler-sdk-att/counters.cpp new file mode 100644 index 0000000000..b13e485cb1 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/counters.cpp @@ -0,0 +1,78 @@ +// MIT License +// +// Copyright (c) 2024 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 "counters.hpp" +#include "outputfile.hpp" + +#include +#include +#include +#include +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +CountersFile::CountersFile(const Fspath& _dir, const std::vector& _names) +: dir(_dir) +, names(_names) +{} + +void +CountersFile::AddShaderEngine(int se, const att_perfevent_t* events, size_t num_events) +{ + if(!num_events || !GlobalDefs::get().has_format("json")) return; + + nlohmann::json js; + + for(size_t i = 0; i < num_events; i++) + { + auto& ev = events[i]; + js.emplace_back({ev.time, ev.events0, ev.events1, ev.events2, ev.events3, ev.CU, ev.bank}); + } + + auto filename = dir / ("se" + std::to_string(se) + "_perfcounter.json"); + + OutputFile(filename) << nlohmann::json{"data", js}; + shaders.emplace_back(filename); +} + +CountersFile::~CountersFile() +{ + nlohmann::json counters_names; + for(auto& name : names) + counters_names.emplace_back(name); + + nlohmann::json perfcounter_filenames; + for(auto& name : shaders) + perfcounter_filenames.emplace_back(name); + + nlohmann::json js; + js["counters"] = counters_names; + js["shaders"] = perfcounter_filenames; + + OutputFile(dir / "graph_options.json") << nlohmann::json{"data", js}; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/counters.hpp b/source/lib/rocprofiler-sdk-att/counters.hpp new file mode 100644 index 0000000000..252c5b6c03 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/counters.hpp @@ -0,0 +1,50 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" + +#include +#include +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +class CountersFile +{ +public: + CountersFile(const Fspath& dir, const std::vector& names); + ~CountersFile(); + + void AddShaderEngine(int se, const att_perfevent_t* events, size_t num_events); + +private: + const Fspath dir; + std::vector names{}; + std::vector shaders{}; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/dl.cpp b/source/lib/rocprofiler-sdk-att/dl.cpp new file mode 100644 index 0000000000..294ad29ecb --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/dl.cpp @@ -0,0 +1,55 @@ +// MIT License +// +// Copyright (c) 2024 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 "dl.hpp" +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +DL::DL(const char* dlname) +{ + handle = dlopen(dlname, RTLD_NOW | RTLD_LOCAL); + if(!handle) return; + + att_parse_data_fn = + reinterpret_cast(dlsym(handle, "rocprofiler_att_decoder_parse_data")); + att_info_fn = + reinterpret_cast(dlsym(handle, "rocprofiler_att_decoder_get_info_string")); + att_status_fn = + reinterpret_cast(dlsym(handle, "rocprofiler_att_decoder_get_status_string")); +}; + +DL::~DL() +{ + if(handle) dlclose(handle); +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/dl.hpp b/source/lib/rocprofiler-sdk-att/dl.hpp new file mode 100644 index 0000000000..8fafbe2b78 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/dl.hpp @@ -0,0 +1,48 @@ +// MIT License +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once +#include +#include "att_decoder.h" + +namespace rocprofiler +{ +namespace att_wrapper +{ +class DL +{ + using ParseFn = decltype(rocprofiler_att_decoder_parse_data); + using InfoFn = decltype(rocprofiler_att_decoder_get_info_string); + using StatusFn = decltype(rocprofiler_att_decoder_get_status_string); + +public: + DL(const char* dlname); + ~DL(); + + ParseFn* att_parse_data_fn = nullptr; + InfoFn* att_info_fn = nullptr; + StatusFn* att_status_fn = nullptr; + void* handle = nullptr; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/filenames.cpp b/source/lib/rocprofiler-sdk-att/filenames.cpp new file mode 100644 index 0000000000..872ffc3200 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/filenames.cpp @@ -0,0 +1,63 @@ +// MIT License +// +// Copyright (c) 2024 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 "filenames.hpp" +#include "outputfile.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +void +FilenameMgr::addwave(const Fspath& name, Coord coord, size_t begint, size_t endt) +{ + streams.emplace(coord, WaveName{name.filename(), begint, endt}); +} + +FilenameMgr::~FilenameMgr() +{ + if(!GlobalDefs::get().has_format("json")) return; + + using std::to_string; + nlohmann::json namelist; + for(auto& [coord, data] : streams) + { + namelist[to_string(coord.se)][to_string(coord.sm)][to_string(coord.sl)] + [to_string(coord.id)] = {data.name, data.begin, data.end}; + } + + nlohmann::json metadata = {{"global_begin_time", 0}, + {"gfxv", (gfxip > 9) ? "navi" : "vega"}, + {"gfxip", gfxip}, + {"version", TOOL_VERSION}}; + + metadata["wave_filenames"] = namelist; + OutputFile(filename) << metadata; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/filenames.hpp b/source/lib/rocprofiler-sdk-att/filenames.hpp new file mode 100644 index 0000000000..ed7f87306e --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/filenames.hpp @@ -0,0 +1,79 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" + +#include +#include +#include "att_decoder.h" +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +class FilenameMgr +{ +public: + struct Coord + { + int se{0}; + int sm{0}; + int sl{0}; + int id{0}; + bool operator==(const Coord& other) const + { + return se == other.se && sm == other.sm && sl == other.sl && id == other.id; + } + bool operator<(const Coord& other) const + { + if(se != other.se) return se < other.se; + if(sm != other.sm) return sm < other.sm; + if(sl != other.sl) return sl < other.sl; + return id < other.id; + } + }; + struct WaveName + { + std::string name{}; + size_t begin{0}; + size_t end{0}; + }; + + FilenameMgr(const Fspath& _dir) + : dir(_dir) + , filename(_dir / "filenames.json") + {} + ~FilenameMgr(); + + void addwave(const Fspath& file, Coord coord, size_t start, size_t end); + + Fspath dir{}; + Fspath filename{}; + std::map streams{}; + int gfxip = 9; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/occupancy.cpp b/source/lib/rocprofiler-sdk-att/occupancy.cpp new file mode 100644 index 0000000000..5ea60acf0a --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/occupancy.cpp @@ -0,0 +1,109 @@ +// MIT License +// +// Copyright (c) 2024 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 "occupancy.hpp" +#include +#include "outputfile.hpp" + +#include +#include +#include +#include + +#define OCCUPANCY_RES 8 + +namespace rocprofiler +{ +namespace att_wrapper +{ +union occupancy_data_v1 +{ + struct + { + uint64_t kernel_id : 12; + uint64_t simd : 2; + uint64_t slot : 4; + uint64_t enable : 1; + uint64_t cu : 4; + uint64_t time : 41; + }; + uint64_t raw; +}; + +std::map kernel_ids{{pcinfo_t{0, 0}, 0}}; +std::atomic current_id{1}; + +static int +get_kernel_id(pcinfo_t pc) +{ + if(kernel_ids.find(pc) != kernel_ids.end()) return kernel_ids.at(pc); + + return kernel_ids.emplace(pc, current_id.fetch_add(1)).first->second; +} + +static uint64_t +convert(const att_occupancy_info_v2_t& v2) +{ + occupancy_data_v1 v1{}; + v1.time = v2.time / OCCUPANCY_RES; + v1.simd = v2.simd; + v1.slot = v2.slot; + v1.enable = v2.start; + v1.cu = v2.cu; + v1.kernel_id = get_kernel_id(v2.pc); + return v1.raw; +}; + +void +OccupancyFile::OccupancyFile(const Fspath& dir, + std::shared_ptr table, + const std::map>& occ) +{ + if(!GlobalDefs::get().has_format("json")) return; + nlohmann::json jocc; + + for(auto& [se, eventlist] : occ) + { + nlohmann::json list; + for(auto& event : eventlist) + list.push_back(convert(event)); + jocc[std::to_string(se)] = list; + } + + for(auto& [pc, id] : kernel_ids) + { + std::stringstream ss; + try + { + ss << table->getSymbolMap(pc.marker_id).at(pc.addr).name; + } catch(std::exception& e) + { + ss << pc.marker_id << " / 0x" << std::hex << pc.addr << std::dec; + } + jocc["dispatches"][std::to_string(id)] = ss.str(); + } + + OutputFile(dir / "occupancy.json") << jocc; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/occupancy.hpp b/source/lib/rocprofiler-sdk-att/occupancy.hpp new file mode 100644 index 0000000000..574cd9d738 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/occupancy.hpp @@ -0,0 +1,46 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" + +#include +#include +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +namespace OccupancyFile +{ +using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; + +void +OccupancyFile(const Fspath& dir, + std::shared_ptr table, + const std::map>& occ); +}; // namespace OccupancyFile + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/outputfile.hpp b/source/lib/rocprofiler-sdk-att/outputfile.hpp new file mode 100644 index 0000000000..4c1d45dbeb --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/outputfile.hpp @@ -0,0 +1,66 @@ +// MIT License +// +// Copyright (c) 2024 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. + +// output_directory: "@CMAKE_CURRENT_BINARY_DIR@/roctracer-roctx-trace" + +#pragma once + +#include "lib/common/logging.hpp" +#include "lib/rocprofiler-sdk-att/att_lib_wrapper.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +class OutputFile +{ +public: + OutputFile(const std::string& str) + { + if(!Enabled()) return; + ofs = std::ofstream(str, std::ofstream::out); + ROCP_FATAL_IF(!ofs.is_open()) << "Could not open output file " << str << std::endl; + } + + template + OutputFile& operator<<(const T& v) + { + if(Enabled()) ofs << v; + return *this; + } + + static bool& Enabled() + { + static bool _enabled = true; + return _enabled; + } + + std::ofstream ofs; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/profile_interface.cpp b/source/lib/rocprofiler-sdk-att/profile_interface.cpp new file mode 100644 index 0000000000..a8ee575240 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/profile_interface.cpp @@ -0,0 +1,253 @@ +// MIT License +// +// Copyright (c) 2024 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. +// +// undefine NDEBUG so asserts are implemented +#ifdef NDEBUG +# undef NDEBUG +#endif + +#include "profile_interface.hpp" +#include +#include +#include +#include +#include "att_decoder.h" +#include "dl.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +struct trace_data_t +{ + int64_t id{0}; + uint8_t* data{nullptr}; + uint64_t size{0}; + ToolData* tool{nullptr}; +}; + +rocprofiler_att_decoder_status_t +get_trace_data(rocprofiler_att_decoder_record_type_t trace_id, + int /* shader_id */, + void* trace_events, + size_t trace_size, + void* userdata) +{ + C_API_BEGIN + + CHECK_NOTNULL(userdata); + trace_data_t& trace_data = *reinterpret_cast(userdata); + CHECK_NOTNULL(trace_data.tool); + ToolData& tool = *reinterpret_cast(trace_data.tool); + + if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_INFO) + { + auto* infos = (rocprofiler_att_decoder_info_t*) trace_events; + for(size_t i = 0; i < trace_size; i++) + ROCP_WARNING << tool.dl->att_info_fn(infos[i]) << std::endl; + + return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; + } + + if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_GFXIP) + { + tool.config.filemgr->gfxip = reinterpret_cast(trace_events); + } + else if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_OCCUPANCY) + { + for(size_t i = 0; i < trace_size; i++) + tool.config.occupancy.push_back( + reinterpret_cast(trace_events)[i]); + } + + if(trace_id != ROCPROFILER_ATT_DECODER_TYPE_WAVE) return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; + + bool bInvalid = false; + for(size_t wave_n = 0; wave_n < trace_size; wave_n++) + { + auto& wave = reinterpret_cast(trace_events)[wave_n]; + + WaveFile(tool.config, wave); + + for(size_t j = 0; j < wave.instructions_size; j++) + { + auto& inst = wave.instructions_array[j]; + if(inst.pc.marker_id == 0 && inst.pc.addr == 0) + continue; + else if(inst.category >= att_wave_inst_category_t::ATT_INST_LAST) + continue; + + try + { + auto& line = tool.get(inst.pc); + line.hitcount.fetch_add(1, std::memory_order_relaxed); + line.latency.fetch_add(inst.duration, std::memory_order_relaxed); + } catch(...) + { + bInvalid = true; + } + } + } + if(bInvalid) ROCP_WARNING << "Could not fetch some instructions!"; + + return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; + + C_API_END + + return ROCPROFILER_ATT_DECODER_STATUS_ERROR; +} + +uint64_t +copy_trace_data(int* seid, uint8_t** buffer, uint64_t* buffer_size, void* userdata) +{ + trace_data_t& data = *reinterpret_cast(userdata); + *seid = data.id; + *buffer_size = data.size; + *buffer = data.data; + data.size = 0; + return *buffer_size; +} + +rocprofiler_att_decoder_status_t +isa_callback(char* isa_instruction, + uint64_t* isa_memory_size, + uint64_t* isa_size, + pcinfo_t pc, + void* userdata) +{ + C_API_BEGIN + CHECK_NOTNULL(userdata); + trace_data_t& trace_data = *reinterpret_cast(userdata); + CHECK_NOTNULL(trace_data.tool); + ToolData& tool = *reinterpret_cast(trace_data.tool); + + std::shared_ptr instruction{nullptr}; + + try + { + CodeLine& line = tool.get(pc); + instruction = line.code_line; + } catch(std::exception& e) + { + ROCP_WARNING << pc.marker_id << ":" << pc.addr << ' ' << e.what(); + return ROCPROFILER_ATT_DECODER_STATUS_ERROR; + } + + if(!instruction.get()) return ROCPROFILER_ATT_DECODER_STATUS_ERROR_INVALID_ARGUMENT; + + { + size_t tmp_isa_size = *isa_size; + *isa_size = instruction->inst.size(); + + if(*isa_size > tmp_isa_size) return ROCPROFILER_ATT_DECODER_STATUS_ERROR_OUT_OF_RESOURCES; + } + + memcpy(isa_instruction, instruction->inst.data(), *isa_size); + *isa_memory_size = instruction->size; + + C_API_END + return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; +} + +ToolData::ToolData(const std::vector& _data, WaveConfig& _config, std::shared_ptr
_dl) +: cfile(_config.code) +, config(_config) +, dl(std::move(_dl)) +{ + trace_data_t data{.id = config.shader_engine, + .data = (uint8_t*) _data.data(), + .size = _data.size(), + .tool = this}; + + auto status = dl->att_parse_data_fn(copy_trace_data, get_trace_data, isa_callback, &data); + if(status != ROCPROFILER_ATT_DECODER_STATUS_SUCCESS) + ROCP_ERROR << "Callback failed with status " << dl->att_status_fn(status); +} + +ToolData::~ToolData() {} + +std::string +demangle(std::string_view line) +{ + int status; + char* c_name = abi::__cxa_demangle(line.data(), 0, 0, &status); + + if(c_name == nullptr) return ""; + + std::string str = c_name; + free(c_name); + return str; +} + +CodeLine& +ToolData::get(pcinfo_t _pc) +{ + auto& isa_map = cfile->isa_map; + if(isa_map.find(_pc) != isa_map.end()) return *isa_map.at(_pc); + + // Attempt to disassemble full kernel + if(_pc.marker_id) try + { + rocprofiler::sdk::codeobj::segment::CodeobjTableTranslator symbol_table; + for(auto& [vaddr, symbol] : cfile->table->getSymbolMap(_pc.marker_id)) + symbol_table.insert({symbol.vaddr, symbol.mem_size, _pc.marker_id}); + + auto addr_range = symbol_table.find_codeobj_in_range(_pc.addr); + try + { + auto symbol = cfile->table->getSymbolMap(_pc.marker_id).at(addr_range.addr); + auto pair = KernelName{symbol.name, demangle(symbol.name)}; + cfile->kernel_names.emplace(pcinfo_t{addr_range.addr, _pc.marker_id}, pair); + } catch(...) + { + ROCP_INFO << "Missing kernelSymbol at " << _pc.marker_id << ':' << addr_range.addr; + } + + for(auto addr = addr_range.addr; addr < addr_range.addr + addr_range.size;) + { + pcinfo_t info{.addr = addr, .marker_id = addr_range.id}; + auto& cline = *(isa_map.emplace(info, std::make_unique()).first->second); + + cline.line_number = isa_map.size() + cfile->kernel_names.size() - 1; + cfile->line_numbers[info] = cline.line_number; + + cline.code_line = cfile->table->get(addr_range.id, addr); + addr += cline.code_line->size; + if(!cline.code_line->size) throw std::invalid_argument("Line has 0 bytes!"); + } + + if(isa_map.find(_pc) != isa_map.end()) return *isa_map.at(_pc); + } catch(std::exception& e) + {} + + auto& cline = *(isa_map.emplace(_pc, std::make_unique()).first->second); + + cline.line_number = isa_map.size(); + cfile->line_numbers[_pc] = cline.line_number; + + cline.code_line = cfile->table->get(_pc.marker_id, _pc.addr); + + return cline; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/profile_interface.hpp b/source/lib/rocprofiler-sdk-att/profile_interface.hpp new file mode 100644 index 0000000000..dd65a07941 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/profile_interface.hpp @@ -0,0 +1,77 @@ +// MIT License +// +// Copyright (c) 2024 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. +// +// undefine NDEBUG so asserts are implemented +#ifdef NDEBUG +# undef NDEBUG +#endif +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "att_decoder.h" +#include "code.hpp" +#include "wave.hpp" + +#define C_API_BEGIN \ + try \ + { +#define C_API_END \ + } \ + catch(std::exception & e) \ + { \ + std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl; \ + } \ + catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; } + +namespace rocprofiler +{ +namespace att_wrapper +{ +using Instruction = rocprofiler::sdk::codeobj::disassembly::Instruction; +using SymbolInfo = rocprofiler::sdk::codeobj::disassembly::SymbolInfo; + +struct ToolData +{ + ToolData(const std::vector& data, WaveConfig& config, std::shared_ptr _dl); + ~ToolData(); + + CodeLine& get(pcinfo_t pc); + + std::shared_ptr cfile{}; + WaveConfig& config; + std::shared_ptr
dl{}; + + std::vector shader_data{}; + size_t num_waves = 0; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt new file mode 100644 index 0000000000..832dc851ac --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt @@ -0,0 +1,36 @@ +# +# +# +project(rocprofiler-att-parser-tests LANGUAGES CXX) + +add_executable(att-parser-tool-v3) +target_link_libraries( + att-parser-tool-v3 + PRIVATE rocprofiler-sdk-att-parser rocprofiler-sdk::rocprofiler-sdk-json + rocprofiler-sdk::rocprofiler-sdk-common-library) +target_sources(att-parser-tool-v3 PRIVATE standalone_tool_main.cpp) + +add_executable(att-decoder-test) +target_link_libraries( + att-decoder-test + PRIVATE rocprofiler-sdk-att-parser + rocprofiler-sdk::rocprofiler-sdk-json + rocprofiler-sdk::rocprofiler-sdk-common-library + rocprofiler-sdk::rocprofiler-sdk-glog + rocprofiler-sdk::rocprofiler-sdk-static-library + GTest::gtest + GTest::gtest_main) +target_sources(att-decoder-test PRIVATE att_decoder_test.cpp) + +add_library(att_decoder_testing SHARED) +target_sources(att_decoder_testing PRIVATE dummy_decoder.cpp) + +if(NOT ROCPROFILER_MEMCHECK) + gtest_add_tests( + TARGET att-decoder-test + SOURCES att_decoder_test.cpp + TEST_LIST att-decoder-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + set_tests_properties(${att-decoder-test_TESTS} PROPERTIES TIMEOUT 10 LABELS + "unittests") +endif() diff --git a/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp b/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp new file mode 100644 index 0000000000..d10f555f74 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp @@ -0,0 +1,138 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/common/logging.hpp" +#include "lib/rocprofiler-sdk-att/att_lib_wrapper.hpp" +#include "lib/rocprofiler-sdk-att/code.hpp" +#include "lib/rocprofiler-sdk-att/outputfile.hpp" +#include "lib/rocprofiler-sdk/registration.hpp" + +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +class ATTDecoderTest : public ATTDecoder +{ +public: + ATTDecoderTest() + : ATTDecoder(rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TESTING) + { + rocprofiler::att_wrapper::OutputFile::Enabled() = false; + GlobalDefs::get().output_formats = "json,csv"; + registration::init_logging(); + }; + + void test_parse() + { + ATTFileMgr mgr("out/", dl); + + auto append_isa = [&](const char* line) { + // matches addresses in dummy_decoder.cpp + pcinfo_t pc{}; + pc.addr = mgr.codefile->isa_map.size() * 8; + pc.marker_id = 1; + + auto code = std::make_unique(); + code->code_line = std::make_shared(); + code->code_line->inst = line; + code->line_number = mgr.codefile->isa_map.size(); + + mgr.codefile->isa_map.emplace(pc, std::move(code)); + }; + + mgr.codefile->kernel_names[pcinfo_t{}] = KernelName{"_Kernel", "Kernel"}; + + append_isa("s_load_"); + append_isa("s_store_"); + append_isa("s_waitcnt vmcnt(0) lgkmcnt(0)"); + + std::vector dummy_data; + dummy_data.resize(128); + + mgr.parseShader(0, dummy_data); + mgr.parseShader(1, dummy_data); + } +}; + +TEST(att_decoder_test, dlopen) +{ + registration::init_logging(); + auto query = query_att_decode_capability(); + ROCP_FATAL_IF(query.empty()) << "No decoder capability available!"; +} + +TEST(att_decoder_test, filewrite) +{ + ATTDecoderTest decoder; + ROCP_FATAL_IF(!decoder.valid()) << "Failed to initialize decoder library!"; + + decoder.test_parse(); +} + +TEST(att_decoder_test, warn_failures) +{ + std::vector codeobjs; + codeobjs.resize(5); + codeobjs.at(0).name = "memory://unknown"; + codeobjs.at(1).name = "memory://unknown&offset=123&size=123"; + codeobjs.at(2).name = "file://nofile"; + codeobjs.at(3).name = "file://nofile&offset=123&size=123"; + codeobjs.at(4).name = "myfile123.out"; + + std::vector att_files; + att_files.emplace_back(std::string("file123.att")); + + ATTDecoderTest decoder; + ROCP_FATAL_IF(!decoder.valid()) << "Failed to initialize decoder library!"; + + decoder.parse(".", ".", att_files, codeobjs, "csv,json"); +} + +TEST(att_decoder_test, code_write) +{ + registration::init_logging(); + rocprofiler::att_wrapper::OutputFile::Enabled() = false; + GlobalDefs::get().output_formats = "json,csv"; + + CodeFile file(".", nullptr); + + pcinfo_t addr{}; + addr.marker_id = 0; + addr.addr = 0x1000; + file.kernel_names[addr] = KernelName{"_Kernel", "Kernel"}; + + for(size_t i = 0; i < 4; i++) + { + auto line = std::make_unique(); + line->line_number = i; + + line->code_line = std::make_shared(); + line->code_line->inst = "v_add"; + file.isa_map[addr] = std::move(line); + } +} + +}; // namespace att_wrapper +}; // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/tests/dummy_decoder.cpp b/source/lib/rocprofiler-sdk-att/tests/dummy_decoder.cpp new file mode 100644 index 0000000000..ac903902eb --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/tests/dummy_decoder.cpp @@ -0,0 +1,142 @@ +// MIT License +// +// Copyright (c) 2024 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 "../att_decoder.h" + +#include +#include + +__attribute__((visibility("default"))) rocprofiler_att_decoder_status_t +rocprofiler_att_decoder_parse_data(rocprofiler_att_decoder_se_data_callback_t se_data_callback, + rocprofiler_att_decoder_trace_callback_t trace_callback, + rocprofiler_att_decoder_isa_callback_t isa_callback, + void* userdata) +{ + const int gfxip = 9; + + trace_callback(ROCPROFILER_ATT_DECODER_TYPE_GFXIP, 0, (void*) gfxip, 0, userdata); + { + std::vector infos{}; + for(size_t i = 1; i < ROCPROFILER_ATT_DECODER_INFO_LAST; i++) + infos.emplace_back(static_cast(i)); + + trace_callback(ROCPROFILER_ATT_DECODER_TYPE_INFO, 0, infos.data(), infos.size(), userdata); + } + { + uint64_t memory_size = 0, size = 16; + std::array inst; + + isa_callback(inst.data(), &memory_size, &size, pcinfo_t{0, 0}, userdata); + } + { + int se_id = 0; + uint8_t* buffer = nullptr; + size_t buffer_size = 0; + + while(se_data_callback(&se_id, &buffer, &buffer_size, userdata)) + {}; + } + + { + std::vector vec; + att_occupancy_info_v2_t occ{}; + occ.cu = occ.se = occ.simd = occ.slot = 1; + occ.pc.marker_id = 0; + occ.pc.addr = 0; + + occ.time = 0; + occ.start = 1; + vec.push_back(occ); + occ.simd = 0; + vec.push_back(occ); + + occ.time = 1024; + occ.start = 0; + vec.push_back(occ); + occ.simd = 1; + vec.push_back(occ); + + trace_callback(ROCPROFILER_ATT_DECODER_TYPE_OCCUPANCY, 0, vec.data(), vec.size(), userdata); + } + + { + std::vector waves{}; + + att_wave_data_t wave{}; + wave.cu = wave.simd = wave.wave_id = wave.traceID = 1; + + wave.begin_time = 0; + wave.end_time = 1024; + + std::vector states; + for(int j = 0; j < 2; j++) + for(int i = 1; i < ATT_WAVE_STATE_LAST; i++) + states.emplace_back(att_wave_state_t{i, 128}); + + std::vector insts; + for(int i = 1; i < ATT_INST_LAST; i++) + { + att_wave_instruction_t inst{}; + inst.category = i; + inst.duration = 48; + inst.time = i * 64 - 32; + inst.pc.marker_id = 1; + inst.pc.addr = 8 * i; + insts.emplace_back(std::move(inst)); + } + + wave.instructions_array = insts.data(); + wave.instructions_size = insts.size(); + wave.timeline_array = states.data(); + wave.timeline_size = states.size(); + + waves.push_back(wave); + wave.simd = 2; + waves.push_back(wave); + + trace_callback(ROCPROFILER_ATT_DECODER_TYPE_WAVE, 0, waves.data(), waves.size(), userdata); + } + + return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; +} + +__attribute__((visibility("default"))) const char* +rocprofiler_att_decoder_get_info_string(rocprofiler_att_decoder_info_t info) +{ + return std::vector{"ROCPROFILER_ATT_DECODER_INFO_NONE", + "ROCPROFILER_ATT_DECODER_INFO_DATA_LOST", + "ROCPROFILER_ATT_DECODER_INFO_STITCH_INCOMPLETE", + "ROCPROFILER_ATT_DECODER_INFO_LAST"} + .at((size_t) info); +} + +__attribute__((visibility("default"))) const char* +rocprofiler_att_decoder_get_status_string(rocprofiler_att_decoder_status_t status) +{ + return std::vector{"ROCPROFILER_ATT_DECODER_STATUS_SUCCESS", + "ROCPROFILER_ATT_DECODER_STATUS_ERROR", + "ROCPROFILER_ATT_DECODER_STATUS_ERROR_OUT_OF_RESOURCES", + "ROCPROFILER_ATT_DECODER_STATUS_ERROR_INVALID_ARGUMENT", + "ROCPROFILER_ATT_DECODER_STATUS_ERROR_INVALID_SHADER_DATA", + "ROCPROFILER_ATT_DECODER_STATUS_LAST"} + .at((size_t) status); +} diff --git a/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp b/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp new file mode 100644 index 0000000000..982418c4de --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp @@ -0,0 +1,167 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/common/logging.hpp" +#include "lib/rocprofiler-sdk-att/att_lib_wrapper.hpp" +#include "lib/rocprofiler-sdk-att/outputfile.hpp" +#include "lib/rocprofiler-sdk/registration.hpp" + +#include +#include +#include + +using rocprofiler::att_wrapper::ATTDecoder; +using Fspath = rocprofiler::att_wrapper::Fspath; + +int +main(int argc, char** argv) +{ + if(argc < 2 || std::string_view(argv[1]) == "--help" || std::string_view(argv[1]) == "-h") + { + std::cout + << "Usage: att-parser-tool json_filepath [output_dir] [output_formats]\nParameters:\n" + << "\tjson_filepath: Path of rocprofv3's generated results.json\n" + << "\toutput_dir: Optional output directory. Default: json_filepath parent dir\n" + << "\toutput_formats: json, perfetto, csv. Default: all\n" + << std::endl; + exit(0); + } + + static auto flag = std::once_flag{}; + std::call_once(flag, []() { + auto logging_cfg = rocprofiler::common::logging_config{.install_failure_handler = true}; + rocprofiler::common::init_logging("ROCPROF", logging_cfg); + FLAGS_colorlogtostderr = true; + }); + + auto cap = rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_SUMMARY; + { + auto query = rocprofiler::att_wrapper::query_att_decode_capability(); + ROCP_FATAL_IF(query.empty()) << "No decoder capability available!"; + + for(auto& avail_cap : query) + cap = std::max(cap, avail_cap); + } + + ATTDecoder decoder(cap); + ROCP_FATAL_IF(!decoder.valid()) << "Failed to initialize decoder library!"; + + auto get_run_number = [](const std::string& path) -> int { + auto get_filename = [](const std::string& _path) -> std::string { + return Fspath(_path).filename().c_str(); + }; + + auto name = get_filename(path); + auto run_pos = name.rfind('_'); + auto extension_pos = name.rfind(".att"); + + if(extension_pos == std::string::npos || extension_pos <= run_pos) throw std::exception(); + + return std::stoi(name.substr(run_pos + 1, extension_pos - run_pos)); + }; + + Fspath input_path = rocprofiler::common::filesystem::absolute(argv[1]); + + std::string formats = "json,csv"; + if(argc >= 4) formats = argv[3]; + + Fspath output_path = input_path.parent_path(); + if(argc >= 3) output_path = argv[2]; + + nlohmann::json sdk_json; + { + nlohmann::json full_json; + std::ifstream ifile(input_path); + ROCP_FATAL_IF(!ifile.is_open()) << "Failed to open json file!"; + ifile >> full_json; + sdk_json = full_json["rocprofiler-sdk-tool"][0]; + } + + std::unordered_map> all_runs; + + for(auto& file : sdk_json["strings"]["att_files"]) + { + try + { + int n = get_run_number(file); + if(all_runs.find(n) == all_runs.end()) all_runs[n] = {}; + + all_runs[n].push_back(file); + + } catch(std::exception&) + { + ROCP_WARNING << "Invalid ATT filename " << file; + } + } + + for(auto& [run_number, att_filenames] : all_runs) + { + std::vector att_files{}; + std::vector codeobj_files{}; + + std::map snapshot_files{}; + for(auto elem : sdk_json["strings"]["code_object_snapshot_files"]) + snapshot_files[elem["key"]] = elem["value"]; + + for(auto& codeobj : sdk_json["code_objects"]) + if(std::string(codeobj["uri"]).size()) + { + std::string filename = codeobj["uri"]; + size_t id = size_t(codeobj["code_object_id"]); + if(filename.empty()) continue; + + try + { + filename = snapshot_files.at(id); + } catch(...) + { + ROCP_WARNING << "codeobject id " << id << " not found " << filename; + } + + try + { + codeobj_files.push_back({filename, + id, + size_t(codeobj["load_delta"]), + size_t(codeobj["load_size"])}); + } catch(std::exception& e) + { + ROCP_WARNING << "Could not load " << filename << ": " << e.what(); + } catch(std::string& r) + { + ROCP_WARNING << "Could not load " << filename << ": " << r; + } catch(...) + { + ROCP_WARNING << "Could not load " << filename; + } + } + + std::string run_name = input_path.filename().c_str(); + std::string ui_name = run_name.substr(0, run_name.find(".json")); + auto output_dir = output_path / ("ui_output_" + ui_name + std::to_string(run_number)); + decoder.parse(input_path.parent_path(), output_dir, att_filenames, codeobj_files, formats); + } + + ROCP_INFO << "Finalizing ATT Tool"; + + return 0; +} diff --git a/source/lib/rocprofiler-sdk-att/util.hpp b/source/lib/rocprofiler-sdk-att/util.hpp new file mode 100644 index 0000000000..dd7ce317ca --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/util.hpp @@ -0,0 +1,88 @@ +// MIT License +// +// Copyright (c) 2024 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 + +#define TOOL_VERSION_MAJOR 3 +#define TOOL_VERSION_MINOR 0 +#define TOOL_VERSION_REV 0 +#define TOOL_VERSION "3.0.0" + +#include +#include "att_decoder.h" +#include "lib/common/logging.hpp" + +#include +#include +#include + +template <> +struct std::hash +{ +public: + size_t operator()(const pcinfo_t& a) const + { + return (a.marker_id << 32) ^ (a.marker_id >> 32) ^ a.addr; + } +}; + +inline bool +operator==(const pcinfo_t& a, const pcinfo_t& b) +{ + return a.addr == b.addr && a.marker_id == b.marker_id; +}; + +inline bool +operator<(const pcinfo_t& a, const pcinfo_t& b) +{ + if(a.marker_id == b.marker_id) return a.addr < b.addr; + return a.marker_id < b.marker_id; +}; + +namespace rocprofiler +{ +namespace att_wrapper +{ +class GlobalDefs +{ +public: + static GlobalDefs& get() + { + static GlobalDefs def; + return def; + } + bool has_format(std::string_view fmt) const + { + return output_formats.find(fmt) != std::string::npos; + } + + std::string output_formats; +}; + +struct KernelName +{ + std::string name{}; + std::string demangled{}; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/waitcnt/CMakeLists.txt new file mode 100644 index 0000000000..35828f8c25 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/CMakeLists.txt @@ -0,0 +1,6 @@ +# +# ATT waitcnt ISA +# +if(ROCPROFILER_BUILD_TESTS) + add_subdirectory(tests) +endif() diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/analysis.cpp b/source/lib/rocprofiler-sdk-att/waitcnt/analysis.cpp new file mode 100644 index 0000000000..da10fd789b --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/analysis.cpp @@ -0,0 +1,85 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/rocprofiler-sdk-att/waitcnt/analysis.hpp" +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +std::map> WaitcntList::_cache; + +int64_t +MemoryCounter::extract_waitcnt(const std::string& str) const +{ + // Find the position of the number in the string + size_t pos = str.find(name); + size_t iter = pos + name.size() + 1; + + if(pos == std::string::npos || iter >= str.size()) return 0; + + // If no hex preamble found, assume it's decimal + if(str.find("0x") == std::string::npos) return atoi(str.data() + iter); + + // Skip any additional spaces + while(iter < str.size() && str.at(iter) == ' ') + iter++; + + // Check if the number is in hexadecimal format (starts with "0x"), otherwise return decimal + if(iter + 2 < str.size() && str.at(iter) == '0' && str.at(iter + 1) == 'x') + return std::stoi(str.substr(iter + 2), nullptr, 16); + else if(iter < str.size()) + return atoi(str.data() + iter); + else + return 0; +} + +std::vector +MemoryCounter::join_and_reset(int64_t offset, std::vector& flats) +{ + std::vector ret = std::move(flats); + flats.clear(); + + offset = std::max(std::min(offset, list.size()), 0l); + ret.insert(ret.end(), list.begin(), list.begin() + offset); + list.erase(list.begin(), list.begin() + offset); + return ret; +} + +std::optional> +MemoryCounter::handle_mem_op(const std::string& inst, std::vector& flat_list) +{ + int64_t wait_n = extract_waitcnt(inst); + + if(wait_n == 0) order = Ordering::MEMORY_SEQUENTIAL; + + if(order == Ordering::MEMORY_SEQUENTIAL) + { + auto joined = join_and_reset(list.size() - wait_n, flat_list); + if(!joined.empty()) return joined; + } + return std::nullopt; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/analysis.hpp b/source/lib/rocprofiler-sdk-att/waitcnt/analysis.hpp new file mode 100644 index 0000000000..c8a35fa770 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/analysis.hpp @@ -0,0 +1,107 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/rocprofiler-sdk-att/att_decoder.h" +#include "lib/rocprofiler-sdk-att/att_lib_wrapper.hpp" +#include "lib/rocprofiler-sdk-att/code.hpp" + +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +struct LineWaitcnt +{ + int line_number{0}; + std::vector dependencies{}; +}; + +struct WaitcntList +{ + using isa_map_t = std::map>; + using wave_t = att_wave_data_t; + + WaitcntList() = default; + + static const WaitcntList& Get(int gfxip, const wave_t& wave, isa_map_t& isa_map) + { + auto it = _cache.find(wave.traceID); + if(it != _cache.end()) return *it->second; + + auto ptr = std::make_unique(); + + if(gfxip == 9) + ptr->mem_unroll = gfx9_construct(wave, isa_map); + else if(gfxip == 10 || gfxip == 11) + ptr->mem_unroll = gfx10_construct(wave, isa_map); + else if(gfxip == 12) + ptr->mem_unroll = gfx12_construct(wave, isa_map); + else + throw std::runtime_error("Invalid gfxip: " + std::to_string(gfxip)); + + return *_cache.emplace(wave.traceID, std::move(ptr)).first->second; + } + + static std::vector gfx9_construct(const wave_t& wave, isa_map_t& isa_map); + static std::vector gfx10_construct(const wave_t& wave, isa_map_t& isa_map); + static std::vector gfx12_construct(const wave_t& wave, isa_map_t& isa_map); + + std::vector mem_unroll{}; + +private: + static std::map> _cache; +}; + +class MemoryCounter +{ +public: + enum Ordering + { + MEMORY_SEQUENTIAL = 0, + MEMORY_PARALLEL + }; + + MemoryCounter(std::string _name) + : name(_name) + {} + + int64_t extract_waitcnt(const std::string& str) const; + + std::vector join_and_reset(int64_t offset, std::vector& flats); + + std::optional> handle_mem_op(const std::string& inst, + std::vector& flat_list); + + const std::string name; + Ordering order = Ordering::MEMORY_SEQUENTIAL; + std::vector list{}; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/gfx10.cpp b/source/lib/rocprofiler-sdk-att/waitcnt/gfx10.cpp new file mode 100644 index 0000000000..907e3a04dc --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/gfx10.cpp @@ -0,0 +1,216 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/rocprofiler-sdk-att/waitcnt/analysis.hpp" + +#include +#include + +#define CLASS_BITS (0xF) + +namespace rocprofiler +{ +namespace att_wrapper +{ +enum MemoryInstType +{ + TYPE_UNCLASSIFIED = 0, + TYPE_NOT_MEM, + TYPE_SCALAR, + TYPE_LDS_MSG, + TYPE_MSG_RTN, + TYPE_GLOBAL_LOAD, + TYPE_GLOBAL_STOR, + TYPE_FLAT_LOAD, + TYPE_FLAT_STOR, + TYPE_WAITCNT +}; + +union MemoryInst +{ + MemoryInst() = default; + MemoryInst(int type) + : raw(type){}; + + struct + { + int inst : 8; // One of MemoryInstType + int lgkm : 1; + int vmcn : 1; + int vscn : 1; + }; + int raw = 0; +}; + +static_assert(sizeof(MemoryInst) == sizeof(int)); + +static MemoryInst +classify(const std::string& inst) +{ + constexpr size_t npos = std::string::npos; + + if(inst.find("s_wait") == 0) + { + if(inst.find("s_wait_alu") != npos) return MemoryInstType::TYPE_NOT_MEM; + + MemoryInst type = MemoryInstType::TYPE_WAITCNT; + if(inst.find("lgkm") != npos) type.lgkm = true; + if(inst.find("vscnt") != npos) type.vscn = true; + if(inst.find("vmcnt") != npos) type.vmcn = true; + + return type; + } + + if(inst.find("v_") == 0) return MemoryInstType::TYPE_NOT_MEM; + + if(inst.find("s_") == 0) + { + if(inst.find("s_load") == 0 || inst.find("s_store") == 0) + return MemoryInstType::TYPE_SCALAR; + + if(inst.find("s_sendmsg") == 0) + { + if(inst.find("s_sendmsg_rtn") == 0) + return MemoryInstType::TYPE_MSG_RTN; + else + return MemoryInstType::TYPE_LDS_MSG; + } + + return MemoryInstType::TYPE_NOT_MEM; + } + + bool bStore = inst.find("store") != npos; + + if(inst.find("global_") == 0 || inst.find("buffer_") <= 1 || inst.find("scratch_") == 0) + return bStore ? MemoryInstType::TYPE_GLOBAL_STOR : MemoryInstType::TYPE_GLOBAL_LOAD; + else if(inst.find("flat_") == 0) + return bStore ? MemoryInstType::TYPE_FLAT_STOR : MemoryInstType::TYPE_FLAT_LOAD; + else if(inst.find("ds_") == 0) + return MemoryInstType::TYPE_LDS_MSG; + else + return MemoryInstType::TYPE_NOT_MEM; +} + +std::vector +WaitcntList::gfx10_construct(const wave_t& wave, isa_map_t& isa_map) +{ + std::vector mem_unroll; + MemoryCounter lgkm("lgkmcnt"); + MemoryCounter vmcnt("vmcnt"); + MemoryCounter vscnt("vscnt"); + + std::vector flat_load{}; + std::vector flat_stor{}; + + for(size_t i = 0; i < wave.instructions_size; i++) + { + auto& event = wave.instructions_array[i]; + auto it = isa_map.find(event.pc); + if(it == isa_map.end() || !it->second->code_line || it->second->code_line->inst.empty()) + { + static thread_local std::once_flag failed_flag{}; + std::call_once(failed_flag, []() { ROCP_WARNING << "Analysis failed!"; }); + break; + } + + MemoryInst type = it->second->type; + const std::string& inst_str = it->second->code_line->inst; + + if(type.inst == MemoryInstType::TYPE_UNCLASSIFIED) + { + type = classify(inst_str); + it->second->type = type.raw; + } + + int line_number = it->second->line_number; + + if(type.inst == MemoryInstType::TYPE_NOT_MEM) continue; + + if(type.inst == MemoryInstType::TYPE_LDS_MSG) + { + lgkm.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_MSG_RTN) + { + lgkm.list.push_back(line_number); // Increments by 2 + lgkm.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_SCALAR) + { + lgkm.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + lgkm.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_GLOBAL_LOAD) + { + vmcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_GLOBAL_STOR) + { + vscnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_FLAT_LOAD) + { + lgkm.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + vmcnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + flat_load.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_FLAT_STOR) + { + lgkm.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + vscnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + flat_stor.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_WAITCNT) + { + if(type.vscn) + { + if(auto joined = vscnt.handle_mem_op(inst_str, flat_stor)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.vmcn) + { + if(auto joined = vmcnt.handle_mem_op(inst_str, flat_load)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.lgkm) + { + if(auto joined = lgkm.handle_mem_op(inst_str, flat_load)) + { + if(!flat_stor.empty()) + { + if(auto stor = lgkm.handle_mem_op(inst_str, flat_stor)) + joined->insert(joined->end(), stor->begin(), stor->end()); + } + + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + } + } + } + + return mem_unroll; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/gfx12.cpp b/source/lib/rocprofiler-sdk-att/waitcnt/gfx12.cpp new file mode 100644 index 0000000000..6895c5c752 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/gfx12.cpp @@ -0,0 +1,276 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/rocprofiler-sdk-att/waitcnt/analysis.hpp" + +#include +#include + +#define CLASS_BITS (0xF) + +namespace rocprofiler +{ +namespace att_wrapper +{ +enum MemoryInstType +{ + TYPE_UNCLASSIFIED = 0, + TYPE_NOT_MEM, + TYPE_SCALAR, + TYPE_LDS, + TYPE_LDS_DIRECT, + TYPE_MSG, + TYPE_MSG_RTN, + TYPE_GLOBAL_LOAD, + TYPE_GLOBAL_STOR, + TYPE_FLAT_LOAD, + TYPE_FLAT_STOR, + TYPE_BVH, + TYPE_WAITCNT +}; + +union MemoryInst +{ + MemoryInst() = default; + MemoryInst(int type) + : raw(type){}; + + struct + { + int inst : 8; // One of MemoryInstType + int dscnt : 1; + int bvhcn : 1; + int expcn : 1; + int kmcnt : 1; + int ldcnt : 1; + int stcnt : 1; + int sampl : 1; + }; + int raw = 0; +}; + +static_assert(sizeof(MemoryInst) == sizeof(int)); + +static MemoryInst +classify(const std::string& inst) +{ + constexpr size_t npos = std::string::npos; + + if(inst.find("s_wait") == 0) + { + if(inst.find("s_wait_alu") != npos) return MemoryInstType::TYPE_NOT_MEM; + + MemoryInst type = MemoryInstType::TYPE_WAITCNT; + if(inst.find("dscnt") != npos) type.dscnt = true; + if(inst.find("bvhcnt") != npos) type.bvhcn = true; + if(inst.find("expcnt") != npos) type.expcn = true; + if(inst.find("kmcnt") != npos) type.kmcnt = true; + if(inst.find("loadcnt") != npos) type.ldcnt = true; + if(inst.find("storecnt") != npos) type.stcnt = true; + if(inst.find("samplecnt") != npos) type.sampl = true; + + return type; + } + + if(inst.find("v_") == 0) return MemoryInstType::TYPE_NOT_MEM; + + if(inst.find("s_") == 0) + { + if(inst.find("s_load") == 0 || inst.find("s_store") == 0) + return MemoryInstType::TYPE_SCALAR; + + if(inst.find("s_sendmsg") == 0) + { + if(inst.find("s_sendmsg_rtn") == 0) + return MemoryInstType::TYPE_MSG_RTN; + else + return MemoryInstType::TYPE_MSG; + } + + return MemoryInstType::TYPE_NOT_MEM; + } + + if(inst.find("ds_") <= 1) + { + if(inst.find("param_load") != npos || inst.find("ds_direct") != npos) + return MemoryInstType::TYPE_LDS_DIRECT; + else + return MemoryInstType::TYPE_LDS; + } + + bool bStore = inst.find("store") != npos || inst.find("global_wb") == 0; + + if(inst.find("global_") == 0 || inst.find("buffer_") <= 1 || inst.find("scratch_") == 0) + return bStore ? MemoryInstType::TYPE_GLOBAL_STOR : MemoryInstType::TYPE_GLOBAL_LOAD; + else if(inst.find("flat_") == 0) + return bStore ? MemoryInstType::TYPE_FLAT_STOR : MemoryInstType::TYPE_FLAT_LOAD; + else if(inst.find("bvh_") != npos) + return MemoryInstType::TYPE_BVH; + else + return MemoryInstType::TYPE_NOT_MEM; +} + +std::vector +WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map) +{ + static std::vector empty_list{}; + std::vector mem_unroll{}; + + MemoryCounter loadcnt("loadcnt"); + MemoryCounter storecnt("storecnt"); + MemoryCounter samplecnt("samplecnt"); + MemoryCounter dscnt("dscnt"); + MemoryCounter kmcnt("kmcnt"); + MemoryCounter expcnt("expcnt"); + MemoryCounter bvhcnt("bvhcnt"); + + std::vector flat_load{}; + std::vector flat_stor{}; + + for(size_t i = 0; i < wave.instructions_size; i++) + { + auto& event = wave.instructions_array[i]; + auto it = isa_map.find(event.pc); + if(it == isa_map.end() || !it->second->code_line || it->second->code_line->inst.empty()) + { + static thread_local std::once_flag failed_flag{}; + std::call_once(failed_flag, []() { ROCP_WARNING << "Analysis failed!"; }); + break; + } + + MemoryInst type = it->second->type; + const std::string& inst_str = it->second->code_line->inst; + + if(type.inst == MemoryInstType::TYPE_UNCLASSIFIED) + { + type = classify(inst_str); + it->second->type = type.raw; + } + + int line_number = it->second->line_number; + + if(type.inst == MemoryInstType::TYPE_NOT_MEM) continue; + + if(type.inst == MemoryInstType::TYPE_LDS) + { + dscnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_LDS_DIRECT) + { + expcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_MSG) + { + kmcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_MSG_RTN) + { + kmcnt.list.push_back(line_number); + kmcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_SCALAR) + { + kmcnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + kmcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_GLOBAL_LOAD) + { + loadcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_GLOBAL_STOR) + { + storecnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_FLAT_LOAD) + { + dscnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + loadcnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + flat_load.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_FLAT_STOR) + { + dscnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + storecnt.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + flat_stor.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_BVH) + { + bvhcnt.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_WAITCNT) + { + if(type.bvhcn) + { + if(auto joined = bvhcnt.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + if(type.expcn) + { + if(auto joined = expcnt.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.sampl) + { + if(auto joined = samplecnt.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.kmcnt) + { + if(auto joined = kmcnt.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.stcnt) + { + if(auto joined = storecnt.handle_mem_op(inst_str, flat_stor)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.ldcnt) + { + if(auto joined = loadcnt.handle_mem_op(inst_str, flat_load)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type.dscnt) + { + if(auto joined = dscnt.handle_mem_op(inst_str, flat_load)) + { + if(!flat_stor.empty()) + { + if(auto stor = dscnt.handle_mem_op(inst_str, flat_stor)) + joined->insert(joined->end(), stor->begin(), stor->end()); + } + + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + } + } + } + + return mem_unroll; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/gfx9.cpp b/source/lib/rocprofiler-sdk-att/waitcnt/gfx9.cpp new file mode 100644 index 0000000000..19d7670198 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/gfx9.cpp @@ -0,0 +1,150 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/rocprofiler-sdk-att/waitcnt/analysis.hpp" + +#include +#include + +#define CLASS_BITS (0xF) +#define LGK_BIT (1 << 4) +#define VMEM_BIT (1 << 5) + +namespace rocprofiler +{ +namespace att_wrapper +{ +enum MemoryInstType +{ + TYPE_UNCLASSIFIED = 0, + TYPE_NOT_MEM, + TYPE_SCALAR, + TYPE_LDS_MSG, + TYPE_GLOBAL, + TYPE_FLAT, + TYPE_WAITCNT +}; + +int +classify(const std::string& inst) +{ + if(inst.find("s_waitcnt") == 0) + { + int type = TYPE_WAITCNT; + if(inst.find("lgk") != std::string::npos) type |= LGK_BIT; + if(inst.find("vmcnt") != std::string::npos) type |= VMEM_BIT; + + return type; + } + + if(inst.find("v_") == 0) return MemoryInstType::TYPE_NOT_MEM; + + if(inst.find("s_") == 0) + { + if(inst.find("s_load") == 0 || inst.find("s_store") == 0) + return MemoryInstType::TYPE_SCALAR; + else if(inst.find("s_sendmsg") == 0) + return MemoryInstType::TYPE_LDS_MSG; + else + return MemoryInstType::TYPE_NOT_MEM; + } + + if(inst.find("global_") == 0 || inst.find("buffer_") <= 1 || inst.find("scratch_") == 0) + return MemoryInstType::TYPE_GLOBAL; + else if(inst.find("flat_") == 0) + return MemoryInstType::TYPE_FLAT; + else if(inst.find("ds_") == 0) + return MemoryInstType::TYPE_LDS_MSG; + else + return MemoryInstType::TYPE_NOT_MEM; +} + +std::vector +WaitcntList::gfx9_construct(const wave_t& wave, isa_map_t& isa_map) +{ + std::vector mem_unroll; + + MemoryCounter lgkm("lgkmcnt"); + MemoryCounter vmem("vmcnt"); + std::vector vflat_list; + + for(size_t i = 0; i < wave.instructions_size; i++) + { + auto& event = wave.instructions_array[i]; + auto it = isa_map.find(event.pc); + if(it == isa_map.end() || !it->second->code_line || it->second->code_line->inst.empty()) + { + static thread_local std::once_flag failed_flag{}; + std::call_once(failed_flag, []() { ROCP_WARNING << "Analysis failed!"; }); + break; + } + + int& type = it->second->type; + const std::string& inst_str = it->second->code_line->inst; + + if(type == TYPE_UNCLASSIFIED) type = classify(inst_str); + + int line_number = it->second->line_number; + int typeclass = type & CLASS_BITS; + + if(typeclass == MemoryInstType::TYPE_NOT_MEM) continue; + + if(typeclass == MemoryInstType::TYPE_LDS_MSG) + { + lgkm.list.push_back(line_number); + } + else if(typeclass == MemoryInstType::TYPE_SCALAR) + { + lgkm.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + lgkm.list.push_back(line_number); + } + else if(typeclass == MemoryInstType::TYPE_GLOBAL) + { + vmem.list.push_back(line_number); + } + else if(typeclass == MemoryInstType::TYPE_FLAT) + { + lgkm.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + vmem.order = MemoryCounter::Ordering::MEMORY_PARALLEL; + vflat_list.push_back(line_number); + } + else if(typeclass == MemoryInstType::TYPE_WAITCNT) + { + if(type & LGK_BIT) + { + if(auto joined = lgkm.handle_mem_op(inst_str, vflat_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + + if(type & VMEM_BIT) + { + if(auto joined = vmem.handle_mem_op(inst_str, vflat_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + } + } + + return mem_unroll; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/waitcnt/tests/CMakeLists.txt new file mode 100644 index 0000000000..372e8786e1 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/tests/CMakeLists.txt @@ -0,0 +1,24 @@ +# +# Waitcnt ISA tests +# +project(rocprofiler-att-parser-waitcnt-tests LANGUAGES CXX) + +add_executable(att-decoder-waitcnt-test) +target_link_libraries( + att-decoder-waitcnt-test + PRIVATE rocprofiler-sdk-att-parser + rocprofiler-sdk::rocprofiler-sdk-json + rocprofiler-sdk::rocprofiler-sdk-common-library + rocprofiler-sdk::rocprofiler-sdk-glog + rocprofiler-sdk::rocprofiler-sdk-static-library + GTest::gtest + GTest::gtest_main) +target_sources(att-decoder-waitcnt-test PRIVATE att_decoder_waitcnt_test.cpp) + +gtest_add_tests( + TARGET att-decoder-waitcnt-test + SOURCES att_decoder_waitcnt_test.cpp + TEST_LIST att-decoder-waitcnt-test_TESTS + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) +set_tests_properties(${att-decoder-waitcnt-test_TESTS} PROPERTIES TIMEOUT 10 LABELS + "unittests") diff --git a/source/lib/rocprofiler-sdk-att/waitcnt/tests/att_decoder_waitcnt_test.cpp b/source/lib/rocprofiler-sdk-att/waitcnt/tests/att_decoder_waitcnt_test.cpp new file mode 100644 index 0000000000..9f28b586b5 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/waitcnt/tests/att_decoder_waitcnt_test.cpp @@ -0,0 +1,375 @@ +// MIT License +// +// Copyright (c) 2024 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 "lib/common/logging.hpp" +#include "lib/rocprofiler-sdk-att/waitcnt/analysis.hpp" +#include "lib/rocprofiler-sdk/registration.hpp" + +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +TEST(att_decoder_waitcnt_test, gfx9) +{ + registration::init_logging(); + + constexpr size_t LOOP_CNT = 4; + + WaitcntList::isa_map_t isa_map{}; + + auto append_isa = [&](size_t line_number, const char* line) { + pcinfo_t pc{}; + pc.addr = line_number; + pc.marker_id = 0; + + auto code = std::make_unique(); + code->code_line = std::make_shared(); + code->code_line->inst = line; + code->line_number = line_number; + + isa_map.emplace(pc, std::move(code)); + }; + + append_isa(0, "s_nop 0"); + append_isa(1, "v_add_ 0"); + append_isa(2, "s_waitcnt vmcnt(0) lgkmcnt(0)"); + append_isa(3, "s_add_ 0"); + + append_isa(4, "global_load_"); + append_isa(5, "buffer_store_"); + append_isa(6, "scratch_load_"); + append_isa(7, "s_waitcnt lkgmcnt(0)"); + append_isa(8, "s_waitcnt vmcnt(2)"); + append_isa(9, "s_waitcnt vmcnt(1)"); + append_isa(10, "s_load_"); + append_isa(11, "s_store_"); + append_isa(12, "s_sendmsg "); + append_isa(13, "s_waitcnt vmcnt(0) lgkmcnt(0)"); + append_isa(14, "ds_load_"); + append_isa(15, "ds_store_"); + append_isa(16, "ds_load_"); + append_isa(17, "s_waitcnt lgkmcnt(2)"); + append_isa(18, "s_waitcnt lgkmcnt(1)"); + append_isa(19, "flat_load_"); + append_isa(20, "s_waitcnt vmcnt( 0) lgkmcnt(0x0)"); // some weird strings + append_isa(21, "invalid"); + + std::vector insts{}; + + for(size_t j = 0; j < LOOP_CNT; j++) + { + for(size_t i = 0; i < isa_map.size(); i++) + { + att_wave_instruction_t inst{}; + inst.pc.addr = i; + insts.push_back(inst); + } + } + + WaitcntList::wave_t wave{}; + wave.traceID = 1; + wave.instructions_array = insts.data(); + wave.instructions_size = insts.size(); + + auto& data = WaitcntList::Get(9, wave, isa_map); + + std::map> dependencies{}; + + for(auto& [dst, src] : data.mem_unroll) + { + auto& dep = dependencies[dst]; + for(auto& p : src) + dep.insert(p); + } + + auto set_equal = [&](int dep, const std::set& set) { + for(int s : set) + ASSERT_NE(dependencies.at(dep).find(s), dependencies.at(dep).end()); + ASSERT_EQ(dependencies.at(dep).size(), set.size()); + }; + + ASSERT_EQ(dependencies.size(), 6); + set_equal(8, {4}); + set_equal(9, {5}); + set_equal(13, {6, 10, 11, 12}); + set_equal(17, {14}); + set_equal(18, {15}); + set_equal(20, {16, 19}); +} + +TEST(att_decoder_waitcnt_test, gfx10) +{ + registration::init_logging(); + + WaitcntList::isa_map_t isa_map{}; + + auto append_isa = [&](size_t line_number, const char* line) { + pcinfo_t pc{}; + pc.addr = line_number; + pc.marker_id = 0; + + auto code = std::make_unique(); + code->code_line = std::make_shared(); + code->code_line->inst = line; + code->line_number = line_number; + + isa_map.emplace(pc, std::move(code)); + }; + + append_isa(0, "buffer_load_"); + append_isa(1, "global_load_"); + append_isa(2, "v_add_ 0"); + append_isa(3, "s_add_ 0"); + append_isa(4, "buffer_store_"); + append_isa(5, "s_waitcnt vmcnt(1)"); + + append_isa(6, "scratch_load_"); + append_isa(7, "scratch_store_"); + append_isa(8, "s_wait_alu "); + append_isa(9, "s_waitcnt vmcnt 0x2"); + append_isa(10, "s_waitcnt vmcnt(1)"); + append_isa(11, "s_waitcnt vscnt(1)"); + append_isa(12, "s_waitcnt vmcnt(0)"); + append_isa(13, "s_waitcnt vscnt(0)"); + + append_isa(14, "s_load"); + append_isa(15, "s_store"); + append_isa(16, "s_waitcnt lgkmcnt 0"); + append_isa(17, "s_sendmsg"); + append_isa(18, "s_sendmsg_rtn"); + append_isa(19, "s_waitcnt lgkmcnt 0x2"); // waits on sendmsg_rtn + append_isa(20, "flat_load_"); + append_isa(21, "flat_store_"); + append_isa(22, "s_waitcnt vmcnt(0) lgkmcnt(0) vscnt(0)"); + + append_isa(23, "ds_load"); + append_isa(24, "ds_store"); + append_isa(25, "s_waitcnt lgkmcnt 0x1"); + append_isa(26, "s_waitcnt lgkmcnt 0"); + append_isa(27, "invalid"); + + std::vector insts{}; + for(size_t i = 0; i < isa_map.size(); i++) + { + att_wave_instruction_t inst{}; + inst.pc.addr = i; + insts.push_back(inst); + } + + WaitcntList::wave_t wave{}; + wave.traceID = 2; + wave.instructions_array = insts.data(); + wave.instructions_size = insts.size(); + + auto& data = WaitcntList::Get(10, wave, isa_map); + + std::map> dependencies{}; + + for(auto& [dst, src] : data.mem_unroll) + { + auto& dep = dependencies[dst]; + for(auto& p : src) + dep.insert(p); + } + + auto set_equal = [&](int dep, const std::set& set) { + for(int s : set) + ASSERT_NE(dependencies.at(dep).find(s), dependencies.at(dep).end()); + ASSERT_EQ(dependencies.at(dep).size(), set.size()); + }; + + ASSERT_EQ(dependencies.size(), 10); + set_equal(5, {0}); + set_equal(10, {1}); + set_equal(11, {4}); + set_equal(12, {6}); + set_equal(13, {7}); + set_equal(16, {14, 15}); + set_equal(19, {17}); + set_equal(22, {18, 20, 21}); + set_equal(25, {23}); + set_equal(26, {24}); +} + +TEST(att_decoder_waitcnt_test, gfx12) +{ + registration::init_logging(); + + WaitcntList::isa_map_t isa_map{}; + + auto append_isa = [&](size_t line_number, const char* line) { + pcinfo_t pc{}; + pc.addr = line_number; + pc.marker_id = 0; + + auto code = std::make_unique(); + code->code_line = std::make_shared(); + code->code_line->inst = line; + code->line_number = line_number; + + isa_map.emplace(pc, std::move(code)); + }; + + // messages + append_isa(0, "s_wait_alu "); + append_isa(1, "s_waitcnt samplecnt(0)"); + append_isa(2, "s_sendmsg "); + append_isa(3, "s_sendmsg "); + append_isa(4, "s_sendmsg_rtn"); + append_isa(5, "s_sendmsg_rtn"); + append_isa(6, "s_waitcnt kmcnt(4)"); + append_isa(7, "s_waitcnt kmcnt(2)"); + append_isa(8, "s_waitcnt kmcnt(0)"); + + // scalar + append_isa(9, "s_load_"); + append_isa(10, "s_store_"); + append_isa(11, "s_sendmsg "); + append_isa(12, "s_waitcnt kmcnt(0)"); + + // flat + append_isa(13, "flat_load_"); + append_isa(14, "flat_store_"); + append_isa(15, "global_load_"); + append_isa(16, "ds_load"); + append_isa(17, + "s_waitcnt bvhcnt(0) expcnt(0) kmcnt(0) kmcnt(0) loadcnt(0) storecnt(0) " + "samplecnt(0) dscnt(0)"); + + // load/store + append_isa(18, "global_load"); + append_isa(19, "buffer_load"); + append_isa(20, "global_store"); + append_isa(21, "global_wb"); + append_isa(22, "buffer_store"); + append_isa(23, "scratch_load"); + append_isa(24, "scratch_store"); + append_isa(25, "s_waitcnt loadcnt(2)"); + append_isa(26, "s_waitcnt storecnt(2)"); + append_isa(27, "s_waitcnt storecnt(0) loadcnt(0)"); + + // skipped + append_isa(28, "s_wait_alu"); + append_isa(29, "s_mul "); + append_isa(30, "v_mul "); + + // ds vs exp vs bvh + append_isa(31, "ds_store"); + append_isa(32, "bvh_"); + append_isa(33, "bvh_"); + append_isa(34, "ds_param_load"); + append_isa(35, "ds_direct"); + append_isa(36, "ds_load"); + append_isa(37, "s_waitcnt dscnt(1)"); + append_isa(38, "s_waitcnt expcnt(0) bvhcnt(0)"); + append_isa(39, "s_waitcnt dscnt(0)"); + append_isa(40, "invalid"); + + std::vector insts{}; + for(size_t i = 0; i < isa_map.size(); i++) + { + att_wave_instruction_t inst{}; + inst.pc.addr = i; + insts.push_back(inst); + } + + WaitcntList::wave_t wave{}; + wave.traceID = 3; + wave.instructions_array = insts.data(); + wave.instructions_size = insts.size(); + + auto& data = WaitcntList::Get(12, wave, isa_map); + + std::map> dependencies{}; + + for(auto& [dst, src] : data.mem_unroll) + { + auto& dep = dependencies[dst]; + for(auto& p : src) + dep.insert(p); + } + + auto set_equal = [&](int dep, const std::set& set) { + for(int s : set) + ASSERT_NE(dependencies.at(dep).find(s), dependencies.at(dep).end()); + ASSERT_EQ(dependencies.at(dep).size(), set.size()); + }; + + ASSERT_EQ(dependencies.size(), 11); + set_equal(6, {2, 3}); + set_equal(7, {4}); + set_equal(8, {5}); + set_equal(12, {9, 10, 11}); + set_equal(17, {13, 14, 15, 16}); + set_equal(25, {18}); + set_equal(26, {20, 21}); + set_equal(27, {19, 22, 23, 24}); + set_equal(37, {31}); + set_equal(38, {32, 33, 34, 35}); + set_equal(39, {36}); +} + +TEST(att_decoder_waitcnt_test, fail_conditions) +{ + registration::init_logging(); + + WaitcntList::isa_map_t isa_map{}; + + std::vector insts{}; + + for(size_t i = 0; i < 10; i++) + { + att_wave_instruction_t inst{}; + inst.pc.addr = i; + insts.push_back(inst); + } + + WaitcntList::wave_t wave{}; + wave.traceID = 4; + wave.instructions_array = insts.data(); + wave.instructions_size = insts.size(); + + // It should give warning and return + ASSERT_TRUE(WaitcntList::Get(9, wave, isa_map).mem_unroll.empty()); + wave.traceID++; + ASSERT_TRUE(WaitcntList::Get(10, wave, isa_map).mem_unroll.empty()); + wave.traceID++; + ASSERT_TRUE(WaitcntList::Get(12, wave, isa_map).mem_unroll.empty()); + wave.traceID++; + + // it cant operate on invalid gfxip + try + { + WaitcntList::Get(-1, wave, isa_map); + // fail + ASSERT_TRUE(false); + } catch(std::runtime_error& e) + { + // pass + } +} + +}; // namespace att_wrapper +}; // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/wave.cpp b/source/lib/rocprofiler-sdk-att/wave.cpp new file mode 100644 index 0000000000..b729580aa0 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/wave.cpp @@ -0,0 +1,133 @@ +// MIT License +// +// Copyright (c) 2024 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 "wave.hpp" +#include +#include "outputfile.hpp" + +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +WaveFile::WaveFile(WaveConfig& config, const att_wave_data_t& wave) +{ + if(wave.contexts) ROCP_WARNING << "Wave had " << wave.contexts << " context save-restores"; + + if(!GlobalDefs::get().has_format("json")) return; + if(wave.instructions_size == 0 && wave.timeline_size < 3) return; + + assert(config.filemgr); + + int assigned_id = config.id_count.at(wave.simd).at(wave.wave_id).fetch_add(1); + { + std::stringstream namess; + namess << "se" << config.shader_engine << "_sm" << (int) wave.simd << "_sl" + << (int) wave.wave_id << "_wv" << assigned_id << ".json"; + + filename = config.filemgr->dir / namess.str(); + config.filemgr->addwave( + filename, + FilenameMgr::Coord{ + config.shader_engine, (int) wave.simd, (int) wave.wave_id, assigned_id}, + wave.begin_time, + wave.end_time); + } + + nlohmann::json instructions; + + for(size_t i = 0; i < wave.instructions_size; i++) + { + auto& inst = wave.instructions_array[i]; + instructions.push_back({(int64_t) inst.time, + (int) inst.category, + (int) inst.stall, + (int64_t) inst.duration, + config.code->line_numbers[inst.pc]}); + } + + nlohmann::json timeline; + int64_t acc_time = wave.begin_time; + + for(size_t i = 0; i < wave.timeline_size; i++) + { + int type = wave.timeline_array[i].type; + int duration = wave.timeline_array[i].duration; + + config.wstates.at(type)->add(acc_time, duration); + timeline.push_back({type, duration}); + + acc_time += duration; + } + + nlohmann::json waitcnt; + + try + { + const WaitcntList& wait_list = + WaitcntList::Get(config.filemgr->gfxip, wave, config.code->isa_map); + + for(auto& line : wait_list.mem_unroll) + if(line.dependencies.size()) + { + nlohmann::json json_line; + for(int dep : line.dependencies) + json_line.push_back({dep, 0}); + waitcnt.push_back({line.line_number, json_line}); + } + } catch(std::exception& e) + { + ROCP_ERROR << e.what(); + } catch(...) + { + ROCP_ERROR << "Generic error"; + } + + nlohmann::json wave_entry = { + {"cu", wave.cu}, + {"id", assigned_id}, + {"simd", wave.simd}, + {"slot", wave.wave_id}, + {"begin", wave.begin_time}, + {"end", wave.end_time}, + + {"instructions", instructions}, + {"timeline", timeline}, + {"waitcnt", waitcnt}, + }; + + nlohmann::json metadata = { + {"name", "SE" + std::to_string(config.shader_engine)}, + {"duration", wave.end_time - wave.begin_time}, + {"wave", wave_entry}, + {"num_stitched", wave.instructions_size}, + {"num_insts", wave.instructions_size}, + }; + + OutputFile(filename) << metadata; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/wave.hpp b/source/lib/rocprofiler-sdk-att/wave.hpp new file mode 100644 index 0000000000..a836ae325f --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/wave.hpp @@ -0,0 +1,80 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_decoder.h" +#include "code.hpp" +#include "filenames.hpp" +#include "waitcnt/analysis.hpp" +#include "wstates.hpp" + +#include "att_lib_wrapper.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +constexpr size_t SIMD_NUM = 4; +constexpr size_t SIMD_SIZE = 32; + +class WaveConfig +{ + using WavestateArray = std::array, ATT_WAVE_STATE_LAST>; + using SIMD = std::array, SIMD_SIZE>; + +public: + WaveConfig(int se_id, + std::shared_ptr& _mgr, + std::shared_ptr& _code, + WavestateArray& _wstates) + : shader_engine(se_id) + , wstates(_wstates) + , code(_code) + , filemgr(_mgr) + {} + + const int shader_engine; + WavestateArray wstates; + + std::array id_count{}; + std::shared_ptr code; + std::shared_ptr filemgr; + + std::map kernel_names{}; + std::vector occupancy{}; +}; + +class WaveFile +{ +public: + WaveFile(WaveConfig& config, const att_wave_data_t& wave); + Fspath filename{}; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/wstates.cpp b/source/lib/rocprofiler-sdk-att/wstates.cpp new file mode 100644 index 0000000000..ca46d3bd89 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/wstates.cpp @@ -0,0 +1,76 @@ +// MIT License +// +// Copyright (c) 2024 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 "wstates.hpp" +#include "outputfile.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +WstatesFile::WstatesFile(int state, const Fspath& dir) +: filename(dir / ("wstates" + std::to_string(state) + ".json")) +{} + +WstatesFile::~WstatesFile() +{ + if(!events.size() || !GlobalDefs::get().has_format("json")) return; + + std::sort(events.begin(), events.end(), [](const event_t& a, const event_t& b) { + return a.first < b.first; + }); + + int64_t accum = 0; + int64_t prev_time = INT64_MIN; + nlohmann::json jtime; + nlohmann::json jstate; + + for(auto& [time, value] : events) + { + accum += value; + if(!jtime.size() || time != prev_time) + { + jtime.push_back(time); + jstate.push_back(accum); + } + else + { + jstate.back() = accum; + } + prev_time = time; + }; + + nlohmann::json jfile; + jfile["time"] = jtime; + jfile["state"] = jstate; + jfile["name"] = filename.filename(); + + OutputFile(filename) << jfile; +} + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/wstates.hpp b/source/lib/rocprofiler-sdk-att/wstates.hpp new file mode 100644 index 0000000000..6b8d835593 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/wstates.hpp @@ -0,0 +1,57 @@ +// MIT License +// +// Copyright (c) 2024 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 "att_lib_wrapper.hpp" + +#include +#include +#include +#include "util.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +class WstatesFile +{ + using event_t = std::pair; + +public: + WstatesFile(int state, const Fspath& dir); + ~WstatesFile(); + + void add(int64_t time, int64_t duration) + { + events.emplace_back(time, 1); + events.emplace_back(time + duration, -1); + }; + + Fspath filename{}; + std::vector events{}; + + static constexpr size_t NUM_WSTATES = 5; +}; + +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h b/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h index 4692072e4f..e67239b477 100644 --- a/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h +++ b/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h @@ -411,244 +411,6 @@ aqlprofile_iterate_event_coord(aqlprofile_agent_handle_t agent, aqlprofile_coordinate_callback_t callback, void* userdata); -typedef union -{ - uint64_t raw; - struct - { - uint64_t isValid : 1; - uint64_t isNavi : 1; - uint64_t npiWaveData : 1; - uint64_t version : 13; - }; -} att_output_flags_t; - -typedef struct -{ - int64_t time; - uint16_t events0; - uint16_t events1; - uint16_t events2; - uint16_t events3; - uint8_t CU; - uint8_t bank; -} att_perfevent_t; - -typedef struct -{ - uint64_t kernel_id : 12; - uint64_t simd : 2; - uint64_t slot : 4; - uint64_t enable : 1; - uint64_t cu : 4; - uint64_t time : 41; // Time_value/8 -} att_occupancy_info_t; - -typedef struct -{ - int32_t type; - int32_t duration; -} wave_state_t; - -typedef struct -{ - int64_t time; - int64_t duration; -} wave_instruction_t; - -enum WAVESLOT_STATE -{ - WS_EMPTY = 0, - WS_IDLE = 1, - WS_EXEC = 2, - WS_WAIT = 3, - WS_STALL = 4, - WS_UNKNOWN = 5, -}; - -enum WaveInstCategory -{ - NONE = 0, - SMEM = 1, - SALU = 2, - VMEM = 3, - FLAT = 4, - LDS = 5, - VALU = 6, - JUMP = 7, - NEXT = 8, - IMMED = 9, - TRAP = 10, - PCINFO = 15, - WAVE_NOT_FINISHED, -}; - -enum WaveTrapStatus -{ - TRAP_RESTORED = 0, - TRAP_REQUEST = 1, - TRAP_SAVED = 1, - TRAP_STANDBY = 2 -}; - -typedef struct -{ - size_t addr; - size_t marker_id; -} pcinfo_t; - -typedef struct __attribute__((packed)) -{ - uint64_t category : 8; - uint64_t hitcount : 56; - uint64_t latency; - pcinfo_t pc; -} att_trace_event_t; - -typedef struct -{ - uint8_t simd; - uint8_t wave_id; - uint8_t trap_status; - uint8_t reserved; - - // VMEM Pipeline: instrs and stalls - int num_vmem_instrs; - int num_vmem_stalls; - // FLAT instrs and stalls - int num_flat_instrs; - int num_flat_stalls; - - // LDS instr and stalls - int num_lds_instrs; - int num_lds_stalls; - - // SCA instrs stalls - int num_salu_instrs; - int num_smem_instrs; - int num_salu_stalls; - int num_smem_stalls; - - // Branch - int num_branch_instrs; - int num_branch_taken_instrs; - int num_branch_stalls; - - // total VMEM/FLAT/LDS/SMEM instructions issued - int num_mem_instrs; // total issued memory instructions - int num_valu_stalls; - size_t num_valu_instrs; - size_t num_issued_instrs; // total issued instructions (compute + memory) - - int64_t begin_time; // Begin and end cycle - int64_t end_time; - int64_t traceID; - - size_t timeline_size; - size_t instructions_size; - wave_state_t* timeline_array; - wave_instruction_t* instructions_array; -} wave_data_t; - -/** - * @brief Callback for rocprofiler to return ISA to aqlprofile ATT parser. - * The caller must copy a desired instruction on isa_instruction and source_reference, - * while obeying the max length passed by the caller. - * If the caller's length is insufficient, then this function writes the minimum sizes to isa_size - * and source_size and returns HSA_STATUS_ERROR_OUT_OF_RESOURCES. - * If call returns HSA_STATUS_SUCCESS, isa_size and source_size are written with bytes used. - * @param[out] isa_instruction Where to copy the ISA line to. - * @param[out] source_reference Reference to source line and/or additional comments in the binary. - * @param[out] isa_memory_size (Auto) The number of bytes to next instruction. 0 for custom ISA. - * @param[inout] isa_size Size of returned ISA string. - * @param[inout] source_size Size of returned reference/comment string. - * @param[in] marker_id The generated ATT marker for given codeobject ID. - * @param[in] offset The offset from base vaddr for given codeobj ID. - * If marker_id == 0, this parameter is raw virtual address with no codeobj ID information. - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - * @retval HSA_STATUS_SUCCESS on success. - * @retval HSA_STATUS_ERROR on generic error. - * @retval HSA_STATUS_ERROR_INVALID_ARGUMENT for invalid offset or invalid marker_id. - * @retval HSA_STATUS_ERROR_OUT_OF_RESOURCES for insufficient isa_size or source_size. - */ -typedef hsa_status_t (*aqlprofile_att_isa_callback_t)(char* isa_instruction, - char* source_reference, - uint64_t* isa_memory_size, - uint64_t* isa_size, - uint64_t* source_size, - uint64_t marker_id, - uint64_t offset, - void* userdata); - -/** - * @brief Callback for rocprofiler to return traces back to rocprofiler. - * @param[in] trace_type_id The type of this trace as in _iterate_event_ids(). - * @param[in] correlation_id The ID of shader engine or trace callback number. - * @param[in] trace_events A pointer to sequence of events, of size trace_size. - * @param[in] trace_size The number of events in the trace. - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - */ -typedef hsa_status_t (*aqlprofile_att_trace_callback_t)(int trace_type_id, - int correlation_id, - void* trace_events, - uint64_t trace_size, - void* userdata); - -/** - * @brief Callback for the ATT parser to retrieve Shader Engine data. - * Returns the amount of data filled. If no more data is available, then callback return 0 - * If the space available in the buffer is less than required for parsing the full data, - * the full data is transfered over multiple calls. - * When all data has been transfered from current shader_engine_id, the caller has the option to - * 1) Return -1 on shader_engine ID and parsing terminates - * 2) Move to the next shader engine. - * @param[out] shader_engine_id The ID of given shader engine. - * @param[out] buffer The buffer to fill up with SE data. - * @param[out] buffer_size The space available in the buffer. - * @param[in] userdata Arbitrary data pointer to be sent back to the user via callback. - * @returns Number of bytes remaining in shader engine. - * @retval 0 if no more SE data is available. Parsing will stop. - * @retval buffer_size if the buffer does not hold enough data for the current shader engine. - * @retval 0 > ret > buffer_size for partially filled buffer, and caller moves over to next SE. - */ -typedef uint64_t (*aqlprofile_att_se_data_callback_t)(int* shader_engine_id, - uint8_t** buffer, - uint64_t* buffer_size, - void* userdata); - -/** - * @brief Callback returning from aqlprofile_att_parser_iterate_event_list - * @param[in] trace_event_id ID of the event. - * @param[in] trace_event_metadata Null-terminated string, entries separated by ';' - * @param[in] userdata userdata. - */ -typedef void (*aqlprofile_att_parser_iterate_event_cb_t)(int trace_event_id, - const char* trace_event_metadata, - void* userdata); - -/** - * @brief Iterate over all available event types. - * @param[in] callback Callback where events are returned to. - * @param[in] userdata userdata. - */ -void -aqlprofile_att_parser_iterate_event_list(aqlprofile_att_parser_iterate_event_cb_t callback, - void* userdata); - -/** - * @brief Iterate over all event coordinates for a given agent_t and event_t. - * @param[in] se_data_callback Callback to return shader engine data from. - * @param[in] trace_callback Callback where the trace data is returned to. - * Each trace will be marked by the ID returned on aqlprofile_att_parser_iterate_event_list. - * @param[in] isa_callback Callback to return ISA lines. - * @param[in] userdata Userdata passed back to caller via callback. - */ -hsa_status_t -aqlprofile_att_parse_data(aqlprofile_att_se_data_callback_t se_data_callback, - aqlprofile_att_trace_callback_t trace_callback, - aqlprofile_att_isa_callback_t isa_callback, - void* userdata); - typedef struct { uint64_t id; diff --git a/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt b/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt index 4f68e9da05..b7e1eb107b 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt @@ -1,5 +1,4 @@ -set(ROCPROFILER_LIB_THREAD_TRACE_SOURCES att_core.cpp att_service.cpp att_parser.cpp - code_object.cpp) +set(ROCPROFILER_LIB_THREAD_TRACE_SOURCES att_core.cpp att_service.cpp code_object.cpp) set(ROCPROFILER_LIB_THREAD_TRACE_HEADERS att_core.hpp code_object.hpp) target_sources( rocprofiler-sdk-object-library PRIVATE ${ROCPROFILER_LIB_THREAD_TRACE_SOURCES} diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index ae000d2764..c84b169a1e 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -61,8 +61,9 @@ constexpr uint64_t MAX_BUFFER_SIZE = std::numeric_limits::max(); // aq struct cbdata_t { + rocprofiler_agent_id_t agent; rocprofiler_att_shader_data_callback_t cb_fn; - const rocprofiler_user_data_t* dispatch_userdata; + const rocprofiler_user_data_t* userdata; }; common::Synchronized> client; @@ -238,7 +239,7 @@ thread_trace_callback(uint32_t shader, void* buffer, uint64_t size, void* callba { auto& cb_data = *static_cast(callback_data); - cb_data.cb_fn(shader, buffer, size, *cb_data.dispatch_userdata); + cb_data.cb_fn(cb_data.agent, shader, buffer, size, *cb_data.userdata); return HSA_STATUS_SUCCESS; } @@ -247,8 +248,9 @@ ThreadTracerQueue::iterate_data(aqlprofile_handle_t handle, rocprofiler_user_dat { cbdata_t cb_dt{}; - cb_dt.cb_fn = params.shader_cb_fn; - cb_dt.dispatch_userdata = &data; + cb_dt.agent = agent_id; + cb_dt.cb_fn = params.shader_cb_fn; + cb_dt.userdata = &data; auto status = aqlprofile_att_iterate_data(handle, thread_trace_callback, &cb_dt); CHECK_HSA(status, "Failed to iterate ATT data"); @@ -342,13 +344,13 @@ DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue, }); }; - auto control_flags = params.dispatch_cb_fn(queue.get_id(), - queue.get_agent().get_rocp_agent(), + auto control_flags = params.dispatch_cb_fn(queue.get_agent().get_rocp_agent()->id, + queue.get_id(), rocprof_corr_id, kernel_id, dispatch_id, - user_data, - params.callback_userdata); + params.callback_userdata.ptr, + user_data); if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) { @@ -514,6 +516,7 @@ AgentThreadTracer::start_context() void AgentThreadTracer::stop_context() { + using wait_t = std::tuple>; std::unique_lock lk(agent_mut); if(tracers.empty()) @@ -522,10 +525,9 @@ AgentThreadTracer::stop_context() return; } - std::vector>> - wait_list{}; + std::vector wait_list{}; - for(auto& [id, tracer] : tracers) + for(auto& [_, tracer] : tracers) { auto packet = tracer->get_control(false); packet->populate_after(); @@ -537,8 +539,7 @@ AgentThreadTracer::stop_context() for(auto& [tracer, handle, signal] : wait_list) { signal->WaitOn(); - rocprofiler_user_data_t userdata{.ptr = tracer->params.callback_userdata}; - tracer->iterate_data(handle, userdata); + tracer->iterate_data(handle, tracer->params.callback_userdata); } } diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp index f671c4f025..b8491aada9 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp @@ -58,7 +58,7 @@ struct thread_trace_parameter_pack rocprofiler_context_id_t context_id{0}; rocprofiler_att_dispatch_callback_t dispatch_cb_fn{nullptr}; rocprofiler_att_shader_data_callback_t shader_cb_fn{nullptr}; - void* callback_userdata{nullptr}; + rocprofiler_user_data_t callback_userdata{}; // Parameters uint8_t target_cu = 1; diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_parser.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_parser.cpp deleted file mode 100644 index 4d7ec7b3df..0000000000 --- a/source/lib/rocprofiler-sdk/thread_trace/att_parser.cpp +++ /dev/null @@ -1,213 +0,0 @@ -// MIT License -// -// Copyright (c) 2024 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 -#include -#include -#include -#include -#include -#include "lib/rocprofiler-sdk/aql/aql_profile_v2.h" - -#define AQLPROFILE_OCCUPANCY_RESOLUTION 8 - -namespace rocprofiler -{ -namespace att_parser -{ -hsa_status_t -forward_hsa_error(rocprofiler_status_t error_code) -{ - static thread_local std::unordered_map error_fwd = { - {ROCPROFILER_STATUS_SUCCESS, HSA_STATUS_SUCCESS}, - {ROCPROFILER_STATUS_ERROR, HSA_STATUS_ERROR}, - {ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT, HSA_STATUS_ERROR_INVALID_ARGUMENT}, - {ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES, HSA_STATUS_ERROR_OUT_OF_RESOURCES}, - }; - - try - { - return error_fwd.at(error_code); - } catch(std::exception& e) - {} - - return HSA_STATUS_ERROR; -} - -rocprofiler_status_t -forward_hsa_error(hsa_status_t error_code) -{ - static thread_local std::unordered_map error_fwd = { - {HSA_STATUS_SUCCESS, ROCPROFILER_STATUS_SUCCESS}, - {HSA_STATUS_ERROR, ROCPROFILER_STATUS_ERROR}, - {HSA_STATUS_ERROR_INVALID_ARGUMENT, ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT}, - {HSA_STATUS_ERROR_OUT_OF_RESOURCES, ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES}, - }; - - try - { - return error_fwd.at(error_code); - } catch(std::exception& e) - {} - - return ROCPROFILER_STATUS_ERROR; -} - -struct userdata_callback_table_t -{ - rocprofiler_att_parser_trace_callback_t trace; - rocprofiler_att_parser_isa_callback_t isa; - rocprofiler_att_parser_se_data_callback_t se_data; - void* user; - - std::vector kernel_id_map; -}; - -thread_local int TRACE_DATA_ID{-1}; -thread_local int KERNEL_ADDR_ID{-1}; -thread_local int OCCUPANCY_ID{-1}; - -void -iterate_trace_type(int id, const char* metadata, void*) -{ - if(std::string_view(metadata).find("occupancy") == 0) - OCCUPANCY_ID = id; - else if(std::string_view(metadata).find("kernel_ids_addr") == 0) - KERNEL_ADDR_ID = id; - else if(std::string_view(metadata).find("tracedata") == 0) - TRACE_DATA_ID = id; -} - -hsa_status_t -trace_callback(int trace_type_id, - int /* correlation_id */, - void* trace_events, - uint64_t trace_size, - void* userdata) -{ - assert(userdata); - auto& table = *reinterpret_cast(userdata); - - if(trace_type_id == KERNEL_ADDR_ID) - { - table.kernel_id_map.resize(trace_size); - const auto* events = reinterpret_cast(trace_events); - - for(size_t i = 0; i < trace_size; i++) - table.kernel_id_map.at(i) = events[i]; - } - else if(trace_type_id == OCCUPANCY_ID) - { - const auto* events = reinterpret_cast(trace_events); - for(size_t i = 0; i < trace_size; i++) - { - rocprofiler_att_data_type_occupancy_t occ{}; - occ.timestamp = events[i].time * AQLPROFILE_OCCUPANCY_RESOLUTION; - occ.enabled = events[i].enable; - try - { - pcinfo_t kernel_id_addr = table.kernel_id_map.at(events[i].kernel_id); - occ.marker_id = kernel_id_addr.marker_id; - occ.offset = kernel_id_addr.addr; - } catch(...) - {} // Not having a kernel_id_map entry is unexpected, but valid - table.trace(ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY, (void*) &occ, table.user); - } - } - else if(trace_type_id == TRACE_DATA_ID) - { - const auto* events = reinterpret_cast(trace_events); - for(size_t i = 0; i < trace_size; i++) - { - rocprofiler_att_data_type_isa_t isa{}; - isa.marker_id = events[i].pc.marker_id; - isa.offset = events[i].pc.addr; - isa.hitcount = events[i].hitcount; - isa.latency = events[i].latency; - table.trace(ROCPROFILER_ATT_PARSER_DATA_TYPE_ISA, (void*) &isa, table.user); - } - } - - return HSA_STATUS_SUCCESS; -} - -hsa_status_t -isa_callback(char* isa, - char* /* source_reference */, - uint64_t* memory_size, - uint64_t* isa_size, - uint64_t* source_size, - uint64_t marker, - uint64_t offset, - void* userdata) -{ - assert(userdata); - assert(source_size); - *source_size = 0; - const auto& table = *reinterpret_cast(userdata); - rocprofiler_status_t status = table.isa(isa, memory_size, isa_size, marker, offset, table.user); - - if(status != ROCPROFILER_STATUS_SUCCESS) - return rocprofiler::att_parser::forward_hsa_error(status); - return HSA_STATUS_SUCCESS; -} - -uint64_t -se_data_callback(int* seid, uint8_t** buffer, uint64_t* buffer_size, void* userdata) -{ - assert(userdata); - auto& table = *reinterpret_cast(userdata); - return table.se_data(seid, buffer, buffer_size, table.user); -} - -} // namespace att_parser -} // namespace rocprofiler - -extern "C" { -rocprofiler_status_t -rocprofiler_att_parse_data(rocprofiler_att_parser_se_data_callback_t user_se_data_callback, - rocprofiler_att_parser_trace_callback_t user_trace_callback, - rocprofiler_att_parser_isa_callback_t user_isa_callback, - void* userdata) -{ - static thread_local bool bInit = []() { - aqlprofile_att_parser_iterate_event_list(rocprofiler::att_parser::iterate_trace_type, - nullptr); - return true; - }(); - (void) bInit; - - rocprofiler::att_parser::userdata_callback_table_t table; - table.trace = user_trace_callback; - table.isa = user_isa_callback; - table.se_data = user_se_data_callback; - table.user = userdata; - - hsa_status_t status = aqlprofile_att_parse_data(rocprofiler::att_parser::se_data_callback, - rocprofiler::att_parser::trace_callback, - rocprofiler::att_parser::isa_callback, - (void*) &table); - - if(status != HSA_STATUS_SUCCESS) return rocprofiler::att_parser::forward_hsa_error(status); - return ROCPROFILER_STATUS_SUCCESS; -} -} diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp index 6bde7087b2..cb864bd190 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp @@ -47,16 +47,16 @@ rocprofiler_configure_dispatch_thread_trace_service( return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); - if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED; + if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; if(ctx->agent_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{}; - pack.context_id = context_id; - pack.dispatch_cb_fn = dispatch_callback; - pack.shader_cb_fn = shader_callback; - pack.callback_userdata = callback_userdata; + pack.context_id = context_id; + pack.dispatch_cb_fn = dispatch_callback; + pack.shader_cb_fn = shader_callback; + pack.callback_userdata.ptr = callback_userdata; if(pack.dispatch_cb_fn == nullptr) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; @@ -102,13 +102,13 @@ rocprofiler_configure_agent_thread_trace_service( size_t num_parameters, rocprofiler_agent_id_t agent, rocprofiler_att_shader_data_callback_t shader_callback, - void* callback_userdata) + rocprofiler_user_data_t userdata) { if(rocprofiler::registration::get_init_status() > -1) return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); - if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED; + if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; if(!ctx->agent_thread_trace) ctx->agent_thread_trace = std::make_unique(); @@ -117,7 +117,7 @@ rocprofiler_configure_agent_thread_trace_service( pack.context_id = context_id; pack.shader_cb_fn = shader_callback; - pack.callback_userdata = callback_userdata; + pack.callback_userdata = userdata; auto id_map = rocprofiler::counters::getPerfCountersIdMap(); for(size_t p = 0; p < num_parameters; p++) diff --git a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index 9d49c7f899..4ad2e3d262 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -147,14 +147,14 @@ TEST(thread_trace, configure_test) ctx, params.data(), params.size(), - [](rocprofiler_queue_id_t, - const rocprofiler_agent_t*, + [](rocprofiler_agent_id_t, + rocprofiler_queue_id_t, rocprofiler_correlation_id_t, rocprofiler_kernel_id_t, rocprofiler_dispatch_id_t, - rocprofiler_user_data_t*, - void*) { return ROCPROFILER_ATT_CONTROL_NONE; }, - [](int64_t, void*, size_t, rocprofiler_user_data_t) {}, + void*, + rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, + [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, nullptr); ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); @@ -198,14 +198,14 @@ TEST(thread_trace, perfcounters_configure_test) ctx, params.data(), params.size(), - [](rocprofiler_queue_id_t, - const rocprofiler_agent_t*, + [](rocprofiler_agent_id_t, + rocprofiler_queue_id_t, rocprofiler_correlation_id_t, rocprofiler_kernel_id_t, rocprofiler_dispatch_id_t, - rocprofiler_user_data_t*, - void*) { return ROCPROFILER_ATT_CONTROL_NONE; }, - [](int64_t, void*, size_t, rocprofiler_user_data_t) {}, + void*, + rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, + [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, nullptr); auto* context = rocprofiler::context::get_mutable_registered_context(ctx); @@ -286,8 +286,8 @@ query_available_agents(rocprofiler_agent_version_t /* version */, params.data(), params.size(), agent->id, - [](int64_t, void*, size_t, rocprofiler_user_data_t) {}, - nullptr); + [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, + rocprofiler_user_data_t{}); } return ROCPROFILER_STATUS_SUCCESS; } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index e2d4704019..292e5e8933 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -58,8 +58,8 @@ add_subdirectory(hsa-memory-allocation) add_subdirectory(scratch-memory-tracing) add_subdirectory(c-tool) add_subdirectory(page-migration) -add_subdirectory(pc_sampling) add_subdirectory(thread-trace) +add_subdirectory(pc_sampling) add_subdirectory(hip-graph-tracing) add_subdirectory(counter-collection) if(ROCPROFILER_BUILD_OPENMP_TESTS) diff --git a/tests/thread-trace/CMakeLists.txt b/tests/thread-trace/CMakeLists.txt index a44e42aa2f..be2aee55ee 100644 --- a/tests/thread-trace/CMakeLists.txt +++ b/tests/thread-trace/CMakeLists.txt @@ -108,7 +108,6 @@ set_source_files_properties(kernel_branch.cpp PROPERTIES COMPILE_FLAGS "-g -O2") set_source_files_properties(kernel_branch.cpp PROPERTIES LANGUAGE HIP) set_source_files_properties(kernel_lds.cpp PROPERTIES COMPILE_FLAGS "-g -O2") set_source_files_properties(kernel_lds.cpp PROPERTIES LANGUAGE HIP) -set_source_files_properties(agent_test.cpp PROPERTIES LANGUAGE HIP) set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP) # Single dispatch test @@ -117,7 +116,7 @@ target_sources( thread-trace-api-single-test PRIVATE main.cpp trace_callbacks.cpp single_dispatch.cpp kernel_branch.cpp kernel_lds.cpp) target_link_libraries(thread-trace-api-single-test - PRIVATE rocprofiler-sdk::rocprofiler-sdk amd_comgr dw) + PRIVATE rocprofiler-sdk::rocprofiler-sdk) add_test(NAME thread-trace-api-single-test COMMAND $) @@ -133,8 +132,8 @@ add_executable(thread-trace-api-multi-test) target_sources( thread-trace-api-multi-test PRIVATE main.cpp trace_callbacks.cpp multi_dispatch.cpp kernel_branch.cpp kernel_lds.cpp) -target_link_libraries(thread-trace-api-multi-test PRIVATE rocprofiler-sdk::rocprofiler-sdk - amd_comgr dw) +target_link_libraries(thread-trace-api-multi-test + PRIVATE rocprofiler-sdk::rocprofiler-sdk) add_test(NAME thread-trace-api-multi-test COMMAND $) @@ -147,8 +146,8 @@ set_tests_properties( # Agent profiling test add_executable(thread-trace-api-agent-test) -target_sources(thread-trace-api-agent-test PRIVATE agent_test.cpp) - +target_sources(thread-trace-api-agent-test PRIVATE main.cpp trace_callbacks.cpp agent.cpp + kernel_branch.cpp kernel_lds.cpp) target_link_libraries(thread-trace-api-agent-test PRIVATE rocprofiler-sdk::rocprofiler-sdk) diff --git a/tests/thread-trace/agent.cpp b/tests/thread-trace/agent.cpp new file mode 100644 index 0000000000..d81c31ecbe --- /dev/null +++ b/tests/thread-trace/agent.cpp @@ -0,0 +1,213 @@ +// MIT License +// +// Copyright (c) 2024 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. +// +// undefine NDEBUG so asserts are implemented +#ifdef NDEBUG +# undef NDEBUG +#endif + +#include "trace_callbacks.hpp" + +#include + +#define C_API_BEGIN \ + try \ + { +#define C_API_END \ + } \ + catch(std::exception & e) \ + { \ + std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl; \ + } \ + catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; } + +namespace ATTTest +{ +namespace Agent +{ +rocprofiler_client_id_t* client_id = nullptr; +rocprofiler_context_id_t agent_ctx = {}; +rocprofiler_context_id_t tracing_ctx = {}; + +void +dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* /* user_data */, + void* /* userdata */) +{ + if(record.kind != ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH) return; + if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) return; + + assert(record.payload); + auto* rdata = static_cast(record.payload); + int dispatch_id = (int) rdata->dispatch_info.dispatch_id; + + auto get_int_var = [](const char* var_name, int def) { + const char* var = getenv(var_name); + if(var) return atoi(var); + return def; + }; + static int begin_dispatch = get_int_var("ROCPROFILER_ATT_BEGIN", 1); + static int end_dispatch = get_int_var("ROCPROFILER_ATT_END", 4); + static std::atomic isprofiling{false}; + + static std::mutex mut; + static std::set captured_ids; + + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) + { + if(dispatch_id == begin_dispatch) + { + ROCPROFILER_CALL(rocprofiler_start_context(agent_ctx), "context start"); + isprofiling.store(true); + } + if(isprofiling && dispatch_id <= end_dispatch) + { + std::unique_lock lk(mut); + captured_ids.insert(dispatch_id); + } + return; + } + + assert(record.phase == ROCPROFILER_CALLBACK_PHASE_NONE); + + if(!isprofiling) return; + + std::unique_lock lk(mut); + captured_ids.erase(dispatch_id); + if(!captured_ids.empty()) return; + + bool _exp = true; + if(!isprofiling.compare_exchange_strong(_exp, false, std::memory_order_relaxed)) return; + + ROCPROFILER_CALL(rocprofiler_stop_context(agent_ctx), "context stop"); +} + +rocprofiler_status_t +query_available_agents(rocprofiler_agent_version_t /* version */, + const void** agents, + size_t num_agents, + void* user_data) +{ + rocprofiler_user_data_t user{}; + user.ptr = user_data; + + for(size_t idx = 0; idx < num_agents; idx++) + { + const auto* agent = static_cast(agents[idx]); + if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; + + std::vector parameters; + parameters.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, 1}); + parameters.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, 0xF}); + parameters.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, 0x6000000}); + parameters.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, 0x11}); + + ROCPROFILER_CALL( + rocprofiler_configure_agent_thread_trace_service(agent_ctx, + parameters.data(), + parameters.size(), + agent->id, + Callbacks::shader_data_callback, + user), + "thread trace service configure"); + } + return ROCPROFILER_STATUS_SUCCESS; +} + +int +tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) +{ + (void) fini_func; + ROCPROFILER_CALL(rocprofiler_create_context(&tracing_ctx), "context creation"); + ROCPROFILER_CALL(rocprofiler_create_context(&agent_ctx), "context creation"); + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(tracing_ctx, + ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, + nullptr, + 0, + Callbacks::tool_codeobj_tracing_callback, + tool_data), + "code object tracing service configure"); + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(tracing_ctx, + ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + nullptr, + 0, + dispatch_tracing_callback, + tool_data), + "dispatch tracing service configure"); + + ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, + &query_available_agents, + sizeof(rocprofiler_agent_t), + tool_data), + "Failed to find GPU agents"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(agent_ctx, &valid_ctx), "validity check"); + assert(valid_ctx != 0); + ROCPROFILER_CALL(rocprofiler_context_is_valid(tracing_ctx, &valid_ctx), "validity check"); + assert(valid_ctx != 0); + + ROCPROFILER_CALL(rocprofiler_start_context(tracing_ctx), "context start"); + + // no errors + return 0; +} + +void +tool_fini(void* tool_data) +{ + Callbacks::finalize_json(tool_data); + delete static_cast(tool_data); +} + +} // namespace Agent +} // namespace ATTTest + +extern "C" rocprofiler_tool_configure_result_t* +rocprofiler_configure(uint32_t /* version */, + const char* /* runtime_version */, + uint32_t priority, + rocprofiler_client_id_t* id) +{ + // only activate if main tool + if(priority > 0) return nullptr; + + // set the client name + id->name = "ATT_test_agent"; + + // store client info + ATTTest::Agent::client_id = id; + + // create configure data + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Agent::tool_init, + &ATTTest::Agent::tool_fini, + new Callbacks::ToolData{"att_agent_test/"}}; + + // return pointer to configure data + return &cfg; +} diff --git a/tests/thread-trace/agent_test.cpp b/tests/thread-trace/agent_test.cpp deleted file mode 100644 index 3f7fb4f7e0..0000000000 --- a/tests/thread-trace/agent_test.cpp +++ /dev/null @@ -1,168 +0,0 @@ -// 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. -// -// undefine NDEBUG so asserts are implemented -#ifdef NDEBUG -# undef NDEBUG -#endif - -#include -#include -#include -#include "common.hpp" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#define HIP_API_CALL(CALL) assert((CALL) == hipSuccess) - -namespace ATTTest -{ -namespace Agent -{ -rocprofiler_context_id_t client_ctx = {0}; -rocprofiler_client_id_t* client_id = nullptr; -std::atomic valid_data{false}; - -void -shader_data_callback(int64_t /* se_id */, - void* se_data, - size_t data_size, - rocprofiler_user_data_t /* userdata */) -{ - if(se_data && data_size) valid_data.store(true); -} - -rocprofiler_status_t -query_available_agents(rocprofiler_agent_version_t /* version */, - const void** agents, - size_t num_agents, - void* /* user_data */) -{ - for(size_t idx = 0; idx < num_agents; idx++) - { - const auto* agent = static_cast(agents[idx]); - if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; - - ROCPROFILER_CALL(rocprofiler_configure_agent_thread_trace_service( - client_ctx, nullptr, 0, agent->id, shader_data_callback, nullptr), - "thread trace service configure"); - - return ROCPROFILER_STATUS_SUCCESS; - } - return ROCPROFILER_STATUS_ERROR; -} - -int -tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) -{ - ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); - - ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, - query_available_agents, - sizeof(rocprofiler_agent_t), - nullptr), - ""); - - int valid = 0; - ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid), "context validity check"); - return (valid == 0) ? -1 : 0; -} - -void -tool_fini(void* /* tool_data */) -{ - assert(valid_data.load()); -} - -} // namespace Agent -} // namespace ATTTest - -extern "C" rocprofiler_tool_configure_result_t* -rocprofiler_configure(uint32_t /* version */, - const char* /* runtime_version */, - uint32_t priority, - rocprofiler_client_id_t* id) -{ - // only activate if main tool - if(priority > 0) return nullptr; - - // set the client name - id->name = "ATT_test_agent_api"; - - // store client info - ATTTest::Agent::client_id = id; - - // create configure data - static auto cfg = - rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), - &ATTTest::Agent::tool_init, - &ATTTest::Agent::tool_fini, - nullptr}; - - // return pointer to configure data - return &cfg; -} - -void -run(int dev) -{ - constexpr size_t size = 0x1000; - float* ptr = nullptr; - - HIP_API_CALL(hipSetDevice(dev)); - - HIP_API_CALL(hipMalloc(&ptr, size * sizeof(float))); - HIP_API_CALL(hipMemset(ptr, 0x55, size * sizeof(float))); - HIP_API_CALL(hipFree(ptr)); -} - -int -main() -{ - int ndev = 0; - HIP_API_CALL(hipGetDeviceCount(&ndev)); - - for(int dev = 0; dev < ndev; dev++) - run(dev); - - ROCPROFILER_CALL(rocprofiler_start_context(ATTTest::Agent::client_ctx), "context start"); - - for(int dev = 0; dev < ndev; dev++) - run(dev); - usleep(100); - - ROCPROFILER_CALL(rocprofiler_stop_context(ATTTest::Agent::client_ctx), "context stop"); - return 0; -} diff --git a/tests/thread-trace/multi_dispatch.cpp b/tests/thread-trace/multi_dispatch.cpp index 857351557e..dc2522ce32 100644 --- a/tests/thread-trace/multi_dispatch.cpp +++ b/tests/thread-trace/multi_dispatch.cpp @@ -25,22 +25,7 @@ # undef NDEBUG #endif -#include -#include -#include -#include "common.hpp" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include +#include "trace_callbacks.hpp" constexpr double WAVE_RATIO_TOLERANCE = 0.05; constexpr size_t NUM_KERNELS = 5; @@ -52,13 +37,13 @@ namespace Multi rocprofiler_client_id_t* client_id = nullptr; rocprofiler_att_control_flags_t -dispatch_callback(rocprofiler_queue_id_t /* queue_id */, - const rocprofiler_agent_t* /* agent */, +dispatch_callback(rocprofiler_agent_id_t /* agent */, + rocprofiler_queue_id_t /* queue_id */, rocprofiler_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t /* kernel_id */, rocprofiler_dispatch_id_t /* dispatch_id */, - rocprofiler_user_data_t* dispatch_userdata, - void* userdata) + void* userdata, + rocprofiler_user_data_t* dispatch_userdata) { static std::atomic count{0}; if(count.fetch_add(1) > NUM_KERNELS) return ROCPROFILER_ATT_CONTROL_NONE; @@ -72,7 +57,6 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */, int tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) { - Callbacks::callbacks_init(); static rocprofiler_context_id_t client_ctx = {0}; ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); @@ -117,25 +101,8 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) void tool_fini(void* tool_data) { - assert(tool_data && "tool_fini callback passed null!"); - ToolData& tool = *reinterpret_cast(tool_data); - - double wave_started = (double) tool.waves_started.load(); - double wave_event_ratio = wave_started / (wave_started + (double) tool.waves_ended.load()); - assert(wave_event_ratio > 0.5 - WAVE_RATIO_TOLERANCE); - assert(wave_event_ratio < 0.5 + WAVE_RATIO_TOLERANCE); - - // Expected: Two kernels in kernel_run.cpp - assert(tool.wave_start_locations.size() >= 2); - - // Expected at least one known code object ID - bool bHasMarkerId = false; - for(auto& pc : tool.wave_start_locations) - bHasMarkerId |= pc.marker_id != 0; - - assert(bHasMarkerId); - - Callbacks::callbacks_fini(); + Callbacks::finalize_json(tool_data); + delete static_cast(tool_data); } } // namespace Multi @@ -156,14 +123,12 @@ rocprofiler_configure(uint32_t /* version */, // store client info ATTTest::Multi::client_id = id; - auto* data = new ATTTest::ToolData{}; - // create configure data - static auto cfg = - rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), - &ATTTest::Multi::tool_init, - &ATTTest::Multi::tool_fini, - reinterpret_cast(data)}; + static auto cfg = rocprofiler_tool_configure_result_t{ + sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Multi::tool_init, + &ATTTest::Multi::tool_fini, + reinterpret_cast(new Callbacks::ToolData{"att_multi_test/"})}; // return pointer to configure data return &cfg; diff --git a/tests/thread-trace/single_dispatch.cpp b/tests/thread-trace/single_dispatch.cpp index abec806751..4f68855c9b 100644 --- a/tests/thread-trace/single_dispatch.cpp +++ b/tests/thread-trace/single_dispatch.cpp @@ -25,23 +25,7 @@ # undef NDEBUG #endif -#include -#include -#include -#include "common.hpp" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include +#include "trace_callbacks.hpp" constexpr double WAVE_RATIO_TOLERANCE = 0.05; @@ -52,17 +36,17 @@ namespace Single rocprofiler_client_id_t* client_id = nullptr; rocprofiler_att_control_flags_t -dispatch_callback(rocprofiler_queue_id_t /* queue_id */, - const rocprofiler_agent_t* /* agent */, - rocprofiler_correlation_id_t /* correlation_id */, +dispatch_callback(rocprofiler_agent_id_t /* agent */, + rocprofiler_queue_id_t /* queue_id */, + rocprofiler_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t /* dispatch_id */, - rocprofiler_user_data_t* dispatch_userdata, - void* userdata) + void* userdata, + rocprofiler_user_data_t* dispatch_userdata) { C_API_BEGIN assert(userdata && "Dispatch callback passed null!"); - ToolData& tool = *reinterpret_cast(userdata); + auto& tool = *reinterpret_cast(userdata); dispatch_userdata->ptr = userdata; static std::string_view desired_func_name = "branching_kernel"; @@ -86,7 +70,6 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */, int tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) { - Callbacks::callbacks_init(); static rocprofiler_context_id_t client_ctx = {0}; ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); @@ -125,34 +108,8 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) void tool_fini(void* tool_data) { - assert(tool_data && "tool_fini callback passed null!"); - ToolData& tool = *reinterpret_cast(tool_data); - - std::unique_lock isa_lk(tool.isa_map_mut); - - // Find largest instruction - size_t max_inst_size = 0; - for(auto& [addr, lines] : tool.isa_map) - max_inst_size = std::max(max_inst_size, lines->inst.size()); - - size_t total_hit = 0; - size_t total_cycles = 0; - - for(auto& [addr, line] : tool.isa_map) - { - total_hit += line->hitcount.load(std::memory_order_relaxed); - total_cycles += line->latency.load(std::memory_order_relaxed); - } - - assert(total_cycles > 0); - assert(total_hit > 0); - - double wave_started = (double) tool.waves_started.load(); - double wave_event_ratio = wave_started / (wave_started + (double) tool.waves_ended.load()); - assert(wave_event_ratio > 0.5 - WAVE_RATIO_TOLERANCE); - assert(wave_event_ratio < 0.5 + WAVE_RATIO_TOLERANCE); - - Callbacks::callbacks_fini(); + Callbacks::finalize_json(tool_data); + delete static_cast(tool_data); } } // namespace Single @@ -173,14 +130,12 @@ rocprofiler_configure(uint32_t /* version */, // store client info ATTTest::Single::client_id = id; - auto* data = new ATTTest::ToolData{}; - // create configure data - static auto cfg = - rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), - &ATTTest::Single::tool_init, - &ATTTest::Single::tool_fini, - reinterpret_cast(data)}; + static auto cfg = rocprofiler_tool_configure_result_t{ + sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Single::tool_init, + &ATTTest::Single::tool_fini, + reinterpret_cast(new Callbacks::ToolData{"att_single_test/"})}; // return pointer to configure data return &cfg; diff --git a/tests/thread-trace/trace_callbacks.cpp b/tests/thread-trace/trace_callbacks.cpp index 71812e8e2e..ef807920b0 100644 --- a/tests/thread-trace/trace_callbacks.cpp +++ b/tests/thread-trace/trace_callbacks.cpp @@ -25,217 +25,148 @@ # undef NDEBUG #endif -#include -#include -#include -#include +#include "trace_callbacks.hpp" #include -#include "common.hpp" -#include +#ifdef ENABLE_ATT_FILES +# include +#endif + +#include #include -#include -#include -#include -#include -#include -#include -#include -#include -#include +#include -namespace ATTTest -{ namespace Callbacks { using code_obj_load_data_t = rocprofiler_callback_tracing_code_object_load_data_t; using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; -using CodeobjAddressTranslate = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; -using Instruction = rocprofiler::sdk::codeobj::disassembly::Instruction; - -CodeobjAddressTranslate* codeobjTranslate = nullptr; - -struct trace_data_t -{ - int64_t id; - uint8_t* data; - uint64_t size; - ToolData* tool; -}; void tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* /* user_data */, - void* callback_data) + void* userdata) { C_API_BEGIN if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return; if(record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) return; - assert(callback_data && "Shader callback passed null!"); - ToolData& tool = *reinterpret_cast(callback_data); + assert(userdata && "Dispatch callback passed null!"); + auto& tool = *reinterpret_cast(userdata); if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) { - std::unique_lock lg(tool.isa_map_mut); - auto* data = static_cast(record.payload); + auto* data = static_cast(record.payload); tool.kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name); } if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return; auto* data = static_cast(record.payload); - if(!data || !data->uri) return; - std::unique_lock lg(tool.isa_map_mut); + static std::atomic filecnt{0}; + std::string name = "codeobj_" + std::to_string(filecnt.fetch_add(1)) + ".out"; +#ifdef ENABLE_ATT_FILES if(std::string_view(data->uri).find("file:///") == 0) { - codeobjTranslate->addDecoder( - data->uri, data->code_object_id, data->load_delta, data->load_size); + rocprofiler::sdk::codeobj::disassembly::CodeObjectBinary binary(data->uri); + + std::ofstream file(tool.out_dir + name, std::ios::binary); + assert(file.is_open() && "Could not open codeobj file for writing"); + file.write((char*) binary.buffer.data(), binary.buffer.size()); } else { - codeobjTranslate->addDecoder(reinterpret_cast(data->memory_base), - data->memory_size, - data->code_object_id, - data->load_delta, - data->load_size); + std::ofstream file(tool.out_dir + name, std::ios::binary); + file.write((char*) data->memory_base, data->memory_size); } +#endif + + auto _lk = std::unique_lock{tool.mut}; + tool.codeobjs.push_back( + {data->load_delta, data->load_size, data->code_object_id, name, data->uri}); C_API_END } void -get_trace_data(rocprofiler_att_parser_data_type_t type, void* att_data, void* userdata) -{ - C_API_BEGIN - assert(userdata && "ISA callback passed null!"); - trace_data_t& trace_data = *reinterpret_cast(userdata); - assert(trace_data.tool && "ISA callback passed null!"); - ToolData& tool = *reinterpret_cast(trace_data.tool); - - std::unique_lock lk(tool.isa_map_mut); - - if(type == ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY) - { - const auto& ev = *reinterpret_cast(att_data); - tool.wave_start_locations.insert({ev.offset, ev.marker_id}); - if(ev.enabled) - tool.waves_started.fetch_add(1); - else - tool.waves_ended.fetch_add(1); - } - - if(type != ROCPROFILER_ATT_PARSER_DATA_TYPE_ISA) return; - - auto& event = *reinterpret_cast(att_data); - - pcInfo pc{event.offset, event.marker_id}; - auto it = tool.isa_map.find(pc); - if(it == tool.isa_map.end()) - { - auto ptr = std::make_unique(); - try - { - auto unique_inst = codeobjTranslate->get(pc.marker_id, pc.addr); - if(unique_inst == nullptr) return; - ptr->inst = unique_inst->inst; - } catch(...) - { - return; - } - it = tool.isa_map.emplace(pc, std::move(ptr)).first; - } - - it->second->hitcount.fetch_add(event.hitcount, std::memory_order_relaxed); - it->second->latency.fetch_add(event.latency, std::memory_order_relaxed); - C_API_END -} - -uint64_t -copy_trace_data(int* seid, uint8_t** buffer, uint64_t* buffer_size, void* userdata) -{ - trace_data_t& data = *reinterpret_cast(userdata); - *seid = data.id; - *buffer_size = data.size; - *buffer = data.data; - data.size = 0; - return *buffer_size; -} - -rocprofiler_status_t -isa_callback(char* isa_instruction, - uint64_t* isa_memory_size, - uint64_t* isa_size, - uint64_t marker_id, - uint64_t offset, - void* userdata) -{ - C_API_BEGIN - assert(userdata && "ISA callback passed null!"); - trace_data_t& trace_data = *reinterpret_cast(userdata); - assert(trace_data.tool && "ISA callback passed null!"); - ToolData& tool = *reinterpret_cast(trace_data.tool); - - std::unique_ptr instruction; - - try - { - std::unique_lock unique_lock(tool.isa_map_mut); - instruction = codeobjTranslate->get(marker_id, offset); - } catch(...) - { - return ROCPROFILER_STATUS_ERROR; - } - - if(!instruction.get()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - - { - size_t tmp_isa_size = *isa_size; - *isa_size = instruction->inst.size(); - - if(*isa_size > tmp_isa_size) return ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES; - } - - memcpy(isa_instruction, instruction->inst.data(), *isa_size); - *isa_memory_size = instruction->size; - - auto ptr = std::make_unique(); - ptr->inst = instruction->inst; - tool.isa_map.emplace(pcInfo{offset, marker_id}, std::move(ptr)); - return ROCPROFILER_STATUS_SUCCESS; - C_API_END - return ROCPROFILER_STATUS_ERROR; -} - -void -shader_data_callback(int64_t se_id, +shader_data_callback(rocprofiler_agent_id_t agent, + int64_t se_id, void* se_data, size_t data_size, rocprofiler_user_data_t userdata) { C_API_BEGIN - assert(userdata.ptr && "Shader callback passed null!"); - ToolData& tool = *reinterpret_cast(userdata.ptr); - trace_data_t data{.id = se_id, .data = (uint8_t*) se_data, .size = data_size, .tool = &tool}; - auto status = rocprofiler_att_parse_data(copy_trace_data, get_trace_data, isa_callback, &data); - if(status != ROCPROFILER_STATUS_SUCCESS) - std::cerr << "shader_data_callback failed with status " << status << std::endl; + assert(userdata.ptr && "Dispatch callback passed null!"); + auto& tool = *reinterpret_cast(userdata.ptr); + + std::string name = "agent_" + std::to_string(agent.handle) + "_shader_engine_" + + std::to_string(se_id) + "_" + std::to_string(agent.handle) + ".att"; + +#ifdef ENABLE_ATT_FILES + { + std::ofstream file(tool.out_dir + name, std::ios::binary); + assert(file.is_open() && "Could not open ATT file for writing"); + file.write((char*) se_data, data_size); + } +#endif + + assert(se_data); + assert(data_size); + + auto _lk = std::unique_lock{tool.mut}; + tool.att_files.push_back(name); + C_API_END } void -callbacks_init() +finalize_json(void* userdata) { - codeobjTranslate = new CodeobjAddressTranslate(); -} -void -callbacks_fini() -{ - delete codeobjTranslate; + assert(userdata && "Dispatch callback passed null!"); + + auto& tool = *reinterpret_cast(userdata); + auto _lk = std::unique_lock{tool.mut}; + assert(!tool.att_files.empty()); + +#ifdef ENABLE_ATT_FILES + nlohmann::json att_json; + for(auto& file : tool.att_files) + att_json.push_back(file); + + nlohmann::json codeobj_json; + nlohmann::json snapshot_json; + for(auto& file : tool.codeobjs) + { + nlohmann::json codeobj; + codeobj["code_object_id"] = file.id; + codeobj["load_delta"] = file.addr; + codeobj["load_size"] = file.size; + codeobj["uri"] = file.uri; + codeobj["filename"] = file.filename; + codeobj_json.push_back(codeobj); + + nlohmann::json pair_json; + pair_json["key"] = file.id; + pair_json["value"] = file.filename; + snapshot_json.push_back(pair_json); + } + nlohmann::json tool_json; + tool_json["strings"]["att_files"] = att_json; + tool_json["code_objects"] = codeobj_json; + tool_json["strings"]["code_object_snapshot_files"] = snapshot_json; + + nlohmann::json array; + array.push_back(tool_json); + + nlohmann::json sdk_json; + sdk_json["rocprofiler-sdk-tool"] = array; + + std::ofstream json_file(tool.out_dir + (std::to_string(getpid()) + "_results.json")); + assert(json_file.is_open() && "Could not open json file for writing!"); + json_file << sdk_json; +#endif } } // namespace Callbacks -} // namespace ATTTest diff --git a/tests/thread-trace/common.hpp b/tests/thread-trace/trace_callbacks.hpp similarity index 72% rename from tests/thread-trace/common.hpp rename to tests/thread-trace/trace_callbacks.hpp index 66d21650ea..1ae6b5332c 100644 --- a/tests/thread-trace/common.hpp +++ b/tests/thread-trace/trace_callbacks.hpp @@ -22,39 +22,27 @@ #pragma once +#include +#include +#include #include +#include #include #include #include #include #include -#include +#include #include -#include +#include #include #include +#include #define ROCPROFILER_VAR_NAME_COMBINE(X, Y) X##Y #define ROCPROFILER_VARIABLE(X, Y) ROCPROFILER_VAR_NAME_COMBINE(X, Y) -#define ROCPROFILER_CALL(result, msg) \ - { \ - rocprofiler_status_t ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) = result; \ - if(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) != ROCPROFILER_STATUS_SUCCESS) \ - { \ - std::string status_msg = \ - rocprofiler_get_status_string(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__)); \ - std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ - << " failed with error code " << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \ - << ": " << status_msg << std::endl; \ - std::stringstream errmsg{}; \ - errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ - << status_msg << ")"; \ - throw std::runtime_error(errmsg.str()); \ - } \ - } - #define C_API_BEGIN \ try \ { @@ -66,60 +54,59 @@ } \ catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; } -namespace ATTTest -{ -struct TrackedIsa -{ - std::atomic hitcount{0}; - std::atomic latency{0}; - std::string inst{}; -}; - -struct pcInfo -{ - size_t addr; - size_t marker_id; - - bool operator==(const pcInfo& other) const - { - return addr == other.addr && marker_id == other.marker_id; - } - bool operator<(const pcInfo& other) const - { - if(marker_id == other.marker_id) return addr < other.addr; - return marker_id < other.marker_id; +#define ROCPROFILER_CALL(result, msg) \ + { \ + rocprofiler_status_t CHECKSTATUS = result; \ + if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS); \ + std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ + << " failed with error code " << CHECKSTATUS << ": " << status_msg \ + << std::endl; \ + std::stringstream errmsg{}; \ + errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ + << status_msg << ")"; \ + throw std::runtime_error(errmsg.str()); \ + } \ } + +namespace Callbacks +{ +struct CodeobjInfo +{ + int64_t addr = 0; + size_t size = 0; + size_t id = 0; + std::string filename{}; + std::string uri{}; }; struct ToolData { - std::unordered_map kernel_id_to_kernel_name = {}; - std::map> isa_map; + ToolData(const char* out) + : out_dir(out){}; - std::atomic waves_started = 0; - std::atomic waves_ended = 0; - std::mutex isa_map_mut; - std::set wave_start_locations{}; + std::string out_dir{}; + std::mutex mut{}; + std::vector codeobjs{}; + std::vector att_files{}; + + std::unordered_map kernel_id_to_kernel_name = {}; }; -namespace Callbacks -{ void tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t*, void* callback_data); void -shader_data_callback(int64_t se_id, +shader_data_callback(rocprofiler_agent_id_t agent, + int64_t se_id, void* se_data, size_t data_size, rocprofiler_user_data_t userdata); void -callbacks_init(); - -void -callbacks_fini(); +finalize_json(void* userdata); } // namespace Callbacks -} // namespace ATTTest