diff --git a/samples/advanced_thread_trace/CMakeLists.txt b/samples/advanced_thread_trace/CMakeLists.txt deleted file mode 100644 index 694bbabde2..0000000000 --- a/samples/advanced_thread_trace/CMakeLists.txt +++ /dev/null @@ -1,59 +0,0 @@ -# -# -# -cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) - -if(NOT CMAKE_HIP_COMPILER) - find_program( - amdclangpp_EXECUTABLE - NAMES amdclang++ - HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm - PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm - PATH_SUFFIXES bin llvm/bin NO_CACHE) - mark_as_advanced(amdclangpp_EXECUTABLE) - - if(amdclangpp_EXECUTABLE) - set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") - endif() -endif() - -project(rocprofiler-sdk-samples-advanced-thread-trace LANGUAGES CXX HIP) - -foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) - if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") - set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") - endif() -endforeach() - -find_package(rocprofiler-sdk REQUIRED) -find_package( - amd_comgr - REQUIRED - CONFIG - HINTS - ${rocm_version_DIR} - ${ROCM_PATH} - PATHS - ${rocm_version_DIR} - ${ROCM_PATH} - PATH_SUFFIXES - lib/cmake/amd_comgr) - -set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP) -set_source_files_properties(main.cpp PROPERTIES COMPILE_FLAGS "-g") - -add_executable(advanced-thread-trace) -target_sources(advanced-thread-trace PRIVATE main.cpp client.cpp) -target_link_libraries( - advanced-thread-trace - PRIVATE rocprofiler-sdk::rocprofiler-sdk amd_comgr - rocprofiler-sdk::samples-common-library rocprofiler-sdk::samples-build-flags) - -rocprofiler_samples_get_preload_env(PRELOAD_ENV) - -add_test(NAME advanced-thread-trace COMMAND $) - -set_tests_properties( - advanced-thread-trace - PROPERTIES TIMEOUT 45 LABELS "samples" ENVIRONMENT "${PRELOAD_ENV}" - FAIL_REGULAR_EXPRESSION "threw an exception") diff --git a/samples/advanced_thread_trace/client.cpp b/samples/advanced_thread_trace/client.cpp deleted file mode 100644 index 37724f5741..0000000000 --- a/samples/advanced_thread_trace/client.cpp +++ /dev/null @@ -1,541 +0,0 @@ -// MIT License -// -// Copyright (c) 2023-2025 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 - -/** - * @file samples/advanced_thread_trace/client.cpp - * - * @brief Example rocprofiler client (tool) - */ - -#include -#include -#include -#include -#include -#include -#include -#include - -#include "common/defines.hpp" -#include "common/filesystem.hpp" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#define OUTPUT_OFSTREAM "advanced_thread_trace.log" -#define TARGET_CU 1 -#define SIMD_SELECT 0x3 -#define BUFFER_SIZE 0x6000000 -#define SE_MASK 0x11 -constexpr bool COPY_MEMORY_CODEOBJ = true; - -#define C_API_BEGIN \ - try \ - { -#define C_API_END \ - } \ - catch(std::exception & e) \ - { \ - std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl; \ - } \ - catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; } - -struct pcinfo_t -{ - uint64_t marker_id; - uint64_t addr; -}; - -bool -operator==(const pcinfo_t& a, const pcinfo_t& b) -{ - return a.addr == b.addr && a.marker_id == b.marker_id; -}; - -bool -operator<(const pcinfo_t& a, const pcinfo_t& b) -{ - if(a.marker_id == b.marker_id) return a.addr < b.addr; - return a.marker_id < b.marker_id; -}; - -namespace client -{ -using code_obj_load_data_t = rocprofiler_callback_tracing_code_object_load_data_t; -using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; - -using Instruction = rocprofiler::sdk::codeobj::disassembly::Instruction; -using CodeobjAddressTranslate = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; -using SymbolInfo = rocprofiler::sdk::codeobj::disassembly::SymbolInfo; - -rocprofiler_client_id_t* client_id = nullptr; -rocprofiler_context_id_t client_ctx = {0}; - -struct isa_map_elem_t -{ - std::atomic hitcount{0}; - std::atomic latency{0}; - std::unique_ptr code_line{nullptr}; -}; - -struct ToolData -{ - ToolData() - { - try - { - output_file.open(OUTPUT_OFSTREAM); - } catch(...) - {} - - if(output_file.is_open()) - std::cout << "Writing code-object-isa-decode log to: " << OUTPUT_OFSTREAM << std::endl; - else - std::cout << "Could not open log file: " << OUTPUT_OFSTREAM << ", writing to stdout\n"; - }; - - std::shared_mutex isa_map_mut; - std::mutex output_mut; - CodeobjAddressTranslate codeobjTranslate; - std::map> isa_map; - std::unordered_map kernels_in_codeobj = {}; - std::unordered_map kernel_id_to_kernel_name = {}; - int num_waves = 0; - - std::ostream& output() - { - if(output_file.is_open()) - return output_file; - else - return std::cout; - } - - std::stringstream printKernel(uint64_t vaddr) - { - std::stringstream ss; - try - { - ss << '\n' << std::hex; - SymbolInfo& info = kernels_in_codeobj.at(vaddr); - - ss << std::hex << "Found: " << info.name << " at addr: 0x" << vaddr << " with offset 0x" - << info.faddr << " vaddr 0x" << info.vaddr << std::dec << '\n'; - } catch(std::exception& e) - { - ss << e.what() << '\n'; - } - return ss; - } - -private: - std::ofstream output_file; -}; - -struct source_location -{ - std::string function = {}; - std::string file = {}; - uint32_t line = 0; - std::string context = {}; -}; - -struct trace_data_t -{ - 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 */) -{ - C_API_BEGIN - if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return; - if(record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) return; - - if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) - { - 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); - } - - if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return; - - 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( - data->uri, 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; - } - 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); - for(auto& [vaddr, symbol] : symbolmap) - tool->kernels_in_codeobj[vaddr] = symbol; - } - C_API_END -} - -rocprofiler_att_control_flags_t -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, - rocprofiler_dispatch_id_t /* dispatch_id */, - rocprofiler_user_data_t* /* userdata */, - void* /* userdata */) -{ - C_API_BEGIN - - 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); - if(kernel_name.find(desired_func_name) == std::string::npos) - return ROCPROFILER_ATT_CONTROL_NONE; - - int id = call_id.fetch_add(1); - if(id == 1) return ROCPROFILER_ATT_CONTROL_START_AND_STOP; - } catch(...) - { - std::cerr << "Could not find kernel id: " << kernel_id << std::endl; - } - - C_API_END - return ROCPROFILER_ATT_CONTROL_NONE; -} - -void -get_trace_data(rocprofiler_att_parser_data_type_t type, void* att_data, void* userdata) -{ - C_API_BEGIN - assert(userdata && "ISA callback passed null!"); - - 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_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()) - { - shared_lock.unlock(); - { - 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); - } catch(std::exception& e) - { - std::cerr << pc.marker_id << ":" << pc.addr << ' ' << e.what() << std::endl; - return; - } catch(...) - { - std::cerr << "Could not fetch: " << pc.marker_id << ':' << pc.addr << std::endl; - return; - } - it = tool->isa_map.emplace(pc, std::move(ptr)).first; - } - shared_lock.lock(); - } - - it->second->hitcount.fetch_add(event.hitcount, std::memory_order_relaxed); - it->second->latency.fetch_add(event.latency, std::memory_order_relaxed); - C_API_END -} - -uint64_t -copy_trace_data(int* seid, uint8_t** buffer, uint64_t* buffer_size, void* userdata) -{ - trace_data_t& data = *reinterpret_cast(userdata); - *seid = data.id; - *buffer_size = data.size; - *buffer = data.data; - data.size = 0; - return *buffer_size; -} - -rocprofiler_status_t -isa_callback(char* isa_instruction, - uint64_t* isa_memory_size, - uint64_t* isa_size, - uint64_t marker_id, - uint64_t offset, - void* userdata) -{ - C_API_BEGIN - assert(userdata && "ISA callback passed null!"); - - std::unique_ptr instruction; - - { - 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; - - { - size_t tmp_isa_size = *isa_size; - *isa_size = instruction->inst.size(); - - if(*isa_size > tmp_isa_size) return ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES; - } - - memcpy(isa_instruction, instruction->inst.data(), *isa_size); - *isa_memory_size = instruction->size; - - auto ptr = std::make_unique(); - ptr->code_line = std::move(instruction); - 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, - rocprofiler_user_data_t /* userdata */) -{ - C_API_BEGIN - - { - 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}; - 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; - C_API_END -} - -int -tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) -{ - (void) fini_func; - ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); - - ROCPROFILER_CALL( - rocprofiler_configure_callback_tracing_service(client_ctx, - ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, - nullptr, - 0, - tool_codeobj_tracing_callback, - tool_data), - "code object tracing service configure"); - - std::vector parameters = { - {ROCPROFILER_ATT_PARAMETER_TARGET_CU, {TARGET_CU}}, - {ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {SIMD_SELECT}}, - {ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {BUFFER_SIZE}}, - {ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {SE_MASK}}}; - - 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; - ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx), - "context validity check"); - if(valid_ctx == 0) - { - // notify rocprofiler that initialization failed - // and all the contexts, buffers, etc. created - // should be ignored - return -1; - } - - ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "context start"); - - // no errors - return 0; -} - -void -tool_fini(void* /* data */) -{ - 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] : client::tool->isa_map) - if(lines.get()) max_inst_size = std::max(max_inst_size, lines->code_line->inst.size()); - - std::string empty_space; - empty_space.resize(max_inst_size, ' '); - - size_t vmc_latency = 0; - size_t lgk_latency = 0; - size_t scalar_latency = 0; - size_t vector_latency = 0; - size_t other_latency = 0; - - size_t scalar_exec = 0; - size_t vector_exec = 0; - size_t other_exec = 0; - - 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; - - 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) - { - other_exec += hitcount; - if(code_line.find("lgkmcnt") != std::string::npos) - lgk_latency += latency; - else - vmc_latency += latency; - } - else if(code_line.find("v_") == 0) - { - vector_exec += hitcount; - vector_latency += latency; - } - else if(code_line.find("s_") == 0) - { - scalar_exec += hitcount; - scalar_latency += latency; - } - else - { - other_exec += hitcount; - other_latency += latency; - } - } - - size_t total_exec = vector_exec + scalar_exec + other_exec; - size_t memory_latency = vmc_latency + lgk_latency; - size_t total_latency = memory_latency + vector_latency + scalar_latency + other_latency; - float vmc_fraction = 100 * vmc_latency / float(total_latency); - float lgk_fraction = 100 * lgk_latency / float(total_latency); - - 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 - -extern "C" rocprofiler_tool_configure_result_t* -rocprofiler_configure(uint32_t version, - const char* runtime_version, - uint32_t priority, - rocprofiler_client_id_t* id) -{ - // set the client name - id->name = "Adv_Thread_Trace_Sample"; - - // store client info - client::client_id = id; - - // compute major/minor/patch version info - uint32_t major = version / 10000; - uint32_t minor = (version % 10000) / 100; - uint32_t patch = version % 100; - - // generate info string - auto info = std::stringstream{}; - info << id->name << " (priority=" << priority << ") is using rocprofiler-sdk v" << major << "." - << minor << "." << patch << " (" << runtime_version << ")"; - - std::clog << info.str() << std::endl; - - // create configure data - static auto cfg = - rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), - &client::tool_init, - &client::tool_fini, - nullptr}; - - // return pointer to configure data - return &cfg; -} diff --git a/samples/advanced_thread_trace/main.cpp b/samples/advanced_thread_trace/main.cpp deleted file mode 100644 index 1266f52ec3..0000000000 --- a/samples/advanced_thread_trace/main.cpp +++ /dev/null @@ -1,245 +0,0 @@ -// MIT License -// -// Copyright (c) 2024-2025 Advanced Micro Devices, Inc. All rights reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining a copy -// of this software and associated documentation files (the "Software"), to deal -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in all -// copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#include "hip/hip_runtime.h" - -#include -#include -#include -#include -#include -#include -#include -#include "transpose_kernels.hpp" - -#define PRINT_ALIGN 36 - -namespace -{ -using lock_guard_t = std::lock_guard; -auto print_lock = std::mutex{}; -} // namespace - -enum TransposeType -{ - TRANSPOSE_NAIVE, - TRANSPOSE_INPLACE_LDS, - TRANSPOSE_NO_BANK_CONFLICTS -}; - -class ITranspose -{ -public: - virtual void run(TransposeType ttype, int numThreadsY, int num_iter) = 0; - virtual ~ITranspose(){}; -}; - -template -class Transpose : public ITranspose -{ -public: - Transpose(int dev, size_t _M) - : devID(dev) - , M(_M) - , databytes(_M * _M * sizeof(T)) - { - HIP_API_CALL(hipSetDevice(devID)); - HIP_API_CALL(hipStreamCreate(&stream)); - - std::default_random_engine _engine{std::random_device{}() * rand()}; - std::uniform_int_distribution _dist{0, 1000}; - - inp_matrix = new T[M * M]; - out_matrix = new T[M * M]; - - for(size_t i = 0; i < M * M; i++) - inp_matrix[i] = static_cast(_dist(_engine)); - memset(out_matrix, 0, databytes); - - HIP_API_CALL(hipMalloc(&in, databytes)); - HIP_API_CALL(hipMalloc(&out, databytes)); - HIP_API_CALL(hipMemsetAsync(in, 0, databytes, stream)); - HIP_API_CALL(hipMemsetAsync(out, 0, databytes, stream)); - HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, databytes, hipMemcpyDefault, stream)); - - HIP_API_CALL(hipEventCreate(&start)); - HIP_API_CALL(hipEventCreate(&stop)); - } - - void run(TransposeType ttype, int numThreadsY, int num_iter) override - { - HIP_API_CALL(hipSetDevice(devID)); - dim3 grid(M / TILE_DIM, M / TILE_DIM, 1); - dim3 block(TILE_DIM, numThreadsY, 1); - - auto Kernel = transposeNaive; - std::string KernelName = "transposeNaive"; - if(ttype == TransposeType::TRANSPOSE_NO_BANK_CONFLICTS) - { - Kernel = transposeLdsNoBankConflicts; - KernelName = "transposeLdsNoBankConflicts"; - } - else if(ttype == TransposeType::TRANSPOSE_INPLACE_LDS) - { - Kernel = transposeLdsSwapInplace; - KernelName = "transposeLdsSwapInplace"; - } - - { - std::string functypeid = __PRETTY_FUNCTION__; - auto it_beg = functypeid.rfind("[T = "); - auto it_end = functypeid.rfind(']'); - - if(it_beg != std::string::npos) it_beg += std::string("[T = ").size(); - - if(it_beg < it_end && it_end != std::string::npos) - KernelName += '<' + functypeid.substr(it_beg, it_end - it_beg) + '>'; - } - - HIP_API_CALL(hipStreamSynchronize(stream)); - HIP_API_CALL(hipEventRecord(start, stream)); - - for(int i = 0; i < num_iter; i++) - { - Kernel<<>>(out, in, M); - HIP_API_CALL(hipGetLastError()); - } - - HIP_API_CALL(hipEventRecord(stop, stream)); - HIP_API_CALL(hipMemcpyAsync(out_matrix, out, databytes, hipMemcpyDefault, stream)); - HIP_API_CALL(hipEventSynchronize(stop)); - - float time; - HIP_API_CALL(hipEventElapsedTime(&time, start, stop)); - float GB = databytes * num_iter * 2 / float(1 << 30); - - { - lock_guard_t _lk{print_lock}; - std::cout << "The average performance of " << std::setw(38) << KernelName << " : " - << (1000 * GB / time) << " GB/s" << std::endl; - } - - verify(); - } - - void verify() const - { - HIP_API_CALL(hipStreamSynchronize(stream)); - for(int i = 0; i < 10; i++) - { - int row = rand() % M; - int col = rand() % M; - if(inp_matrix[row * M + col] != out_matrix[col * M + row]) - { - lock_guard_t _lk{print_lock}; - std::cout << "mismatch: " << row << ", " << col << " : " - << inp_matrix[row * M + col] << " | " << out_matrix[col * M + row] - << std::endl; - } - } - } - - virtual ~Transpose() - { - HIP_API_CALL(hipSetDevice(devID)); - HIP_API_CALL(hipEventDestroy(start)); - HIP_API_CALL(hipEventDestroy(stop)); - - HIP_API_CALL(hipFree(in)); - HIP_API_CALL(hipFree(out)); - HIP_API_CALL(hipStreamDestroy(stream)); - - delete[] inp_matrix; - delete[] out_matrix; - } - - const int devID; - const size_t M; - const size_t databytes; - - hipStream_t stream; - hipEvent_t start, stop; - - T* inp_matrix = nullptr; - T* out_matrix = nullptr; - - T* in = nullptr; - T* out = nullptr; -}; - -int -main(int argc, char** argv) -{ - int deviceId = 0; - int blockDimY = 8; - int num_iter = 1; - int mat_size = 8192; - - for(int i = 1; i < argc; ++i) - { - auto _arg = std::string{argv[i]}; - if(_arg == "?" || _arg == "-h" || _arg == "--help") - { - std::cout << "usage: transpose " - << "[MatrixSize (" << mat_size << ")] " - << "[numIter (" << num_iter << ")] " - << "[blockDimY (" << blockDimY << ")] " - << "[DEVICE_ID (" << deviceId << ")] " << std::endl; - exit(EXIT_SUCCESS); - } - } - if(argc > 1) mat_size = atoll(argv[1]); - if(argc > 2) num_iter = atoll(argv[2]); - if(argc > 3) blockDimY = atoll(argv[3]); - if(argc > 4) deviceId = atoll(argv[4]); - - printf("[transpose] Matrix size: %d, device ID: %d, num iter: %d, blockDimY: %d\n", - mat_size, - deviceId, - num_iter, - blockDimY); - - int ndevice = 0; - HIP_API_CALL(hipGetDeviceCount(&ndevice)); - printf("[transpose] Number of devices found: %i\n", ndevice); - assert(ndevice > 0); - - if(deviceId >= ndevice) exit(EXIT_FAILURE); - - { - std::vector> kernels; - kernels.push_back(std::make_unique>(deviceId, mat_size)); - kernels.push_back(std::make_unique>(deviceId, mat_size)); - kernels.push_back(std::make_unique>(deviceId, mat_size)); - - for(auto& kernel : kernels) - { - kernel->run(TransposeType::TRANSPOSE_NAIVE, blockDimY, num_iter); - kernel->run(TransposeType::TRANSPOSE_INPLACE_LDS, blockDimY, num_iter); - kernel->run(TransposeType::TRANSPOSE_NO_BANK_CONFLICTS, blockDimY, num_iter); - } - } - - HIP_API_CALL(hipDeviceSynchronize()); - - return 0; -} diff --git a/samples/advanced_thread_trace/transpose_kernels.hpp b/samples/advanced_thread_trace/transpose_kernels.hpp deleted file mode 100644 index f7fa00bfe7..0000000000 --- a/samples/advanced_thread_trace/transpose_kernels.hpp +++ /dev/null @@ -1,106 +0,0 @@ -// MIT License -// -// Copyright (c) 2024-2025 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 "hip/hip_runtime.h" - -#define HIP_API_CALL(CALL) \ - { \ - hipError_t error_ = (CALL); \ - if(error_ != hipSuccess) \ - { \ - lock_guard_t _hip_api_print_lk{print_lock}; \ - fprintf(stderr, \ - "%s:%d :: HIP error : %s\n", \ - __FILE__, \ - __LINE__, \ - hipGetErrorString(error_)); \ - exit(EXIT_FAILURE); \ - } \ - } - -#define TILE_DIM 64 - -template -__global__ void -transposeNaive(T* odata, const T* idata, size_t size) -{ - size_t idx = blockIdx.x * TILE_DIM + threadIdx.x; - size_t block_posy = blockIdx.y * TILE_DIM; - - for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y) - odata[size * idx + block_posy + idy] = idata[idx + (block_posy + idy) * size]; -} - -template -__global__ void -transposeLdsNoBankConflicts(T* odata, const T* idata, size_t size) -{ - __shared__ T tile[TILE_DIM][TILE_DIM + 1]; - - size_t idx_in = blockIdx.x * TILE_DIM + threadIdx.x; - size_t idy_in = blockIdx.y * TILE_DIM + threadIdx.y; - size_t index_in = idx_in + idy_in * size; - - size_t idx_out = blockIdx.y * TILE_DIM + threadIdx.x; - size_t idy_out = blockIdx.x * TILE_DIM + threadIdx.y; - size_t index_out = idx_out + idy_out * size; - - for(size_t y = 0; y < TILE_DIM; y += blockDim.y) - tile[threadIdx.y + y][threadIdx.x] = idata[index_in + y * size]; - - __syncthreads(); - - for(size_t y = 0; y < TILE_DIM; y += blockDim.y) - odata[index_out + y * size] = tile[threadIdx.x][threadIdx.y + y]; -} - -// Generates more interesting ISA -template -__global__ void -transposeLdsSwapInplace(T* odata, const T* idata, size_t size) -{ - __shared__ T tile[TILE_DIM][TILE_DIM]; - - const size_t idx_in = blockIdx.x * TILE_DIM + threadIdx.x; - - for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y) - tile[idy][threadIdx.x] = idata[idx_in + (idy + blockIdx.y * TILE_DIM) * size]; - - __syncthreads(); - - for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y) - if(idy < threadIdx.x) - { - T temp = tile[idy][threadIdx.x]; - tile[idy][threadIdx.x] = tile[threadIdx.x][idy]; - tile[threadIdx.x][idy] = temp; - } - - __syncthreads(); - - const size_t idx_out = blockIdx.y * TILE_DIM + threadIdx.x; - - for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y) - odata[(blockIdx.x * TILE_DIM + idy) * size + idx_out] = tile[idy][threadIdx.x]; -} diff --git a/source/bin/rocprofv3.py b/source/bin/rocprofv3.py index 5797cdefe0..95f2df5378 100755 --- a/source/bin/rocprofv3.py +++ b/source/bin/rocprofv3.py @@ -90,7 +90,7 @@ def strtobool(val): def search_path(path_list): supported_option = [] - lib_att_pattern = r"libatt_decoder_(trace|summary|debug|testing)\.so" + lib_att_pattern = r"libatt_decoder_(trace|debug|testing1|testing2)\.so" file_list = [] for path in path_list: @@ -100,7 +100,7 @@ def search_path(path_list): for itr in file_list: _match = re.match(lib_att_pattern, itr) if _match: - lst = re.findall("trace|debug|summary|testing", itr) + lst = re.findall("trace|debug|testing1|testing2", itr) supported_option.extend(lst) return set(supported_option) @@ -682,6 +682,20 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins choices=set(choice_list), ) + att_options.add_argument( + "--att-perfcounters", + help="Set performance counters, and optionally their mask", + default=None, + type=str.upper, + ) + + att_options.add_argument( + "--att-perfcounter-ctrl", + help="Integer in [0,32] range specifying collection period.", + default=None, + type=int, + ) + add_parser_bool_argument( att_options, "--att-serialize-all", @@ -1305,15 +1319,13 @@ def run(app_args, args, **kwargs): f"{type(num_str)} is not supported. {num_str} should be of type integer or string." ) - if args.pmc or ( + if ( args.pc_sampling_beta_enabled or args.pc_sampling_unit or args.pc_sampling_method or args.pc_sampling_interval ): - fatal_error( - "Advanced thread trace cannot be enabled with counter collection or pc sampling" - ) + fatal_error("Advanced thread trace cannot be enabled with pc sampling") if not args.att_parse: fatal_error("provide the parser choice") @@ -1358,12 +1370,24 @@ def run(app_args, args, **kwargs): ":".join(args.att_library_path), overwrite=True, ) - if args.att_percounters: - update_env( - "ROCPROF_ATT_PARAM_PERFCOUNTERS", - " ".join(args.att_perfcounters), - overwrite=True, - ) + if args.att_perfcounters: + if args.pmc: + fatal_error("ATT perfcounters cannot be enabled with PMC") + else: + update_env( + "ROCPROF_ATT_PARAM_PERFCOUNTERS", + args.att_perfcounters, + overwrite=True, + ) + if args.att_perfcounter_ctrl: + if args.pmc: + fatal_error("ATT perfcounters cannot be enabled with PMC") + else: + update_env( + "ROCPROF_ATT_PARAM_PERFCOUNTER_CTRL", + args.att_perfcounter_ctrl, + overwrite=True, + ) if args.log_level in ("info", "trace", "env"): log_config(app_env) diff --git a/source/docs/rocprofv3-schema.json b/source/docs/rocprofv3-schema.json index faa9bcc22e..918ca0a53b 100644 --- a/source/docs/rocprofv3-schema.json +++ b/source/docs/rocprofv3-schema.json @@ -569,22 +569,27 @@ "description": "Comments matching assembly instructions from pc_sample_instructions array. If debug symbols are available, comments provide instructions to source-line mapping. Otherwise, a comment is an empty string." }, "att_filenames" : { - "type": "object", - "properties": { - "key": { - "type": "integer", - "description": "Dispatch id." - }, - "value": { - "type": "array", - "description": "An array of ATT filenames." - } - + "type": "array", + "items": { + "type": "string", + "description": "Filename of saved .att data in the form name_shader_engine_seId_dispatchId.att" } }, "code_object_snapshot_filenames": { "type": "array", - "description": "Code object snapshot filename" + "items": { + "type": "object", + "properties": { + "key": { + "type": "integer", + "description": "Code object load id." + }, + "value": { + "type": "string", + "description": "Saved code object filename" + } + } + } } } }, diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h index c0f59adedf..38d9373532 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h @@ -57,11 +57,11 @@ ROCPROFILER_EXTERN_C_INIT * @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT for invalid rocprofiler_att_parameter_t */ rocprofiler_status_t -rocprofiler_configure_agent_thread_trace_service( +rocprofiler_configure_device_thread_trace_service( rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id, rocprofiler_att_parameter_t* parameters, size_t num_parameters, - rocprofiler_agent_id_t agent_id, rocprofiler_att_shader_data_callback_t shader_callback, rocprofiler_user_data_t callback_userdata) ROCPROFILER_API; diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h index 07e6ad5887..772376e271 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -66,8 +66,9 @@ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)( /** * @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. + * This service serializes all traced kernels, and optionally all non-traced kernels. + * @param [in] context_id id of the context used for start/stop thread_trace. + * @param [in] agent_id rocprofiler_agent_id_t to configure thread trace. * @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. @@ -84,6 +85,7 @@ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)( rocprofiler_status_t rocprofiler_configure_dispatch_thread_trace_service( rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id, rocprofiler_att_parameter_t* parameters, size_t num_parameters, rocprofiler_att_dispatch_callback_t dispatch_callback, diff --git a/source/lib/output/generateJSON.cpp b/source/lib/output/generateJSON.cpp index 1df4b2db35..f88ddfbbab 100644 --- a/source/lib/output/generateJSON.cpp +++ b/source/lib/output/generateJSON.cpp @@ -125,9 +125,10 @@ write_json(json_output& json_ar, auto att_filenames = tool_metadata.get_att_filenames(); auto code_object_snapshot_filenames = std::vector{}; - code_object_snapshot_filenames.reserve(code_object_load_info.size()); + code_object_snapshot_filenames.resize(code_object_load_info.size()); + for(const auto& info : code_object_load_info) - code_object_snapshot_filenames.emplace_back(fs::path(info.name).filename()); + code_object_snapshot_filenames.at(info.id) = fs::path(info.name).filename(); json_ar.setNextName("strings"); json_ar.startNode(); diff --git a/source/lib/output/metadata.cpp b/source/lib/output/metadata.cpp index 9be88638f0..850b685f04 100644 --- a/source/lib/output/metadata.cpp +++ b/source/lib/output/metadata.cpp @@ -23,6 +23,7 @@ #include "metadata.hpp" #include "lib/common/filesystem.hpp" +#include "lib/common/logging.hpp" #include "lib/common/string_entry.hpp" #include "lib/output/agent_info.hpp" #include "lib/output/host_symbol_info.hpp" @@ -31,6 +32,7 @@ #include #include +#include #include #include @@ -246,7 +248,20 @@ metadata::get_code_object_load_info() const return _info; }); - return _data; + uint64_t _sz = 0; + for(const auto& itr : _data) + _sz = std::max(_sz, itr.id); + + ROCP_WARNING_IF((_sz + 1) - _data.size() > 1000) << fmt::format( + "Spares index detected for code object load info: {} < {}", _data.size(), _sz); + + auto _code_obj_data = std::vector{}; + _code_obj_data.resize(_sz + 1, rocprofiler::att_wrapper::CodeobjLoadInfo{}); + // index by the code object id + for(auto& itr : _data) + _code_obj_data.at(itr.id) = itr; + + return _code_obj_data; } std::vector diff --git a/source/lib/rocprofiler-sdk-att/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/CMakeLists.txt index 5126235d1f..d4d050f90f 100755 --- a/source/lib/rocprofiler-sdk-att/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-att/CMakeLists.txt @@ -12,6 +12,7 @@ set(ATT_TOOL_SOURCE_FILES filenames.cpp occupancy.cpp wstates.cpp + perfcounter.cpp profile_interface.cpp dl.cpp) diff --git a/source/lib/rocprofiler-sdk-att/att_decoder.h b/source/lib/rocprofiler-sdk-att/att_decoder.h index 9daab12978..ec37e17fe7 100644 --- a/source/lib/rocprofiler-sdk-att/att_decoder.h +++ b/source/lib/rocprofiler-sdk-att/att_decoder.h @@ -124,6 +124,17 @@ typedef struct att_wave_instruction_t* instructions_array; } att_wave_data_t; +typedef struct att_perfevent_t +{ + int64_t time; + uint16_t events0; + uint16_t events1; + uint16_t events2; + uint16_t events3; + uint8_t CU; + uint8_t bank; +} att_perfevent_t; + typedef rocprofiler_att_decoder_status_t (*rocprofiler_att_decoder_isa_callback_t)( char* instruction, uint64_t* memory_size, diff --git a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp index bfd3c0ba15..5fbbef77f0 100644 --- a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp +++ b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp @@ -48,13 +48,13 @@ get_lib_names() std::vector> lib_names = { {tool_att_capability_t::ATT_CAPABILITIES_DEBUG, "libatt_decoder_debug.so"}, {tool_att_capability_t::ATT_CAPABILITIES_TRACE, "libatt_decoder_trace.so"}, - {tool_att_capability_t::ATT_CAPABILITIES_SUMMARY, "libatt_decoder_summary.so"}, - {tool_att_capability_t::ATT_CAPABILITIES_TESTING, "libatt_decoder_testing.so"}, + {tool_att_capability_t::ATT_CAPABILITIES_TESTING1, "libatt_decoder_testing1.so"}, + {tool_att_capability_t::ATT_CAPABILITIES_TESTING2, "libatt_decoder_testing2.so"}, }; return lib_names; } -ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr
_dl) +ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr
_dl, std::vector _counters) : dir(std::move(_dir)) , dl(std::move(_dl)) { @@ -62,8 +62,10 @@ ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr
_dl) table = std::make_shared(); codefile = std::make_shared(dir, table); filenames = std::make_shared(dir); + for(size_t i = 0; i < ATT_WAVE_STATE_LAST; i++) wstates.at(i) = std::make_shared(i, dir); + filenames->perfcounters = std::move(_counters); } ATTFileMgr::~ATTFileMgr() { OccupancyFile::OccupancyFile(dir, table, occupancy); } @@ -124,6 +126,7 @@ ATTDecoder::parse(const Fspath& input_dir, const Fspath& output_dir, const std::vector& att_files, const std::vector& codeobj_files, + const std::vector& counters_names, const std::string& output_formats) { auto& formats = GlobalDefs::get().output_formats; @@ -132,7 +135,7 @@ ATTDecoder::parse(const Fspath& input_dir, return std::tolower(c); }); - ATTFileMgr mgr(output_dir, dl); + ATTFileMgr mgr(output_dir, dl, counters_names); for(const auto& file : codeobj_files) { diff --git a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp index 033db690b2..3506bc213b 100644 --- a/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp +++ b/source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp @@ -47,9 +47,9 @@ struct CodeobjLoadInfo enum tool_att_capability_t { - ATT_CAPABILITIES_TESTING = 0, // used for code coverage testing - ATT_CAPABILITIES_SUMMARY, // used for CSV output only - ATT_CAPABILITIES_TRACE, // used for all outputs + ATT_CAPABILITIES_TESTING1 = 0, // used for code coverage testing + ATT_CAPABILITIES_TESTING2, // used for code coverage testing + ATT_CAPABILITIES_TRACE, // used for all outputs ATT_CAPABILITIES_DEBUG, ATT_CAPABILITIES_LAST = ATT_CAPABILITIES_DEBUG, }; @@ -77,6 +77,7 @@ public: const Fspath& output_dir, const std::vector& att_files, const std::vector& codeobj_files, + const std::vector& counters_names, const std::string& output_formats); bool valid() const; @@ -90,7 +91,7 @@ class ATTFileMgr using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; public: - ATTFileMgr(Fspath _dir, std::shared_ptr
_dl); + ATTFileMgr(Fspath _dir, std::shared_ptr
_dl, std::vector _counters); ~ATTFileMgr(); void parseShader(int se_id, const std::vector& data); diff --git a/source/lib/rocprofiler-sdk-att/filenames.cpp b/source/lib/rocprofiler-sdk-att/filenames.cpp index 6af3f502a5..d11f0d8e83 100644 --- a/source/lib/rocprofiler-sdk-att/filenames.cpp +++ b/source/lib/rocprofiler-sdk-att/filenames.cpp @@ -50,12 +50,13 @@ FilenameMgr::~FilenameMgr() [to_string(coord.id)] = {data.name, data.begin, data.end}; } - nlohmann::json metadata = {{"global_begin_time", 0}, - {"gfxv", (gfxip > 9) ? "navi" : "vega"}, - {"gfxip", gfxip}, - {"version", TOOL_VERSION}}; + const nlohmann::json metadata = {{"global_begin_time", 0}, + {"gfxv", (gfxip > 9) ? "navi" : "vega"}, + {"gfxip", gfxip}, + {"version", TOOL_VERSION}, + {"counter_names", perfcounters}, + {"wave_filenames", namelist}}; - metadata["wave_filenames"] = namelist; OutputFile(filename) << metadata; } diff --git a/source/lib/rocprofiler-sdk-att/filenames.hpp b/source/lib/rocprofiler-sdk-att/filenames.hpp index 0430ee63e3..0849097042 100644 --- a/source/lib/rocprofiler-sdk-att/filenames.hpp +++ b/source/lib/rocprofiler-sdk-att/filenames.hpp @@ -72,6 +72,7 @@ public: Fspath dir{}; Fspath filename{}; std::map streams{}; + std::vector perfcounters{}; int gfxip = 9; }; diff --git a/source/lib/rocprofiler-sdk-att/perfcounter.cpp b/source/lib/rocprofiler-sdk-att/perfcounter.cpp new file mode 100644 index 0000000000..9682dce406 --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/perfcounter.cpp @@ -0,0 +1,62 @@ +// MIT License +// +// Copyright (c) 2024-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "perfcounter.hpp" +#include "outputfile.hpp" +#include "wave.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace att_wrapper +{ +void +PerfcounterFile(WaveConfig& config, const att_perfevent_t* events, size_t event_count) +{ + nlohmann::json data; + for(size_t i = 0; i < event_count; i++) + { + const auto& event = events[i]; + + nlohmann::json json_event; + json_event.push_back(event.time); + json_event.push_back(event.events0); + json_event.push_back(event.events1); + json_event.push_back(event.events2); + json_event.push_back(event.events3); + json_event.push_back(event.CU); + json_event.push_back(event.bank); + + data.push_back(json_event); + } + + const nlohmann::json json{{"data", data}}; + + std::string filename = "se" + std::to_string(config.shader_engine) + "_perfcounter.json"; + OutputFile(config.filemgr->dir / filename) << json; +} +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/perfcounter.hpp b/source/lib/rocprofiler-sdk-att/perfcounter.hpp new file mode 100644 index 0000000000..1ceb13e71d --- /dev/null +++ b/source/lib/rocprofiler-sdk-att/perfcounter.hpp @@ -0,0 +1,34 @@ +// MIT License +// +// Copyright (c) 2024-2025 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "att_lib_wrapper.hpp" + +namespace rocprofiler +{ +namespace att_wrapper +{ +void +PerfcounterFile(class WaveConfig& config, const att_perfevent_t* events, size_t event_count); +} // namespace att_wrapper +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-att/profile_interface.cpp b/source/lib/rocprofiler-sdk-att/profile_interface.cpp index 5687619a43..7ad5ea1a25 100644 --- a/source/lib/rocprofiler-sdk-att/profile_interface.cpp +++ b/source/lib/rocprofiler-sdk-att/profile_interface.cpp @@ -26,12 +26,14 @@ #endif #include "profile_interface.hpp" +#include "att_decoder.h" +#include "dl.hpp" +#include "perfcounter.hpp" + #include #include #include #include -#include "att_decoder.h" -#include "dl.hpp" namespace rocprofiler { @@ -64,11 +66,8 @@ get_trace_data(rocprofiler_att_decoder_record_type_t trace_id, auto* infos = (rocprofiler_att_decoder_info_t*) trace_events; for(size_t i = 0; i < trace_size; i++) ROCP_WARNING << tool.dl->att_info_fn(infos[i]); - - return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; } - - if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_GFXIP) + else if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_GFXIP) { tool.config.filemgr->gfxip = reinterpret_cast(trace_events); } @@ -78,6 +77,10 @@ get_trace_data(rocprofiler_att_decoder_record_type_t trace_id, tool.config.occupancy.push_back( reinterpret_cast(trace_events)[i]); } + else if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_PERFEVENT) + { + PerfcounterFile(tool.config, reinterpret_cast(trace_events), trace_size); + } if(trace_id != ROCPROFILER_ATT_DECODER_TYPE_WAVE) return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS; diff --git a/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt index 55d64f5617..ed5555b8de 100644 --- a/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt @@ -24,17 +24,17 @@ target_link_libraries( GTest::gtest GTest::gtest_main) -add_library(att_decoder_testing SHARED) -add_library(rocprofiler-sdk::att-decoder-testing ALIAS att_decoder_testing) -target_sources(att_decoder_testing PRIVATE dummy_decoder.cpp) -set_target_properties(att_decoder_testing PROPERTIES VERSION ${PROJECT_VERSION} - SOVERSION ${PROJECT_VERSION_MAJOR}) +add_library(att_decoder_testing1 SHARED) +add_library(rocprofiler-sdk::att-decoder-testing1 ALIAS att_decoder_testing1) +target_sources(att_decoder_testing1 PRIVATE dummy_decoder.cpp) +set_target_properties(att_decoder_testing1 PROPERTIES VERSION ${PROJECT_VERSION} + SOVERSION ${PROJECT_VERSION_MAJOR}) -add_library(att_decoder_summary SHARED) -add_library(rocprofiler-sdk::att-decoder-summary ALIAS att_decoder_summary) -target_sources(att_decoder_summary PRIVATE dummy_decoder.cpp) +add_library(att_decoder_testing2 SHARED) +add_library(rocprofiler-sdk::att-decoder-testing2 ALIAS att_decoder_testing2) +target_sources(att_decoder_testing2 PRIVATE dummy_decoder.cpp) set_target_properties( - att_decoder_summary + att_decoder_testing2 PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}/lib/att VERSION ${PROJECT_VERSION} SOVERSION ${PROJECT_VERSION_MAJOR}) diff --git a/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp b/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp index 055766d12a..acf998ac3e 100644 --- a/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp +++ b/source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp @@ -37,7 +37,7 @@ class ATTDecoderTest : public ATTDecoder { public: ATTDecoderTest() - : ATTDecoder(rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TESTING) + : ATTDecoder(rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TESTING1) { rocprofiler::att_wrapper::OutputFile::Enabled() = false; GlobalDefs::get().output_formats = "json,csv"; @@ -46,7 +46,7 @@ public: void test_parse() { - ATTFileMgr mgr("out/", dl); + ATTFileMgr mgr("out/", dl, {}); auto append_isa = [&](const char* line) { // matches addresses in dummy_decoder.cpp @@ -107,7 +107,7 @@ TEST(att_decoder_test, warn_failures) ATTDecoderTest decoder; ROCP_FATAL_IF(!decoder.valid()) << "Failed to initialize decoder library!"; - decoder.parse(".", ".", att_files, codeobjs, "csv,json"); + decoder.parse(".", ".", att_files, codeobjs, {}, "csv,json"); } TEST(att_decoder_test, code_write) diff --git a/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp b/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp index 9f6f423297..1c41e7e95e 100644 --- a/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp +++ b/source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp @@ -53,7 +53,7 @@ main(int argc, char** argv) FLAGS_colorlogtostderr = true; }); - auto cap = rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_SUMMARY; + auto cap = rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TRACE; { auto query = rocprofiler::att_wrapper::query_att_decode_capability(); ROCP_FATAL_IF(query.empty()) << "No decoder capability available!"; @@ -98,7 +98,7 @@ main(int argc, char** argv) std::unordered_map> all_runs; - for(auto& file : sdk_json["strings"]["att_files"]) + for(auto& file : sdk_json["strings"]["att_filenames"]) { try { @@ -119,7 +119,7 @@ main(int argc, char** argv) std::vector codeobj_files{}; std::map snapshot_files{}; - for(auto elem : sdk_json["strings"]["code_object_snapshot_files"]) + for(auto elem : sdk_json["strings"]["code_object_snapshot_filenames"]) snapshot_files[elem["key"]] = elem["value"]; for(auto& codeobj : sdk_json["code_objects"]) @@ -158,7 +158,8 @@ main(int argc, char** argv) std::string run_name = input_path.filename().c_str(); std::string ui_name = run_name.substr(0, run_name.find(".json")); auto output_dir = output_path / ("ui_output_" + ui_name + std::to_string(run_number)); - decoder.parse(input_path.parent_path(), output_dir, att_filenames, codeobj_files, formats); + decoder.parse( + input_path.parent_path(), output_dir, att_filenames, codeobj_files, {}, formats); } ROCP_INFO << "Finalizing ATT Tool"; diff --git a/source/lib/rocprofiler-sdk-tool/config.cpp b/source/lib/rocprofiler-sdk-tool/config.cpp index 2ef53a4053..15e839d15d 100644 --- a/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/source/lib/rocprofiler-sdk-tool/config.cpp @@ -113,13 +113,13 @@ has_counter_format(std::string const& str) } // validate kernel names -std::unordered_set +std::unordered_set get_kernel_filter_range(const std::string& kernel_filter) { if(kernel_filter.empty()) return {}; auto delim = rocprofiler::sdk::parse::tokenize(kernel_filter, "[], "); - auto range_set = std::unordered_set{}; + auto range_set = std::unordered_set{}; for(const auto& itr : delim) { if(itr.find('-') != std::string::npos) @@ -129,8 +129,8 @@ get_kernel_filter_range(const std::string& kernel_filter) ROCP_FATAL_IF(drange.size() != 2) << "bad range format for '" << itr << "'. Expected [A-B] where A and B are numbers"; - uint32_t start_range = std::stoul(drange.front()); - uint32_t end_range = std::stoul(drange.back()); + size_t start_range = std::stoul(drange.front()); + size_t end_range = std::stoul(drange.back()); for(auto i = start_range; i <= end_range; i++) range_set.emplace(i); } @@ -180,7 +180,6 @@ parse_att_counters(std::string line) }; // regex to check if string is of the form "counter_name:simd_mask" - std::regex pattern(R"([a-zA-Z0-9_]+(:0x[0-9a-fA-F]+)?)"); std::set unique_counters; auto input_ss = std::stringstream{line}; @@ -190,13 +189,6 @@ parse_att_counters(std::string line) input_ss >> counter; if(counter.empty()) break; - // check if the counter string matches the pattern - if(!std::regex_match(counter, pattern)) - { - ROCP_FATAL << "Invalid counter format for ATT: " << counter - << ". Expected format : Counter_name:optional_simd_mask(hexadecimal)"; - } - // Consider only those counters where combination of counter name and simd mask is unique if(unique_counters.insert(counter).second == false) continue; diff --git a/source/lib/rocprofiler-sdk-tool/config.hpp b/source/lib/rocprofiler-sdk-tool/config.hpp index dda1ce5c32..3fe697cd18 100644 --- a/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/source/lib/rocprofiler-sdk-tool/config.hpp @@ -128,6 +128,7 @@ struct config : output_config uint64_t att_param_buffer_size = get_env("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x6000000); uint64_t att_param_simd_select = get_env("ROCPROF_ATT_PARAM_SIMD_SELECT", 0xF); uint64_t att_param_target_cu = get_env("ROCPROF_ATT_PARAM_TARGET_CU", 1); + uint64_t att_param_perf_ctrl = get_env("ROCPROF_ATT_PARAM_PERFCOUNTER_CTRL", 0); std::string kernel_filter_include = get_env("ROCPROF_KERNEL_FILTER_INCLUDE_REGEX", ".*"); std::string kernel_filter_exclude = get_env("ROCPROF_KERNEL_FILTER_EXCLUDE_REGEX", ""); @@ -135,7 +136,7 @@ struct config : output_config std::string pc_sampling_unit = get_env("ROCPROF_PC_SAMPLING_UNIT", "none"); std::string extra_counters_contents = get_env("ROCPROF_EXTRA_COUNTERS_CONTENTS", ""); - std::unordered_set kernel_filter_range = {}; + std::unordered_set kernel_filter_range = {}; std::vector> counters = {}; std::string att_capability = get_env("ROCPROF_ATT_CAPABILITY", ""); std::vector att_param_perfcounters = {}; diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index 99892c8aa1..a3e93efc97 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -193,17 +193,16 @@ as_pointer() } using targeted_kernels_map_t = - std::unordered_map>; + std::unordered_map>; using counter_dimension_info_map_t = std::unordered_map>; using agent_info_map_t = std::unordered_map; -using kernel_iteration_t = std::unordered_map; +using kernel_iteration_t = std::unordered_map; using kernel_rename_map_t = std::unordered_map; using kernel_rename_stack_t = std::stack; -auto* tool_metadata = as_pointer(tool::metadata::inprocess{}); -auto target_kernels = common::Synchronized{}; -auto kernel_iteration = common::Synchronized{}; +auto* tool_metadata = as_pointer(tool::metadata::inprocess{}); +auto target_kernels = common::Synchronized{}; std::mutex att_shader_data; thread_local auto thread_dispatch_rename = as_pointer(); @@ -239,13 +238,13 @@ add_kernel_rename_and_stream_display_pairs(kernel_rename_and_stream_display_pair } bool -add_kernel_target(uint64_t _kern_id, const std::unordered_set& range) +add_kernel_target(uint64_t _kern_id, const std::unordered_set& range) { return target_kernels .wlock( - [](targeted_kernels_map_t& _targets_v, - uint64_t _kern_id_v, - const std::unordered_set& _range) { + [](targeted_kernels_map_t& _targets_v, + uint64_t _kern_id_v, + const std::unordered_set& _range) { return _targets_v.emplace(_kern_id_v, _range); }, _kern_id, @@ -254,10 +253,11 @@ add_kernel_target(uint64_t _kern_id, const std::unordered_set& range) } bool -is_targeted_kernel(uint64_t _kern_id) +is_targeted_kernel(uint64_t _kern_id, + common::Synchronized& _kernel_iteration) { - const std::unordered_set* range = target_kernels.rlock( - [](const auto& _targets_v, uint64_t _kern_id_v) -> const std::unordered_set* { + const std::unordered_set* range = target_kernels.rlock( + [](const auto& _targets_v, uint64_t _kern_id_v) -> const std::unordered_set* { if(_targets_v.find(_kern_id_v) != _targets_v.end()) return &_targets_v.at(_kern_id_v); return nullptr; }, @@ -265,21 +265,28 @@ is_targeted_kernel(uint64_t _kern_id) if(range) { - return kernel_iteration.rlock( - [](const auto& _kernel_iter, - uint64_t _kernel_id, - const std::unordered_set& _range) { - auto itr = _kernel_iter.at(_kernel_id); + _kernel_iteration.wlock( + [](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) { + auto itr = _kernel_iter.find(_kernel_id); + if(itr == _kernel_iter.end()) + _kernel_iter.emplace(_kernel_id, 1); + else + itr->second++; + }, + _kern_id); + return _kernel_iteration.rlock( + [](const auto& _kernel_iter, + uint64_t _kernel_id, + const std::unordered_set& _range) { + auto itr = _kernel_iter.at(_kernel_id); // If the iteration range is not given then all iterations of the kernel is profiled if(_range.empty()) { if(!tool::get_config().advanced_thread_trace) return true; - else - { - if(itr == 1) return true; - } + else if(itr == 1) + return true; } else if(_range.find(itr) != _range.end()) return true; @@ -1129,32 +1136,33 @@ get_instruction_index(rocprofiler_pc_t pc) } // namespace std::vector -get_att_perfcounter_params(std::vector& att_perf_counters) +get_att_perfcounter_params(rocprofiler_agent_id_t agent, + std::vector& att_perf_counters) { - std::vector _data; + std::vector _data{}; if(att_perf_counters.empty()) return _data; - static const auto gpu_agents = get_gpu_agents(); - static const auto gpu_agents_counter_info = get_agent_counter_info(); + static const auto agent_counter_info = get_agent_counter_info(); - for(const auto& [agent_, tool_counter_info_] : gpu_agents_counter_info) + for(const auto& att_perf_counter : att_perf_counters) { - for(const auto& counter_info_ : tool_counter_info_) - { - if(std::string_view(counter_info_.block) != "SQ") continue; + bool counter_found = false; - for(const auto& att_perf_counter : att_perf_counters) - { - if(std::string_view(counter_info_.name) == att_perf_counter.counter_name) - { - auto param = rocprofiler_att_parameter_t{}; - param.type = ROCPROFILER_ATT_PARAMETER_PERFCOUNTER, - param.counter_id = counter_info_.id, - param.simd_mask = att_perf_counter.simd_mask; - _data.emplace_back(param); - } - } + for(const auto& counter_info_ : agent_counter_info.at(agent)) + { + if(std::string_view(counter_info_.name) != att_perf_counter.counter_name) continue; + + auto param = rocprofiler_att_parameter_t{}; + param.type = ROCPROFILER_ATT_PARAMETER_PERFCOUNTER; + param.counter_id = counter_info_.id; + param.simd_mask = att_perf_counter.simd_mask; + _data.emplace_back(param); + counter_found = true; + break; } + + ROCP_WARNING_IF(!counter_found) + << "Agent " << agent.handle << " counter not found: " << att_perf_counter.counter_name; } return _data; @@ -1232,23 +1240,12 @@ att_dispatch_callback(rocprofiler_agent_id_t /* agent_id */, void* /*userdata_config*/, rocprofiler_user_data_t* userdata_shader) { - userdata_shader->value = dispatch_id; - kernel_iteration.wlock( - [](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) { - auto itr = _kernel_iter.find(_kernel_id); - if(itr == _kernel_iter.end()) - _kernel_iter.emplace(_kernel_id, 1); - else - { - itr->second++; - } - }, - kernel_id); + static auto kernel_iteration = common::Synchronized{}; - if(is_targeted_kernel(kernel_id)) - { + userdata_shader->value = dispatch_id; + + if(is_targeted_kernel(kernel_id, kernel_iteration)) return ROCPROFILER_ATT_CONTROL_START_AND_STOP; - } return ROCPROFILER_ATT_CONTROL_NONE; } @@ -1258,22 +1255,12 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_user_data_t* user_data, void* /*callback_data_args*/) { + static auto kernel_iteration = common::Synchronized{}; + auto kernel_id = dispatch_data.dispatch_info.kernel_id; auto agent_id = dispatch_data.dispatch_info.agent_id; - kernel_iteration.wlock( - [](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) { - auto itr = _kernel_iter.find(_kernel_id); - if(itr == _kernel_iter.end()) - _kernel_iter.emplace(_kernel_id, 1); - else - { - itr->second++; - } - }, - kernel_id); - - if(!is_targeted_kernel(kernel_id)) + if(!is_targeted_kernel(kernel_id, kernel_iteration)) { return; } @@ -1555,34 +1542,50 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) if(tool::get_config().advanced_thread_trace) { - auto parameters = std::vector{}; + auto global_parameters = std::vector{}; uint64_t target_cu = tool::get_config().att_param_target_cu; uint64_t simd_select = tool::get_config().att_param_simd_select; uint64_t buffer_sz = tool::get_config().att_param_buffer_size; uint64_t shader_mask = tool::get_config().att_param_shader_engine_mask; + uint64_t perfcounter_ctrl = tool::get_config().att_param_perf_ctrl; auto& att_perf = tool::get_config().att_param_perfcounters; bool att_serialize_all = tool::get_config().att_serialize_all; - auto att_perf_params = get_att_perfcounter_params(att_perf); - parameters.insert(parameters.end(), att_perf_params.begin(), att_perf_params.end()); - // TODO: att params could be different for different devices. How to support? - // Input file schema might also need to change to support multiple ATT params - - parameters.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, {target_cu}}); - parameters.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {simd_select}}); - parameters.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {buffer_sz}}); - parameters.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {shader_mask}}); - parameters.push_back( + global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, {target_cu}}); + global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {simd_select}}); + global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {buffer_sz}}); + global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {shader_mask}}); + global_parameters.push_back( {ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, {static_cast(att_serialize_all)}}); - ROCPROFILER_CALL( - rocprofiler_configure_dispatch_thread_trace_service(get_client_ctx(), - parameters.data(), - parameters.size(), - att_dispatch_callback, - att_shader_data_callback, - tool_data), - "thread trace service configure"); + if(perfcounter_ctrl != 0 && !att_perf.empty()) + { + global_parameters.push_back( + {ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL, {perfcounter_ctrl}}); + } + else if(perfcounter_ctrl != 0 || !att_perf.empty()) + { + ROCP_FATAL << "ATT Perf requires setting both perfcounter_ctrl and perfcounter list!"; + } + + for(auto& [id, agent] : tool_metadata->agents_map) + { + if(agent.type != ROCPROFILER_AGENT_TYPE_GPU) continue; + + auto agent_params = global_parameters; + for(auto& counter : get_att_perfcounter_params(id, att_perf)) + agent_params.push_back(counter); + + ROCPROFILER_CALL( + rocprofiler_configure_dispatch_thread_trace_service(get_client_ctx(), + id, + agent_params.data(), + agent_params.size(), + att_dispatch_callback, + att_shader_data_callback, + tool_data), + "thread trace service configure"); + } } if(tool::get_config().hip_runtime_api_trace || tool::get_config().hip_compiler_api_trace) @@ -1933,40 +1936,6 @@ tool_fini(void* /*tool_data*/) tool::generate_csv(tool::get_config(), *tool_metadata, contributions); } - if(tool::get_config().advanced_thread_trace) - { - const std::unordered_map - tool_att_capability_map = { - {"testing", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING}, - {"summary", rocprofiler::att_wrapper::ATT_CAPABILITIES_SUMMARY}, - {"trace", rocprofiler::att_wrapper::ATT_CAPABILITIES_TRACE}, - {"debug", rocprofiler::att_wrapper::ATT_CAPABILITIES_DEBUG}}; - - ROCP_FATAL_IF(tool::get_config().att_capability.empty()) - << "Provide the decoder parser method as input"; - - auto att_capability_value = tool_att_capability_map.at(tool::get_config().att_capability); - auto decoder = rocprofiler::att_wrapper::ATTDecoder(att_capability_value); - ROCP_FATAL_IF(!decoder.valid()) << "Decoder library not found at ROCPROF_ATT_LIBRARY_PATH"; - auto codeobj = tool_metadata->get_code_object_load_info(); - auto output_path = tool::format_path(tool::get_config().output_path); - for(auto& [dispatch_id, att_filename_data] : tool_metadata->att_filenames) - { - std::string formats = "json,csv"; - // if(tool::get_config().json_output) formats += "json,"; - // if(tool::get_config().csv_output) formats += "csv,"; - - std::stringstream ui_name; - ui_name << fmt::format("ui_output_agent_{}_dispatch_{}", - std::to_string(att_filename_data.first.handle), - dispatch_id); - auto out_path = fmt::format("{}/{}", output_path, ui_name.str()); - std::string in_path = "."; - - decoder.parse(in_path, out_path, att_filename_data.second, codeobj, formats); - } - } - if(tool::get_config().json_output && num_output > 0) { auto json_ar = tool::open_json(tool::get_config()); @@ -2044,6 +2013,51 @@ tool_fini(void* /*tool_data*/) { tool::generate_stats(tool::get_config(), *tool_metadata, contributions); } + + if(tool::get_config().advanced_thread_trace) + { + const std::unordered_map + tool_att_capability_map = { + {"testing1", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING1}, + {"testing2", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING2}, + {"trace", rocprofiler::att_wrapper::ATT_CAPABILITIES_TRACE}, + {"debug", rocprofiler::att_wrapper::ATT_CAPABILITIES_DEBUG}}; + + ROCP_FATAL_IF(tool::get_config().att_capability.empty()) + << "Provide the decoder parser method as input"; + + auto att_capability_value = tool_att_capability_map.at(tool::get_config().att_capability); + auto decoder = rocprofiler::att_wrapper::ATTDecoder(att_capability_value); + ROCP_FATAL_IF(!decoder.valid()) << "Decoder library not found at ROCPROF_ATT_LIBRARY_PATH"; + auto codeobj = tool_metadata->get_code_object_load_info(); + auto output_path = tool::format_path(tool::get_config().output_path); + + std::vector perf{}; + for(auto& counter : tool::get_config().att_param_perfcounters) + { + std::stringstream ss; + ss << counter.counter_name; + + if(counter.simd_mask != 0xF) ss << ':' << std::hex << counter.simd_mask; + + perf.emplace_back(ss.str()); + } + + for(auto& [dispatch_id, att_filename_data] : tool_metadata->att_filenames) + { + std::string formats = "json,csv"; + + auto ui_name = std::stringstream{}; + ui_name << fmt::format("ui_output_agent_{}_dispatch_{}", + std::to_string(att_filename_data.first.handle), + dispatch_id); + auto out_path = fmt::format("{}/{}", output_path, ui_name.str()); + auto in_path = std::string("."); + + decoder.parse(in_path, out_path, att_filename_data.second, codeobj, perf, formats); + } + } + auto destroy_output = [](auto& _buffered_output_v) { _buffered_output_v.destroy(); }; destroy_output(kernel_dispatch_with_stream_output); diff --git a/source/lib/rocprofiler-sdk/context/context.cpp b/source/lib/rocprofiler-sdk/context/context.cpp index 1a11be653e..8d5872e7ab 100644 --- a/source/lib/rocprofiler-sdk/context/context.cpp +++ b/source/lib/rocprofiler-sdk/context/context.cpp @@ -339,7 +339,7 @@ start_context(rocprofiler_context_id_t context_id) auto status = ROCPROFILER_STATUS_SUCCESS; if(cfg->counter_collection) rocprofiler::counters::start_context(cfg); - if(cfg->agent_thread_trace) cfg->agent_thread_trace->start_context(); + if(cfg->device_thread_trace) cfg->device_thread_trace->start_context(); if(cfg->dispatch_thread_trace) cfg->dispatch_thread_trace->start_context(); if(cfg->device_counter_collection) status = rocprofiler::counters::start_agent_ctx(cfg); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 @@ -374,7 +374,7 @@ stop_context(rocprofiler_context_id_t idx) rocprofiler::counters::stop_context(const_cast(_expected)); } - if(_expected->agent_thread_trace) _expected->agent_thread_trace->stop_context(); + if(_expected->device_thread_trace) _expected->device_thread_trace->stop_context(); if(_expected->dispatch_thread_trace) _expected->dispatch_thread_trace->stop_context(); diff --git a/source/lib/rocprofiler-sdk/context/context.hpp b/source/lib/rocprofiler-sdk/context/context.hpp index e7bcf5160c..13075e0ce5 100644 --- a/source/lib/rocprofiler-sdk/context/context.hpp +++ b/source/lib/rocprofiler-sdk/context/context.hpp @@ -131,7 +131,7 @@ struct context std::unique_ptr pc_sampler = {}; std::unique_ptr dispatch_thread_trace = {}; - std::unique_ptr agent_thread_trace = {}; + std::unique_ptr device_thread_trace = {}; template bool is_tracing(KindT _kind) const; diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 0a44962da1..cfb3a85fe8 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -46,7 +46,7 @@ namespace counters * * We return an AQLPacket containing the start/stop/read packets for injection. */ -std::unique_ptr +hsa::Queue::pkt_and_serialize_t queue_cb(const context::context* ctx, const std::shared_ptr& info, const hsa::Queue& queue, @@ -59,43 +59,25 @@ queue_cb(const context::context* ctx, { CHECK(info && ctx); - // 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(&queue) - .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); - } - }); - }; - // Packet generated when no instrumentation is performed. May contain serialization // packets/barrier packets (and can be empty). auto no_instrumentation = [&]() { auto ret_pkt = std::make_unique(); + info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); }); // If we have a counter collection context but it is not enabled, we still might need // to add barrier packets to transition from serialized -> unserialized execution. This // transition is coordinated by the serializer. - maybe_add_serialization(ret_pkt); - info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); }); return ret_pkt; }; - if(!ctx || !ctx->counter_collection) return nullptr; + if(!ctx || !ctx->counter_collection) return {nullptr, false}; bool is_enabled = false; ctx->counter_collection->enabled.rlock( [&](const auto& collect_ctx) { is_enabled = collect_ctx; }); - if(!is_enabled || !info->user_cb) - { - return no_instrumentation(); - } + if(!is_enabled || !info->user_cb) return {no_instrumentation(), true}; auto _corr_id_v = rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; @@ -133,10 +115,7 @@ queue_cb(const context::context* ctx, info->user_cb(dispatch_data, &req_profile, user_data, info->callback_args); - if(req_profile.handle == 0) - { - return no_instrumentation(); - } + if(req_profile.handle == 0) return {no_instrumentation(), true}; auto prof_config = get_controller().get_profile_cfg(req_profile); CHECK(prof_config); @@ -145,18 +124,15 @@ queue_cb(const context::context* ctx, auto status = info->get_packet(ret_pkt, prof_config); CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) << rocprofiler_get_status_string(status); - maybe_add_serialization(ret_pkt); - if(ret_pkt->empty) + if(!ret_pkt->empty) { - return ret_pkt; + ret_pkt->populate_before(); + ret_pkt->populate_after(); + for(auto& aql_pkt : ret_pkt->after_krn_pkt) + aql_pkt.completion_signal.handle = 0; } - ret_pkt->populate_before(); - ret_pkt->populate_after(); - for(auto& aql_pkt : ret_pkt->after_krn_pkt) - aql_pkt.completion_signal.handle = 0; - - return ret_pkt; + return {std::move(ret_pkt), true}; } /** @@ -169,7 +145,6 @@ completed_cb(const context::context* ctx, inst_pkt_t& pkts, kernel_dispatch::profiling_time dispatch_time) { - auto& session = *ptr_session; CHECK(info && ctx); std::shared_ptr prof_config; @@ -189,14 +164,8 @@ completed_cb(const context::context* ctx, } }); - if(!pkt) return; - - CHECK_NOTNULL(hsa::get_queue_controller()) - ->serializer(&session.queue) - .wlock([&](auto& serializer) { serializer.kernel_completion_signal(session.queue); }); - // We have no profile config, nothing to output. - if(!prof_config) return; + if(!pkt || !prof_config) return; completed_cb_params_t params{info, ptr_session, dispatch_time, prof_config, std::move(pkt)}; process_callback_data(std::move(params)); diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp index 3fa5c7505c..f03a424135 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp @@ -34,7 +34,7 @@ using ClientID = int64_t; using inst_pkt_t = common::container:: small_vector, ClientID>, 4>; -std::unique_ptr +hsa::Queue::pkt_and_serialize_t queue_cb(const context::context* ctx, const std::shared_ptr& info, const hsa::Queue& queue, diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index d909757ae5..896b88c397 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -444,8 +444,8 @@ TEST(core, check_callbacks) extern_ids, &corr_id); - ASSERT_TRUE(ret_pkt) << fmt::format("Expected a packet to be generated for - {}", - metric.name()); + ASSERT_TRUE(ret_pkt.pkt) + << fmt::format("Expected a packet to be generated for - {}", metric.name()); /** * Create the buffer and run test @@ -469,7 +469,7 @@ TEST(core, check_callbacks) counters::inst_pkt_t pkts; pkts.emplace_back( - std::make_pair(std::move(ret_pkt), static_cast(0))); + std::make_pair(std::move(ret_pkt.pkt), static_cast(0))); completed_cb(&ctx, cb_info, sess, pkts, kernel_dispatch::profiling_time{}); rocprofiler_flush_buffer(opt_buff_id); rocprofiler_destroy_buffer(opt_buff_id); diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index b43b1c5754..b196c58f19 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -137,6 +137,15 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) } }); + if(queue_info_session.is_serialized) + { + CHECK_NOTNULL(hsa::get_queue_controller()) + ->serializer(&queue_info_session.queue) + .wlock([&](auto& serializer) { + serializer.kernel_completion_signal(queue_info_session.queue); + }); + } + // Delete signals and packets, signal we have completed. if(queue_info_session.interrupt_signal.handle != 0u) { @@ -367,24 +376,36 @@ WriteInterceptor(const void* packets, // completed_cb_t) auto inst_pkt = inst_pkt_t{}; + // True if any service (ATT,SPM,CC) requests this dispatch to be serialized + bool bRequest_Serialize = false; + // Signal callbacks that a kernel_pkt is being enqueued queue.signal_callback([&](const auto& map) { for(const auto& [client_id, cb_pair] : map) { - if(auto maybe_pkt = cb_pair.first(queue, - kernel_pkt, - kernel_id, - dispatch_id, - &user_data, - tracing_data_v.external_correlation_ids, - corr_id)) - { - inst_pkt.push_back(std::make_pair(std::move(maybe_pkt), client_id)); - } + auto [packet, bSerial] = cb_pair.first(queue, + kernel_pkt, + kernel_id, + dispatch_id, + &user_data, + tracing_data_v.external_correlation_ids, + corr_id); + bRequest_Serialize |= bSerial; + if(packet) inst_pkt.push_back(std::make_pair(std::move(packet), client_id)); } }); bool inserted_before = false; + if(bRequest_Serialize) + { + inserted_before = true; + CHECK_NOTNULL(hsa::get_queue_controller()) + ->serializer(&queue) + .rlock([&](const auto& serializer) { + for(auto& s_pkt : serializer.kernel_dispatch(queue)) + transformed_packets.emplace_back(s_pkt.ext_amd_aql_pm4); + }); + } for(const auto& pkt_injection : inst_pkt) { for(const auto& pkt : pkt_injection.first->before_krn_pkt) @@ -460,7 +481,8 @@ WriteInterceptor(const void* packets, .correlation_id = corr_id, .kernel_pkt = kernel_pkt, .callback_record = callback_record, - .tracing_data = tracing_data_v}; + .tracing_data = tracing_data_v, + .is_serialized = bRequest_Serialize}; auto shared = std::make_shared(std::move(info_session)); @@ -521,7 +543,8 @@ Queue::Queue(const AgentCache& agent, << "Could not setup intercept profiler"; if(!context::get_registered_contexts([](const context::context* ctx) { - return (ctx->counter_collection || ctx->device_counter_collection); + return (ctx->counter_collection || ctx->device_counter_collection || + ctx->dispatch_thread_trace || ctx->device_thread_trace); }).empty()) { CHECK(_agent.cpu_pool().handle != 0); diff --git a/source/lib/rocprofiler-sdk/hsa/queue.hpp b/source/lib/rocprofiler-sdk/hsa/queue.hpp index ea9ff9972a..ed336f9263 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -73,17 +73,23 @@ public: using callback_t = void (*)(hsa_status_t status, hsa_queue_t* source, void* data); using queue_info_session_t = queue_info_session; - // Function prototype used to notify consumers that a kernel has been - // enqueued. An AQL packet can be returned that will be injected into - // the queue. - using queue_cb_t = std::function( - const Queue&, - const rocprofiler_packet&, - rocprofiler_kernel_id_t, - rocprofiler_dispatch_id_t, - rocprofiler_user_data_t*, - const queue_info_session_t::external_corr_id_map_t&, - const context::correlation_id*)>; + struct pkt_and_serialize_t + { + std::unique_ptr pkt{nullptr}; + bool request_serialize{false}; + }; + + // Function prototype used to notify consumers that a kernel has been enqueued. + // Pair first: An AQL packet can be returned that will be injected into the queue. + // Pair second: Boolean flag indicating the dispatch needs to be serialized. + using queue_cb_t = + std::function; // Signals the completion of the kernel packet. using completed_cb_t = std::function queue) { - for(const auto& itr : context::get_registered_contexts()) - { - if(auto* trace = itr->dispatch_thread_trace.get()) - trace->resource_init(queue->get_agent(), get_core_table(), get_ext_table()); - } - CHECK(queue); _callback_cache.wlock([&](auto& callbacks) { _queues.wlock([&](auto& map) { @@ -176,16 +170,6 @@ QueueController::destroy_queue(hsa_queue_t* id) { if(!id) return; - for(const auto& itr : context::get_registered_contexts()) - { - if(!itr->dispatch_thread_trace) continue; - - _queues.wlock([&](auto& map) { - if(map.find(id) != map.end()) - itr->dispatch_thread_trace->resource_deinit(map.at(id)->get_agent()); - }); - } - const auto* queue = get_queue(*id); // return if queue does not exist @@ -469,7 +453,7 @@ enable_queue_intercept() itr->is_tracing(ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY); if(itr->counter_collection || itr->pc_sampler || has_kernel_tracing || - has_scratch_reporting || itr->device_counter_collection || itr->agent_thread_trace || + has_scratch_reporting || itr->device_counter_collection || itr->device_thread_trace || itr->dispatch_thread_trace) return true; } diff --git a/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp b/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp index c5e28b0a1b..9cc0259068 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp @@ -64,6 +64,7 @@ struct queue_info_session rocprofiler_packet kernel_pkt = {}; callback_record_t callback_record = {}; tracing::tracing_data tracing_data = {}; + bool is_serialized = false; }; } // namespace hsa } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp b/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp index 93b2711569..ab943b4347 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp @@ -347,7 +347,9 @@ pc_sampling_service_finish_configuration(context::pc_sampling_service* service) rocprofiler_dispatch_id_t /*dispatch_id*/, rocprofiler_user_data_t*, const rocprofiler::hsa::Queue::queue_info_session_t::external_corr_id_map_t&, - const context::correlation_id*) { return nullptr; }, + const context::correlation_id*) { + return rocprofiler::hsa::Queue::pkt_and_serialize_t{}; + }, // Completion CB [](const rocprofiler::hsa::Queue& q, rocprofiler::hsa::rocprofiler_packet kern_pkt, diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index fe75c824b5..1be92e4e6c 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -890,6 +890,7 @@ rocprofiler_set_api_table(const char* name, // need to construct agent mappings before initializing the queue controller rocprofiler::agent::construct_agent_cache(hsa_api_table); + rocprofiler::thread_trace::initialize(hsa_api_table); rocprofiler::hsa::queue_controller_init(hsa_api_table); // Process agent ctx's that were started prior to HSA init rocprofiler::counters::device_counting_service_hsa_registration(); @@ -898,7 +899,6 @@ rocprofiler_set_api_table(const char* name, rocprofiler::hsa::memory_allocation_init(hsa_api_table->core_, lib_instance); rocprofiler::hsa::memory_allocation_init(hsa_api_table->amd_ext_, lib_instance); rocprofiler::code_object::initialize(hsa_api_table); - rocprofiler::thread_trace::initialize(hsa_api_table); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 if(runtime_pc_sampling_table) rocprofiler::pc_sampling::code_object::initialize(hsa_api_table); diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index 8f24f97f16..494c573d38 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -26,6 +26,7 @@ #include "lib/common/container/stable_vector.hpp" #include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" @@ -68,20 +69,6 @@ struct cbdata_t common::Synchronized> client; -CoreApiTable& -get_core() -{ - static CoreApiTable api{}; - return api; -} - -AmdExtTable& -get_ext() -{ - static AmdExtTable api{}; - return api; -} - bool thread_trace_parameter_pack::are_params_valid() const { @@ -110,14 +97,16 @@ class Signal public: Signal(hsa_ext_amd_aql_pm4_packet_t* packet) { - get_ext().hsa_amd_signal_create_fn(0, 0, nullptr, 0, &signal); + auto& core = *hsa::get_core_table(); + auto& ext = *hsa::get_amd_ext_table(); + ext.hsa_amd_signal_create_fn(0, 0, nullptr, 0, &signal); packet->completion_signal = signal; - get_core().hsa_signal_store_screlease_fn(signal, 1); + core.hsa_signal_store_screlease_fn(signal, 1); } ~Signal() { WaitOn(); - get_core().hsa_signal_destroy_fn(signal); + hsa::get_core_table()->hsa_signal_destroy_fn(signal); } Signal(Signal& other) = delete; Signal(const Signal& other) = delete; @@ -126,7 +115,7 @@ public: void WaitOn() const { - auto* wait_fn = get_core().hsa_signal_wait_scacquire_fn; + auto wait_fn = hsa::get_core_table()->hsa_signal_wait_scacquire_fn; while(wait_fn(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED) != 0) {} } @@ -136,10 +125,12 @@ public: }; std::unique_ptr -ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait) +ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait) const { + auto* core = hsa::get_core_table(); + std::unique_ptr signal{}; - const uint64_t write_idx = add_write_index_relaxed_fn(queue, 1); + const uint64_t write_idx = core->hsa_queue_add_write_index_relaxed_fn(queue, 1); size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t); // NOLINTNEXTLINE(performance-no-int-to-ptr) @@ -154,40 +145,44 @@ ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait) auto* header = reinterpret_cast*>(queue_slot); header->store(slot_data[0], std::memory_order_release); - signal_store_screlease_fn(queue->doorbell_signal, write_idx); + core->hsa_signal_store_screlease_fn(queue->doorbell_signal, write_idx); return signal; } ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params, - const hsa::AgentCache& cache, - const CoreApiTable& coreapi, - const AmdExtTable& ext) + rocprofiler_agent_id_t cache) : params(std::move(_params)) -, agent_id(cache.get_rocp_agent()->id) +, agent_id(cache) { - factory = std::make_unique(cache, this->params, coreapi, ext); + ROCP_TRACE << "Constructing ATT instance for agent " << agent_id.handle; + auto* core = hsa::get_core_table(); + auto* ext = hsa::get_amd_ext_table(); + + factory = std::make_unique( + *rocprofiler::agent::get_agent_cache(rocprofiler::agent::get_agent(agent_id)), + this->params, + *core, + *ext); control_packet = factory->construct_control_packet(); - auto status = coreapi.hsa_queue_create_fn(cache.get_hsa_agent(), - QUEUE_SIZE, - HSA_QUEUE_TYPE_SINGLE, - nullptr, - nullptr, - UINT32_MAX, - UINT32_MAX, - &this->queue); + auto hsa_agent = rocprofiler::agent::get_hsa_agent(agent_id); + CHECK(hsa_agent.has_value()); + + auto status = core->hsa_queue_create_fn(*hsa_agent, + QUEUE_SIZE, + HSA_QUEUE_TYPE_SINGLE, + nullptr, + nullptr, + UINT32_MAX, + UINT32_MAX, + &this->queue); if(status != HSA_STATUS_SUCCESS) { ROCP_ERROR << "Failed to create thread trace async queue"; this->queue = nullptr; } - queue_destroy_fn = coreapi.hsa_queue_destroy_fn; - signal_store_screlease_fn = coreapi.hsa_signal_store_screlease_fn; - add_write_index_relaxed_fn = coreapi.hsa_queue_add_write_index_relaxed_fn; - load_read_index_relaxed_fn = coreapi.hsa_queue_load_read_index_relaxed_fn; - codeobj_reg = std::make_unique( [this](rocprofiler_agent_id_t agent, uint64_t codeobj_id, uint64_t addr, uint64_t size) { if(agent == this->agent_id) this->load_codeobj(codeobj_id, addr, size); @@ -199,14 +194,15 @@ ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params, ThreadTracerQueue::~ThreadTracerQueue() { + ROCP_TRACE << "Destroying ATT Queue..."; std::unique_lock lk(trace_resources_mut); if(active_traces.load() < 1) { - if(queue_destroy_fn) queue_destroy_fn(this->queue); + hsa::get_core_table()->hsa_queue_destroy_fn(this->queue); return; } - ROCP_WARNING << "Thread tracer being destroyed with thread trace active"; + ROCP_CI_LOG(WARNING) << "Thread tracer being destroyed with thread trace active"; control_packet->clear(); control_packet->populate_after(); @@ -253,7 +249,10 @@ ThreadTracerQueue::iterate_data(aqlprofile_handle_t handle, rocprofiler_user_dat cb_dt.userdata = &data; auto status = aqlprofile_att_iterate_data(handle, thread_trace_callback, &cb_dt); - CHECK_HSA(status, "Failed to iterate ATT data"); + if(status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) + ROCP_WARNING << "Thread trace buffer full!"; + else + CHECK_HSA(status, "Failed to iterate ATT data"); active_traces.fetch_sub(1); } @@ -284,42 +283,36 @@ ThreadTracerQueue::unload_codeobj(code_object_id_t id) } void -DispatchThreadTracer::resource_init(const hsa::AgentCache& cache, - const CoreApiTable& coreapi, - const AmdExtTable& ext) +DispatchThreadTracer::resource_init() { - auto agent = cache.get_hsa_agent(); - std::unique_lock lk(agents_map_mut); + auto rocp_agents = rocprofiler::agent::get_agents(); - auto agent_it = agents.find(agent); - if(agent_it != agents.end()) + auto lk = std::unique_lock{agents_map_mut}; + + for(const auto* rocp_agent : rocp_agents) { - agent_it->second->active_queues.fetch_add(1); - return; - } + auto it = params.find(rocp_agent->id); + if(it == params.end()) continue; - auto new_tracer = std::make_unique(this->params, cache, coreapi, ext); - agents.emplace(agent, std::move(new_tracer)); + auto cache = rocprofiler::agent::get_hsa_agent(rocp_agent); + CHECK(cache.has_value()); + agents[*cache] = std::make_unique(it->second, rocp_agent->id); + } } void -DispatchThreadTracer::resource_deinit(const hsa::AgentCache& cache) +DispatchThreadTracer::resource_deinit() { - std::unique_lock lk(agents_map_mut); - - auto agent_it = agents.find(cache.get_hsa_agent()); - if(agent_it == agents.end()) return; - - if(agent_it->second->active_queues.fetch_sub(1) > 1) return; - - agents.erase(cache.get_hsa_agent()); + ROCP_TRACE << "Clearing agents"; + auto lk = std::unique_lock{agents_map_mut}; + agents.clear(); } /** * Callback we get from HSA interceptor when a kernel packet is being enqueued. * We return an AQLPacket containing the start/stop/read packets for injection. */ -std::unique_ptr +hsa::Queue::pkt_and_serialize_t DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t dispatch_id, @@ -332,73 +325,36 @@ DispatchThreadTracer::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(&queue) - .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_agent().get_rocp_agent()->id, - queue.get_id(), - rocprof_corr_id, - kernel_id, - dispatch_id, - params.callback_userdata.ptr, - user_data); - - if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) - { - auto empty = std::make_unique(); - if(params.bSerialize) 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->get_control(true); + if(it == agents.end()) return {nullptr, false}; + auto& agent = *CHECK_NOTNULL(it->second); + const auto& parameters = agent.params; + + auto control_flags = parameters.dispatch_cb_fn(queue.get_agent().get_rocp_agent()->id, + queue.get_id(), + rocprof_corr_id, + kernel_id, + dispatch_id, + parameters.callback_userdata.ptr, + user_data); + + if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) return {nullptr, parameters.bSerialize}; + + auto packet = agent.get_control(true); post_move_data.fetch_add(1); - maybe_add_serialization(packet); - packet->populate_before(); packet->populate_after(); - return packet; + return {std::move(packet), true}; } -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(&session.queue).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) { - std::unique_ptr signal{nullptr}; - if(params.bSerialize) signal = std::make_unique(session); - if(post_move_data.load() < 1) return; for(auto& aql_pkt : aql) @@ -414,8 +370,6 @@ DispatchThreadTracer::post_kernel_call(DispatchThreadTracer::inst_pkt_t& a auto it = agents.find(pkt->GetAgent()); if(it != agents.end() && it->second != nullptr) it->second->iterate_data(pkt->GetHandle(), session.user_data); - - if(!signal) std::make_unique(session); } } @@ -456,19 +410,22 @@ DispatchThreadTracer::start_context() void DispatchThreadTracer::stop_context() // NOLINT(readability-convert-member-functions-to-static) { - CHECK_NOTNULL(hsa::get_queue_controller())->disable_serialization(); + auto* controller = hsa::get_queue_controller(); + if(!controller) return; client.wlock([&](auto& client_id) { if(!client_id) return; // Remove our callbacks from HSA's queue controller - CHECK_NOTNULL(hsa::get_queue_controller())->remove_callback(*client_id); + controller->remove_callback(*client_id); client_id = std::nullopt; }); + + controller->disable_serialization(); } void -AgentThreadTracer::resource_init(const CoreApiTable& coreapi, const AmdExtTable& ext) +DeviceThreadTracer::resource_init() { auto rocp_agents = rocprofiler::agent::get_agents(); @@ -476,41 +433,36 @@ AgentThreadTracer::resource_init(const CoreApiTable& coreapi, const AmdExtTable& for(const auto* rocp_agent : rocp_agents) { - auto id = rocp_agent->id; - const auto* cache = rocprofiler::agent::get_agent_cache(rocp_agent); + auto it = params.find(CHECK_NOTNULL(rocp_agent)->id); + if(it == params.end()) continue; - if(tracers.find(id) != tracers.end()) - ROCP_WARNING << "Agent configured twice: " << id.handle; - else if(params.find(id) == params.end()) - ROCP_INFO << "Skipping agent " << id.handle; - else if(cache == nullptr) - ROCP_WARNING << "Invalid agent id: " << id.handle; - else - tracers[id] = std::make_unique(params.at(id), *cache, coreapi, ext); + agents[it->first] = std::make_unique(it->second, rocp_agent->id); } } void -AgentThreadTracer::resource_deinit() +DeviceThreadTracer::resource_deinit() { + ROCP_TRACE << "Clearing agents"; std::unique_lock lk(agent_mut); - tracers.clear(); + agents.clear(); } void -AgentThreadTracer::start_context() +DeviceThreadTracer::start_context() { + ROCP_TRACE << "Start device context"; std::unique_lock lk(agent_mut); - if(tracers.empty()) + if(agents.empty()) { - ROCP_FATAL << "Thread trace context not present for agent!"; + ROCP_WARNING << "Thread trace context not present for agent!"; return; } std::vector> wait_list{}; - for(auto& [_, tracer] : tracers) + for(auto& [_, tracer] : agents) { auto packet = tracer->get_control(true); packet->populate_before(); @@ -521,20 +473,20 @@ AgentThreadTracer::start_context() } void -AgentThreadTracer::stop_context() +DeviceThreadTracer::stop_context() { using wait_t = std::tuple>; std::unique_lock lk(agent_mut); - if(tracers.empty()) + if(agents.empty()) { - ROCP_FATAL << "Thread trace context not present for agent!"; + ROCP_WARNING << "Thread trace context not present for agent!"; return; } std::vector wait_list{}; - for(auto& [_, tracer] : tracers) + for(auto& [_, tracer] : agents) { auto packet = tracer->get_control(false); packet->populate_after(); @@ -553,25 +505,25 @@ AgentThreadTracer::stop_context() void initialize(HsaApiTable* table) { - assert(table->core_ && table->amd_ext_); - get_core() = *table->core_; - get_ext() = *table->amd_ext_; + ROCP_FATAL_IF(!table->core_ || !table->amd_ext_); code_object::initialize(table); for(auto& ctx : context::get_registered_contexts()) { - if(ctx->agent_thread_trace) - ctx->agent_thread_trace->resource_init(*table->core_, *table->amd_ext_); + if(ctx->device_thread_trace) ctx->device_thread_trace->resource_init(); + if(ctx->dispatch_thread_trace) ctx->dispatch_thread_trace->resource_init(); } } void finalize() { + ROCP_TRACE << "Finalize called"; for(auto& ctx : context::get_registered_contexts()) { - if(ctx->agent_thread_trace) ctx->agent_thread_trace->resource_deinit(); + if(ctx->device_thread_trace) ctx->device_thread_trace->resource_deinit(); + if(ctx->dispatch_thread_trace) ctx->dispatch_thread_trace->resource_deinit(); } code_object::finalize(); diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp index efb007b673..36f490358c 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp @@ -25,6 +25,7 @@ #include "lib/rocprofiler-sdk/context/correlation_id.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" #include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" +#include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp" #include "lib/rocprofiler-sdk/thread_trace/code_object.hpp" @@ -72,11 +73,9 @@ struct thread_trace_parameter_pack // GFX9 Only std::vector> perfcounters; - static constexpr size_t DEFAULT_SIMD = 0x7; - static constexpr size_t DEFAULT_PERFCOUNTER_SIMD_MASK = 0xF; - static constexpr size_t DEFAULT_SE_MASK = 0x21; - static constexpr size_t DEFAULT_BUFFER_SIZE = 0x8000000; - static constexpr size_t PERFCOUNTER_SIMD_MASK_SHIFT = 28; + static constexpr size_t DEFAULT_SIMD = 0xF; + static constexpr size_t DEFAULT_SE_MASK = 0x1; + static constexpr size_t DEFAULT_BUFFER_SIZE = 0x8000000; bool are_params_valid() const; }; @@ -86,10 +85,7 @@ class ThreadTracerQueue using code_object_id_t = uint64_t; public: - ThreadTracerQueue(thread_trace_parameter_pack _params, - const hsa::AgentCache&, - const CoreApiTable&, - const AmdExtTable&); + ThreadTracerQueue(thread_trace_parameter_pack _params, rocprofiler_agent_id_t); virtual ~ThreadTracerQueue(); void load_codeobj(code_object_id_t id, uint64_t addr, uint64_t size); @@ -98,17 +94,17 @@ public: std::unique_ptr get_control(bool bStart); void iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data); - hsa_queue_t* queue = nullptr; + 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; [[nodiscard]] std::unique_ptr Submit(hsa_ext_amd_aql_pm4_packet_t* packet, - bool bWait); + bool bWait) const; template [[nodiscard]] std::unique_ptr SubmitAndSignalLast(VecType vec) @@ -125,11 +121,6 @@ private: std::unique_ptr codeobj_reg{nullptr}; rocprofiler_agent_id_t agent_id; - - 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}; }; class DispatchThreadTracer @@ -139,17 +130,21 @@ class DispatchThreadTracer using inst_pkt_t = common::container::small_vector, 4>; public: - DispatchThreadTracer(thread_trace_parameter_pack _params) - : params(std::move(_params)) - {} + DispatchThreadTracer() = default; ~DispatchThreadTracer() = default; void start_context(); void stop_context(); - void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&); - void resource_deinit(const hsa::AgentCache&); + void resource_init(); + void resource_deinit(); - std::unique_ptr pre_kernel_call(const hsa::Queue& queue, + void add_agent(rocprofiler_agent_id_t agent, thread_trace_parameter_pack pack) + { + auto lk = std::unique_lock{agents_map_mut}; + params[agent] = std::move(pack); + } + + hsa::Queue::pkt_and_serialize_t pre_kernel_call(const hsa::Queue& queue, uint64_t kernel_id, rocprofiler_dispatch_id_t dispatch_id, rocprofiler_user_data_t* user_data, @@ -157,23 +152,22 @@ public: void post_kernel_call(inst_pkt_t& aql, const hsa::queue_info_session& session); - std::unordered_map> agents{}; + std::unordered_map> agents{}; + std::unordered_map params{}; std::shared_mutex agents_map_mut{}; std::atomic post_move_data{0}; - - thread_trace_parameter_pack params{}; }; -class AgentThreadTracer +class DeviceThreadTracer { public: - AgentThreadTracer() = default; - ~AgentThreadTracer() = default; + DeviceThreadTracer() = default; + ~DeviceThreadTracer() = default; void start_context(); void stop_context(); - void resource_init(const CoreApiTable&, const AmdExtTable&); + void resource_init(); void resource_deinit(); void add_agent(rocprofiler_agent_id_t id, thread_trace_parameter_pack _params) @@ -187,7 +181,7 @@ public: return params.find(id) != params.end(); } - std::map> tracers{}; + std::map> agents{}; std::map params{}; std::mutex agent_mut; diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp index affd174a99..71d5835a17 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp @@ -31,25 +31,30 @@ #include "rocprofiler-sdk/amd_detail/thread_trace.h" using DispatchThreadTracer = rocprofiler::thread_trace::DispatchThreadTracer; -using AgentThreadTracer = rocprofiler::thread_trace::AgentThreadTracer; +using DeviceThreadTracer = rocprofiler::thread_trace::DeviceThreadTracer; extern "C" { rocprofiler_status_t rocprofiler_configure_dispatch_thread_trace_service( rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_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) { + ROCP_TRACE << "Configuring Dispatch ATT for agent " << agent_id.handle; + 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_FOUND; - if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; - if(ctx->agent_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; + if(ctx->device_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; + + if(!ctx->dispatch_thread_trace) + ctx->dispatch_thread_trace = std::make_unique(); auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{}; @@ -92,19 +97,21 @@ rocprofiler_configure_dispatch_thread_trace_service( if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - ctx->dispatch_thread_trace = std::make_unique(pack); + ctx->dispatch_thread_trace->add_agent(agent_id, pack); return ROCPROFILER_STATUS_SUCCESS; } rocprofiler_status_t -rocprofiler_configure_agent_thread_trace_service( +rocprofiler_configure_device_thread_trace_service( rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id, rocprofiler_att_parameter_t* parameters, size_t num_parameters, - rocprofiler_agent_id_t agent, rocprofiler_att_shader_data_callback_t shader_callback, rocprofiler_user_data_t userdata) { + ROCP_TRACE << "Configuring Device ATT for agent " << agent_id.handle; + if(rocprofiler::registration::get_init_status() > -1) return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; @@ -112,7 +119,7 @@ rocprofiler_configure_agent_thread_trace_service( if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; - if(!ctx->agent_thread_trace) ctx->agent_thread_trace = std::make_unique(); + if(!ctx->device_thread_trace) ctx->device_thread_trace = std::make_unique(); auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{}; @@ -154,8 +161,7 @@ rocprofiler_configure_agent_thread_trace_service( if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - assert(ctx->agent_thread_trace); - ctx->agent_thread_trace->add_agent(agent, pack); + ctx->device_thread_trace->add_agent(agent_id, pack); return ROCPROFILER_STATUS_SUCCESS; } } diff --git a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index b514ea3f6c..1b7cd7cb82 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -27,6 +27,7 @@ #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include "lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" #include "lib/rocprofiler-sdk/registration.hpp" @@ -64,11 +65,13 @@ test_init() HsaApiTable table; table.amd_ext_ = &get_ext_table(); table.core_ = &get_api_table(); + rocprofiler::hsa::copy_table(table.core_, 0); + rocprofiler::hsa::copy_table(table.amd_ext_, 0); agent::construct_agent_cache(&table); hsa::get_queue_controller()->init(get_api_table(), get_ext_table()); return true; }; - [[maybe_unused]] static bool run_ince = init(); + [[maybe_unused]] static bool run_once = init(); } } // namespace rocprofiler @@ -85,6 +88,7 @@ TEST(thread_trace, resource_creation) auto agents = hsa::get_queue_controller()->get_supported_agents(); ASSERT_GT(agents.size(), 0); + for(const auto& [_, agent] : agents) { auto params = thread_trace::thread_trace_parameter_pack{}; @@ -101,14 +105,12 @@ TEST(thread_trace, resource_creation) { thread_trace::thread_trace_parameter_pack params{}; - thread_trace::DispatchThreadTracer tracer(std::move(params)); + thread_trace::DispatchThreadTracer tracer{}; for(const auto& [_, agent] : agents) - { - // Init twice to simulate two queues - tracer.resource_init(agent, get_api_table(), get_ext_table()); - tracer.resource_init(agent, get_api_table(), get_ext_table()); - } + tracer.add_agent(agent.get_rocp_agent()->id, params); + + tracer.resource_init(); for(auto& [_, agenttracer] : tracer.agents) { @@ -117,18 +119,13 @@ TEST(thread_trace, resource_creation) agenttracer->unload_codeobj(1); } - for(const auto& [_, agent] : agents) - { - // Deinit twice to remove both queues - tracer.resource_deinit(agent); - tracer.resource_deinit(agent); - } + tracer.resource_deinit(); } - hsa_shut_down(); } TEST(thread_trace, configure_test) { + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); registration::init_logging(); @@ -143,29 +140,35 @@ 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_dispatch_thread_trace_service( - ctx, - params.data(), - params.size(), - [](rocprofiler_agent_id_t, - rocprofiler_queue_id_t, - rocprofiler_correlation_id_t, - rocprofiler_kernel_id_t, - rocprofiler_dispatch_id_t, - void*, - rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, - [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, - nullptr); + auto agents = hsa::get_queue_controller()->get_supported_agents(); + ASSERT_GT(agents.size(), 0); + + for(auto& [_, agent] : agents) + { + rocprofiler_configure_dispatch_thread_trace_service( + ctx, + agent.get_rocp_agent()->id, + params.data(), + params.size(), + [](rocprofiler_agent_id_t, + rocprofiler_queue_id_t, + rocprofiler_correlation_id_t, + rocprofiler_kernel_id_t, + rocprofiler_dispatch_id_t, + void*, + rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, + [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, + nullptr); + } - ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); ROCPROFILER_CALL(rocprofiler_start_context(ctx), "context start failed"); ROCPROFILER_CALL(rocprofiler_stop_context(ctx), "context stop failed"); context::pop_client(1); - hsa_shut_down(); } TEST(thread_trace, perfcounters_configure_test) { + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); registration::init_logging(); @@ -194,36 +197,45 @@ TEST(thread_trace, perfcounters_configure_test) expected.insert({std::atoi(metric.event().c_str()), simd_mask}); } - rocprofiler_configure_dispatch_thread_trace_service( - ctx, - params.data(), - params.size(), - [](rocprofiler_agent_id_t, - rocprofiler_queue_id_t, - rocprofiler_correlation_id_t, - rocprofiler_kernel_id_t, - rocprofiler_dispatch_id_t, - void*, - rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, - [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, - nullptr); + auto agents = hsa::get_queue_controller()->get_supported_agents(); + ASSERT_GT(agents.size(), 0); + + for(auto& [_, agent] : agents) + { + rocprofiler_configure_dispatch_thread_trace_service( + ctx, + agent.get_rocp_agent()->id, + params.data(), + params.size(), + [](rocprofiler_agent_id_t, + rocprofiler_queue_id_t, + rocprofiler_correlation_id_t, + rocprofiler_kernel_id_t, + rocprofiler_dispatch_id_t, + void*, + rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; }, + [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, + nullptr); + } auto* context = rocprofiler::context::get_mutable_registered_context(ctx); auto* tracer = context->dispatch_thread_trace.get(); ASSERT_NE(tracer, nullptr); - ASSERT_EQ(tracer->params.perfcounter_ctrl, 1); - ASSERT_EQ(tracer->params.perfcounters.size(), 3); - for(const auto& param : tracer->params.perfcounters) - EXPECT_TRUE(expected.find(param) != expected.end()) - << "valid AQLprofile mask not generated for perfcounters"; + for(auto& [id, agent] : tracer->agents) + { + ASSERT_EQ(agent->params.perfcounter_ctrl, 1); + ASSERT_EQ(agent->params.perfcounters.size(), 3); + for(const auto& param : agent->params.perfcounters) + EXPECT_TRUE(expected.find(param) != expected.end()) + << "valid AQLprofile mask not generated for perfcounters"; + } context::pop_client(1); - hsa_shut_down(); } TEST(thread_trace, perfcounters_aql_options_test) { - hsa_init(); + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); registration::init_logging(); @@ -243,12 +255,11 @@ TEST(thread_trace, perfcounters_aql_options_test) _params.perfcounters.push_back({std::atoi(metric.event().c_str()), simd_mask}); _params.perfcounter_ctrl = 2; auto new_tracer = std::make_unique( - _params, begin(agents)->second, get_api_table(), get_ext_table()); + _params, begin(agents)->second.get_rocp_agent()->id); ASSERT_EQ(new_tracer->factory->aql_params.size(), sqtt_default_num_options + perf_counters.size()); context::pop_client(1); - hsa_shut_down(); } rocprofiler_status_t @@ -281,11 +292,11 @@ query_available_agents(rocprofiler_agent_version_t /* version */, params.push_back(att_param); } - rocprofiler_configure_agent_thread_trace_service( + rocprofiler_configure_device_thread_trace_service( *reinterpret_cast(ctx_ptr), + agent->id, params.data(), params.size(), - agent->id, [](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {}, rocprofiler_user_data_t{}); } @@ -294,6 +305,7 @@ query_available_agents(rocprofiler_agent_version_t /* version */, TEST(thread_trace, agent_configure_test) { + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); registration::init_logging(); diff --git a/tests/rocprofv3/CMakeLists.txt b/tests/rocprofv3/CMakeLists.txt index 6f3406c1cf..bfcf3a73e8 100644 --- a/tests/rocprofv3/CMakeLists.txt +++ b/tests/rocprofv3/CMakeLists.txt @@ -38,7 +38,7 @@ add_subdirectory(pc-sampling) add_subdirectory(collection-period) add_subdirectory(rocdecode-trace) add_subdirectory(rocjpeg-trace) -if(TARGET att_decoder_testing) +if(TARGET att_decoder_testing1) add_subdirectory(advanced-thread-trace) endif() add_subdirectory(hip-stream-display) diff --git a/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt b/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt index 7365ec2ec1..4d224c94be 100644 --- a/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt +++ b/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt @@ -18,12 +18,14 @@ find_package(rocprofiler-sdk REQUIRED) # hsa multiqueue dependency test +message("Add rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute") + add_test( NAME rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute COMMAND $ --log-level env --advanced-thread-trace 1 --att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD - --att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing + --att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing1 --att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o out --output-format json ${PRELOAD_ARGS} -- $) @@ -38,7 +40,7 @@ add_test( COMMAND $ --log-level env --advanced-thread-trace 1 --att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD - --att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing + --att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing1 --att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o out --output-format json ${PRELOAD_ARGS} -- $) @@ -96,8 +98,8 @@ set_tests_properties( "rocprofv3-test-hsa-multiqueue-att-json-execute" FAIL_REGULAR_EXPRESSION "AssertionError") -if(TARGET rocprofiler-sdk::att-decoder-testing AND TARGET - rocprofiler-sdk::att-decoder-summary) +if(TARGET rocprofiler-sdk::att-decoder-testing1 AND TARGET + rocprofiler-sdk::att-decoder-testing2) set(MISSING_TEST_DECODER_LIBS OFF) else() set(MISSING_TEST_DECODER_LIBS ON) @@ -144,7 +146,7 @@ add_test( NAME rocprofv3-test-att-library-path-cmd-line COMMAND $ --att --att-library-path - ${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse summary --log-level env --echo + ${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing2 --log-level env --echo -- sleep 0) set_tests_properties( @@ -156,7 +158,7 @@ add_test( NAME rocprofv3-test-att-library-path-cmd-line-will-fail COMMAND $ --att --att-library-path - ${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing --log-level env --echo + ${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing1 --log-level env --echo -- sleep 0) set_tests_properties( @@ -168,7 +170,7 @@ set_tests_properties( # Uses ROCPROF_ATT_LIBRARY_PATH to specify ATT library path # add_test(NAME rocprofv3-test-att-library-path-env-var - COMMAND $ --att --att-parse summary + COMMAND $ --att --att-parse testing2 --log-level env --echo -- sleep 0) set_tests_properties( @@ -178,7 +180,7 @@ set_tests_properties( "${MISSING_TEST_DECODER_LIBS}") add_test(NAME rocprofv3-test-att-library-path-env-var-will-fail - COMMAND $ --att --att-parse testing + COMMAND $ --att --att-parse testing1 --log-level env --echo -- sleep 0) set_tests_properties( @@ -193,3 +195,19 @@ set_tests_properties( ON DISABLED "${MISSING_TEST_DECODER_LIBS}") + +# +# Uses ATT and Counter Collection at the same time +# +add_test( + NAME rocprofv3-test-hsa-multiqueue-att-plus-pmc-execute + COMMAND + $ --log-level env --pmc SQ_WAVES + --advanced-thread-trace --att-parse testing1 -d + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o out --output-format json + ${PRELOAD_ARGS} -- $) + +set_tests_properties( + rocprofv3-test-hsa-multiqueue-att-plus-pmc-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT + LD_LIBRARY_PATH=${CMAKE_LIBRARY_OUTPUT_DIRECTORY}:$ENV{LD_LIBRARY_PATH}) diff --git a/tests/rocprofv3/advanced-thread-trace/att_input.json b/tests/rocprofv3/advanced-thread-trace/att_input.json index 03dea9e053..7cfa71d2ae 100644 --- a/tests/rocprofv3/advanced-thread-trace/att_input.json +++ b/tests/rocprofv3/advanced-thread-trace/att_input.json @@ -9,12 +9,13 @@ ], "truncate_kernels": true, "advanced_thread_trace": true, - "att_parse": "testing", + "att_parse": "testing1", "att_target_cu": 1, "att_shader_engine_mask": "0x11", "att_simd_select": "0x3", "att_buffer_size": "0x6000000", - "att_perfcounters": "SQ_WAVES:0x1 SQ_INSTS_VALU:0x3 SQ_INSTS_SALU:0xF" + "att_perfcounters": "SQ_WAVES:0x1 SQ_INSTS_VALU:0x3 SQ_INSTS_SALU", + "att_perfcounter_ctrl": 3 } ] } diff --git a/tests/rocprofv3/advanced-thread-trace/att_input.yml.in b/tests/rocprofv3/advanced-thread-trace/att_input.yml.in index c33de7b07d..5a494e7276 100644 --- a/tests/rocprofv3/advanced-thread-trace/att_input.yml.in +++ b/tests/rocprofv3/advanced-thread-trace/att_input.yml.in @@ -1,30 +1,30 @@ jobs: - advanced_thread_trace: True - att_parse: summary + att_parse: testing2 att_library_path: - @LIBRARY_OUTPUT_DIR@/att - advanced_thread_trace: True - att_parse: summary + att_parse: testing2 att_library_path: - @LIBRARY_OUTPUT_DIR@/att - @LIBRARY_OUTPUT_DIR@ - advanced_thread_trace: True - att_parse: testing + att_parse: testing1 att_library_path: - @LIBRARY_OUTPUT_DIR@/att - @LIBRARY_OUTPUT_DIR@ - advanced_thread_trace: True - att_parse: testing + att_parse: testing1 - advanced_thread_trace: True - att_parse: testing + att_parse: testing1 att_library_path: - @LIBRARY_OUTPUT_DIR@ - advanced_thread_trace: True - att_parse: testing + att_parse: testing1 att_library_path: - @LIBRARY_OUTPUT_DIR@ diff --git a/tests/rocprofv3/advanced-thread-trace/validate.py b/tests/rocprofv3/advanced-thread-trace/validate.py index 873c1c392c..e8c6c06498 100644 --- a/tests/rocprofv3/advanced-thread-trace/validate.py +++ b/tests/rocprofv3/advanced-thread-trace/validate.py @@ -19,11 +19,14 @@ def test_code_object_memory(code_object_file_path, json_data, output_path): data = json_data["rocprofiler-sdk-tool"] tool_memory_load = data["strings"]["code_object_snapshot_filenames"] gfx_pattern = "gfx[a-z0-9]+" - match = re.search(gfx_pattern, tool_memory_load[0]) + match = re.search(gfx_pattern, tool_memory_load[1]) assert match != None gpu_name = match.group(0) - tool_memory_load_1 = open(os.path.join(output_path, tool_memory_load[0]), "rb") - tool_memory_load_2 = open(os.path.join(output_path, tool_memory_load[1]), "rb") + + read_bytes = lambda filename: open(os.path.join(output_path, filename), "rb").read() + # Loads all saved code objects + tool_memory = [read_bytes(saved) for saved in tool_memory_load[1:]] + found = False for hsa_file in code_object_file_path["hsa_memory_load"]: @@ -33,11 +36,9 @@ def test_code_object_memory(code_object_file_path, json_data, output_path): if gpu == gpu_name: found = True - hsa_memory_load = open(hsa_file, "rb") - hsa_memory_fs = hsa_memory_load.read() - tool_memory_fs_1 = tool_memory_load_1.read() - tool_memory_fs_2 = tool_memory_load_2.read() - assert hsa_memory_fs == tool_memory_fs_2 or hsa_memory_fs == tool_memory_fs_1 + hsa_memory_bytes = open(hsa_file, "rb").read() + # Checks if hsa_file is one of the saved code objects + assert any([hsa_memory_bytes == fs for fs in tool_memory]) break assert found == True diff --git a/tests/thread-trace/agent.cpp b/tests/thread-trace/agent.cpp index aab994ec3f..3e45cd8f2d 100644 --- a/tests/thread-trace/agent.cpp +++ b/tests/thread-trace/agent.cpp @@ -123,12 +123,12 @@ query_available_agents(rocprofiler_agent_version_t /* version */, parameters.push_back({ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, 0}); ROCPROFILER_CALL( - rocprofiler_configure_agent_thread_trace_service(agent_ctx, - parameters.data(), - parameters.size(), - agent->id, - Callbacks::shader_data_callback, - user), + rocprofiler_configure_device_thread_trace_service(agent_ctx, + agent->id, + parameters.data(), + parameters.size(), + Callbacks::shader_data_callback, + user), "thread trace service configure"); } return ROCPROFILER_STATUS_SUCCESS; diff --git a/tests/thread-trace/multi_dispatch.cpp b/tests/thread-trace/multi_dispatch.cpp index 1277a19498..fcdc99a860 100644 --- a/tests/thread-trace/multi_dispatch.cpp +++ b/tests/thread-trace/multi_dispatch.cpp @@ -73,14 +73,36 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) std::vector params{}; params.push_back({ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, 1}); + std::vector agents{}; + 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"); + rocprofiler_query_available_agents( + ROCPROFILER_AGENT_INFO_VERSION_0, + [](rocprofiler_agent_version_t, const void** _agents, size_t _num_agents, void* _data) { + auto* agent_v = static_cast*>(_data); + for(size_t i = 0; i < _num_agents; ++i) + { + auto* agent = static_cast(_agents[i]); + if(agent->type == ROCPROFILER_AGENT_TYPE_GPU) agent_v->emplace_back(agent->id); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + sizeof(rocprofiler_agent_v0_t), + &agents), + "Failed to iterate agents"); + + for(auto id : agents) + { + ROCPROFILER_CALL( + rocprofiler_configure_dispatch_thread_trace_service(client_ctx, + id, + 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/tests/thread-trace/single_dispatch.cpp b/tests/thread-trace/single_dispatch.cpp index 36df4fde3c..c5432bc325 100644 --- a/tests/thread-trace/single_dispatch.cpp +++ b/tests/thread-trace/single_dispatch.cpp @@ -83,10 +83,36 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) tool_data), "code object tracing service configure"); + std::vector agents{}; + ROCPROFILER_CALL( - rocprofiler_configure_dispatch_thread_trace_service( - client_ctx, nullptr, 0, dispatch_callback, Callbacks::shader_data_callback, tool_data), - "thread trace service configure"); + rocprofiler_query_available_agents( + ROCPROFILER_AGENT_INFO_VERSION_0, + [](rocprofiler_agent_version_t, const void** _agents, size_t _num_agents, void* _data) { + auto* agent_v = static_cast*>(_data); + for(size_t i = 0; i < _num_agents; ++i) + { + auto* agent = static_cast(_agents[i]); + if(agent->type == ROCPROFILER_AGENT_TYPE_GPU) agent_v->emplace_back(agent->id); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + sizeof(rocprofiler_agent_v0_t), + &agents), + "Failed to iterate agents"); + + for(auto id : agents) + { + ROCPROFILER_CALL( + rocprofiler_configure_dispatch_thread_trace_service(client_ctx, + id, + nullptr, + 0, + 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),