From 37b0cfd7dc16c9df56dc0e10e4b2450dafbe684a Mon Sep 17 00:00:00 2001 From: Giovanni Lenzi Baraldi Date: Thu, 13 Jun 2024 15:29:29 -0300 Subject: [PATCH] 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 Co-authored-by: Vladimir Indic <139573562+vlaindic@users.noreply.github.com> [ROCm/rocprofiler-sdk commit: 9676295d3d118327dd68b7e405fd057fc6921256] --- .../samples/advanced_thread_trace/client.cpp | 151 ++++---- .../rocprofiler-sdk/amd_detail/CMakeLists.txt | 3 +- .../rocprofiler-sdk/amd_detail/thread_trace.h | 186 +-------- .../amd_detail/thread_trace_agent.h | 64 ++++ .../amd_detail/thread_trace_core.h | 171 +++++++++ .../amd_detail/thread_trace_dispatch.h | 82 ++++ .../lib/rocprofiler-sdk/aql/aql_profile_v2.h | 21 +- .../rocprofiler-sdk/aql/packet_construct.cpp | 33 +- .../rocprofiler-sdk/aql/packet_construct.hpp | 18 +- .../lib/rocprofiler-sdk/context/context.hpp | 4 +- .../lib/rocprofiler-sdk/hsa/aql_packet.cpp | 43 ++- .../lib/rocprofiler-sdk/hsa/aql_packet.hpp | 88 +++-- .../rocprofiler-sdk/hsa/queue_controller.cpp | 38 +- .../rocprofiler-sdk/hsa/queue_controller.hpp | 4 - .../rocprofiler-sdk/thread_trace/att_core.cpp | 354 ++++++++++-------- .../rocprofiler-sdk/thread_trace/att_core.hpp | 145 ++++--- .../thread_trace/att_service.cpp | 106 ++++-- .../thread_trace/tests/att_packet_test.cpp | 39 +- .../tests/thread-trace/CMakeLists.txt | 23 ++ .../tests/thread-trace/agent_test.cpp | 168 +++++++++ .../tests/thread-trace/common.hpp | 5 +- .../tests/thread-trace/multi_dispatch.cpp | 23 +- .../tests/thread-trace/single_dispatch.cpp | 11 +- .../tests/thread-trace/trace_callbacks.cpp | 9 +- 24 files changed, 1154 insertions(+), 635 deletions(-) create mode 100644 projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h create mode 100644 projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h create mode 100644 projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h create mode 100644 projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp diff --git a/projects/rocprofiler-sdk/samples/advanced_thread_trace/client.cpp b/projects/rocprofiler-sdk/samples/advanced_thread_trace/client.cpp index 22c0d0d0bc..c93d19e50a 100644 --- a/projects/rocprofiler-sdk/samples/advanced_thread_trace/client.cpp +++ b/projects/rocprofiler-sdk/samples/advanced_thread_trace/client.cpp @@ -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(callback_data); - if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) { - std::unique_lock lg(tool.isa_map_mut); + std::unique_lock lg(tool->isa_map_mut); auto* data = static_cast(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 lg(tool.isa_map_mut); + std::unique_lock lg(tool->isa_map_mut); auto* data = static_cast(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(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(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(userdata); - std::shared_lock lg(tool.isa_map_mut); + std::shared_lock lg(tool->isa_map_mut); static std::atomic 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(userdata); - assert(trace_data.tool && "ISA callback passed null!"); - ToolData& tool = *reinterpret_cast(trace_data.tool); - std::shared_lock shared_lock(tool.isa_map_mut); + std::shared_lock 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(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 unique_lock(tool.isa_map_mut); + std::unique_lock unique_lock(tool->isa_map_mut); auto ptr = std::make_unique(); 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(userdata); - assert(trace_data.tool && "ISA callback passed null!"); - ToolData& tool = *reinterpret_cast(trace_data.tool); std::shared_ptr instruction; { - std::unique_lock unique_lock(tool.isa_map_mut); - instruction = tool.codeobjTranslate.get(marker_id, offset); + std::unique_lock 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(); 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(userdata); { - std::unique_lock lk(tool.output_mut); - tool.output() << "SE ID: " << se_id << " with size " << data_size << std::hex << '\n'; + std::unique_lock 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(tool_data); - - std::unique_lock isa_lk(tool.isa_map_mut); - std::unique_lock out_lk(tool.output_mut); + std::unique_lock isa_lk(client::tool->isa_map_mut); + std::unique_lock 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(data)}; + nullptr}; // return pointer to configure data return &cfg; diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/CMakeLists.txt b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/CMakeLists.txt index 8bfe7fe5cb..1e0bdd706a 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/CMakeLists.txt @@ -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} diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace.h index 4ece91300c..e089bc5e3a 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace.h @@ -22,186 +22,6 @@ #pragma once -#include -#include -#include -#include -#include - -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 +#include +#include diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h new file mode 100644 index 0000000000..f1db19dfbd --- /dev/null +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h @@ -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 +#include +#include +#include +#include + +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 diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h new file mode 100644 index 0000000000..21265ee43f --- /dev/null +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h @@ -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 +#include +#include +#include + +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 diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h new file mode 100644 index 0000000000..686a77d1dd --- /dev/null +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -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 +#include +#include +#include +#include + +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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h index 6de750beb2..4692072e4f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h @@ -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; /** diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index 4298edd776..67ba50ea0a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -219,31 +219,28 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache& uint32_t perf_ctrl = static_cast(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& -ThreadTraceAQLPacketFactory::get_aql_params() -{ - return aql_params; -} - std::unique_ptr -ThreadTraceAQLPacketFactory::construct_packet() +ThreadTraceAQLPacketFactory::construct_control_packet() { - uint32_t num_params = static_cast(get_aql_params().size()); - auto profile = - aqlprofile_att_profile_t{tracepool.gpu_agent, get_aql_params().data(), num_params}; - auto packet = std::make_unique(this->tracepool, profile); + auto num_params = static_cast(aql_params.size()); + auto profile = aqlprofile_att_profile_t{tracepool.gpu_agent, aql_params.data(), num_params}; + auto packet = std::make_unique(this->tracepool, profile); packet->clear(); return packet; } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp index 9c1a1b50d9..ac31e9892a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp @@ -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& get_aql_params(); - std::unique_ptr construct_packet(); - std::unique_ptr construct_load_marker_packet(uint64_t id, - uint64_t addr, - uint64_t size); + + std::unique_ptr construct_control_packet(); + std::unique_ptr construct_load_marker_packet(uint64_t id, + uint64_t addr, + uint64_t size); + std::unique_ptr construct_unload_marker_packet(uint64_t id); + std::vector aql_params; + private: - hsa::TraceMemoryPool tracepool; - std::vector aql_params; + hsa::TraceMemoryPool tracepool; }; } // namespace aql diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp index a18749b985..80626a8cbe 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp @@ -138,8 +138,8 @@ struct context std::unique_ptr counter_collection = {}; std::unique_ptr agent_counter_collection = {}; std::unique_ptr pc_sampler = {}; - // TODO: Make a unique pointer instead - std::shared_ptr thread_trace = {}; + + std::unique_ptr thread_trace = {}; }; // set the client index needs to be called before allocate_context() diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp index dea20a56a2..e8a21ab4cc 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp @@ -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(data)->tracepool; + auto& pool = *reinterpret_cast(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(data)->tracepool; + auto& pool = *reinterpret_cast(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(data)->tracepool; + auto& pool = *reinterpret_cast(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(_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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp index 8136e87697..10004f12f8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp @@ -25,6 +25,8 @@ #include "lib/common/container/small_vector.hpp" #include "lib/rocprofiler-sdk/aql/aql_profile_v2.h" +#include + #include #include @@ -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& _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 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(tracepool, id, addr, size, true, false); + std::make_shared(*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 tracepool; + aqlprofile_att_control_aql_packets_t packets; -private: - aqlprofile_att_control_aql_packets_t packets; - std::unordered_map> loaded_codeobj; + std::unordered_map> loaded_codeobj; }; } // namespace hsa diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index f81c9ea7be..b3d284ce9d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -143,8 +143,11 @@ constexpr rocprofiler_agent_t default_agent = void QueueController::add_queue(hsa_queue_t* id, std::unique_ptr 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), + 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 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; } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp index d719390413..59a14f5c6c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp @@ -103,7 +103,6 @@ public: private: using client_id_map_t = std::unordered_map; using agent_cache_map_t = std::unordered_map; - using resource_alloc_t = void(const AgentCache&, const CoreApiTable&, const AmdExtTable&); CoreApiTable _core_table = {}; AmdExtTable _ext_table = {}; @@ -111,9 +110,6 @@ private: common::Synchronized _callback_cache = {}; agent_cache_map_t _supported_agents = {}; common::Synchronized _profiler_serializer; - - std::vector> pre_initialize; - std::vector> pre_deinitialize; }; QueueController* diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index 431059d83d..a5d1c9e7f9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -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> 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(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 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(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 -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 +ThreadTracerQueue::get_control(bool bStart) { - if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) return nullptr; - std::unique_lock 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(*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(callback_data)->tool_userdata; - auto callback_fn = *static_cast(callback_data)->cb_fn; + auto& cb_data = *static_cast(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&& aql) +ThreadTracerQueue::iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data) { - std::unique_lock lk(trace_resources_mut); - - active_resources.packet = std::move(aql); - - if(!active_resources.packet || data_is_ready.load() < 1) return; - auto* pkt = static_cast(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 lk(trace_resources_mut); - if(auto* pkt = static_cast(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 lk(trace_resources_mut); - if(auto* pkt = static_cast(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(record.payload); assert(rec); - GlobalThreadTracer& tracer = *static_cast(callback_data); - auto agent = rec->hsa_agent; + DispatchThreadTracer& tracer = *static_cast(callback_data); + auto agent = rec->hsa_agent; std::shared_lock 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 lk(agents_map_mut); @@ -344,13 +256,12 @@ GlobalThreadTracer::resource_init(const hsa::AgentCache& cache, return; } - auto new_tracer = std::make_unique(this->params, cache, coreapi, ext); - new_tracer->active_queues.store(1); + auto new_tracer = std::make_unique(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 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 -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(); + maybe_add_serialization(empty); + return empty; + } std::shared_lock 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 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 lk(mut); + + if(tracer != nullptr) + { + tracer->active_queues.fetch_add(1); + return; + } + + tracer = std::make_unique(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 lk(mut); + if(tracer == nullptr) return; + + if(tracer->active_queues.fetch_sub(1) == 1) tracer.reset(); +} + +void +AgentThreadTracer::start_context() +{ + std::unique_lock 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 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(record.payload); + assert(rec); + + AgentThreadTracer& tracer = *static_cast(callback_data); + std::unique_lock 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 diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp index 204caa910b..caee1fb45a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp @@ -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 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 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 get_control(bool bStart); + void iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data); - void post_kernel_call(std::unique_ptr&& aql); - - hsa_queue_t* queue = nullptr; - std::mutex trace_resources_mut; - thread_trace_parameter_pack params; - std::unique_ptr cached_resources; - ThreadTraceActiveResource active_resources; - std::atomic data_is_ready{0}; - std::atomic active_queues{1}; - std::vector remaining_codeobj_record; + hsa_queue_t* queue = nullptr; + std::mutex trace_resources_mut; + thread_trace_parameter_pack params; + std::atomic active_traces{0}; + std::atomic active_queues{1}; + std::unique_ptr control_packet; std::unique_ptr 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; using inst_pkt_t = common::container::small_vector, 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 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> loaded_codeobjs; - std::unordered_map> agents; + std::unordered_map> agents; + + std::shared_mutex agents_map_mut; + std::atomic post_move_data{0}; - std::atomic 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 tracer{nullptr}; + + thread_trace_parameter_pack params; + rocprofiler_context_id_t codeobj_client_ctx{0}; +}; + }; // namespace thread_trace } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp index dd6aea9deb..60a7fe9e49 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp @@ -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(param_pack); + auto tracer = std::make_unique(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(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; } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index fec6333ca9..e8f814cc2e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -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(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> 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( + auto new_tracer = std::make_unique( _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(); diff --git a/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt b/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt index 70d38730dd..7c86c14227 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt @@ -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}:$") +else() + set(PRELOAD_ENV "LD_PRELOAD=$") +endif() + +add_test(NAME thread-trace-api-agent-test + COMMAND $) + +set_tests_properties( + thread-trace-api-agent-test + PROPERTIES TIMEOUT 10 LABELS "integration-tests" ENVIRONMENT "${PRELOAD_ENV}" + FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") diff --git a/projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp b/projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp new file mode 100644 index 0000000000..aeb9a2d4c2 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/thread-trace/agent_test.cpp @@ -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 +#include +#include +#include "common.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#define HIP_API_CALL(CALL) assert((CALL) == hipSuccess) + +namespace ATTTest +{ +namespace Agent +{ +rocprofiler_context_id_t client_ctx = {}; +rocprofiler_client_id_t* client_id = nullptr; +std::atomic valid_data{false}; + +void +shader_data_callback(int64_t /* se_id */, + void* se_data, + size_t data_size, + rocprofiler_user_data_t /* userdata */) +{ + if(se_data && data_size) valid_data.store(true); +} + +rocprofiler_status_t +query_available_agents(rocprofiler_agent_version_t /* version */, + const void** agents, + size_t num_agents, + void* /* user_data */) +{ + for(size_t idx = 0; idx < num_agents; idx++) + { + const auto* agent = static_cast(agents[idx]); + if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; + + ROCPROFILER_CALL(rocprofiler_configure_agent_thread_trace_service( + client_ctx, nullptr, 0, agent->id, shader_data_callback, nullptr), + "thread trace service configure"); + + return ROCPROFILER_STATUS_SUCCESS; + } + return ROCPROFILER_STATUS_ERROR; +} + +int +tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) +{ + ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); + + ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, + query_available_agents, + sizeof(rocprofiler_agent_t), + nullptr), + ""); + + int valid = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid), "context validity check"); + return (valid == 0) ? -1 : 0; +} + +void +tool_fini(void* /* tool_data */) +{ + assert(valid_data.load()); +} + +} // namespace Agent +} // namespace ATTTest + +extern "C" rocprofiler_tool_configure_result_t* +rocprofiler_configure(uint32_t /* version */, + const char* /* runtime_version */, + uint32_t priority, + rocprofiler_client_id_t* id) +{ + // only activate if main tool + if(priority > 0) return nullptr; + + // set the client name + id->name = "ATT_test_agent_api"; + + // store client info + ATTTest::Agent::client_id = id; + + // create configure data + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Agent::tool_init, + &ATTTest::Agent::tool_fini, + nullptr}; + + // return pointer to configure data + return &cfg; +} + +void +run(int dev) +{ + constexpr size_t size = 0x1000; + float* ptr = nullptr; + + HIP_API_CALL(hipSetDevice(dev)); + + HIP_API_CALL(hipMalloc(&ptr, size * sizeof(float))); + HIP_API_CALL(hipMemset(ptr, 0x55, size * sizeof(float))); + HIP_API_CALL(hipFree(ptr)); +} + +int +main() +{ + int ndev = 0; + HIP_API_CALL(hipGetDeviceCount(&ndev)); + + for(int dev = 0; dev < ndev; dev++) + run(dev); + + ROCPROFILER_CALL(rocprofiler_start_context(ATTTest::Agent::client_ctx), "context start"); + + for(int dev = 0; dev < ndev; dev++) + run(dev); + usleep(100); + + ROCPROFILER_CALL(rocprofiler_stop_context(ATTTest::Agent::client_ctx), "context stop"); + return 0; +} diff --git a/projects/rocprofiler-sdk/tests/thread-trace/common.hpp b/projects/rocprofiler-sdk/tests/thread-trace/common.hpp index c16c0081a8..10d4e5fcde 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/common.hpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/common.hpp @@ -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(); diff --git a/projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp b/projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp index 0a49362997..6a5dc4577f 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp @@ -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(userdata); + ToolData& tool = *reinterpret_cast(userdata); + dispatch_userdata->ptr = userdata; static std::atomic 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 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), diff --git a/projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp b/projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp index 79a9e90ccb..a911e2d300 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/single_dispatch.cpp @@ -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(userdata); + ToolData& tool = *reinterpret_cast(userdata); + dispatch_userdata->ptr = userdata; static std::atomic 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"); diff --git a/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp b/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp index 0c0b56b018..bd0d5a296c 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp @@ -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(userdata); + assert(userdata.ptr && "Shader callback passed null!"); + ToolData& tool = *reinterpret_cast(userdata.ptr); trace_data_t data{.id = se_id, .data = (uint8_t*) se_data, .size = data_size, .tool = &tool}; auto status = rocprofiler_att_parse_data(copy_trace_data, get_trace_data, isa_callback, &data);