AI/rocm-systems
2
0
Fork 0
Κώδικας Ζητήματα Pull Requests Δράσεις Πακέτα Έργα Κυκλοφορίες Wiki Δραστηριότητα

ATT API changes - add user_data field and separation of dispatch vs agent profiling (#893)

Περιήγηση Στο Πηγαίο 
* DRM Issue Fix for SLES 15 (#897)

* DRM Issue Fix

* Formatting Fix

* PC sampling: CID manager unit test (#898)

* Adding per-dispatch userdata field to ATT

* Clang tidy

* Formatting

* Update source/lib/rocprofiler-sdk/hsa/aql_packet.hpp

Co-authored-by: Vladimir Indic <139573562+vlaindic@users.noreply.github.com>

* Adding dispatch_id, fixing user_data and update aql_profile_v2

* Formatting

* Tidy fixes

* Second fix for userdata

* removing assert for union

* Adding serialization. Created agent profiling-like thread trace

* Implemented agent thread trace

* Update source/lib/rocprofiler-sdk/hsa/aql_packet.hpp

Co-authored-by: Vladimir Indic <139573562+vlaindic@users.noreply.github.com>

* Restructured thread trace packets

* Added agent API tests

* Fixing multigpu for agent test

* Formatting

* Formatting

* Improving header locations

* Fixing merge conflicts

* Tidy

* Tidy

* Tidy

---------

Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com>
Co-authored-by: Vladimir Indic <139573562+vlaindic@users.noreply.github.com>

[ROCm/rocprofiler-sdk commit: 9676295d3d]
Αυτή η υποβολή περιλαμβάνεται σε:
Giovanni Lenzi Baraldi
2024-06-13 15:29:29 -03:00
υποβλήθηκε από GitHub
γονέας f7c804f916
υποβολή 37b0cfd7dc
24 αρχεία άλλαξαν με 1154 προσθήκες και 635 διαγραφές
Λήψη Αρχείου Patch Λήψη Αρχείου Diff Ανάπτυξη όλων των αρχείων Σύμπτυξη όλων των αρχείων
projects/rocprofiler-sdk/samples/advanced_thread_trace/client.cpp
+70 -81
Προβολή Αρχείου
@@ -183,58 +183,53 @@ struct source_location
struct trace_data_t
{
int64_t id;
uint8_t* data;
uint64_t size;
ToolData* tool;
int64_t id;
uint8_t* data;
uint64_t size;
};
auto* tool = new ToolData{};
void
tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* user_data,
void* callback_data)
rocprofiler_user_data_t* /* user_data */,
void* /* callback_data */)
{
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);
if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
{
std::unique_lock<std::shared_mutex> lg(tool.isa_map_mut);
std::unique_lock<std::shared_mutex> lg(tool->isa_map_mut);
auto* data = static_cast<kernel_symbol_data_t*>(record.payload);
tool.kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name);
tool->kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name);
}
if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return;
std::unique_lock<std::shared_mutex> lg(tool.isa_map_mut);
std::unique_lock<std::shared_mutex> lg(tool->isa_map_mut);
auto* data = static_cast<code_obj_load_data_t*>(record.payload);
if(std::string_view(data->uri).find("file:///") == 0)
{
tool.codeobjTranslate.addDecoder(
tool->codeobjTranslate.addDecoder(
data->uri, data->code_object_id, data->load_delta, data->load_size);
auto symbolmap = tool.codeobjTranslate.getSymbolMap(data->code_object_id);
auto symbolmap = tool->codeobjTranslate.getSymbolMap(data->code_object_id);
for(auto& [vaddr, symbol] : symbolmap)
tool.kernels_in_codeobj[vaddr] = symbol;
tool->kernels_in_codeobj[vaddr] = symbol;
}
else if(COPY_MEMORY_CODEOBJ)
{
tool.codeobjTranslate.addDecoder(reinterpret_cast<const void*>(data->memory_base),
data->memory_size,
data->code_object_id,
data->load_delta,
data->load_size);
auto symbolmap = tool.codeobjTranslate.getSymbolMap(data->code_object_id);
tool->codeobjTranslate.addDecoder(reinterpret_cast<const void*>(data->memory_base),
data->memory_size,
data->code_object_id,
data->load_delta,
data->load_size);
auto symbolmap = tool->codeobjTranslate.getSymbolMap(data->code_object_id);
for(auto& [vaddr, symbol] : symbolmap)
tool.kernels_in_codeobj[vaddr] = symbol;
tool->kernels_in_codeobj[vaddr] = symbol;
}
(void) user_data;
(void) callback_data;
C_API_END
}
@@ -243,20 +238,20 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */,
const rocprofiler_agent_t* /* agent */,
rocprofiler_correlation_id_t /* correlation_id */,
rocprofiler_kernel_id_t kernel_id,
void* userdata)
rocprofiler_dispatch_id_t /* dispatch_id */,
rocprofiler_user_data_t* /* userdata */,
void* /* userdata */)
{
C_API_BEGIN
assert(userdata && "Dispatch callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
std::shared_lock<std::shared_mutex> lg(tool.isa_map_mut);
std::shared_lock<std::shared_mutex> lg(tool->isa_map_mut);
static std::atomic<int> call_id{0};
static std::string_view desired_func_name = "transposeLds";
try
{
auto& kernel_name = tool.kernel_id_to_kernel_name.at(kernel_id);
auto& kernel_name = tool->kernel_id_to_kernel_name.at(kernel_id);
if(kernel_name.find(desired_func_name) == std::string::npos)
return ROCPROFILER_ATT_CONTROL_NONE;
@@ -276,29 +271,26 @@ get_trace_data(rocprofiler_att_parser_data_type_t type, void* att_data, void* us
{
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::shared_lock<std::shared_mutex> shared_lock(tool.isa_map_mut);
std::shared_lock<std::shared_mutex> shared_lock(tool->isa_map_mut);
if(type == ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY) tool.num_waves++;
if(type == ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY) tool->num_waves++;
if(type != ROCPROFILER_ATT_PARSER_DATA_TYPE_ISA) return;
auto& event = *reinterpret_cast<rocprofiler_att_data_type_isa_t*>(att_data);
pcinfo_t pc{event.marker_id, event.offset};
auto it = tool.isa_map.find(pc);
if(it == tool.isa_map.end())
auto it = tool->isa_map.find(pc);
if(it == tool->isa_map.end())
{
shared_lock.unlock();
{
std::unique_lock<std::shared_mutex> unique_lock(tool.isa_map_mut);
std::unique_lock<std::shared_mutex> unique_lock(tool->isa_map_mut);
auto ptr = std::make_unique<isa_map_elem_t>();
try
{
ptr->code_line = tool.codeobjTranslate.get(pc.marker_id, pc.addr);
ptr->code_line = tool->codeobjTranslate.get(pc.marker_id, pc.addr);
} catch(std::exception& e)
{
std::cerr << pc.marker_id << ":" << pc.addr << ' ' << e.what() << std::endl;
@@ -308,7 +300,7 @@ get_trace_data(rocprofiler_att_parser_data_type_t type, void* att_data, void* us
std::cerr << "Could not fetch: " << pc.marker_id << ':' << pc.addr << std::endl;
return;
}
it = tool.isa_map.emplace(pc, std::move(ptr)).first;
it = tool->isa_map.emplace(pc, std::move(ptr)).first;
}
shared_lock.lock();
}
@@ -339,15 +331,12 @@ isa_callback(char* isa_instruction,
{
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::shared_ptr<Instruction> instruction;
{
std::unique_lock<std::shared_mutex> unique_lock(tool.isa_map_mut);
instruction = tool.codeobjTranslate.get(marker_id, offset);
std::unique_lock<std::shared_mutex> unique_lock(tool->isa_map_mut);
instruction = tool->codeobjTranslate.get(marker_id, offset);
}
if(!instruction.get()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
@@ -364,23 +353,25 @@ isa_callback(char* isa_instruction,
auto ptr = std::make_unique<isa_map_elem_t>();
ptr->code_line = std::move(instruction);
tool.isa_map.emplace(pcinfo_t{marker_id, offset}, std::move(ptr));
C_API_END
tool->isa_map.emplace(pcinfo_t{marker_id, offset}, std::move(ptr));
return ROCPROFILER_STATUS_SUCCESS;
C_API_END
return ROCPROFILER_STATUS_ERROR;
}
void
shader_data_callback(int64_t se_id, void* se_data, size_t data_size, void* userdata)
shader_data_callback(int64_t se_id,
void* se_data,
size_t data_size,
rocprofiler_user_data_t /* userdata */)
{
C_API_BEGIN
assert(userdata && "Shader callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
{
std::unique_lock<std::mutex> lk(tool.output_mut);
tool.output() << "SE ID: " << se_id << " with size " << data_size << std::hex << '\n';
std::unique_lock<std::mutex> lk(tool->output_mut);
tool->output() << "SE ID: " << se_id << " with size " << data_size << std::hex << '\n';
}
trace_data_t data{.id = se_id, .data = (uint8_t*) se_data, .size = data_size, .tool = &tool};
trace_data_t data{.id = se_id, .data = (uint8_t*) se_data, .size = data_size};
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;
@@ -408,12 +399,12 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
{ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {BUFFER_SIZE}},
{ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {SE_MASK}}};
ROCPROFILER_CALL(rocprofiler_configure_thread_trace_service(client_ctx,
parameters.data(),
parameters.size(),
dispatch_callback,
shader_data_callback,
tool_data),
ROCPROFILER_CALL(rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
parameters.data(),
parameters.size(),
dispatch_callback,
shader_data_callback,
tool_data),
"thread trace service configure");
int valid_ctx = 0;
@@ -434,17 +425,14 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
}
void
tool_fini(void* tool_data)
tool_fini(void* /* data */)
{
assert(tool_data && "tool_fini callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(tool_data);
std::unique_lock<std::shared_mutex> isa_lk(tool.isa_map_mut);
std::unique_lock<std::mutex> out_lk(tool.output_mut);
std::unique_lock<std::shared_mutex> isa_lk(client::tool->isa_map_mut);
std::unique_lock<std::mutex> out_lk(client::tool->output_mut);
// Find largest instruction
size_t max_inst_size = 0;
for(auto& [addr, lines] : tool.isa_map)
for(auto& [addr, lines] : client::tool->isa_map)
if(lines.get()) max_inst_size = std::max(max_inst_size, lines->code_line->inst.size());
std::string empty_space;
@@ -460,16 +448,16 @@ tool_fini(void* tool_data)
size_t vector_exec = 0;
size_t other_exec = 0;
for(auto& [addr, line] : tool.isa_map)
for(auto& [addr, line] : client::tool->isa_map)
if(line.get())
{
size_t hitcount = line->hitcount.load(std::memory_order_relaxed);
size_t latency = line->latency.load(std::memory_order_relaxed);
auto& code_line = line->code_line->inst;
tool.output() << std::hex << "0x" << addr.addr << std::dec << ' ' << code_line
<< empty_space.substr(0, max_inst_size - code_line.size())
<< " Hit: " << hitcount << " - Latency: " << latency << '\n';
client::tool->output() << std::hex << "0x" << addr.addr << std::dec << ' ' << code_line
<< empty_space.substr(0, max_inst_size - code_line.size())
<< " Hit: " << hitcount << " - Latency: " << latency << '\n';
if(code_line.find("s_waitcnt") == 0)
{
@@ -502,14 +490,17 @@ tool_fini(void* tool_data)
float vmc_fraction = 100 * vmc_latency / float(total_latency);
float lgk_fraction = 100 * lgk_latency / float(total_latency);
tool.output() << "Total executed instructions: " << total_exec << '\n'
<< "Total executed vector instructions: " << vector_exec << " with average "
<< vector_latency / float(vector_exec) << " cycles.\n"
<< "Total executed scalar instructions: " << scalar_exec << " with average "
<< scalar_latency / float(scalar_exec) << " cycles.\n"
<< "Vector memory ops occupied: " << vmc_fraction << "% of cycles.\n"
<< "Scalar and LDS memory ops occupied: " << lgk_fraction << "% of cycles.\n"
<< "Num waves created: " << (tool.num_waves / 2) << std::endl;
client::tool->output() << "Total executed instructions: " << total_exec << '\n'
<< "Total executed vector instructions: " << vector_exec
<< " with average " << vector_latency / float(vector_exec)
<< " cycles.\n"
<< "Total executed scalar instructions: " << scalar_exec
<< " with average " << scalar_latency / float(scalar_exec)
<< " cycles.\n"
<< "Vector memory ops occupied: " << vmc_fraction << "% of cycles.\n"
<< "Scalar and LDS memory ops occupied: " << lgk_fraction
<< "% of cycles.\n"
<< "Num waves created: " << (client::tool->num_waves / 2) << std::endl;
}
} // namespace client
@@ -538,14 +529,12 @@ rocprofiler_configure(uint32_t version,
std::clog << info.str() << std::endl;
auto* data = new client::ToolData{};
// create configure data
static auto cfg =
rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t),
&client::tool_init,
&client::tool_fini,
reinterpret_cast<void*>(data)};
nullptr};
// return pointer to configure data
return &cfg;
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/CMakeLists.txt
+2 -1
Προβολή Αρχείου
@@ -3,7 +3,8 @@
# Installation of amd_detail headers
#
#
set(ROCPROFILER_AMD_DETAIL_HEADER_FILES thread_trace.h)
set(ROCPROFILER_AMD_DETAIL_HEADER_FILES thread_trace.h thread_trace_core.h
thread_trace_dispatch.h thread_trace_agent.h)
install(
FILES ${ROCPROFILER_AMD_DETAIL_HEADER_FILES}
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace.h
+3 -183
Προβολή Αρχείου
@@ -22,186 +22,6 @@
#pragma once
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/defines.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/hsa.h>
#include <cstdint>
ROCPROFILER_EXTERN_C_INIT
/**
* @defgroup THREAD_TRACE Thread Trace Service
* @brief Provides API calls to enable and handle thread trace data
*
* @{
*/
typedef enum
{
ROCPROFILER_ATT_PARAMETER_TARGET_CU = 0, ///< Select the Target CU or WGP
ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, ///< Bitmask of shader engines.
ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, ///< Size of combined GPU buffer for ATT
ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, ///< Bitmask (GFX9) or ID (Navi) of SIMDs
ROCPROFILER_ATT_PARAMETER_CODE_OBJECT_TRACE_ENABLE, ///< Enables Codeobj Markers IDs into ATT
ROCPROFILER_ATT_PARAMETER_PERFCOUNTER, ///< Enables Perfcounter with simd mask (GFX9 only)
ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL, ///< Defines the update period (GFX9 only)
ROCPROFILER_ATT_PARAMETER_LAST
} rocprofiler_att_parameter_type_t;
typedef struct
{
rocprofiler_att_parameter_type_t type;
union
{
uint64_t value;
struct
{
rocprofiler_counter_id_t counter_id;
uint64_t simd_mask : 4;
};
};
} rocprofiler_att_parameter_t;
typedef enum
{
ROCPROFILER_ATT_CONTROL_NONE = 0,
ROCPROFILER_ATT_CONTROL_START = 1,
ROCPROFILER_ATT_CONTROL_STOP = 2,
ROCPROFILER_ATT_CONTROL_START_AND_STOP = 3
} rocprofiler_att_control_flags_t;
/**
* @brief Callback to be triggered every kernel dispatch, indicating to start and/or stop ATT
*/
typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)(
rocprofiler_queue_id_t queue_id,
const rocprofiler_agent_t* agent,
rocprofiler_correlation_id_t correlation_id,
rocprofiler_kernel_id_t kernel_id,
void* userdata);
/**
* @brief Callback to be triggered every time some ATT data is generated by the device
* @param [in] shader_engine_id ID of shader engine, as enabled by SE_MASK
* @param [in] data Pointer to the buffer containing the ATT data
* @param [in] data_size Number of bytes in "data"
* @param [in] userdata Passed back to user
*/
typedef void (*rocprofiler_att_shader_data_callback_t)(int64_t shader_engine_id,
void* data,
size_t data_size,
void* userdata);
/**
* @brief Enables the advanced thread trace service.
* @param [in] context_id context_id.
* @param [in] parameters List of ATT-specific parameters.
* @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.
*/
rocprofiler_status_t
rocprofiler_configure_thread_trace_service(rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata) ROCPROFILER_API;
/**
* @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 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 (*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_EXTERN_C_FINI
#include <rocprofiler-sdk/amd_detail/thread_trace_agent.h>
#include <rocprofiler-sdk/amd_detail/thread_trace_core.h>
#include <rocprofiler-sdk/amd_detail/thread_trace_dispatch.h>
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h
+64
Προβολή Αρχείου
@@ -0,0 +1,64 @@
// 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 <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/amd_detail/thread_trace_core.h>
#include <rocprofiler-sdk/defines.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/hsa.h>
ROCPROFILER_EXTERN_C_INIT
/**
* @defgroup THREAD_TRACE Thread Trace Service
* @brief Provides API calls to enable and handle thread trace data
*
* @{
*/
/**
* @brief Configure Thread Trace Service for agent. There may only be one agent profile
* configured per context and can be only one active context that is profiling a single agent
* at a time. Multiple agent contexts can be started at the same time if they are profiling
* different agents.
*
* @param [in] context_id context id
* @param [in] parameters List of ATT-specific parameters.
* @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.
*/
rocprofiler_status_t
rocprofiler_configure_agent_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata) ROCPROFILER_API;
/** @} */
ROCPROFILER_EXTERN_C_FINI
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h
+171
Προβολή Αρχείου
@@ -0,0 +1,171 @@
// 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 <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/defines.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/hsa.h>
ROCPROFILER_EXTERN_C_INIT
/**
* @defgroup THREAD_TRACE Thread Trace Service
* @brief Provides API calls to enable and handle thread trace data
*
* @{
*/
typedef enum
{
ROCPROFILER_ATT_PARAMETER_TARGET_CU = 0, ///< Select the Target CU or WGP
ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, ///< Bitmask of shader engines.
ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, ///< Size of combined GPU buffer for ATT
ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, ///< Bitmask (GFX9) or ID (Navi) of SIMDs
ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL,
ROCPROFILER_ATT_PARAMETER_PERFCOUNTER,
ROCPROFILER_ATT_PARAMETER_LAST
} rocprofiler_att_parameter_type_t;
typedef struct
{
rocprofiler_att_parameter_type_t type;
union
{
uint64_t value;
struct
{
rocprofiler_counter_id_t counter_id;
uint64_t simd_mask : 4;
};
};
} rocprofiler_att_parameter_t;
/**
* @brief Callback to be triggered every time some ATT data is generated by the device
* @param [in] shader_engine_id ID of shader engine, as enabled by SE_MASK
* @param [in] data Pointer to the buffer containing the ATT data
* @param [in] data_size Number of bytes in "data"
* @param [in] userdata_dispatch Passed back to user from rocprofiler_att_dispatch_callback_t()
* @param [in] userdata_config Passed back to user from configure_[...]_service()
*/
typedef void (*rocprofiler_att_shader_data_callback_t)(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 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 (*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_EXTERN_C_FINI
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h
+82
Προβολή Αρχείου
@@ -0,0 +1,82 @@
// 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 <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/amd_detail/thread_trace_core.h>
#include <rocprofiler-sdk/defines.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/hsa.h>
ROCPROFILER_EXTERN_C_INIT
/**
* @defgroup THREAD_TRACE Thread Trace Service
* @brief Provides API calls to enable and handle thread trace data
*
* @{
*/
typedef enum
{
ROCPROFILER_ATT_CONTROL_NONE = 0,
ROCPROFILER_ATT_CONTROL_START = 1,
ROCPROFILER_ATT_CONTROL_STOP = 2,
ROCPROFILER_ATT_CONTROL_START_AND_STOP = 3
} rocprofiler_att_control_flags_t;
/**
* @brief Callback to be triggered every kernel dispatch, indicating to start and/or stop ATT
*/
typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)(
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);
/**
* @brief Enables the advanced thread trace service for dispatch-based tracing.
* The tool has an option to enable/disable thread trace on every dispatch callback.
* This service enables kernel serialization.
* @param [in] context_id context_id.
* @param [in] parameters List of ATT-specific parameters.
* @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.
*/
rocprofiler_status_t
rocprofiler_configure_dispatch_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata) ROCPROFILER_API;
/** @} */
ROCPROFILER_EXTERN_C_FINI
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h
+18 -3
Προβολή Αρχείου
@@ -184,14 +184,29 @@ aqlprofile_get_pmc_info(const aqlprofile_pmc_profile_t* profile,
aqlprofile_pmc_info_type_t attribute,
void* value);
// Profile parameter object
typedef struct
{
hsa_ven_amd_aqlprofile_parameter_name_t parameter_name;
union
{
uint32_t value;
struct
{
uint32_t counter_id : 28;
uint32_t simd_mask : 4;
};
};
} aqlprofile_att_parameter_t;
/**
* @brief AQLprofile struct containing information for Advanced Thread Trace
*/
typedef struct
{
hsa_agent_t agent;
const hsa_ven_amd_aqlprofile_parameter_t* parameters;
uint32_t parameter_count;
hsa_agent_t agent;
const aqlprofile_att_parameter_t* parameters;
uint32_t parameter_count;
} aqlprofile_att_profile_t;
/**
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp
+15 -18
Προβολή Αρχείου
@@ -219,31 +219,28 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache&
uint32_t perf_ctrl = static_cast<uint32_t>(params.perfcounter_ctrl);
aql_params.clear();
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET, cu});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SE_MASK, shader_engine_mask});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SIMD_SELECTION, simd});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE, buffer_size});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_PERFCOUNTER_CTRL, perf_ctrl});
for(uint32_t perf_counter : params.perfcounters)
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET, {cu}});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SE_MASK, {shader_engine_mask}});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SIMD_SELECTION, {simd}});
aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE, {buffer_size}});
if(perf_ctrl != 0 && !params.perfcounters.empty())
{
aql_params.push_back(
{HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_PERFCOUNTER_NAME, perf_counter});
{HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_PERFCOUNTER_CTRL, {perf_ctrl - 1}});
auto perf_param = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_PERFCOUNTER_NAME;
for(uint32_t perf_counter : params.perfcounters)
aql_params.push_back({perf_param, {perf_counter}});
}
}
const std::vector<hsa_ven_amd_aqlprofile_parameter_t>&
ThreadTraceAQLPacketFactory::get_aql_params()
{
return aql_params;
}
std::unique_ptr<hsa::TraceControlAQLPacket>
ThreadTraceAQLPacketFactory::construct_packet()
ThreadTraceAQLPacketFactory::construct_control_packet()
{
uint32_t num_params = static_cast<uint32_t>(get_aql_params().size());
auto profile =
aqlprofile_att_profile_t{tracepool.gpu_agent, get_aql_params().data(), num_params};
auto packet = std::make_unique<hsa::TraceControlAQLPacket>(this->tracepool, profile);
auto num_params = static_cast<uint32_t>(aql_params.size());
auto profile = aqlprofile_att_profile_t{tracepool.gpu_agent, aql_params.data(), num_params};
auto packet = std::make_unique<hsa::TraceControlAQLPacket>(this->tracepool, profile);
packet->clear();
return packet;
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp
+11 -7
Προβολή Αρχείου
@@ -86,21 +86,25 @@ protected:
class ThreadTraceAQLPacketFactory
{
using thread_trace_parameter_pack = thread_trace::thread_trace_parameter_pack;
public:
ThreadTraceAQLPacketFactory(const hsa::AgentCache& agent,
const thread_trace_parameter_pack& params,
const CoreApiTable& coreapi,
const AmdExtTable& ext);
const std::vector<hsa_ven_amd_aqlprofile_parameter_t>& get_aql_params();
std::unique_ptr<hsa::TraceControlAQLPacket> construct_packet();
std::unique_ptr<hsa::CodeobjMarkerAQLPacket> construct_load_marker_packet(uint64_t id,
uint64_t addr,
uint64_t size);
std::unique_ptr<hsa::TraceControlAQLPacket> construct_control_packet();
std::unique_ptr<hsa::CodeobjMarkerAQLPacket> construct_load_marker_packet(uint64_t id,
uint64_t addr,
uint64_t size);
std::unique_ptr<hsa::CodeobjMarkerAQLPacket> construct_unload_marker_packet(uint64_t id);
std::vector<aqlprofile_att_parameter_t> aql_params;
private:
hsa::TraceMemoryPool tracepool;
std::vector<hsa_ven_amd_aqlprofile_parameter_t> aql_params;
hsa::TraceMemoryPool tracepool;
};
} // namespace aql
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp
+2 -2
Προβολή Αρχείου
@@ -138,8 +138,8 @@ struct context
std::unique_ptr<dispatch_counter_collection_service> counter_collection = {};
std::unique_ptr<agent_counter_collection_service> agent_counter_collection = {};
std::unique_ptr<pc_sampling_service> pc_sampler = {};
// TODO: Make a unique pointer instead
std::shared_ptr<GlobalThreadTracer> thread_trace = {};
std::unique_ptr<thread_trace::ThreadTracerInterface> thread_trace = {};
};
// set the client index needs to be called before allocate_context()
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp
+25 -18
Προβολή Αρχείου
@@ -66,10 +66,10 @@ CounterAQLPacket::~CounterAQLPacket()
}
hsa_status_t
BaseTTAQLPacket::Alloc(void** ptr, size_t size, desc_t flags, void* data)
TraceMemoryPool::Alloc(void** ptr, size_t size, desc_t flags, void* data)
{
if(!data) return HSA_STATUS_ERROR;
auto& pool = reinterpret_cast<BaseTTAQLPacket*>(data)->tracepool;
auto& pool = *reinterpret_cast<TraceMemoryPool*>(data);
if(!pool.allocate_fn || !pool.free_fn || !pool.allow_access_fn) return HSA_STATUS_ERROR;
@@ -91,19 +91,19 @@ BaseTTAQLPacket::Alloc(void** ptr, size_t size, desc_t flags, void* data)
}
void
BaseTTAQLPacket::Free(void* ptr, void* data)
TraceMemoryPool::Free(void* ptr, void* data)
{
assert(data);
auto& pool = reinterpret_cast<BaseTTAQLPacket*>(data)->tracepool;
auto& pool = *reinterpret_cast<TraceMemoryPool*>(data);
if(pool.free_fn) pool.free_fn(ptr);
}
hsa_status_t
BaseTTAQLPacket::Copy(void* dst, const void* src, size_t size, void* data)
TraceMemoryPool::Copy(void* dst, const void* src, size_t size, void* data)
{
if(!data) return HSA_STATUS_ERROR;
auto& pool = reinterpret_cast<BaseTTAQLPacket*>(data)->tracepool;
auto& pool = *reinterpret_cast<TraceMemoryPool*>(data);
if(!pool.api_copy_fn) return HSA_STATUS_ERROR;
@@ -112,9 +112,15 @@ BaseTTAQLPacket::Copy(void* dst, const void* src, size_t size, void* data)
TraceControlAQLPacket::TraceControlAQLPacket(const TraceMemoryPool& _tracepool,
const aqlprofile_att_profile_t& p)
: BaseTTAQLPacket(_tracepool)
: tracepool(std::make_shared<TraceMemoryPool>(_tracepool))
{
auto status = aqlprofile_att_create_packets(&handle, &packets, p, &Alloc, &Free, &Copy, this);
auto status = aqlprofile_att_create_packets(&tracepool->handle,
&packets,
p,
&TraceMemoryPool::Alloc,
&TraceMemoryPool::Free,
&TraceMemoryPool::Copy,
tracepool.get());
CHECK_HSA(status, "failed to create ATT packet");
packets.start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
@@ -122,14 +128,8 @@ TraceControlAQLPacket::TraceControlAQLPacket(const TraceMemoryPool& _tr
packets.start_packet.completion_signal = hsa_signal_t{.handle = 0};
packets.stop_packet.completion_signal = hsa_signal_t{.handle = 0};
this->empty = false;
};
void
TraceControlAQLPacket::populate_before()
{
before_krn_pkt.push_back(packets.start_packet);
for(auto& [_, codeobj] : loaded_codeobj)
if(codeobj) before_krn_pkt.push_back(codeobj->packet);
clear();
};
CodeobjMarkerAQLPacket::CodeobjMarkerAQLPacket(const TraceMemoryPool& _tracepool,
@@ -138,22 +138,29 @@ CodeobjMarkerAQLPacket::CodeobjMarkerAQLPacket(const TraceMemoryPool& _tracepool
uint64_t size,
bool bFromStart,
bool bIsUnload)
: BaseTTAQLPacket(_tracepool)
: tracepool(_tracepool)
{
aqlprofile_att_codeobj_data_t codeobj{};
codeobj.id = id;
codeobj.addr = addr;
codeobj.size = size;
codeobj.agent = _tracepool.gpu_agent;
codeobj.agent = tracepool.gpu_agent;
codeobj.isUnload = bIsUnload;
codeobj.fromStart = bFromStart;
auto status = aqlprofile_att_codeobj_marker(&packet, &handle, codeobj, &Alloc, &Free, this);
auto status = aqlprofile_att_codeobj_marker(&packet,
&tracepool.handle,
codeobj,
&TraceMemoryPool::Alloc,
&TraceMemoryPool::Free,
&tracepool);
CHECK_HSA(status, "failed to create ATT marker");
packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
packet.completion_signal = hsa_signal_t{.handle = 0};
this->empty = false;
clear();
}
} // namespace hsa
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp
+62 -26
Προβολή Αρχείου
@@ -25,6 +25,8 @@
#include "lib/common/container/small_vector.hpp"
#include "lib/rocprofiler-sdk/aql/aql_profile_v2.h"
#include <rocprofiler-sdk/rocprofiler.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
@@ -83,6 +85,16 @@ public:
bool isEmpty() const { return empty; }
};
class EmptyAQLPacket : public AQLPacket
{
public:
EmptyAQLPacket() = default;
~EmptyAQLPacket() override = default;
void populate_before() override{};
void populate_after() override{};
};
class CounterAQLPacket : public AQLPacket
{
friend class rocprofiler::aql::CounterPacketConstruct;
@@ -108,6 +120,8 @@ protected:
struct TraceMemoryPool
{
using desc_t = aqlprofile_buffer_desc_flags_t;
hsa_agent_t gpu_agent;
hsa_amd_memory_pool_t cpu_pool_;
hsa_amd_memory_pool_t gpu_pool_;
@@ -115,33 +129,32 @@ struct TraceMemoryPool
decltype(hsa_amd_agents_allow_access)* allow_access_fn;
decltype(hsa_amd_memory_pool_free)* free_fn;
decltype(hsa_memory_copy)* api_copy_fn;
};
class BaseTTAQLPacket : public AQLPacket
{
friend class rocprofiler::aql::ThreadTraceAQLPacketFactory;
protected:
using desc_t = aqlprofile_buffer_desc_flags_t;
public:
BaseTTAQLPacket(const TraceMemoryPool& _tracepool)
: tracepool(_tracepool){};
~BaseTTAQLPacket() override { aqlprofile_att_delete_packets(this->handle); };
aqlprofile_handle_t GetHandle() const { return handle; }
hsa_agent_t GetAgent() const { return tracepool.gpu_agent; }
protected:
TraceMemoryPool tracepool;
aqlprofile_handle_t handle;
~TraceMemoryPool() { aqlprofile_att_delete_packets(this->handle); };
static hsa_status_t Alloc(void** ptr, size_t size, desc_t flags, void* data);
static void Free(void* ptr, void* data);
static hsa_status_t Copy(void* dst, const void* src, size_t size, void* data);
};
class CodeobjMarkerAQLPacket : public BaseTTAQLPacket
class BaseTTAQLPacket : public AQLPacket
{
friend class rocprofiler::aql::ThreadTraceAQLPacketFactory;
public:
BaseTTAQLPacket(std::shared_ptr<TraceMemoryPool>& _tracepool)
: tracepool(_tracepool){};
~BaseTTAQLPacket() override = default;
aqlprofile_handle_t GetHandle() const { return tracepool->handle; }
hsa_agent_t GetAgent() const { return tracepool->gpu_agent; }
protected:
std::shared_ptr<TraceMemoryPool> tracepool;
};
class CodeobjMarkerAQLPacket : public AQLPacket
{
friend class rocprofiler::aql::ThreadTraceAQLPacketFactory;
@@ -157,10 +170,16 @@ public:
void populate_before() override { before_krn_pkt.push_back(packet); };
void populate_after() override{};
aqlprofile_handle_t GetHandle() const { return tracepool.handle; }
hsa_agent_t GetAgent() const { return tracepool.gpu_agent; }
hsa_ext_amd_aql_pm4_packet_t packet;
protected:
TraceMemoryPool tracepool;
};
class TraceControlAQLPacket : public BaseTTAQLPacket
class TraceControlAQLPacket : public AQLPacket
{
friend class rocprofiler::aql::ThreadTraceAQLPacketFactory;
using code_object_id_t = uint64_t;
@@ -170,19 +189,36 @@ public:
const aqlprofile_att_profile_t& profile);
~TraceControlAQLPacket() override = default;
explicit TraceControlAQLPacket(TraceControlAQLPacket& other)
{
this->tracepool = other.tracepool;
this->packets = other.packets;
this->loaded_codeobj = other.loaded_codeobj;
}
aqlprofile_handle_t GetHandle() const { return tracepool->handle; }
hsa_agent_t GetAgent() const { return tracepool->gpu_agent; }
void populate_before() override
{
before_krn_pkt.push_back(packets.start_packet);
for(auto& [_, codeobj] : loaded_codeobj)
before_krn_pkt.push_back(codeobj->packet);
}
void populate_after() override { after_krn_pkt.push_back(packets.stop_packet); }
void add_codeobj(code_object_id_t id, uint64_t addr, uint64_t size)
{
loaded_codeobj[id] =
std::make_unique<CodeobjMarkerAQLPacket>(tracepool, id, addr, size, true, false);
std::make_shared<CodeobjMarkerAQLPacket>(*tracepool, id, addr, size, true, false);
}
void remove_codeobj(code_object_id_t id) { loaded_codeobj.erase(id); }
void populate_before() override;
void populate_after() override { after_krn_pkt.push_back(packets.stop_packet); }
protected:
std::shared_ptr<TraceMemoryPool> tracepool;
aqlprofile_att_control_aql_packets_t packets;
private:
aqlprofile_att_control_aql_packets_t packets;
std::unordered_map<code_object_id_t, std::unique_ptr<CodeobjMarkerAQLPacket>> loaded_codeobj;
std::unordered_map<code_object_id_t, std::shared_ptr<CodeobjMarkerAQLPacket>> loaded_codeobj;
};
} // namespace hsa
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
+17 -21
Προβολή Αρχείου
@@ -143,8 +143,11 @@ constexpr rocprofiler_agent_t default_agent =
void
QueueController::add_queue(hsa_queue_t* id, std::unique_ptr<Queue> queue)
{
for(auto& pre_initialize_fn : pre_initialize)
pre_initialize_fn(queue->get_agent(), get_core_table(), get_ext_table());
for(const auto& itr : context::get_registered_contexts())
{
if(itr->thread_trace)
itr->thread_trace->resource_init(queue->get_agent(), get_core_table(), get_ext_table());
}
CHECK(queue);
_callback_cache.wlock([&](auto& callbacks) {
@@ -167,11 +170,16 @@ void
QueueController::destroy_queue(hsa_queue_t* id)
{
if(!id) return;
_queues.wlock([&](auto& map) {
for(auto& deinitialize_fn : pre_deinitialize)
for(const auto& itr : context::get_registered_contexts())
{
if(!itr->thread_trace) continue;
_queues.wlock([&](auto& map) {
if(map.find(id) != map.end())
deinitialize_fn(map.at(id)->get_agent(), get_core_table(), get_ext_table());
});
itr->thread_trace->resource_deinit(map.at(id)->get_agent());
});
}
const auto* queue = get_queue(*id);
@@ -254,10 +262,9 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table)
auto enable_intercepter = false;
for(const auto& itr : context::get_registered_contexts())
{
constexpr auto expected_context_size = 200UL;
constexpr auto expected_context_size = 208UL;
static_assert(
sizeof(context::context) ==
expected_context_size + sizeof(std::shared_ptr<rocprofiler::GlobalThreadTracer>),
sizeof(context::context) == expected_context_size,
"If you added a new field to context struct, make sure there is a check here if it "
"requires queue interception. Once you have done so, increment expected_context_size");
@@ -275,18 +282,7 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table)
}
else if(itr->thread_trace)
{
enable_intercepter = true;
std::weak_ptr<rocprofiler::GlobalThreadTracer> trace = itr->thread_trace;
// TODO: Make it wrapper on HSA initialization
pre_initialize.emplace_back(
[trace](const AgentCache& cache, const CoreApiTable& core, const AmdExtTable& ext) {
if(auto locked = trace.lock()) locked->resource_init(cache, core, ext);
});
pre_deinitialize.emplace_back(
[trace](const AgentCache& cache, const CoreApiTable&, const AmdExtTable&) {
if(auto locked = trace.lock()) locked->resource_deinit(cache);
});
enable_intercepter = true;
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp
-4
Προβολή Αρχείου
@@ -103,7 +103,6 @@ public:
private:
using client_id_map_t = std::unordered_map<ClientID, agent_callback_tuple_t>;
using agent_cache_map_t = std::unordered_map<uint32_t, AgentCache>;
using resource_alloc_t = void(const AgentCache&, const CoreApiTable&, const AmdExtTable&);
CoreApiTable _core_table = {};
AmdExtTable _ext_table = {};
@@ -111,9 +110,6 @@ private:
common::Synchronized<client_id_map_t> _callback_cache = {};
agent_cache_map_t _supported_agents = {};
common::Synchronized<hsa::profiler_serializer> _profiler_serializer;
std::vector<std::function<resource_alloc_t>> pre_initialize;
std::vector<std::function<resource_alloc_t>> pre_deinitialize;
};
QueueController*
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp
+202 -152
Προβολή Αρχείου
@@ -55,17 +55,18 @@ constexpr size_t ROCPROFILER_QUEUE_SIZE = 64;
namespace rocprofiler
{
namespace thread_trace
{
struct cbdata_t
{
void* tool_userdata;
rocprofiler_att_shader_data_callback_t cb_fn;
rocprofiler_correlation_id_t corr_id;
const rocprofiler_user_data_t* dispatch_userdata;
};
common::Synchronized<std::optional<int64_t>> client;
bool
AgentThreadTracer::Submit(hsa_ext_amd_aql_pm4_packet_t* packet)
ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet)
{
const uint64_t write_idx = add_write_index_relaxed_fn(queue, 1);
@@ -94,14 +95,14 @@ AgentThreadTracer::Submit(hsa_ext_amd_aql_pm4_packet_t* packet)
return true;
}
AgentThreadTracer::AgentThreadTracer(thread_trace_parameter_pack _params,
ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params,
const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
: params(std::move(_params))
{
factory = std::make_unique<aql::ThreadTraceAQLPacketFactory>(cache, this->params, coreapi, ext);
cached_resources = factory->construct_packet();
control_packet = factory->construct_control_packet();
auto status = coreapi.hsa_queue_create_fn(cache.get_hsa_agent(),
ROCPROFILER_QUEUE_SIZE,
@@ -123,140 +124,72 @@ AgentThreadTracer::AgentThreadTracer(thread_trace_parameter_pack _params,
load_read_index_relaxed_fn = coreapi.hsa_queue_load_read_index_relaxed_fn;
}
AgentThreadTracer::~AgentThreadTracer()
ThreadTracerQueue::~ThreadTracerQueue()
{
std::unique_lock<std::mutex> lk(trace_resources_mut);
if(active_resources.packet != nullptr)
ROCP_WARNING << "Thread tracer being destroyed with thread trace active";
if(!this->queue) return;
auto* packet = static_cast<hsa::TraceControlAQLPacket*>(active_resources.packet.get());
if(packet)
if(active_traces.load() < 1)
{
packet->clear();
packet->populate_after();
for(auto& after_packet : packet->after_krn_pkt)
Submit(&after_packet);
if(queue_destroy_fn) queue_destroy_fn(this->queue);
return;
}
if(queue_destroy_fn) queue_destroy_fn(this->queue);
ROCP_WARNING << "Thread tracer being destroyed with thread trace active";
control_packet->clear();
control_packet->populate_after();
for(auto& after_packet : control_packet->after_krn_pkt)
Submit(&after_packet);
}
/**
* Callback we get from HSA interceptor when a kernel packet is being enqueued.
* We return an AQLPacket containing the start/stop/read packets for injection.
*/
std::unique_ptr<hsa::AQLPacket>
AgentThreadTracer::pre_kernel_call(rocprofiler_att_control_flags_t control_flags,
rocprofiler_queue_id_t queue_id,
rocprofiler_correlation_id_t corr_id)
std::unique_ptr<hsa::TraceControlAQLPacket>
ThreadTracerQueue::get_control(bool bStart)
{
if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) return nullptr;
std::unique_lock<std::mutex> lk(trace_resources_mut);
if(control_flags == ROCPROFILER_ATT_CONTROL_STOP)
{
if(active_resources.packet == nullptr)
{
ROCP_ERROR << "Attempt at stopping a thread trace that has not started!\n";
return nullptr;
}
auto active_resources = std::make_unique<hsa::TraceControlAQLPacket>(*control_packet);
active_resources->clear();
active_resources.packet->clear();
active_resources.packet->populate_after();
data_is_ready.fetch_add(1);
return std::move(active_resources.packet);
}
if(bStart) active_traces.fetch_add(1);
if(active_resources.packet != nullptr)
{
ROCP_ERROR << "Attempt at starting a thread trace while another was active!\n";
return nullptr;
}
else
{
active_resources.corr_id = corr_id;
active_resources.queue_id = queue_id;
}
if(cached_resources == nullptr)
{
ROCP_ERROR << "Attempt to initialize ATT without allocated resources!\n";
return nullptr;
}
cached_resources->clear();
cached_resources->populate_before();
if((control_flags & ROCPROFILER_ATT_CONTROL_STOP) != 0)
{
cached_resources->populate_after();
data_is_ready.fetch_add(1);
}
return std::move(cached_resources);
return active_resources;
}
hsa_status_t
thread_trace_callback(uint32_t shader, void* buffer, uint64_t size, void* callback_data)
{
void* tool_userdata = static_cast<cbdata_t*>(callback_data)->tool_userdata;
auto callback_fn = *static_cast<cbdata_t*>(callback_data)->cb_fn;
auto& cb_data = *static_cast<cbdata_t*>(callback_data);
callback_fn(shader, buffer, size, tool_userdata);
cb_data.cb_fn(shader, buffer, size, *cb_data.dispatch_userdata);
return HSA_STATUS_SUCCESS;
}
void
AgentThreadTracer::post_kernel_call(std::unique_ptr<hsa::AQLPacket>&& aql)
ThreadTracerQueue::iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data)
{
std::unique_lock<std::mutex> lk(trace_resources_mut);
active_resources.packet = std::move(aql);
if(!active_resources.packet || data_is_ready.load() < 1) return;
auto* pkt = static_cast<hsa::TraceControlAQLPacket*>(active_resources.packet.get());
for(auto& record : remaining_codeobj_record)
{
if(!record.bUnload)
pkt->add_codeobj(record.id, record.addr, record.size);
else
pkt->remove_codeobj(record.id);
}
remaining_codeobj_record.clear();
cbdata_t cb_dt{};
cb_dt.corr_id = active_resources.corr_id;
cb_dt.tool_userdata = params.callback_userdata;
cb_dt.cb_fn = params.shader_cb_fn;
cb_dt.cb_fn = params.shader_cb_fn;
cb_dt.dispatch_userdata = &data;
auto status = aqlprofile_att_iterate_data(pkt->GetHandle(), thread_trace_callback, &cb_dt);
auto status = aqlprofile_att_iterate_data(handle, thread_trace_callback, &cb_dt);
CHECK_HSA(status, "Failed to iterate ATT data");
data_is_ready.fetch_sub(1);
cached_resources = std::move(active_resources.packet);
active_traces.fetch_sub(1);
}
void
AgentThreadTracer::load_codeobj(code_object_id_t id, uint64_t addr, uint64_t size)
ThreadTracerQueue::load_codeobj(code_object_id_t id, uint64_t addr, uint64_t size)
{
std::unique_lock<std::mutex> lk(trace_resources_mut);
if(auto* pkt = static_cast<hsa::TraceControlAQLPacket*>(cached_resources.get()))
{
pkt->add_codeobj(id, addr, size);
return;
}
control_packet->add_codeobj(id, addr, size);
remaining_codeobj_record.push_back({id, addr, size, false});
if(!queue) return;
if(!queue || active_traces.load() < 1) return;
auto packet = factory->construct_load_marker_packet(id, addr, size);
bool bSuccess = Submit(&packet->packet);
@@ -266,19 +199,13 @@ AgentThreadTracer::load_codeobj(code_object_id_t id, uint64_t addr, uint64_t siz
}
void
AgentThreadTracer::unload_codeobj(code_object_id_t id)
ThreadTracerQueue::unload_codeobj(code_object_id_t id)
{
std::unique_lock<std::mutex> lk(trace_resources_mut);
if(auto* pkt = static_cast<hsa::TraceControlAQLPacket*>(cached_resources.get()))
{
pkt->remove_codeobj(id);
return;
}
control_packet->remove_codeobj(id);
remaining_codeobj_record.push_back({id, 0, 0, true});
if(!queue) return;
if(!queue || active_traces.load() < 1) return;
auto packet = factory->construct_unload_marker_packet(id);
bool bSuccess = Submit(&packet->packet);
@@ -289,9 +216,9 @@ AgentThreadTracer::unload_codeobj(code_object_id_t id)
// TODO: make this a wrapper on HSA load instead of registering
void
GlobalThreadTracer::codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* /* user_data */,
void* callback_data)
DispatchThreadTracer::codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* /* user_data */,
void* callback_data)
{
if(!callback_data) return;
if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return;
@@ -300,26 +227,11 @@ GlobalThreadTracer::codeobj_tracing_callback(rocprofiler_callback_tracing_record
auto* rec = static_cast<rocprofiler_callback_tracing_code_object_load_data_t*>(record.payload);
assert(rec);
GlobalThreadTracer& tracer = *static_cast<GlobalThreadTracer*>(callback_data);
auto agent = rec->hsa_agent;
DispatchThreadTracer& tracer = *static_cast<DispatchThreadTracer*>(callback_data);
auto agent = rec->hsa_agent;
std::shared_lock<std::shared_mutex> lk(tracer.agents_map_mut);
if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
{
try
{
tracer.loaded_codeobjs.at(rec->hsa_agent).erase(rec->code_object_id);
} catch(std::exception& e)
{
ROCP_WARNING << "Codeobj unload called for invalid ID " << rec->code_object_id;
}
}
else
{
tracer.loaded_codeobjs[agent][rec->code_object_id] = {rec->load_delta, rec->load_size};
}
auto tracer_it = tracer.agents.find(agent);
if(tracer_it == tracer.agents.end()) return;
@@ -330,9 +242,9 @@ GlobalThreadTracer::codeobj_tracing_callback(rocprofiler_callback_tracing_record
}
void
GlobalThreadTracer::resource_init(const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
DispatchThreadTracer::resource_init(const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
{
auto agent = cache.get_hsa_agent();
std::unique_lock<std::shared_mutex> lk(agents_map_mut);
@@ -344,13 +256,12 @@ GlobalThreadTracer::resource_init(const hsa::AgentCache& cache,
return;
}
auto new_tracer = std::make_unique<AgentThreadTracer>(this->params, cache, coreapi, ext);
new_tracer->active_queues.store(1);
auto new_tracer = std::make_unique<ThreadTracerQueue>(this->params, cache, coreapi, ext);
agents.emplace(agent, std::move(new_tracer));
}
void
GlobalThreadTracer::resource_deinit(const hsa::AgentCache& cache)
DispatchThreadTracer::resource_deinit(const hsa::AgentCache& cache)
{
std::unique_lock<std::shared_mutex> lk(agents_map_mut);
@@ -367,9 +278,11 @@ GlobalThreadTracer::resource_deinit(const hsa::AgentCache& cache)
* We return an AQLPacket containing the start/stop/read packets for injection.
*/
std::unique_ptr<hsa::AQLPacket>
GlobalThreadTracer::pre_kernel_call(const hsa::Queue& queue,
rocprofiler_kernel_id_t kernel_id,
const context::correlation_id* corr_id)
DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue,
rocprofiler_kernel_id_t kernel_id,
rocprofiler_dispatch_id_t dispatch_id,
rocprofiler_user_data_t* user_data,
const context::correlation_id* corr_id)
{
rocprofiler_correlation_id_t rocprof_corr_id =
rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data};
@@ -377,27 +290,71 @@ GlobalThreadTracer::pre_kernel_call(const hsa::Queue& queue,
if(corr_id) rocprof_corr_id.internal = corr_id->internal;
// TODO: Get external
// Maybe adds serialization packets to the AQLPacket (if serializer is enabled)
// and maybe adds barrier packets if the state is transitioning from serialized <->
// unserialized
auto maybe_add_serialization = [&](auto& gen_pkt) {
CHECK_NOTNULL(hsa::get_queue_controller())->serializer().rlock([&](const auto& serializer) {
for(auto& s_pkt : serializer.kernel_dispatch(queue))
gen_pkt->before_krn_pkt.push_back(s_pkt.ext_amd_aql_pm4);
});
};
auto control_flags = params.dispatch_cb_fn(queue.get_id(),
queue.get_agent().get_rocp_agent(),
rocprof_corr_id,
kernel_id,
dispatch_id,
user_data,
params.callback_userdata);
if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) return nullptr;
if(control_flags == ROCPROFILER_ATT_CONTROL_NONE)
{
auto empty = std::make_unique<hsa::EmptyAQLPacket>();
maybe_add_serialization(empty);
return empty;
}
std::shared_lock<std::shared_mutex> lk(agents_map_mut);
auto it = agents.find(queue.get_agent().get_hsa_agent());
assert(it != agents.end() && it->second != nullptr);
auto packet = it->second->pre_kernel_call(control_flags, queue.get_id(), rocprof_corr_id);
if(packet != nullptr) post_move_data.fetch_add(1);
auto packet = it->second->get_control(bool(control_flags & ROCPROFILER_ATT_CONTROL_START));
post_move_data.fetch_add(1);
maybe_add_serialization(packet);
if((control_flags & ROCPROFILER_ATT_CONTROL_START) != 0) packet->populate_before();
if((control_flags & ROCPROFILER_ATT_CONTROL_STOP) != 0) packet->populate_after();
return packet;
}
void
GlobalThreadTracer::post_kernel_call(GlobalThreadTracer::inst_pkt_t& aql)
class SignalSerializerExit
{
public:
SignalSerializerExit(const hsa::Queue::queue_info_session_t& _session)
: session(_session)
{}
~SignalSerializerExit()
{
auto* controller = hsa::get_queue_controller();
if(!controller) return;
controller->serializer().wlock(
[&](auto& serializer) { serializer.kernel_completion_signal(session.queue); });
}
const hsa::Queue::queue_info_session_t& session;
};
void
DispatchThreadTracer::post_kernel_call(DispatchThreadTracer::inst_pkt_t& aql,
const hsa::Queue::queue_info_session_t& session)
{
SignalSerializerExit signal(session);
if(post_move_data.load() < 1) return;
for(auto& aql_pkt : aql)
@@ -408,21 +365,26 @@ GlobalThreadTracer::post_kernel_call(GlobalThreadTracer::inst_pkt_t& aql)
std::shared_lock<std::shared_mutex> lk(agents_map_mut);
post_move_data.fetch_sub(1);
if(pkt->after_krn_pkt.empty()) continue;
auto it = agents.find(pkt->GetAgent());
if(it != agents.end() && it->second != nullptr)
it->second->post_kernel_call(std::move(aql_pkt.first));
it->second->iterate_data(pkt->GetHandle(), session.user_data);
}
}
void
GlobalThreadTracer::start_context()
DispatchThreadTracer::start_context()
{
using corr_id_map_t = hsa::Queue::queue_info_session_t::external_corr_id_map_t;
if(codeobj_client_ctx.handle != 0)
{
auto status = rocprofiler_start_context(codeobj_client_ctx);
if(status != ROCPROFILER_STATUS_SUCCESS) throw std::exception();
}
CHECK_NOTNULL(hsa::get_queue_controller())->enable_serialization();
// Only one thread should be attempting to enable/disable this context
client.wlock([&](auto& client_id) {
if(client_id) return;
@@ -431,22 +393,22 @@ GlobalThreadTracer::start_context()
std::nullopt,
[=](const hsa::Queue& q,
const hsa::rocprofiler_packet& /* kern_pkt */,
rocprofiler_kernel_id_t kernel_id,
rocprofiler_dispatch_id_t /* dispatch_id */,
rocprofiler_user_data_t* /* user_data */,
rocprofiler_kernel_id_t kernel_id,
rocprofiler_dispatch_id_t dispatch_id,
rocprofiler_user_data_t* user_data,
const corr_id_map_t& /* extern_corr_ids */,
const context::correlation_id* corr_id) {
return this->pre_kernel_call(q, kernel_id, corr_id);
return this->pre_kernel_call(q, kernel_id, dispatch_id, user_data, corr_id);
},
[=](const hsa::Queue& /* q */,
hsa::rocprofiler_packet /* kern_pkt */,
const hsa::Queue::queue_info_session_t& /* session */,
inst_pkt_t& aql) { this->post_kernel_call(aql); });
const hsa::Queue::queue_info_session_t& session,
inst_pkt_t& aql) { this->post_kernel_call(aql, session); });
});
}
void
GlobalThreadTracer::stop_context()
DispatchThreadTracer::stop_context()
{
client.wlock([&](auto& client_id) {
if(!client_id) return;
@@ -455,6 +417,94 @@ GlobalThreadTracer::stop_context()
hsa::get_queue_controller()->remove_callback(*client_id);
client_id = std::nullopt;
});
auto* controller = hsa::get_queue_controller();
if(controller) controller->disable_serialization();
}
void
AgentThreadTracer::resource_init(const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
{
if(cache.get_rocp_agent()->id != this->agent_id) return;
std::unique_lock<std::mutex> lk(mut);
if(tracer != nullptr)
{
tracer->active_queues.fetch_add(1);
return;
}
tracer = std::make_unique<ThreadTracerQueue>(this->params, cache, coreapi, ext);
}
void
AgentThreadTracer::resource_deinit(const hsa::AgentCache& cache)
{
if(cache.get_rocp_agent()->id != this->agent_id) return;
std::unique_lock<std::mutex> lk(mut);
if(tracer == nullptr) return;
if(tracer->active_queues.fetch_sub(1) == 1) tracer.reset();
}
void
AgentThreadTracer::start_context()
{
std::unique_lock<std::mutex> lk(mut);
if(tracer == nullptr)
{
ROCP_FATAL << "Thread trace context not present for agent!";
return;
}
auto packet = tracer->get_control(true);
packet->populate_before();
for(auto& start : packet->before_krn_pkt)
tracer->Submit(&start);
}
void
AgentThreadTracer::stop_context()
{
std::unique_lock<std::mutex> lk(mut);
auto packet = tracer->get_control(false);
packet->populate_after();
for(auto& stop : packet->after_krn_pkt)
tracer->Submit(&stop);
rocprofiler_user_data_t userdata{.ptr = params.callback_userdata};
tracer->iterate_data(packet->GetHandle(), userdata);
}
void
AgentThreadTracer::codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* /* user_data */,
void* callback_data)
{
if(!callback_data) return;
if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return;
if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return;
auto* rec = static_cast<rocprofiler_callback_tracing_code_object_load_data_t*>(record.payload);
assert(rec);
AgentThreadTracer& tracer = *static_cast<AgentThreadTracer*>(callback_data);
std::unique_lock<std::mutex> lk(tracer.mut);
if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD)
tracer.tracer->load_codeobj(rec->code_object_id, rec->load_delta, rec->load_size);
else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
tracer.tracer->unload_codeobj(rec->code_object_id);
}
} // namespace thread_trace
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp
+83 -62
Προβολή Αρχείου
@@ -40,12 +40,19 @@
namespace rocprofiler
{
namespace hsa
{
class AQLPacket;
};
namespace thread_trace
{
struct thread_trace_parameter_pack
{
rocprofiler_context_id_t context_id;
rocprofiler_att_dispatch_callback_t dispatch_cb_fn;
rocprofiler_att_shader_data_callback_t shader_cb_fn;
void* callback_userdata;
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};
// Parameters
uint8_t target_cu = 1;
@@ -64,85 +71,69 @@ struct thread_trace_parameter_pack
static constexpr size_t PERFCOUNTER_SIMD_MASK_SHIFT = 28;
};
namespace hsa
{
class AQLPacket;
};
struct ThreadTraceActiveResource
{
rocprofiler_correlation_id_t corr_id;
rocprofiler_queue_id_t queue_id;
std::unique_ptr<hsa::AQLPacket> packet{nullptr};
};
class AgentThreadTracer
class ThreadTracerQueue
{
using code_object_id_t = uint64_t;
struct CodeobjRecord
{
code_object_id_t id;
uint64_t addr;
uint64_t size;
bool bUnload;
};
public:
AgentThreadTracer(thread_trace_parameter_pack _params,
ThreadTracerQueue(thread_trace_parameter_pack _params,
const hsa::AgentCache&,
const CoreApiTable&,
const AmdExtTable&);
virtual ~AgentThreadTracer();
virtual ~ThreadTracerQueue();
void load_codeobj(code_object_id_t id, uint64_t addr, uint64_t size);
void unload_codeobj(code_object_id_t id);
std::unique_ptr<hsa::AQLPacket> pre_kernel_call(rocprofiler_att_control_flags_t control_flags,
rocprofiler_queue_id_t queue_id,
rocprofiler_correlation_id_t corr_id);
std::unique_ptr<hsa::TraceControlAQLPacket> get_control(bool bStart);
void iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data);
void post_kernel_call(std::unique_ptr<hsa::AQLPacket>&& aql);
hsa_queue_t* queue = nullptr;
std::mutex trace_resources_mut;
thread_trace_parameter_pack params;
std::unique_ptr<hsa::AQLPacket> cached_resources;
ThreadTraceActiveResource active_resources;
std::atomic<int> data_is_ready{0};
std::atomic<int> active_queues{1};
std::vector<CodeobjRecord> remaining_codeobj_record;
hsa_queue_t* queue = nullptr;
std::mutex trace_resources_mut;
thread_trace_parameter_pack params;
std::atomic<int> active_traces{0};
std::atomic<int> active_queues{1};
std::unique_ptr<hsa::TraceControlAQLPacket> control_packet;
std::unique_ptr<aql::ThreadTraceAQLPacketFactory> factory;
private:
bool Submit(hsa_ext_amd_aql_pm4_packet_t* packet);
private:
decltype(hsa_queue_load_read_index_relaxed)* load_read_index_relaxed_fn{nullptr};
decltype(hsa_queue_add_write_index_relaxed)* add_write_index_relaxed_fn{nullptr};
decltype(hsa_signal_store_screlease)* signal_store_screlease_fn{nullptr};
decltype(hsa_queue_destroy)* queue_destroy_fn{nullptr};
}; // namespace thread_trace
};
class GlobalThreadTracer
class ThreadTracerInterface
{
struct CodeobjAddrRange
{
int64_t addr;
uint64_t size;
};
public:
ThreadTracerInterface() = default;
virtual ~ThreadTracerInterface() = default;
virtual void start_context() = 0;
virtual void stop_context() = 0;
virtual void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) = 0;
virtual void resource_deinit(const hsa::AgentCache&) = 0;
};
class DispatchThreadTracer : public ThreadTracerInterface
{
using code_object_id_t = uint64_t;
using AQLPacketPtr = std::unique_ptr<hsa::AQLPacket>;
using inst_pkt_t = common::container::small_vector<std::pair<AQLPacketPtr, int64_t>, 4>;
using corr_id_map_t = hsa::Queue::queue_info_session_t::external_corr_id_map_t;
using code_object_id_t = uint64_t;
public:
GlobalThreadTracer(thread_trace_parameter_pack _params)
: params(std::move(_params)){};
virtual void start_context();
virtual void stop_context();
virtual void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&);
virtual void resource_deinit(const hsa::AgentCache&);
virtual ~GlobalThreadTracer() = default;
DispatchThreadTracer(thread_trace_parameter_pack _params)
: params(std::move(_params))
{}
~DispatchThreadTracer() override = default;
void start_context() override;
void stop_context() override;
void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) override;
void resource_deinit(const hsa::AgentCache&) override;
static void codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* user_data,
@@ -150,17 +141,47 @@ public:
std::unique_ptr<hsa::AQLPacket> pre_kernel_call(const hsa::Queue& queue,
uint64_t kernel_id,
rocprofiler_dispatch_id_t dispatch_id,
rocprofiler_user_data_t* user_data,
const context::correlation_id* corr_id);
void post_kernel_call(inst_pkt_t& aql);
void post_kernel_call(inst_pkt_t& aql, const hsa::Queue::queue_info_session_t& session);
std::map<hsa_agent_t, std::map<code_object_id_t, CodeobjAddrRange>> loaded_codeobjs;
std::unordered_map<hsa_agent_t, std::unique_ptr<AgentThreadTracer>> agents;
std::unordered_map<hsa_agent_t, std::unique_ptr<ThreadTracerQueue>> agents;
std::shared_mutex agents_map_mut;
std::atomic<int> post_move_data{0};
std::atomic<int> post_move_data{0};
std::shared_mutex agents_map_mut;
rocprofiler_context_id_t codeobj_client_ctx{0};
thread_trace_parameter_pack params;
rocprofiler_context_id_t codeobj_client_ctx{0};
};
class AgentThreadTracer : public ThreadTracerInterface
{
public:
AgentThreadTracer(thread_trace_parameter_pack _params, rocprofiler_agent_id_t _id)
: agent_id(_id)
, params(std::move(_params))
{}
~AgentThreadTracer() override = default;
void start_context() override;
void stop_context() override;
void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) override;
void resource_deinit(const hsa::AgentCache&) override;
static void codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* user_data,
void* callback_data);
rocprofiler_agent_id_t agent_id;
std::mutex mut;
std::unique_ptr<ThreadTracerQueue> tracer{nullptr};
thread_trace_parameter_pack params;
rocprofiler_context_id_t codeobj_client_ctx{0};
};
}; // namespace thread_trace
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp
+84 -22
Προβολή Αρχείου
@@ -37,20 +37,23 @@ get_mask(const rocprofiler::counters::Metric* metric, uint64_t simds_selected)
{
uint32_t mask = std::atoi(metric->event().c_str());
if(simds_selected == 0)
simds_selected = rocprofiler::thread_trace_parameter_pack::DEFAULT_PERFCOUNTER_SIMD_MASK;
mask |= simds_selected << rocprofiler::thread_trace_parameter_pack::PERFCOUNTER_SIMD_MASK_SHIFT;
simds_selected =
rocprofiler::thread_trace::thread_trace_parameter_pack::DEFAULT_PERFCOUNTER_SIMD_MASK;
mask |= simds_selected
<< rocprofiler::thread_trace::thread_trace_parameter_pack::PERFCOUNTER_SIMD_MASK_SHIFT;
return mask;
}
} // namespace
extern "C" {
rocprofiler_status_t ROCPROFILER_API
rocprofiler_configure_thread_trace_service(rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata)
rocprofiler_configure_dispatch_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata)
{
if(rocprofiler::registration::get_init_status() > -1)
return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED;
@@ -59,13 +62,12 @@ rocprofiler_configure_thread_trace_service(rocprofiler_context_id_t
if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED;
if(ctx->thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED;
auto param_pack = rocprofiler::thread_trace_parameter_pack{};
auto param_pack = rocprofiler::thread_trace::thread_trace_parameter_pack{};
param_pack.context_id = context_id;
param_pack.dispatch_cb_fn = dispatch_callback;
param_pack.shader_cb_fn = shader_callback;
param_pack.callback_userdata = callback_userdata;
bool bEnableCodeobj = false;
const auto& id_map = *CHECK_NOTNULL(rocprofiler::counters::getPerfCountersIdMap());
for(size_t p = 0; p < num_parameters; p++)
@@ -82,9 +84,6 @@ rocprofiler_configure_thread_trace_service(rocprofiler_context_id_t
break;
case ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE: param_pack.buffer_size = param.value; break;
case ROCPROFILER_ATT_PARAMETER_SIMD_SELECT: param_pack.simd_select = param.value; break;
case ROCPROFILER_ATT_PARAMETER_CODE_OBJECT_TRACE_ENABLE:
bEnableCodeobj = param.value != 0;
break;
case ROCPROFILER_ATT_PARAMETER_PERFCOUNTER:
if(const auto* metric_ptr =
rocprofiler::common::get_val(id_map, param.counter_id.handle))
@@ -97,23 +96,86 @@ rocprofiler_configure_thread_trace_service(rocprofiler_context_id_t
}
}
ctx->thread_trace = std::make_shared<rocprofiler::GlobalThreadTracer>(param_pack);
auto tracer = std::make_unique<rocprofiler::thread_trace::DispatchThreadTracer>(param_pack);
if(!bEnableCodeobj) return ROCPROFILER_STATUS_SUCCESS; // Skip TRACING_CODE_OBJECT setup
auto& client_ctx = ctx->thread_trace->codeobj_client_ctx;
rocprofiler_status_t status = rocprofiler_create_context(&client_ctx);
rocprofiler_status_t status = rocprofiler_create_context(&tracer->codeobj_client_ctx);
if(status != ROCPROFILER_STATUS_SUCCESS) return status;
status = rocprofiler_configure_callback_tracing_service(
client_ctx,
tracer->codeobj_client_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
nullptr,
0,
rocprofiler::GlobalThreadTracer::codeobj_tracing_callback,
ctx->thread_trace.get());
rocprofiler::thread_trace::DispatchThreadTracer::codeobj_tracing_callback,
tracer.get());
ctx->thread_trace = std::move(tracer);
return status;
}
rocprofiler_status_t ROCPROFILER_API
rocprofiler_configure_agent_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_agent_id_t agent,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_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->thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED;
auto param_pack = rocprofiler::thread_trace::thread_trace_parameter_pack{};
param_pack.context_id = context_id;
param_pack.shader_cb_fn = shader_callback;
param_pack.callback_userdata = callback_userdata;
const auto& id_map = *CHECK_NOTNULL(rocprofiler::counters::getPerfCountersIdMap());
for(size_t p = 0; p < num_parameters; p++)
{
const rocprofiler_att_parameter_t& param = parameters[p];
if(param.type > ROCPROFILER_ATT_PARAMETER_LAST)
return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
switch(param.type)
{
case ROCPROFILER_ATT_PARAMETER_TARGET_CU: param_pack.target_cu = param.value; break;
case ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK:
param_pack.shader_engine_mask = param.value;
break;
case ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE: param_pack.buffer_size = param.value; break;
case ROCPROFILER_ATT_PARAMETER_SIMD_SELECT: param_pack.simd_select = param.value; break;
case ROCPROFILER_ATT_PARAMETER_PERFCOUNTER:
if(const auto* metric_ptr =
rocprofiler::common::get_val(id_map, param.counter_id.handle))
param_pack.perfcounters.push_back(get_mask(metric_ptr, param.simd_mask));
break;
case ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL:
param_pack.perfcounter_ctrl = param.value;
break;
case ROCPROFILER_ATT_PARAMETER_LAST: return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
}
}
auto tracer = std::make_unique<rocprofiler::thread_trace::AgentThreadTracer>(param_pack, agent);
rocprofiler_status_t status = rocprofiler_create_context(&tracer->codeobj_client_ctx);
if(status != ROCPROFILER_STATUS_SUCCESS) return status;
status = rocprofiler_configure_callback_tracing_service(
tracer->codeobj_client_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
nullptr,
0,
rocprofiler::thread_trace::AgentThreadTracer::codeobj_tracing_callback,
tracer.get());
ctx->thread_trace = std::move(tracer);
return status;
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp
+22 -17
Προβολή Αρχείου
@@ -86,11 +86,11 @@ TEST(thread_trace, resource_creation)
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
auto params = thread_trace_parameter_pack{};
auto params = thread_trace::thread_trace_parameter_pack{};
aql::ThreadTraceAQLPacketFactory factory(agent, params, get_api_table(), get_ext_table());
auto packet = factory.construct_packet();
auto packet = factory.construct_control_packet();
packet->populate_before();
packet->populate_after();
@@ -102,8 +102,8 @@ TEST(thread_trace, resource_creation)
}
{
thread_trace_parameter_pack params{};
GlobalThreadTracer tracer(std::move(params));
thread_trace::thread_trace_parameter_pack params{};
thread_trace::DispatchThreadTracer tracer(std::move(params));
for(const auto& [_, agent] : agents)
{
@@ -145,7 +145,7 @@ TEST(thread_trace, configure_test)
params.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {0x1000000}});
params.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {0xF}});
rocprofiler_configure_thread_trace_service(
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
params.data(),
params.size(),
@@ -153,8 +153,10 @@ TEST(thread_trace, configure_test)
const rocprofiler_agent_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, void*) {},
[](int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
@@ -194,7 +196,7 @@ TEST(thread_trace, perfcounters_configure_test)
expected.insert(std::atoi(metric.event().c_str()) | (simd_mask << 28));
}
rocprofiler_configure_thread_trace_service(
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
params.data(),
params.size(),
@@ -202,16 +204,19 @@ TEST(thread_trace, perfcounters_configure_test)
const rocprofiler_agent_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, void*) {},
[](int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
auto* context = rocprofiler::context::get_mutable_registered_context(ctx);
thread_trace_parameter_pack _params = context->thread_trace->params;
auto* context = rocprofiler::context::get_mutable_registered_context(ctx);
auto* tracer = dynamic_cast<thread_trace::DispatchThreadTracer*>(context->thread_trace.get());
ASSERT_EQ(_params.perfcounter_ctrl, 1);
ASSERT_EQ(_params.perfcounters.size(), 3);
for(uint32_t param : _params.perfcounters)
ASSERT_NE(tracer, nullptr);
ASSERT_EQ(tracer->params.perfcounter_ctrl, 1);
ASSERT_EQ(tracer->params.perfcounters.size(), 3);
for(uint32_t param : tracer->params.perfcounters)
EXPECT_TRUE(expected.find(param) != expected.end())
<< "valid AQLprofile mask not generated for perfcounters";
context::pop_client(1);
@@ -230,8 +235,8 @@ TEST(thread_trace, perfcounters_aql_options_test)
const std::uint8_t sqtt_default_num_options = 5;
auto agents = hsa::get_queue_controller()->get_supported_agents();
thread_trace_parameter_pack _params = {};
auto metrics = rocprofiler::counters::getMetricsForAgent("gfx90a");
thread_trace::thread_trace_parameter_pack _params = {};
auto metrics = rocprofiler::counters::getMetricsForAgent("gfx90a");
std::vector<std::pair<std::string, uint64_t>> perf_counters = {
{"SQ_WAVES", 0x1}, {"SQ_WAVES", 0x2}, {"GRBM_COUNT", 0x3}};
for(auto& [counter_name, simd_mask] : perf_counters)
@@ -240,10 +245,10 @@ TEST(thread_trace, perfcounters_aql_options_test)
_params.perfcounters.push_back(std::atoi(metric.event().c_str()) |
(simd_mask << 28));
_params.perfcounter_ctrl = 2;
auto new_tracer = std::make_unique<AgentThreadTracer>(
auto new_tracer = std::make_unique<thread_trace::ThreadTracerQueue>(
_params, begin(agents)->second, get_api_table(), get_ext_table());
ASSERT_EQ(new_tracer->factory->get_aql_params().size(),
ASSERT_EQ(new_tracer->factory->aql_params.size(),
sqtt_default_num_options + perf_counters.size());
context::pop_client(1);
hsa_shut_down();
projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt
+23
Προβολή Αρχείου
@@ -108,6 +108,7 @@ 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
@@ -155,3 +156,25 @@ set_tests_properties(
thread-trace-api-multi-test
PROPERTIES TIMEOUT 10 LABELS "integration-tests" ENVIRONMENT "${PRELOAD_ENV}"
FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}")
# Agent profiling test
add_executable(thread-trace-api-agent-test)
target_sources(thread-trace-api-agent-test PRIVATE agent_test.cpp)
target_link_libraries(thread-trace-api-agent-test
PRIVATE rocprofiler-sdk::rocprofiler-sdk)
if(ROCPROFILER_MEMCHECK_PRELOAD_ENV)
set(PRELOAD_ENV
"${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$<TARGET_FILE:thread-trace-api-agent-test>")
else()
set(PRELOAD_ENV "LD_PRELOAD=$<TARGET_FILE:thread-trace-api-agent-test>")
endif()
add_test(NAME thread-trace-api-agent-test
COMMAND $<TARGET_FILE:thread-trace-api-agent-test>)
set_tests_properties(
thread-trace-api-agent-test
PROPERTIES TIMEOUT 10 LABELS "integration-tests" ENVIRONMENT "${PRELOAD_ENV}"
FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}")
projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp
+168
Προβολή Αρχείου
@@ -0,0 +1,168 @@
// 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 = {};
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/common.hpp
+4 -1
Προβολή Αρχείου
@@ -87,7 +87,10 @@ tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
void* callback_data);
void
shader_data_callback(int64_t se_id, void* se_data, size_t data_size, void* userdata);
shader_data_callback(int64_t se_id,
void* se_data,
size_t data_size,
rocprofiler_user_data_t userdata);
void
callbacks_init();
projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp
+13 -10
Προβολή Αρχείου
@@ -55,11 +55,14 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */,
const rocprofiler_agent_t* /* agent */,
rocprofiler_correlation_id_t /* correlation_id */,
rocprofiler_kernel_id_t kernel_id,
void* userdata)
rocprofiler_dispatch_id_t /* dispatch_id */,
rocprofiler_user_data_t* dispatch_userdata,
void* userdata)
{
C_API_BEGIN
assert(userdata && "Dispatch callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
dispatch_userdata->ptr = userdata;
static std::atomic<int> call_id{0};
static std::string_view desired_func_name = "branching_kernel";
@@ -102,15 +105,15 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data)
"code object tracing service configure");
std::vector<rocprofiler_att_parameter_t> params{};
params.push_back({ROCPROFILER_ATT_PARAMETER_CODE_OBJECT_TRACE_ENABLE, {1}});
ROCPROFILER_CALL(rocprofiler_configure_thread_trace_service(client_ctx,
params.data(),
params.size(),
dispatch_callback,
Callbacks::shader_data_callback,
tool_data),
"thread trace service configure");
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
params.data(),
params.size(),
dispatch_callback,
Callbacks::shader_data_callback,
tool_data),
"thread trace service configure");
int valid_ctx = 0;
ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),
projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp
+7 -4
Προβολή Αρχείου
@@ -56,11 +56,14 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */,
const rocprofiler_agent_t* /* agent */,
rocprofiler_correlation_id_t /* correlation_id */,
rocprofiler_kernel_id_t kernel_id,
void* userdata)
rocprofiler_dispatch_id_t /* dispatch_id */,
rocprofiler_user_data_t* dispatch_userdata,
void* userdata)
{
C_API_BEGIN
assert(userdata && "Dispatch callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
dispatch_userdata->ptr = userdata;
static std::atomic<int> call_id{0};
static std::string_view desired_func_name = "branching_kernel";
@@ -71,7 +74,7 @@ dispatch_callback(rocprofiler_queue_id_t /* queue_id */,
if(kernel_name.find(desired_func_name) == std::string::npos)
return ROCPROFILER_ATT_CONTROL_NONE;
if(call_id.fetch_add(1) == 0) return ROCPROFILER_ATT_CONTROL_START_AND_STOP;
return ROCPROFILER_ATT_CONTROL_START_AND_STOP;
} catch(...)
{
std::cerr << "Could not find kernel id: " << kernel_id << std::endl;
@@ -99,7 +102,7 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data)
"code object tracing service configure");
ROCPROFILER_CALL(
rocprofiler_configure_thread_trace_service(
rocprofiler_configure_dispatch_thread_trace_service(
client_ctx, nullptr, 0, dispatch_callback, Callbacks::shader_data_callback, tool_data),
"thread trace service configure");
projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp
+6 -3
Προβολή Αρχείου
@@ -210,11 +210,14 @@ isa_callback(char* isa_instruction,
}
void
shader_data_callback(int64_t se_id, void* se_data, size_t data_size, void* userdata)
shader_data_callback(int64_t se_id,
void* se_data,
size_t data_size,
rocprofiler_user_data_t userdata)
{
C_API_BEGIN
assert(userdata && "Shader callback passed null!");
ToolData& tool = *reinterpret_cast<ToolData*>(userdata);
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);