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

SWDEV-492607: Adding ATT wrapper (#40)

Browse Source 
* Adding att parser wrapper

* Adding ATT tests as optional

* Adding decoder API for query capability

* Removed samples

* Formatting

* adding new line

* Removed perfetto and moved to static library

* using default search for lib

* Updated to SDK

* Namespace changes

* Added tests

* Small refactor

* Updated API to receive agent_id

* Fixing tests

* Tidy fixes

* Not write to file

* Switch to filesystem.hpp

* Compilation fixes

* Formatting

* Tidy fix

* Removed likely

* Adding tests

* Added gfx9 test

* Adding gfx12 tests

* Formatting

* Enable tidy

* Fix tests

* Fix deadlock on agent test

* Workaround ASAN

* Moving query outside class.

* Fix standalone tool

* Addressing comments

* Formatting

* Change query name

* Fixed some tests. Updated PR comments.

* Formatting

* Improved coverage

* Formatting

* Fix for comments

* Formatting

* Adding some description. Fix error type.

---------

Co-authored-by: Giovanni Baraldi <gbaraldi@amd.com>

[ROCm/rocprofiler-sdk commit: 2c8e88a76b]
This commit is contained in:
Baraldi, Giovanni
2024-12-19 03:53:32 +01:00
committed by GitHub
parent b2aa29eb0a
commit f8442415f8
59 changed files with 4250 additions and 1086 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/.gitignore
+1 -1
View File
@@ -35,7 +35,7 @@
/compile_commands.json
/build*
/.cache
/.vscode
*.vscode
# Github Workflows
/.github
projects/rocprofiler-sdk/.gitmodules
+3
View File
@@ -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
projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake
+1
View File
@@ -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)
projects/rocprofiler-sdk/external/CMakeLists.txt
Vendored
+14
View File
@@ -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 $<BUILD_INTERFACE:json::json>)
target_include_directories(
rocprofiler-sdk-json SYSTEM
INTERFACE $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/external/json/include>)
projects/rocprofiler-sdk/external/json
Vendored Submodule
+1
Submodule projects/rocprofiler-sdk/external/json added at e41905fcb0
projects/rocprofiler-sdk/samples/CMakeLists.txt
-1
View File
@@ -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)
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h
+8 -2
View File
@@ -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;
/** @} */
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h
+2 -94
View File
@@ -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
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h
+19 -4
View File
@@ -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(
projects/rocprofiler-sdk/source/lib/CMakeLists.txt
+1
View File
@@ -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)
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/CMakeLists.txt
Executable
+47
View File
@@ -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()
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/att_decoder.h
+160
View File
@@ -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 <cstddef>
# include <cstdint>
extern "C" {
#else
# include <stddef.h>
# include <stdint.h>
#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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp
+197
View File
@@ -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 <dlfcn.h>
#include <algorithm>
#include <fstream>
#include <iostream>
#include <mutex>
#include <nlohmann/json.hpp>
#include <vector>
namespace rocprofiler
{
namespace att_wrapper
{
auto
get_lib_names()
{
std::vector<std::pair<tool_att_capability_t, const char*>> 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> _dl)
: dir(std::move(_dir))
, dl(std::move(_dl))
{
rocprofiler::common::filesystem::create_directories(dir);
table = std::make_shared<AddressTable>();
codefile = std::make_shared<CodeFile>(dir, table);
filenames = std::make_shared<FilenameMgr>(dir);
for(size_t i = 0; i < ATT_WAVE_STATE_LAST; i++)
wstates.at(i) = std::make_shared<WstatesFile>(i, dir);
}
ATTFileMgr::~ATTFileMgr() { OccupancyFile::OccupancyFile(dir, table, occupancy); }
void
ATTFileMgr::parseShader(int se_id, const std::vector<char>& 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<tool_att_capability_t>
query_att_decode_capability()
{
std::vector<tool_att_capability_t> 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<DL>(libname);
return;
}
}
};
void
ATTDecoder::parse(const Fspath& input_dir,
const Fspath& output_dir,
const std::vector<std::string>& att_files,
const std::vector<CodeobjLoadInfo>& 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<char> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp
+110
View File
@@ -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 <algorithm>
#include <fstream>
#include <iostream>
#include <map>
#include <vector>
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<tool_att_capability_t>
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<std::string>& att_files,
const std::vector<CodeobjLoadInfo>& codeobj_files,
const std::string& output_formats);
bool valid() const;
protected:
std::shared_ptr<class DL> dl{nullptr};
};
class ATTFileMgr
{
using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
public:
ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl);
~ATTFileMgr();
void parseShader(int se_id, const std::vector<char>& data);
Fspath dir{};
std::shared_ptr<class DL> dl{nullptr};
std::shared_ptr<class CodeFile> codefile{nullptr};
std::shared_ptr<class FilenameMgr> filenames{nullptr};
std::shared_ptr<AddressTable> table{nullptr};
std::map<size_t, std::vector<att_occupancy_info_v2_t>> occupancy;
std::array<std::shared_ptr<class WstatesFile>, ATT_WAVE_STATE_LAST> wstates;
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/code.cpp
+123
View File
@@ -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 <nlohmann/json.hpp>
#include "outputfile.hpp"
#include <fstream>
#include <iostream>
#include <sstream>
namespace rocprofiler
{
namespace att_wrapper
{
#define ATT_CSV_NAME "att_output.csv"
CodeFile::CodeFile(const Fspath& _dir, std::shared_ptr<AddressTable> _table)
: dir(_dir)
, filename(_dir / "code.json")
, table(_table)
{}
CodeFile::~CodeFile()
{
std::vector<std::pair<pcinfo_t, std::unique_ptr<CodeLine>>> 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<pcinfo_t, std::unique_ptr<CodeLine>>& a,
const std::pair<pcinfo_t, std::unique_ptr<CodeLine>>& 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<pcinfo_t, std::unique_ptr<CodeLine>>& a,
const std::pair<pcinfo_t, std::unique_ptr<CodeLine>>& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/code.hpp
+68
View File
@@ -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 <atomic>
#include <map>
#include <unordered_map>
#include <vector>
#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<size_t> hitcount{0};
std::atomic<size_t> latency{0};
std::shared_ptr<Instruction> code_line{nullptr};
};
class CodeFile
{
using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
public:
CodeFile(const Fspath& dir, std::shared_ptr<AddressTable> table);
~CodeFile();
Fspath dir{};
Fspath filename{};
std::unordered_map<pcinfo_t, int> line_numbers{};
std::map<pcinfo_t, std::unique_ptr<CodeLine>> isa_map{};
std::map<pcinfo_t, KernelName> kernel_names{};
std::shared_ptr<AddressTable> table;
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/counters.cpp
+78
View File
@@ -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 <algorithm>
#include <cassert>
#include <fstream>
#include <nlohmann/json.hpp>
#include "util.hpp"
namespace rocprofiler
{
namespace att_wrapper
{
CountersFile::CountersFile(const Fspath& _dir, const std::vector<std::string>& _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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/counters.hpp
+50
View File
@@ -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 <map>
#include <vector>
#include "util.hpp"
namespace rocprofiler
{
namespace att_wrapper
{
class CountersFile
{
public:
CountersFile(const Fspath& dir, const std::vector<std::string>& names);
~CountersFile();
void AddShaderEngine(int se, const att_perfevent_t* events, size_t num_events);
private:
const Fspath dir;
std::vector<std::string> names{};
std::vector<std::string> shaders{};
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/dl.cpp
+55
View File
@@ -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 <dlfcn.h>
#include <atomic>
#include <cassert>
#include <cstdlib>
#include <mutex>
#include <stdexcept>
#include <string>
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<ParseFn*>(dlsym(handle, "rocprofiler_att_decoder_parse_data"));
att_info_fn =
reinterpret_cast<InfoFn*>(dlsym(handle, "rocprofiler_att_decoder_get_info_string"));
att_status_fn =
reinterpret_cast<StatusFn*>(dlsym(handle, "rocprofiler_att_decoder_get_status_string"));
};
DL::~DL()
{
if(handle) dlclose(handle);
}
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/dl.hpp
+48
View File
@@ -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 <memory>
#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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/filenames.cpp
+63
View File
@@ -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 <fstream>
#include <iostream>
#include <nlohmann/json.hpp>
#include <sstream>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/filenames.hpp
+79
View File
@@ -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 <map>
#include <vector>
#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<Coord, WaveName> streams{};
int gfxip = 9;
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/occupancy.cpp
+109
View File
@@ -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 <nlohmann/json.hpp>
#include "outputfile.hpp"
#include <atomic>
#include <fstream>
#include <iostream>
#include <sstream>
#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<pcinfo_t, int> kernel_ids{{pcinfo_t{0, 0}, 0}};
std::atomic<int> 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<AddressTable> table,
const std::map<size_t, std::vector<att_occupancy_info_v2_t>>& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/occupancy.hpp
+46
View File
@@ -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 <map>
#include <vector>
#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<AddressTable> table,
const std::map<size_t, std::vector<att_occupancy_info_v2_t>>& occ);
}; // namespace OccupancyFile
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/outputfile.hpp
+66
View File
@@ -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 <fstream>
#include <iostream>
#include <nlohmann/json.hpp>
#include <sstream>
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 <typename T>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/profile_interface.cpp
+253
View File
@@ -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 <cxxabi.h>
#include <cstring>
#include <fstream>
#include <shared_mutex>
#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<trace_data_t*>(userdata);
CHECK_NOTNULL(trace_data.tool);
ToolData& tool = *reinterpret_cast<ToolData*>(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<size_t>(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<const att_occupancy_info_v2_t*>(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<att_wave_data_t*>(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<trace_data_t*>(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<trace_data_t*>(userdata);
CHECK_NOTNULL(trace_data.tool);
ToolData& tool = *reinterpret_cast<ToolData*>(trace_data.tool);
std::shared_ptr<Instruction> 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<char>& _data, WaveConfig& _config, std::shared_ptr<DL> _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<CodeLine>()).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<CodeLine>()).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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/profile_interface.hpp
+77
View File
@@ -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 <cxxabi.h>
#include <array>
#include <atomic>
#include <cstdint>
#include <iostream>
#include <map>
#include <string>
#include <string_view>
#include <unordered_map>
#include <vector>
#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<char>& data, WaveConfig& config, std::shared_ptr<class DL> _dl);
~ToolData();
CodeLine& get(pcinfo_t pc);
std::shared_ptr<CodeFile> cfile{};
WaveConfig& config;
std::shared_ptr<DL> dl{};
std::vector<char> shader_data{};
size_t num_waves = 0;
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt
+36
View File
@@ -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()
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp
+138
View File
@@ -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 <gtest/gtest.h>
#include <iostream>
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<CodeLine>();
code->code_line = std::make_shared<CodeLine::Instruction>();
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<char> 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<CodeobjLoadInfo> 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<std::string> 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<CodeLine>();
line->line_number = i;
line->code_line = std::make_shared<CodeLine::Instruction>();
line->code_line->inst = "v_add";
file.isa_map[addr] = std::move(line);
}
}
}; // namespace att_wrapper
}; // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/tests/dummy_decoder.cpp
+142
View File
@@ -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 <array>
#include <vector>
__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<rocprofiler_att_decoder_info_t> infos{};
for(size_t i = 1; i < ROCPROFILER_ATT_DECODER_INFO_LAST; i++)
infos.emplace_back(static_cast<rocprofiler_att_decoder_info_t>(i));
trace_callback(ROCPROFILER_ATT_DECODER_TYPE_INFO, 0, infos.data(), infos.size(), userdata);
}
{
uint64_t memory_size = 0, size = 16;
std::array<char, 16> 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<att_occupancy_info_v2_t> 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<att_wave_data_t> 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<att_wave_state_t> 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<att_wave_instruction_t> 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<const char*>{"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<const char*>{"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);
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp
+167
View File
@@ -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 <iostream>
#include <mutex>
#include <nlohmann/json.hpp>
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<int, std::vector<std::string>> 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<Fspath> att_files{};
std::vector<rocprofiler::att_wrapper::CodeobjLoadInfo> codeobj_files{};
std::map<size_t, std::string> 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;
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/util.hpp
+88
View File
@@ -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 <rocprofiler-sdk/cxx/codeobj/code_printing.hpp>
#include "att_decoder.h"
#include "lib/common/logging.hpp"
#include <memory>
#include <string>
#include <string_view>
template <>
struct std::hash<pcinfo_t>
{
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/CMakeLists.txt
+6
View File
@@ -0,0 +1,6 @@
#
# ATT waitcnt ISA
#
if(ROCPROFILER_BUILD_TESTS)
add_subdirectory(tests)
endif()
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/analysis.cpp
+85
View File
@@ -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 <map>
namespace rocprofiler
{
namespace att_wrapper
{
std::map<size_t, std::unique_ptr<WaitcntList>> 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<int>
MemoryCounter::join_and_reset(int64_t offset, std::vector<int>& flats)
{
std::vector<int> ret = std::move(flats);
flats.clear();
offset = std::max(std::min<int64_t>(offset, list.size()), 0l);
ret.insert(ret.end(), list.begin(), list.begin() + offset);
list.erase(list.begin(), list.begin() + offset);
return ret;
}
std::optional<std::vector<int>>
MemoryCounter::handle_mem_op(const std::string& inst, std::vector<int>& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/analysis.hpp
+107
View File
@@ -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 <atomic>
#include <map>
#include <memory>
#include <unordered_map>
#include <vector>
namespace rocprofiler
{
namespace att_wrapper
{
struct LineWaitcnt
{
int line_number{0};
std::vector<int> dependencies{};
};
struct WaitcntList
{
using isa_map_t = std::map<pcinfo_t, std::unique_ptr<CodeLine>>;
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<WaitcntList>();
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<LineWaitcnt> gfx9_construct(const wave_t& wave, isa_map_t& isa_map);
static std::vector<LineWaitcnt> gfx10_construct(const wave_t& wave, isa_map_t& isa_map);
static std::vector<LineWaitcnt> gfx12_construct(const wave_t& wave, isa_map_t& isa_map);
std::vector<LineWaitcnt> mem_unroll{};
private:
static std::map<size_t, std::unique_ptr<WaitcntList>> _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<int> join_and_reset(int64_t offset, std::vector<int>& flats);
std::optional<std::vector<int>> handle_mem_op(const std::string& inst,
std::vector<int>& flat_list);
const std::string name;
Ordering order = Ordering::MEMORY_SEQUENTIAL;
std::vector<int> list{};
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/gfx10.cpp
+216
View File
@@ -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 <map>
#include <mutex>
#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<LineWaitcnt>
WaitcntList::gfx10_construct(const wave_t& wave, isa_map_t& isa_map)
{
std::vector<LineWaitcnt> mem_unroll;
MemoryCounter lgkm("lgkmcnt");
MemoryCounter vmcnt("vmcnt");
MemoryCounter vscnt("vscnt");
std::vector<int> flat_load{};
std::vector<int> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/gfx12.cpp
+276
View File
@@ -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 <map>
#include <mutex>
#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<LineWaitcnt>
WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map)
{
static std::vector<int> empty_list{};
std::vector<LineWaitcnt> 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<int> flat_load{};
std::vector<int> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/gfx9.cpp
+150
View File
@@ -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 <map>
#include <mutex>
#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<LineWaitcnt>
WaitcntList::gfx9_construct(const wave_t& wave, isa_map_t& isa_map)
{
std::vector<LineWaitcnt> mem_unroll;
MemoryCounter lgkm("lgkmcnt");
MemoryCounter vmem("vmcnt");
std::vector<int> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/tests/CMakeLists.txt
+24
View File
@@ -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")
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/waitcnt/tests/att_decoder_waitcnt_test.cpp
+375
View File
@@ -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 <gtest/gtest.h>
#include <iostream>
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<CodeLine>();
code->code_line = std::make_shared<CodeLine::Instruction>();
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<att_wave_instruction_t> 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<int, std::set<int>> 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<int>& 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<CodeLine>();
code->code_line = std::make_shared<CodeLine::Instruction>();
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<att_wave_instruction_t> 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<int, std::set<int>> 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<int>& 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<CodeLine>();
code->code_line = std::make_shared<CodeLine::Instruction>();
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<att_wave_instruction_t> 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<int, std::set<int>> 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<int>& 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<att_wave_instruction_t> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/wave.cpp
+133
View File
@@ -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 <nlohmann/json.hpp>
#include "outputfile.hpp"
#include <fstream>
#include <iostream>
#include <sstream>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/wave.hpp
+80
View File
@@ -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 <atomic>
#include <map>
#include <unordered_map>
#include <vector>
namespace rocprofiler
{
namespace att_wrapper
{
constexpr size_t SIMD_NUM = 4;
constexpr size_t SIMD_SIZE = 32;
class WaveConfig
{
using WavestateArray = std::array<std::shared_ptr<WstatesFile>, ATT_WAVE_STATE_LAST>;
using SIMD = std::array<std::atomic<int>, SIMD_SIZE>;
public:
WaveConfig(int se_id,
std::shared_ptr<FilenameMgr>& _mgr,
std::shared_ptr<CodeFile>& _code,
WavestateArray& _wstates)
: shader_engine(se_id)
, wstates(_wstates)
, code(_code)
, filemgr(_mgr)
{}
const int shader_engine;
WavestateArray wstates;
std::array<SIMD, SIMD_NUM> id_count{};
std::shared_ptr<CodeFile> code;
std::shared_ptr<FilenameMgr> filemgr;
std::map<pcinfo_t, KernelName> kernel_names{};
std::vector<att_occupancy_info_v2_t> occupancy{};
};
class WaveFile
{
public:
WaveFile(WaveConfig& config, const att_wave_data_t& wave);
Fspath filename{};
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/wstates.cpp
+76
View File
@@ -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 <algorithm>
#include <cassert>
#include <fstream>
#include <nlohmann/json.hpp>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-att/wstates.hpp
+57
View File
@@ -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 <map>
#include <nlohmann/json.hpp>
#include <vector>
#include "util.hpp"
namespace rocprofiler
{
namespace att_wrapper
{
class WstatesFile
{
using event_t = std::pair<int64_t, int64_t>;
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<event_t> events{};
static constexpr size_t NUM_WSTATES = 5;
};
} // namespace att_wrapper
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h
-238
View File
@@ -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;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/CMakeLists.txt
+1 -2
View File
@@ -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}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp
+14 -13
View File
@@ -61,8 +61,9 @@ constexpr uint64_t MAX_BUFFER_SIZE = std::numeric_limits<int32_t>::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<std::optional<int64_t>> client;
@@ -238,7 +239,7 @@ thread_trace_callback(uint32_t shader, void* buffer, uint64_t size, void* callba
{
auto& cb_data = *static_cast<cbdata_t*>(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<ThreadTracerQueue*, aqlprofile_handle_t, std::unique_ptr<Signal>>;
std::unique_lock<std::mutex> lk(agent_mut);
if(tracers.empty())
@@ -522,10 +525,9 @@ AgentThreadTracer::stop_context()
return;
}
std::vector<std::tuple<ThreadTracerQueue*, aqlprofile_handle_t, std::unique_ptr<Signal>>>
wait_list{};
std::vector<wait_t> 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);
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp
+1 -1
View File
@@ -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;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_parser.cpp
-213
View File
@@ -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 <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <atomic>
#include <string_view>
#include <unordered_map>
#include <vector>
#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<int, hsa_status_t> 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<int, rocprofiler_status_t> 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<pcinfo_t> 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_callback_table_t*>(userdata);
if(trace_type_id == KERNEL_ADDR_ID)
{
table.kernel_id_map.resize(trace_size);
const auto* events = reinterpret_cast<const pcinfo_t*>(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<const att_occupancy_info_t*>(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<const att_trace_event_t*>(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<const userdata_callback_table_t*>(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_callback_table_t*>(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;
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp
+8 -8
View File
@@ -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<AgentThreadTracer>();
@@ -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++)
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp
+12 -12
View File
@@ -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;
}
projects/rocprofiler-sdk/tests/CMakeLists.txt
+1 -1
View File
@@ -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)
projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt
+5 -6
View File
@@ -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 $<TARGET_FILE:thread-trace-api-single-test>)
@@ -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 $<TARGET_FILE:thread-trace-api-multi-test>)
@@ -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)
projects/rocprofiler-sdk/tests/thread-trace/agent.cpp
+213
View File
@@ -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 <set>
#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<rocprofiler_callback_tracing_kernel_dispatch_data_t*>(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<bool> isprofiling{false};
static std::mutex mut;
static std::set<int> 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<std::mutex> lk(mut);
captured_ids.insert(dispatch_id);
}
return;
}
assert(record.phase == ROCPROFILER_CALLBACK_PHASE_NONE);
if(!isprofiling) return;
std::unique_lock<std::mutex> 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<const rocprofiler_agent_v0_t*>(agents[idx]);
if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue;
std::vector<rocprofiler_att_parameter_t> 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<Callbacks::ToolData*>(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;
}
projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp
-168
View File
@@ -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 <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include "common.hpp"
#include <atomic>
#include <cassert>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <map>
#include <mutex>
#include <sstream>
#include <string>
#include <string_view>
#include <unordered_map>
#include <vector>
#include <unistd.h>
#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<bool> 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<const rocprofiler_agent_v0_t*>(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;
}
projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp
+12 -47
View File
@@ -25,22 +25,7 @@
# undef NDEBUG
#endif
#include <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include "common.hpp"
#include <atomic>
#include <cassert>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <map>
#include <mutex>
#include <sstream>
#include <string>
#include <string_view>
#include <vector>
#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<size_t> 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<ToolData*>(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<Callbacks::ToolData*>(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<void*>(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<void*>(new Callbacks::ToolData{"att_multi_test/"})};
// return pointer to configure data
return &cfg;
projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp
+14 -59
View File
@@ -25,23 +25,7 @@
# undef NDEBUG
#endif
#include <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include "common.hpp"
#include <atomic>
#include <cassert>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <map>
#include <mutex>
#include <sstream>
#include <string>
#include <string_view>
#include <unordered_map>
#include <vector>
#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<ToolData*>(userdata);
auto& tool = *reinterpret_cast<Callbacks::ToolData*>(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<ToolData*>(tool_data);
std::unique_lock<std::mutex> 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<Callbacks::ToolData*>(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<void*>(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<void*>(new Callbacks::ToolData{"att_single_test/"})};
// return pointer to configure data
return &cfg;
projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp
+92 -161
View File
@@ -25,217 +25,148 @@
# undef NDEBUG
#endif
#include <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include "trace_callbacks.hpp"
#include <rocprofiler-sdk/cxx/codeobj/code_printing.hpp>
#include "common.hpp"
#include <atomic>
#ifdef ENABLE_ATT_FILES
# include <nlohmann/json.hpp>
#endif
#include <unistd.h>
#include <cassert>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <map>
#include <mutex>
#include <sstream>
#include <string>
#include <string_view>
#include <unordered_map>
#include <fstream>
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<ToolData*>(callback_data);
assert(userdata && "Dispatch callback passed null!");
auto& tool = *reinterpret_cast<Callbacks::ToolData*>(userdata);
if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
{
std::unique_lock<std::mutex> lg(tool.isa_map_mut);
auto* data = static_cast<kernel_symbol_data_t*>(record.payload);
auto* data = static_cast<kernel_symbol_data_t*>(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<code_obj_load_data_t*>(record.payload);
if(!data || !data->uri) return;
std::unique_lock<std::mutex> lg(tool.isa_map_mut);
static std::atomic<int> 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<const void*>(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<trace_data_t*>(userdata);
assert(trace_data.tool && "ISA callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(trace_data.tool);
std::unique_lock<std::mutex> lk(tool.isa_map_mut);
if(type == ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY)
{
const auto& ev = *reinterpret_cast<const rocprofiler_att_data_type_occupancy_t*>(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<rocprofiler_att_data_type_isa_t*>(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<TrackedIsa>();
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<trace_data_t*>(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<trace_data_t*>(userdata);
assert(trace_data.tool && "ISA callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(trace_data.tool);
std::unique_ptr<Instruction> instruction;
try
{
std::unique_lock<std::mutex> 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<TrackedIsa>();
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<ToolData*>(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<Callbacks::ToolData*>(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<Callbacks::ToolData*>(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
projects/rocprofiler-sdk/tests/thread-trace/common.hpp → projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.hpp
+42 -55
View File
@@ -22,39 +22,27 @@
#pragma once
#include <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <atomic>
#include <cassert>
#include <cstdint>
#include <cstdlib>
#include <map>
#include <iostream>
#include <mutex>
#include <set>
#include <sstream>
#include <string>
#include <unordered_map>
#include <vector>
#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<size_t> hitcount{0};
std::atomic<size_t> 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<uint64_t, std::string> kernel_id_to_kernel_name = {};
std::map<pcInfo, std::unique_ptr<TrackedIsa>> isa_map;
ToolData(const char* out)
: out_dir(out){};
std::atomic<int> waves_started = 0;
std::atomic<int> waves_ended = 0;
std::mutex isa_map_mut;
std::set<pcInfo> wave_start_locations{};
std::string out_dir{};
std::mutex mut{};
std::vector<CodeobjInfo> codeobjs{};
std::vector<std::string> att_files{};
std::unordered_map<uint64_t, std::string> 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