From e898079a13d9aab5699dae92365db0eb70711c5e Mon Sep 17 00:00:00 2001 From: "Baraldi, Giovanni" Date: Wed, 23 Jul 2025 03:08:12 +0200 Subject: [PATCH 1/7] Thread trace and Trace Decoder API tests and samples (#416) * Adding test and samples to decoder * Fix sample * Formatting * Fix multi test * Disable sample * Fix tests * Format * Version fix * Locking the decoder * Add atomic * Review comments * Format * Adding readme * merge conflict and adding PCS+ATT test * Review comments * Properly disable PCS test * Update tests/rocprofv3/advanced-thread-trace/CMakeLists.txt * Adding back env var test * Name fix * Preload sample * Addressing review comments * Update docs --------- Co-authored-by: Giovanni Baraldi --- samples/CMakeLists.txt | 1 + samples/thread_trace/CMakeLists.txt | 79 ++++ samples/thread_trace/README.md | 26 ++ samples/thread_trace/agent.cpp | 391 ++++++++++++++++++ samples/thread_trace/main.cpp | 178 ++++++++ source/docs/api-reference/thread_trace.rst | 2 +- source/include/rocprofiler-sdk/cxx/hash.hpp | 2 + .../include/rocprofiler-sdk/cxx/operators.hpp | 6 + .../experimental/thread-trace/trace_decoder.h | 26 +- source/lib/att-tool/att_lib_wrapper.cpp | 6 +- source/lib/att-tool/att_lib_wrapper.hpp | 10 +- source/lib/att-tool/profile_interface.cpp | 6 +- source/lib/att-tool/profile_interface.hpp | 8 +- .../rocprofiler-sdk/thread_trace/decode.cpp | 117 +++--- .../advanced-thread-trace/CMakeLists.txt | 162 ++++---- .../advanced-thread-trace/att_input.yml.in | 6 +- .../att_input_will_fail.yml.in | 26 ++ tests/thread-trace/agent.cpp | 33 +- tests/thread-trace/multi_dispatch.cpp | 27 +- tests/thread-trace/single_dispatch.cpp | 50 +-- tests/thread-trace/trace_callbacks.cpp | 162 +++----- tests/thread-trace/trace_callbacks.hpp | 62 +-- 22 files changed, 982 insertions(+), 404 deletions(-) create mode 100644 samples/thread_trace/CMakeLists.txt create mode 100644 samples/thread_trace/README.md create mode 100644 samples/thread_trace/agent.cpp create mode 100644 samples/thread_trace/main.cpp create mode 100644 tests/rocprofv3/advanced-thread-trace/att_input_will_fail.yml.in diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index b2f2e9633a..1790a85cde 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -47,3 +47,4 @@ add_subdirectory(code_object_isa_decode) add_subdirectory(external_correlation_id_request) add_subdirectory(pc_sampling) add_subdirectory(openmp_target) +add_subdirectory(thread_trace) diff --git a/samples/thread_trace/CMakeLists.txt b/samples/thread_trace/CMakeLists.txt new file mode 100644 index 0000000000..0c9c805631 --- /dev/null +++ b/samples/thread_trace/CMakeLists.txt @@ -0,0 +1,79 @@ +# +# +# +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-thread-trace LANGUAGES CXX HIP) + +find_package(rocprofiler-sdk REQUIRED) +find_package(rocprofiler-sdk-roctx REQUIRED) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set_source_files_properties( + main.cpp PROPERTIES LANGUAGE HIP COMPILE_FLAGS "${CMAKE_HIP_FLAGS_RELWITHDEBINFO}") + +add_executable(thread-trace-sample) +target_sources(thread-trace-sample PRIVATE main.cpp agent.cpp) + +target_link_libraries( + thread-trace-sample + PUBLIC rocprofiler-sdk::samples-build-flags + PRIVATE rocprofiler-sdk::rocprofiler-sdk + rocprofiler-sdk::samples-common-library + rocprofiler-sdk::rocprofiler-sdk-dw + rocprofiler-sdk-roctx::rocprofiler-sdk-roctx + rocprofiler-sdk::rocprofiler-sdk-amd-comgr) + +add_test(NAME thread-trace-sample COMMAND $) + +find_library( + attdecoder_LIBRARY + NAMES rocprof-trace-decoder + HINTS ${ROCM_PATH} + PATHS ${ROCM_PATH} + PATH_SUFFIXES lib) + +if(attdecoder_LIBRARY) + cmake_path(GET attdecoder_LIBRARY PARENT_PATH attdecoder_LIB_DIR) +endif() + +find_package_handle_standard_args(attdecoder REQUIRED_VARS attdecoder_LIB_DIR + attdecoder_LIBRARY) + +set(IS_DISABLED ON) +if(attdecoder_FOUND) + set(IS_DISABLED OFF) +endif() + +rocprofiler_samples_get_preload_env(PRELOAD_ENV) +list(APPEND PRELOAD_ENV "ROCPROFILER_TRACE_DECODER_LIB_PATH=${attdecoder_LIB_DIR}") + +set_tests_properties( + thread-trace-sample PROPERTIES TIMEOUT 60 ENVIRONMENT ${PRELOAD_ENV} LABELS + "samples;thread-trace" DISABLED ${IS_DISABLED}) diff --git a/samples/thread_trace/README.md b/samples/thread_trace/README.md new file mode 100644 index 0000000000..48528a75a0 --- /dev/null +++ b/samples/thread_trace/README.md @@ -0,0 +1,26 @@ +# Thread Trace and ROCprof Trace Decoder + +## Services + +- Thread trace in device profiling mode +- ROCprof Trace Decoder decodes the received thread trace data +- Thread trace start/stop using roctx + +## Properties + +### [agent.cpp](agent.cpp): + +- Configures thread trace in all GPU agents found with `rocprofiler_configure_device_thread_trace_service` +- Waits until `roctxProfilerResume` is called to start thread trace +- Stops tracing at `roctxProfilerPause` +- Receives the trace data in `shader_data_callback` and calls `rocprofiler_trace_decode` to decode the data +- `rocprofiler_trace_decode` calls `parse` (a lambda) +- `parse` receives the dedecoded data and increments hitcount/latencies by pc address +- At application end, `tool_fini` calls `gen_output_stream` to write the top hotspots into `thread_trace.log` + +### [main.cpp](main.cpp): + +- Defines a few different kernels and runs them +- The first loop iteration warms up the kernels +- The second iteration calls `roctxProfilerResume` to start thread trace +- After the loop ends, `roctxProfilerPause` is called to stop tracing diff --git a/samples/thread_trace/agent.cpp b/samples/thread_trace/agent.cpp new file mode 100644 index 0000000000..e8d9240e29 --- /dev/null +++ b/samples/thread_trace/agent.cpp @@ -0,0 +1,391 @@ +// 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. +// +// undefine NDEBUG so asserts are implemented +#ifdef NDEBUG +# undef NDEBUG +#endif + +#include + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#define ROCPROFILER_CALL(result, msg) \ + if(auto ec = (result); ec != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::cerr << "rocprofiler-sdk error at " << __FILE__ << ":" << __LINE__ \ + << " :: " << #result << std::endl; \ + std::cerr << "rocprofiler-sdk error code " << ec << ": " \ + << rocprofiler_get_status_string(ec) << " :: " << msg << std::endl; \ + abort(); \ + } + +#define DECODER_CALL(result) \ + if(auto ec = (result); ec != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::cerr << "Decoder error at " << __FILE__ << ":" << __LINE__ << std::endl; \ + std::cerr << "rocprofiler-sdk error code " << ec << ": " \ + << rocprofiler_get_status_string(ec) << std::endl; \ + } + +#define CHECK_NOTNULL(x) \ + if(!(x)) \ + { \ + abort(); \ + }; + +namespace +{ +constexpr uint64_t TARGET_CU = 1; // CU (gfx9) or WGP (gfx10+) +constexpr uint64_t SHADER_MASK = 0x1; // Only enable SE=0 +constexpr uint64_t BUFFER_SIZE = 0x10000000; // 256MB +}; // namespace + +namespace Results +{ +using pcinfo_t = rocprofiler_thread_trace_decoder_pc_t; + +struct address_sort_t +{ + bool operator()(const pcinfo_t& a, const pcinfo_t& b) const + { + if(a.marker_id == b.marker_id) return a.addr < b.addr; + return a.marker_id < b.marker_id; + } +}; + +struct Latency +{ + uint64_t latency{0}; + uint64_t hitcount{0}; +}; + +// Maps address to latency +using LatencyTable = std::map; +// Used to disassemble instructions at (id, vaddr) pair +using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; + +AddressTable* table{nullptr}; +LatencyTable* latencies{nullptr}; + +// used to calculate mean wave lifetime +int64_t wave_lifetime = 0; +int64_t waves_started = 0; +int64_t waves_ended = 0; + +void +gen_output_stream() +{ + CHECK_NOTNULL(Results::latencies); + CHECK_NOTNULL(Results::table); + + const char* OUTPUT_OFSTREAM = "thread_trace.log"; + std::ofstream file(OUTPUT_OFSTREAM); + + if(!file.is_open()) + std::cout << "Could not open log file: " << OUTPUT_OFSTREAM << ", writing to stdout\n"; + else + std::cout << "Writing log to: " << OUTPUT_OFSTREAM << std::endl; + + std::ostream& output = file.is_open() ? file : std::cout; + + // Sort map by instruction cost + using Element = std::pair; + + std::vector sorted(latencies->begin(), latencies->end()); + std::stable_sort(sorted.begin(), sorted.end(), [](const Element& a, const Element& b) { + return a.second.latency > b.second.latency; + }); + + output << "Top 50 hotspots for trace (cycles):\n"; + for(size_t i = 0; i < sorted.size() && i < 50; i++) + { + auto& addr = sorted.at(i).first; + auto& latency = sorted.at(i).second; + auto inst = table->get(addr.marker_id, addr.addr); + + auto comment = inst->comment; + size_t pos = comment.rfind('/'); + if(pos != std::string::npos && pos + 1 < comment.size()) comment = comment.substr(pos + 1); + + output << "Latency:" << latency.latency << "\tHit:" << latency.hitcount << " \t" + << inst->inst << " [" << comment << "]\n"; + } + + if(waves_started != waves_ended) + std::cerr << "Error: Some waves have not ended!" << std::endl; + else if(waves_started == 0) + std::cerr << "Error: No waves started!" << std::endl; + else + output << "\nMean wave lifetime: " << wave_lifetime / waves_started << " cycles"; + + output << "\nWaves started: " << waves_started << "\nWaves ended: " << waves_ended << "\n"; +}; +} // namespace Results + +namespace Decoder +{ +rocprofiler_thread_trace_decoder_id_t decoder{}; + +void +tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* /* user_data */, + void* /* userdata */) +{ + if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return; + if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return; + + CHECK_NOTNULL(Results::table); + auto* data = static_cast(record.payload); + + if(data->storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) + { + Results::table->addDecoder( + data->uri, data->code_object_id, data->load_delta, data->load_size); + return; + } + + auto* memorybase = reinterpret_cast(data->memory_base); + CHECK_NOTNULL(memorybase); + + DECODER_CALL(rocprofiler_thread_trace_decoder_codeobj_load(decoder, + data->code_object_id, + data->load_delta, + data->load_size, + memorybase, + data->memory_size)); + + Results::table->addDecoder( + memorybase, data->memory_size, data->code_object_id, data->load_delta, data->load_size); +} + +void +shader_data_callback(rocprofiler_agent_id_t /* agent */, + int64_t /* se_id */, + void* se_data, + size_t data_size, + rocprofiler_user_data_t /* userdata */) +{ + CHECK_NOTNULL(Results::latencies); + + auto parse = [](rocprofiler_thread_trace_decoder_record_type_t record_type_id, + void* events, + uint64_t num_events, + void* /* userdata */) { + if(record_type_id == ROCPROFILER_THREAD_TRACE_DECODER_RECORD_OCCUPANCY) + { + for(size_t i = 0; i < num_events; i++) + { + auto& event = static_cast(events)[i]; + + if(event.start) + { + Results::wave_lifetime -= static_cast(event.time); + Results::waves_started++; + } + else + { + Results::wave_lifetime += static_cast(event.time); + Results::waves_ended++; + } + } + } + + if(record_type_id != ROCPROFILER_THREAD_TRACE_DECODER_RECORD_WAVE) return; + + for(size_t w = 0; w < num_events; w++) + { + auto* wave = static_cast(events); + for(size_t i = 0; i < wave->instructions_size; i++) + { + auto& inst = wave->instructions_array[i]; + auto& latency = (*Results::latencies)[inst.pc]; + latency.latency += inst.duration; + latency.hitcount += 1; + } + } + }; + + DECODER_CALL(rocprofiler_trace_decode(decoder, parse, se_data, data_size, nullptr)); +} + +} // namespace Decoder + +namespace ThreadTracer +{ +rocprofiler_client_id_t* client_id = nullptr; +rocprofiler_context_id_t agent_ctx = {}; +rocprofiler_context_id_t tracing_ctx = {}; + +rocprofiler_status_t +query_available_agents(rocprofiler_agent_version_t /* version */, + const void** agents, + size_t num_agents, + void* user_data) +{ + rocprofiler_user_data_t user{}; + user.ptr = user_data; + + for(size_t idx = 0; idx < num_agents; idx++) + { + const auto* agent = static_cast(agents[idx]); + if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; + + auto parameters = std::vector{}; + parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU, TARGET_CU}); + parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, BUFFER_SIZE}); + parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, SHADER_MASK}); + + ROCPROFILER_CALL( + rocprofiler_configure_device_thread_trace_service(agent_ctx, + agent->id, + parameters.data(), + parameters.size(), + Decoder::shader_data_callback, + user), + "thread trace service configure"); + } + return ROCPROFILER_STATUS_SUCCESS; +} + +void +cntrl_tracing_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* /* user_data */, + void* /* cb_data */) +{ + if(record.kind != ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API) return; + + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER && + record.operation == ROCPROFILER_MARKER_CONTROL_API_ID_roctxProfilerPause) + { + ROCPROFILER_CALL(rocprofiler_stop_context(agent_ctx), "stopping context"); + } + else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT && + record.operation == ROCPROFILER_MARKER_CONTROL_API_ID_roctxProfilerResume) + { + ROCPROFILER_CALL(rocprofiler_start_context(agent_ctx), "starting context"); + } +} + +int +tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) +{ + Results::latencies = new Results::LatencyTable{}; + Results::table = new Results::AddressTable{}; + + // This is set by ctests: TODO: move to client.cpp + // If nullptr, searches rocprofiler-sdk install location + const char* lib_path = std::getenv("ROCPROFILER_TRACE_DECODER_LIB_PATH"); + + DECODER_CALL(rocprofiler_thread_trace_decoder_create(&Decoder::decoder, lib_path)); + + ROCPROFILER_CALL(rocprofiler_create_context(&tracing_ctx), "context creation"); + ROCPROFILER_CALL(rocprofiler_create_context(&agent_ctx), "context creation"); + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(tracing_ctx, + ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, + nullptr, + 0, + Decoder::tool_codeobj_tracing_callback, + nullptr), + "code object tracing service configure"); + + ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service( + tracing_ctx, + ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API, + nullptr, + 0, + cntrl_tracing_callback, + nullptr), + "marker tracing callback service configure"); + + ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, + &query_available_agents, + sizeof(rocprofiler_agent_t), + nullptr), + "Failed to find GPU agents"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(agent_ctx, &valid_ctx), "validity check"); + assert(valid_ctx != 0); + ROCPROFILER_CALL(rocprofiler_context_is_valid(tracing_ctx, &valid_ctx), "validity check"); + assert(valid_ctx != 0); + + ROCPROFILER_CALL(rocprofiler_start_context(tracing_ctx), "context start"); + + // no errors + return 0; +} + +void +tool_fini(void* /* tool_data */) +{ + rocprofiler_thread_trace_decoder_destroy(Decoder::decoder); + + Results::gen_output_stream(); + + delete Results::latencies; + delete Results::table; +} + +} // namespace ThreadTracer + +extern "C" rocprofiler_tool_configure_result_t* +rocprofiler_configure(uint32_t /* version */, + const char* /* runtime_version */, + uint32_t priority, + rocprofiler_client_id_t* id) +{ + // only activate if main tool + if(priority > 0) return nullptr; + + // set the client name + id->name = "Thread Trace Sample"; + + // store client info + ThreadTracer::client_id = id; + + // create configure data + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + &ThreadTracer::tool_init, + &ThreadTracer::tool_fini, + nullptr}; + + // return pointer to configure data + return &cfg; +} diff --git a/samples/thread_trace/main.cpp b/samples/thread_trace/main.cpp new file mode 100644 index 0000000000..e189b41472 --- /dev/null +++ b/samples/thread_trace/main.cpp @@ -0,0 +1,178 @@ +// 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 + +#include +#include +#include +#include + +#include + +#include "hip/hip_runtime.h" + +// Two waves per SIMD on MI300 +#define DATA_SIZE (304 * 64 * 4 * 2) +#define HIP_API_CALL(CALL) \ + if((CALL) != hipSuccess) \ + { \ + abort(); \ + } + +#define LDS_SIZE 1024 + +__global__ void +divide_kernel(float* a, const float* b, const float* c, int /* unused */) +{ + int index = blockDim.x * blockIdx.x + threadIdx.x; + + if(index >= DATA_SIZE) return; + + a[index] = (b[index] - c[index]) / abs(c[index] + b[index]) + 1; +} + +__global__ void +looping_lds_kernel(float* a, const float* b, const float* c, int loopcount) +{ + __shared__ float interm[LDS_SIZE]; + + size_t index = blockDim.x * blockIdx.x + threadIdx.x; + + for(size_t i = index; i < DATA_SIZE; i += blockDim.x * gridDim.x) + interm[threadIdx.x % LDS_SIZE] = b[index] + threadIdx.x; + + for(int it = 0; it < loopcount; it++) + { + __syncthreads(); + float value = interm[(it + threadIdx.x + LDS_SIZE / 2) % LDS_SIZE]; + __syncthreads(); + interm[threadIdx.x % LDS_SIZE] += value; + } + + a[index] = interm[threadIdx.x % LDS_SIZE] + c[index]; +} + +__global__ void +fifo_kernel(float* /* a */, const float* /* b */, const float* /* c */, int loops) +{ + using _float4 = __attribute__((__vector_size__(4 * sizeof(float)))) float; + + __shared__ _float4 lds[LDS_SIZE]; + lds[threadIdx.x] = _float4{float(threadIdx.x)}; + lds[threadIdx.x + 512] = _float4{float(threadIdx.x)}; + + __syncthreads(); + + _float4 dst[16]; + + float res1 = 0, res2 = 0; + + for(int l = 0; l < loops; l++) + { +#pragma unroll 16 + for(int i = 0; i < 16; i++) + dst[i] = lds[threadIdx.x + i * 8]; + + __syncthreads(); + +#pragma unroll 16 + for(int i = 0; i < 16; i++) + { + res1 += dst[i][0] + dst[i][1]; + res2 += dst[i][2] + dst[i][3]; + } + asm volatile("v_add_f32 %0, %1, %2" : "=v"(res1) : "v"(res1), "v"(res2)); + } +}; + +class hipMemory +{ +public: + hipMemory(size_t size = DATA_SIZE) + { + HIP_API_CALL(hipMalloc(&ptr, size * sizeof(float))); + HIP_API_CALL(hipMemset(ptr, 0, size * sizeof(float))); + } + ~hipMemory() + { + if(ptr) HIP_API_CALL(hipFree(ptr)); + } + hipMemory(hipMemory&& other) + { + ptr = other.ptr; + other.ptr = nullptr; + } + float* ptr = nullptr; +}; + +class HipStream +{ +public: + HipStream() { HIP_API_CALL(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); } + ~HipStream() { HIP_API_CALL(hipStreamDestroy(stream)); } + + hipStream_t stream; + + hipMemory src1{}; + hipMemory src2{}; + hipMemory dst{}; +}; + +#define Launch(kernel, stream, arglast) \ + hipLaunchKernelGGL( \ + kernel, DATA_SIZE / 512, 512, 0, 0, stream.dst.ptr, stream.src1.ptr, stream.src2.ptr, 6); + +int +main(int /*argc*/, char** /*argv*/) +{ + std::array streams{}; + std::vector kernels{}; + + kernels.push_back(divide_kernel); + kernels.push_back(looping_lds_kernel); + kernels.push_back(fifo_kernel); + + for(size_t i = 0; i < streams.size() * kernels.size(); i++) + { + // Warmup then start + if(i == streams.size()) + { + HIP_API_CALL(hipDeviceSynchronize()); + roctxProfilerResume(0); + } + + auto& stream = streams.at(i % streams.size()); + auto& kernel = kernels.at(i % kernels.size()); + + Launch(kernel, stream, 3); + HIP_API_CALL(hipGetLastError()); + } + + HIP_API_CALL(hipDeviceSynchronize()); + roctxProfilerPause(0); + + return 0; +} diff --git a/source/docs/api-reference/thread_trace.rst b/source/docs/api-reference/thread_trace.rst index 78e1343076..6531fb9fdd 100644 --- a/source/docs/api-reference/thread_trace.rst +++ b/source/docs/api-reference/thread_trace.rst @@ -188,7 +188,7 @@ To decode the raw thread trace data, create and initialize a Trace Decoder: .. code-block:: cpp - rocprofiler_thread_trace_decoder_handle_t decoder{}; + rocprofiler_thread_trace_decoder_id_t decoder{}; // Create the Trace Decoder with the path to the decoder library ROCPROFILER_CALL( diff --git a/source/include/rocprofiler-sdk/cxx/hash.hpp b/source/include/rocprofiler-sdk/cxx/hash.hpp index 6ba6eade8f..200a48f26f 100644 --- a/source/include/rocprofiler-sdk/cxx/hash.hpp +++ b/source/include/rocprofiler-sdk/cxx/hash.hpp @@ -23,6 +23,7 @@ #pragma once +#include #include #include #include @@ -70,6 +71,7 @@ ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_signal_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_executable_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_region_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_thread_trace_decoder_id_t) #undef ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER } // namespace std diff --git a/source/include/rocprofiler-sdk/cxx/operators.hpp b/source/include/rocprofiler-sdk/cxx/operators.hpp index 2eaecd4c49..d09956e124 100644 --- a/source/include/rocprofiler-sdk/cxx/operators.hpp +++ b/source/include/rocprofiler-sdk/cxx/operators.hpp @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -148,6 +149,7 @@ ROCPROFILER_CXX_DECLARE_OPERATORS(const rocprofiler_counter_record_dimension_inf ROCPROFILER_CXX_DECLARE_OPERATORS(const rocprofiler_counter_record_dimension_instance_info_t&) ROCPROFILER_CXX_DECLARE_OPERATORS(const rocprofiler_counter_dimension_info_t&) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_version_triplet_t) +ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_thread_trace_decoder_id_t) // definitions of operator== ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_context_id_t) @@ -164,6 +166,7 @@ ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_signal_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_executable_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_region_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_thread_trace_decoder_id_t) inline bool operator==(const rocprofiler_agent_v0_t& lhs, const rocprofiler_agent_v0_t& rhs) @@ -241,6 +244,7 @@ ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_dim3_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_region_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_amd_memory_pool_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_version_triplet_t) +ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_thread_trace_decoder_id_t) // definitions of operator< ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_context_id_t) @@ -257,6 +261,7 @@ ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_signal_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_executable_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_region_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_thread_trace_decoder_id_t) inline bool operator<(const rocprofiler_counter_record_dimension_info_t& lhs, @@ -334,6 +339,7 @@ ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_dim3_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_region_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_amd_memory_pool_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_version_triplet_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_thread_trace_decoder_id_t) // cleanup defines #undef ROCPROFILER_CXX_DECLARE_OPERATORS diff --git a/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder.h b/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder.h index 746e21278b..b2e34f2c1e 100644 --- a/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder.h +++ b/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder.h @@ -36,10 +36,10 @@ ROCPROFILER_EXTERN_C_INIT /** * @brief Handle containing a loaded rocprof-trace-decoder and a decoder state. */ -typedef struct rocprofiler_thread_trace_decoder_handle_t +typedef struct rocprofiler_thread_trace_decoder_id_t { uint64_t handle; -} rocprofiler_thread_trace_decoder_handle_t; +} rocprofiler_thread_trace_decoder_id_t; /** * @brief Initializes Trace Decoder library with a library search path @@ -51,7 +51,7 @@ typedef struct rocprofiler_thread_trace_decoder_handle_t * @retval ::ROCPROFILER_STATUS_SUCCESS Handle created */ rocprofiler_status_t -rocprofiler_thread_trace_decoder_create(rocprofiler_thread_trace_decoder_handle_t* handle, +rocprofiler_thread_trace_decoder_create(rocprofiler_thread_trace_decoder_id_t* handle, const char* path) ROCPROFILER_API ROCPROFILER_NONNULL(1, 2); /** @@ -59,7 +59,7 @@ rocprofiler_thread_trace_decoder_create(rocprofiler_thread_trace_decoder_handle_ * @param[in] handle Handle to destroy */ void -rocprofiler_thread_trace_decoder_destroy(rocprofiler_thread_trace_decoder_handle_t handle) +rocprofiler_thread_trace_decoder_destroy(rocprofiler_thread_trace_decoder_id_t handle) ROCPROFILER_API; /** @@ -83,11 +83,11 @@ rocprofiler_thread_trace_decoder_destroy(rocprofiler_thread_trace_decoder_handle * @retval ::ROCPROFILER_STATUS_SUCCESS Code object loaded */ rocprofiler_status_t -rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_handle_t handle, - uint64_t load_id, - uint64_t load_addr, - uint64_t load_size, - const void* data, +rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_id_t handle, + uint64_t load_id, + uint64_t load_addr, + uint64_t load_size, + const void* data, uint64_t size) ROCPROFILER_API ROCPROFILER_NONNULL(5); /** @@ -100,7 +100,7 @@ rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_h * @retval ::ROCPROFILER_STATUS_SUCCESS Code object unloaded */ rocprofiler_status_t -rocprofiler_thread_trace_decoder_codeobj_unload(rocprofiler_thread_trace_decoder_handle_t handle, +rocprofiler_thread_trace_decoder_codeobj_unload(rocprofiler_thread_trace_decoder_id_t handle, uint64_t load_id) ROCPROFILER_API; /** @@ -135,7 +135,7 @@ typedef void (*rocprofiler_thread_trace_decoder_callback_t)( * @retval ::ROCPROFILER_STATUS_SUCCESS on success */ rocprofiler_status_t -rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_handle_t handle, +rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_id_t handle, rocprofiler_thread_trace_decoder_callback_t callback, void* data, uint64_t size, @@ -148,8 +148,8 @@ rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_handle_t handle, * @retval null terminated string as description of "info". */ const char* -rocprofiler_thread_trace_decoder_info_string(rocprofiler_thread_trace_decoder_handle_t handle, - rocprofiler_thread_trace_decoder_info_t info) +rocprofiler_thread_trace_decoder_info_string(rocprofiler_thread_trace_decoder_id_t handle, + rocprofiler_thread_trace_decoder_info_t info) ROCPROFILER_API; /** @} */ diff --git a/source/lib/att-tool/att_lib_wrapper.cpp b/source/lib/att-tool/att_lib_wrapper.cpp index e2bf7842c0..d07800b36f 100644 --- a/source/lib/att-tool/att_lib_wrapper.cpp +++ b/source/lib/att-tool/att_lib_wrapper.cpp @@ -41,9 +41,9 @@ namespace rocprofiler { namespace att_wrapper { -ATTFileMgr::ATTFileMgr(Fspath _dir, - std::vector _counters, - rocprofiler_thread_trace_decoder_handle_t _decoder) +ATTFileMgr::ATTFileMgr(Fspath _dir, + std::vector _counters, + rocprofiler_thread_trace_decoder_id_t _decoder) : dir(std::move(_dir)) , decoder(_decoder) { diff --git a/source/lib/att-tool/att_lib_wrapper.hpp b/source/lib/att-tool/att_lib_wrapper.hpp index 58e3f6deb6..b5a8dcb493 100644 --- a/source/lib/att-tool/att_lib_wrapper.hpp +++ b/source/lib/att-tool/att_lib_wrapper.hpp @@ -71,7 +71,7 @@ public: bool valid() const; protected: - rocprofiler_thread_trace_decoder_handle_t decoder{}; + rocprofiler_thread_trace_decoder_id_t decoder{}; }; class ATTFileMgr @@ -79,9 +79,9 @@ class ATTFileMgr using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; public: - ATTFileMgr(Fspath _dir, - std::vector _counters, - rocprofiler_thread_trace_decoder_handle_t _decoder); + ATTFileMgr(Fspath _dir, + std::vector _counters, + rocprofiler_thread_trace_decoder_id_t _decoder); ~ATTFileMgr(); void addDecoder(const char* filepath, uint64_t id, uint64_t load_addr, uint64_t memsize); @@ -95,7 +95,7 @@ public: std::shared_ptr table{nullptr}; std::map> occupancy{}; std::vector codeobjs_to_delete{}; - rocprofiler_thread_trace_decoder_handle_t decoder{}; + rocprofiler_thread_trace_decoder_id_t decoder{}; std::array, ROCPROFILER_THREAD_TRACE_DECODER_WSTATE_LAST> wstates; diff --git a/source/lib/att-tool/profile_interface.cpp b/source/lib/att-tool/profile_interface.cpp index f20348617d..1d50c422df 100644 --- a/source/lib/att-tool/profile_interface.cpp +++ b/source/lib/att-tool/profile_interface.cpp @@ -103,9 +103,9 @@ get_trace_data(rocprofiler_thread_trace_decoder_record_type_t trace_id, C_API_END } -ToolData::ToolData(std::vector& _data, - WaveConfig& _config, - rocprofiler_thread_trace_decoder_handle_t _decoder) +ToolData::ToolData(std::vector& _data, + WaveConfig& _config, + rocprofiler_thread_trace_decoder_id_t _decoder) : cfile(_config.code) , config(_config) , decoder(_decoder) diff --git a/source/lib/att-tool/profile_interface.hpp b/source/lib/att-tool/profile_interface.hpp index dc08f1c28d..54d38835b1 100644 --- a/source/lib/att-tool/profile_interface.hpp +++ b/source/lib/att-tool/profile_interface.hpp @@ -60,9 +60,9 @@ using SymbolInfo = rocprofiler::sdk::codeobj::disassembly::SymbolInfo; struct ToolData { - ToolData(std::vector& data, - WaveConfig& config, - rocprofiler_thread_trace_decoder_handle_t decoder); + ToolData(std::vector& data, + WaveConfig& config, + rocprofiler_thread_trace_decoder_id_t decoder); ~ToolData(); CodeLine& get(pcinfo_t pc); @@ -73,7 +73,7 @@ struct ToolData std::vector shader_data{}; size_t num_waves = 0; - rocprofiler_thread_trace_decoder_handle_t decoder{}; + rocprofiler_thread_trace_decoder_id_t decoder{}; }; } // namespace att_wrapper diff --git a/source/lib/rocprofiler-sdk/thread_trace/decode.cpp b/source/lib/rocprofiler-sdk/thread_trace/decode.cpp index 16ef7bfd15..f0e0980d7e 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/decode.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/decode.cpp @@ -21,6 +21,7 @@ // SOFTWARE. #include "lib/common/static_object.hpp" +#include "lib/common/synchronized.hpp" #include "lib/rocprofiler-sdk/aql/helpers.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" @@ -29,15 +30,19 @@ #include #include +#include +#include #include +#include #include namespace { using DL = rocprofiler::thread_trace::DL; using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate; +using LockedTable = rocprofiler::common::Synchronized; class DecoderInstance { @@ -46,72 +51,72 @@ public: : dl(std::move(_dl)) {} - std::unique_ptr
dl{nullptr}; - AddressTable table{}; + const std::unique_ptr dl{nullptr}; + + LockedTable table{}; }; -std::mutex map_mut; +using DecoderMap = + std::unordered_map>; +using LockedMap = rocprofiler::common::Synchronized; auto& -get_dlopens() +get_dlmap() { - static auto*& _v = rocprofiler::common::static_object< - std::unordered_map>>::construct(); + static auto*& _v = rocprofiler::common::static_object::construct(); return *CHECK_NOTNULL(_v); } std::shared_ptr -get_dl(rocprofiler_thread_trace_decoder_handle_t handle) +get_dl(rocprofiler_thread_trace_decoder_id_t handle) { - auto lk = std::unique_lock{map_mut}; - auto it = get_dlopens().find(handle.handle); - if(it == get_dlopens().end()) return nullptr; - - return it->second; + return get_dlmap().rlock([&](const DecoderMap& map) -> std::shared_ptr { + if(auto it = map.find(handle); it != map.end()) return it->second; + return nullptr; + }); } } // namespace extern "C" { rocprofiler_status_t -rocprofiler_thread_trace_decoder_create(rocprofiler_thread_trace_decoder_handle_t* handle, - const char* path) +rocprofiler_thread_trace_decoder_create(rocprofiler_thread_trace_decoder_id_t* handle, + const char* path) { auto dl = std::make_unique
(path); if(dl->handle == nullptr) return ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE; if(!dl->valid()) return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI; - auto lk = std::unique_lock{map_mut}; - static uint64_t count = 1; + static std::atomic count{1}; + handle->handle = count.fetch_add(1); - auto instance = std::make_shared(std::move(dl)); - - handle->handle = count++; - get_dlopens()[handle->handle] = std::move(instance); + get_dlmap().wlock( + [&](DecoderMap& map) { map[*handle] = std::make_shared(std::move(dl)); }); return ROCPROFILER_STATUS_SUCCESS; } void -rocprofiler_thread_trace_decoder_destroy(rocprofiler_thread_trace_decoder_handle_t handle) +rocprofiler_thread_trace_decoder_destroy(rocprofiler_thread_trace_decoder_id_t handle) { - auto lk = std::unique_lock{map_mut}; - get_dlopens().erase(handle.handle); + get_dlmap().wlock([&](DecoderMap& map) { map.erase(handle); }); } rocprofiler_status_t -rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_handle_t handle, - uint64_t load_id, - uint64_t load_addr, - uint64_t load_size, - const void* data, - uint64_t size) +rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_id_t handle, + uint64_t load_id, + uint64_t load_addr, + uint64_t load_size, + const void* data, + uint64_t size) { auto decoder = get_dl(handle); if(decoder == nullptr) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; try { - decoder->table.addDecoder(data, size, load_id, load_addr, load_size); + decoder->table.wlock([&](AddressTable& table) { + table.addDecoder(data, size, load_id, load_addr, load_size); + }); } catch(...) { return ROCPROFILER_STATUS_ERROR; @@ -120,15 +125,17 @@ rocprofiler_thread_trace_decoder_codeobj_load(rocprofiler_thread_trace_decoder_h } rocprofiler_status_t -rocprofiler_thread_trace_decoder_codeobj_unload(rocprofiler_thread_trace_decoder_handle_t handle, - uint64_t load_id) +rocprofiler_thread_trace_decoder_codeobj_unload(rocprofiler_thread_trace_decoder_id_t handle, + uint64_t load_id) { auto decoder = get_dl(handle); if(decoder == nullptr) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; try { - if(decoder->table.removeDecoder(load_id)) return ROCPROFILER_STATUS_SUCCESS; + bool result = + decoder->table.wlock([&](AddressTable& table) { return table.removeDecoder(load_id); }); + if(result) return ROCPROFILER_STATUS_SUCCESS; } catch(std::exception&) {} @@ -169,32 +176,32 @@ isa_callback(char* isa_instruction, void* userdata) { ROCP_FATAL_IF(userdata == nullptr) << "Userdata is null!"; - auto& table = static_cast(userdata)->decoder->table; - - std::unique_ptr instruction{nullptr}; + auto decoder = static_cast(userdata)->decoder; + ROCP_FATAL_IF(decoder == nullptr) << "decoder is null"; try { - instruction = table.get(pc.marker_id, pc.addr); + auto instruction = decoder->table.wlock( + [&](AddressTable& table) { return table.get(pc.marker_id, pc.addr); }); + + if(!instruction) return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_ERROR_INVALID_ARGUMENT; + + { + size_t tmp_isa_size = *isa_size; + *isa_size = instruction->inst.size(); + + if(*isa_size > tmp_isa_size) + return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_ERROR_OUT_OF_RESOURCES; + } + + memcpy(isa_instruction, instruction->inst.data(), *isa_size); + *isa_memory_size = instruction->size; + } catch(std::exception& e) { - ROCP_WARNING << pc.marker_id << ":" << pc.addr << ' ' << e.what(); + ROCP_CI_LOG(INFO) << pc.marker_id << ":" << pc.addr << ' ' << e.what(); return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_ERROR; } - - if(!instruction) return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_ERROR_INVALID_ARGUMENT; - - { - size_t tmp_isa_size = *isa_size; - *isa_size = instruction->inst.size(); - - if(*isa_size > tmp_isa_size) - return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_ERROR_OUT_OF_RESOURCES; - } - - memcpy(isa_instruction, instruction->inst.data(), *isa_size); - *isa_memory_size = instruction->size; - return ROCPROFILER_THREAD_TRACE_DECODER_STATUS_SUCCESS; } @@ -214,7 +221,7 @@ trace_callback(rocprofiler_thread_trace_decoder_record_type_t record_type_id, extern "C" { rocprofiler_status_t -rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_handle_t handle, +rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_id_t handle, rocprofiler_thread_trace_decoder_callback_t user_callback, void* data, uint64_t size, @@ -249,8 +256,8 @@ rocprofiler_trace_decode(rocprofiler_thread_trace_decoder_handle_t handle, } const char* -rocprofiler_thread_trace_decoder_info_string(rocprofiler_thread_trace_decoder_handle_t handle, - rocprofiler_thread_trace_decoder_info_t info) +rocprofiler_thread_trace_decoder_info_string(rocprofiler_thread_trace_decoder_id_t handle, + rocprofiler_thread_trace_decoder_info_t info) { auto decoder = get_dl(handle); if(decoder == nullptr) return nullptr; diff --git a/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt b/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt index d48dc04315..7a938b5dde 100644 --- a/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt +++ b/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt @@ -42,8 +42,6 @@ rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.p find_package(rocprofiler-sdk REQUIRED) -set(IS_DISABLED ON) - find_library( attdecoder_LIBRARY NAMES rocprof-trace-decoder @@ -58,43 +56,43 @@ endif() find_package_handle_standard_args(attdecoder REQUIRED_VARS attdecoder_LIB_DIR attdecoder_LIBRARY) +set(COMMON_PARAMS_NO_LIB ${PRELOAD_ARGS} --att --output-format json --log-level env -d + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/) + +set(IS_DISABLED ON) if(attdecoder_FOUND) set(IS_DISABLED OFF) - set(LIB_PATH_ENV "ROCPROF_ATT_LIBRARY_PATH=${attdecoder_LIB_DIR}") + set(COMMON_PARAMS --att-library-path ${attdecoder_LIB_DIR} ${COMMON_PARAMS_NO_LIB}) endif() # hsa multiqueue dependency test with lib path add_test( - NAME rocprofv3-test-hsa-multiqueue-att-cmd-env-att-lib-path-execute + NAME rocprofv3-test-att-hsa-multiqueue-cmd-env-att-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-serialize-all 1 -d - ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o out --output-format json - ${PRELOAD_ARGS} --att-library-path ${attdecoder_LIB_DIR} --att-activity 8 -- + $ ${COMMON_PARAMS}/cmd_input + --att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD + --att-buffer-size 0x6000000 --att-simd-select 0x3 --att-serialize-all 1 -o out -- $) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-cmd-env-att-lib-path-execute + rocprofv3-test-att-hsa-multiqueue-cmd-env-att-lib-path-execute PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED}) # hsa multiqueue dependency test with json input add_test( - NAME rocprofv3-test-hsa-multiqueue-att-json-execute + NAME rocprofv3-test-att-hsa-multiqueue-json-execute COMMAND - $ --log-level env --att-library-path - ${attdecoder_LIB_DIR} -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/json_input -i - ${CMAKE_CURRENT_BINARY_DIR}/att_input.json ${PRELOAD_ARGS} -- + $ ${COMMON_PARAMS}/json_input -i + ${CMAKE_CURRENT_BINARY_DIR}/att_input.json -- $) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-json-execute - PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED} ENVIRONMENT - "${LIB_PATH_ENV}") + rocprofv3-test-att-hsa-multiqueue-json-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED}) # validate output add_test( - NAME rocprofv3-test-hsa-multiqueue-att-cmd-validate + NAME rocprofv3-test-att-hsa-multiqueue-cmd-validate COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input ${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/cmd_input/out_results.json @@ -102,7 +100,7 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/cmd_input) add_test( - NAME rocprofv3-test-hsa-multiqueue-att-json-validate + NAME rocprofv3-test-att-hsa-multiqueue-json-validate COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input ${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/json_input/out_results.json @@ -110,46 +108,46 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/json_input) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-cmd-validate + rocprofv3-test-att-hsa-multiqueue-cmd-validate PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS - "rocprofv3-test-hsa-multiqueue-att-cmd-ld-lib-path-execute" + "rocprofv3-test-att-hsa-multiqueue-cmd-ld-lib-path-execute" FAIL_REGULAR_EXPRESSION "AssertionError" DISABLED ${IS_DISABLED}) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-json-validate + rocprofv3-test-att-hsa-multiqueue-json-validate PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS - "rocprofv3-test-hsa-multiqueue-att-json-execute" + "rocprofv3-test-att-hsa-multiqueue-json-execute" FAIL_REGULAR_EXPRESSION "AssertionError" DISABLED ${IS_DISABLED}) -function(configure_att_input _FILENAME _OUTDIR) - set(LIBRARY_OUTPUT_DIR ${_OUTDIR}) - configure_file(${CMAKE_CURRENT_SOURCE_DIR}/att_input.yml.in +function(configure_att_input _FILENAME) + set(LIB_PATH_LOC "${attdecoder_LIB_DIR}") + configure_file(${CMAKE_CURRENT_SOURCE_DIR}//${_FILENAME}.in ${CMAKE_CURRENT_BINARY_DIR}/${_FILENAME} @ONLY) endfunction() -configure_att_input(att_input.yml "${attdecoder_LIB_DIR}") -configure_att_input(att_input_will_fail.yml "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") +configure_att_input(att_input.yml) +configure_att_input(att_input_will_fail.yml) # test yaml input add_test( NAME rocprofv3-test-att-yaml-input COMMAND - $ -i - ${CMAKE_CURRENT_BINARY_DIR}/att_input.yml --log-level env --echo -- + $ ${COMMON_PARAMS_NO_LIB}/yaml -i + ${CMAKE_CURRENT_BINARY_DIR}/att_input.yml --echo -- $) set_tests_properties( @@ -166,7 +164,7 @@ add_test( set_tests_properties( rocprofv3-test-att-yaml-input-will-fail - PROPERTIES TIMEOUT 45 LABELS "integration-tests" WILL_FAIL ON DISABLED True) + PROPERTIES TIMEOUT 45 LABELS "integration-tests" WILL_FAIL ON DISABLED ${IS_DISABLED}) add_test(NAME rocprofv3-test-att-library-path-cmd-line-will-fail COMMAND $ --att --att-library-path . @@ -174,11 +172,12 @@ add_test(NAME rocprofv3-test-att-library-path-cmd-line-will-fail set_tests_properties( rocprofv3-test-att-library-path-cmd-line-will-fail - PROPERTIES TIMEOUT 45 LABELS "integration-tests" WILL_FAIL ON DISABLED True) + PROPERTIES TIMEOUT 45 LABELS "integration-tests" WILL_FAIL ON DISABLED ${IS_DISABLED}) -add_test(NAME rocprofv3-test-att-library-path-env-var-will-fail - COMMAND $ --att --log-level env --echo - -- $) +add_test( + NAME rocprofv3-test-att-library-path-env-var-will-fail + COMMAND $ ${COMMON_PARAMS_NO_LIB}/envfail + --att-library-path . --echo -- $) set_tests_properties( rocprofv3-test-att-library-path-env-var-will-fail @@ -187,65 +186,70 @@ set_tests_properties( LABELS "integration-tests" ENVIRONMENT - "ROCPROF_ATT_LIBRARY_PATH=." + "ROCPROF_ATT_LIBRARY_PATH=${attdecoder_LIB_DIR}" WILL_FAIL ON DISABLED - True) + ${IS_DISABLED}) + +add_test(NAME rocprofv3-test-att-env-var + COMMAND $ ${COMMON_PARAMS_NO_LIB}/envvar + -- $) + +set_tests_properties( + rocprofv3-test-att-env-var + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT + "ROCPROF_ATT_LIBRARY_PATH=${attdecoder_LIB_DIR}" DISABLED ${IS_DISABLED}) # # 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 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o - out --output-format json ${PRELOAD_ARGS} -- $) +add_test(NAME rocprofv3-test-att-hsa-multiqueue-plus-pmc-execute + COMMAND $ ${COMMON_PARAMS}/cmd_input + --pmc SQ_WAVES -o out -- $) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-plus-pmc-execute - PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED} ENVIRONMENT - "${LIB_PATH_ENV}") + rocprofv3-test-att-hsa-multiqueue-plus-pmc-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED}) # Check for conflict PMC + activity -add_test( - NAME rocprofv3-test-hsa-multiqueue-att-activity-pmc-will-fail - COMMAND - $ --log-level env --pmc SQ_WAVES - --advanced-thread-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o - out --att-activity 8 ${PRELOAD_ARGS} -- $) +add_test(NAME rocprofv3-test-att-hsa-multiqueue-activity-pmc-will-fail + COMMAND $ ${COMMON_PARAMS}/cmd_input + --pmc SQ_WAVES -o out --att-activity 8 -- $) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-activity-pmc-will-fail - PROPERTIES TIMEOUT - 45 - LABELS - "integration-tests" - DISABLED - ${IS_DISABLED} - ENVIRONMENT - "${LIB_PATH_ENV}" - WILL_FAIL - ON) + rocprofv3-test-att-hsa-multiqueue-activity-pmc-will-fail + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED} WILL_FAIL ON) # Check for conflict Perfcounters + activity add_test( - NAME rocprofv3-test-hsa-multiqueue-att-activity-perf-will-fail - COMMAND - $ --log-level env --att - --att-perfcounter-ctrl 8 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o - out --att-activity 8 ${PRELOAD_ARGS} -- $) + NAME rocprofv3-test-att-hsa-multiqueue-activity-perf-will-fail + COMMAND $ ${COMMON_PARAMS}/cmd_input + --att-perfcounter-ctrl 8 -o out --att-activity 8 -- $) set_tests_properties( - rocprofv3-test-hsa-multiqueue-att-activity-perf-will-fail - PROPERTIES TIMEOUT - 45 - LABELS - "integration-tests" - DISABLED - ${IS_DISABLED} - ENVIRONMENT - "${LIB_PATH_ENV}" - WILL_FAIL - ON) + rocprofv3-test-att-hsa-multiqueue-activity-perf-will-fail + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DISABLED ${IS_DISABLED} WILL_FAIL ON) + +# ATT plus PC sampling + +rocprofiler_sdk_pc_sampling_stochastic_disabled(IS_PC_SAMPLING_STOCHASTIC_DISABLED) + +set(ATT_PLUS_PCS_DISABLE ${IS_DISABLED}) +if(${IS_PC_SAMPLING_STOCHASTIC_DISABLED}) + set(ATT_PLUS_PCS_DISABLE ON) +endif() +if(${ROCPROFILER_DISABLE_UNSTABLE_CTESTS}) + set(ATT_PLUS_PCS_DISABLE ON) +endif() +set(PCS_ARGS --pc-sampling-beta-enabled 1 --pc-sampling-unit cycles + --pc-sampling-interval 16384 --pc-sampling-method stochastic) + +add_test(NAME rocprofv3-test-att-plus-pc-sampling + COMMAND $ ${COMMON_PARAMS}/cmd_input + ${PCS_ARGS} -o out -- $) + +set_tests_properties( + rocprofv3-test-att-plus-pc-sampling + PROPERTIES TIMEOUT 90 LABELS "integration-tests;thread-trace;pc-sampling" DISABLED + ${ATT_PLUS_PCS_DISABLE}) diff --git a/tests/rocprofv3/advanced-thread-trace/att_input.yml.in b/tests/rocprofv3/advanced-thread-trace/att_input.yml.in index d80dd3e1f6..e46bfc1c48 100644 --- a/tests/rocprofv3/advanced-thread-trace/att_input.yml.in +++ b/tests/rocprofv3/advanced-thread-trace/att_input.yml.in @@ -23,9 +23,9 @@ jobs: - advanced_thread_trace: True att_library_path: - - @LIBRARY_OUTPUT_DIR@/att - - @LIBRARY_OUTPUT_DIR@ + - @LIB_PATH_LOC@/att + - @LIB_PATH_LOC@ - advanced_thread_trace: True att_library_path: - - @LIBRARY_OUTPUT_DIR@ + - @LIB_PATH_LOC@ diff --git a/tests/rocprofv3/advanced-thread-trace/att_input_will_fail.yml.in b/tests/rocprofv3/advanced-thread-trace/att_input_will_fail.yml.in new file mode 100644 index 0000000000..b3bc337db3 --- /dev/null +++ b/tests/rocprofv3/advanced-thread-trace/att_input_will_fail.yml.in @@ -0,0 +1,26 @@ +# 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. + +jobs: + - advanced_thread_trace: True + att_library_path: + - @LIB_PATH_LOC@/att diff --git a/tests/thread-trace/agent.cpp b/tests/thread-trace/agent.cpp index 4d57af5945..708ba44968 100644 --- a/tests/thread-trace/agent.cpp +++ b/tests/thread-trace/agent.cpp @@ -29,17 +29,6 @@ #include -#define C_API_BEGIN \ - try \ - { -#define C_API_END \ - } \ - catch(std::exception & e) \ - { \ - std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl; \ - } \ - catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; } - namespace ATTTest { namespace Agent @@ -135,9 +124,10 @@ query_available_agents(rocprofiler_agent_version_t /* version */, } int -tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) +tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) { - (void) fini_func; + Callbacks::init(); + ROCPROFILER_CALL(rocprofiler_create_context(&tracing_ctx), "context creation"); ROCPROFILER_CALL(rocprofiler_create_context(&agent_ctx), "context creation"); @@ -147,7 +137,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) nullptr, 0, Callbacks::tool_codeobj_tracing_callback, - tool_data), + nullptr), "code object tracing service configure"); ROCPROFILER_CALL( @@ -156,13 +146,13 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) nullptr, 0, dispatch_tracing_callback, - tool_data), + nullptr), "dispatch tracing service configure"); ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, &query_available_agents, sizeof(rocprofiler_agent_t), - tool_data), + nullptr), "Failed to find GPU agents"); int valid_ctx = 0; @@ -177,13 +167,6 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) return 0; } -void -tool_fini(void* tool_data) -{ - Callbacks::finalize_json(tool_data); - delete static_cast(tool_data); -} - } // namespace Agent } // namespace ATTTest @@ -206,8 +189,8 @@ rocprofiler_configure(uint32_t /* version */, static auto cfg = rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), &ATTTest::Agent::tool_init, - &ATTTest::Agent::tool_fini, - new Callbacks::ToolData{"att_agent_test/"}}; + &Callbacks::finalize, + nullptr}; // return pointer to configure data return &cfg; diff --git a/tests/thread-trace/multi_dispatch.cpp b/tests/thread-trace/multi_dispatch.cpp index bf47704b5e..234024e9d7 100644 --- a/tests/thread-trace/multi_dispatch.cpp +++ b/tests/thread-trace/multi_dispatch.cpp @@ -48,15 +48,17 @@ dispatch_callback(rocprofiler_agent_id_t /* agent */, static std::atomic count{0}; if(count.fetch_add(1) > NUM_KERNELS) return ROCPROFILER_THREAD_TRACE_CONTROL_NONE; - assert(userdata && "Dispatch callback passed null!"); + assert(dispatch_userdata && "Dispatch callback passed null!"); dispatch_userdata->ptr = userdata; return ROCPROFILER_THREAD_TRACE_CONTROL_START_AND_STOP; } int -tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) +tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) { + Callbacks::init(); + static rocprofiler_context_id_t client_ctx = {0}; ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); @@ -67,7 +69,7 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) nullptr, 0, Callbacks::tool_codeobj_tracing_callback, - tool_data), + nullptr), "code object tracing service configure"); std::vector params{}; @@ -100,7 +102,7 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) params.size(), dispatch_callback, Callbacks::shader_data_callback, - tool_data), + nullptr), "thread trace service configure"); } @@ -121,13 +123,6 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) return 0; } -void -tool_fini(void* tool_data) -{ - Callbacks::finalize_json(tool_data); - delete static_cast(tool_data); -} - } // namespace Multi } // namespace ATTTest @@ -147,11 +142,11 @@ rocprofiler_configure(uint32_t /* version */, ATTTest::Multi::client_id = id; // create configure data - static auto cfg = rocprofiler_tool_configure_result_t{ - sizeof(rocprofiler_tool_configure_result_t), - &ATTTest::Multi::tool_init, - &ATTTest::Multi::tool_fini, - reinterpret_cast(new Callbacks::ToolData{"att_multi_test/"})}; + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Multi::tool_init, + &Callbacks::finalize, + nullptr}; // return pointer to configure data return &cfg; diff --git a/tests/thread-trace/single_dispatch.cpp b/tests/thread-trace/single_dispatch.cpp index 4feb8b44d1..f09e9c500d 100644 --- a/tests/thread-trace/single_dispatch.cpp +++ b/tests/thread-trace/single_dispatch.cpp @@ -41,35 +41,20 @@ dispatch_callback(rocprofiler_agent_id_t /* agent */, rocprofiler_async_correlation_id_t /* correlation_id */, rocprofiler_kernel_id_t kernel_id, rocprofiler_dispatch_id_t /* dispatch_id */, - void* userdata, - rocprofiler_user_data_t* dispatch_userdata) + void* /* userdata */, + rocprofiler_user_data_t* /* dispatch_userdata */) { - C_API_BEGIN - assert(userdata && "Dispatch callback passed null!"); - auto& tool = *reinterpret_cast(userdata); - dispatch_userdata->ptr = userdata; + static rocprofiler_kernel_id_t target_kernel_id = kernel_id; + if(target_kernel_id == kernel_id) return ROCPROFILER_THREAD_TRACE_CONTROL_START_AND_STOP; - static std::string_view desired_func_name = "branching_kernel"; - - 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_THREAD_TRACE_CONTROL_NONE; - - return ROCPROFILER_THREAD_TRACE_CONTROL_START_AND_STOP; - } catch(...) - { - std::cerr << "Could not find kernel id: " << kernel_id << std::endl; - } - - C_API_END return ROCPROFILER_THREAD_TRACE_CONTROL_NONE; } int -tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) +tool_init(rocprofiler_client_finalize_t /* fini_func */, void* /* tool_data */) { + Callbacks::init(); + static rocprofiler_context_id_t client_ctx = {0}; ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation"); @@ -80,7 +65,7 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) nullptr, 0, Callbacks::tool_codeobj_tracing_callback, - tool_data), + nullptr), "code object tracing service configure"); std::vector agents{}; @@ -110,7 +95,7 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) 0, dispatch_callback, Callbacks::shader_data_callback, - tool_data), + nullptr), "thread trace service configure"); } @@ -131,13 +116,6 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data) return 0; } -void -tool_fini(void* tool_data) -{ - Callbacks::finalize_json(tool_data); - delete static_cast(tool_data); -} - } // namespace Single } // namespace ATTTest @@ -157,11 +135,11 @@ rocprofiler_configure(uint32_t /* version */, ATTTest::Single::client_id = id; // create configure data - static auto cfg = rocprofiler_tool_configure_result_t{ - sizeof(rocprofiler_tool_configure_result_t), - &ATTTest::Single::tool_init, - &ATTTest::Single::tool_fini, - reinterpret_cast(new Callbacks::ToolData{"att_single_test/"})}; + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + &ATTTest::Single::tool_init, + &Callbacks::finalize, + nullptr}; // return pointer to configure data return &cfg; diff --git a/tests/thread-trace/trace_callbacks.cpp b/tests/thread-trace/trace_callbacks.cpp index 3224c57b4c..d40377badb 100644 --- a/tests/thread-trace/trace_callbacks.cpp +++ b/tests/thread-trace/trace_callbacks.cpp @@ -26,11 +26,6 @@ #endif #include "trace_callbacks.hpp" -#include - -#ifdef ENABLE_ATT_FILES -# include -#endif #include #include @@ -38,135 +33,78 @@ namespace Callbacks { -using code_obj_load_data_t = rocprofiler_callback_tracing_code_object_load_data_t; -using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; +rocprofiler_thread_trace_decoder_id_t decoder{}; +std::atomic latency{0}; void tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* /* user_data */, - void* userdata) + void* /* userdata */) { - C_API_BEGIN if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return; - if(record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) return; - - assert(userdata && "Dispatch callback passed null!"); - auto& tool = *reinterpret_cast(userdata); - - if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) - { - auto* data = static_cast(record.payload); - tool.kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name); - } - if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return; - auto* data = static_cast(record.payload); + auto* data = static_cast(record.payload); + if(data->storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) return; - static std::atomic filecnt{0}; - std::string name = "codeobj_" + std::to_string(filecnt.fetch_add(1)) + ".out"; - -#ifdef ENABLE_ATT_FILES - if(std::string_view(data->uri).find("file:///") == 0) + if(record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) { - rocprofiler::sdk::codeobj::disassembly::CodeObjectBinary binary(data->uri); - - std::ofstream file(tool.out_dir + name, std::ios::binary); - assert(file.is_open() && "Could not open codeobj file for writing"); - file.write((char*) binary.buffer.data(), binary.buffer.size()); + DECODER_CALL( + rocprofiler_thread_trace_decoder_codeobj_unload(decoder, data->code_object_id)); + return; } - else - { - std::ofstream file(tool.out_dir + name, std::ios::binary); - file.write((char*) data->memory_base, data->memory_size); - } -#endif - auto _lk = std::unique_lock{tool.mut}; - tool.codeobjs.push_back( - {data->load_delta, data->load_size, data->code_object_id, name, data->uri}); + DECODER_CALL(rocprofiler_thread_trace_decoder_codeobj_load( + decoder, + data->code_object_id, + data->load_delta, + data->load_size, + reinterpret_cast(data->memory_base), + data->memory_size)); +} - C_API_END +typedef void (*rocprofiler_thread_trace_decoder_callback_t)( + rocprofiler_thread_trace_decoder_record_type_t record_type_id, + void* trace_events, + uint64_t trace_size, + void* userdata); + +void +shader_data_callback(rocprofiler_agent_id_t /* agent */, + int64_t /* se_id */, + void* se_data, + size_t data_size, + rocprofiler_user_data_t /* userdata */) +{ + auto parse = [](rocprofiler_thread_trace_decoder_record_type_t record_type_id, + void* trace_events, + uint64_t trace_size, + void* /* userdata */) { + if(record_type_id != ROCPROFILER_THREAD_TRACE_DECODER_RECORD_WAVE) return; + + for(size_t w = 0; w < trace_size; w++) + { + auto* wave = static_cast(trace_events); + for(size_t i = 0; i < wave->instructions_size; i++) + latency += wave->instructions_array[i].duration; + } + }; + DECODER_CALL(rocprofiler_trace_decode(decoder, parse, se_data, data_size, nullptr)); } void -shader_data_callback(rocprofiler_agent_id_t agent, - int64_t se_id, - void* se_data, - size_t data_size, - rocprofiler_user_data_t userdata) +init() { - C_API_BEGIN - - assert(userdata.ptr && "Dispatch callback passed null!"); - auto& tool = *reinterpret_cast(userdata.ptr); - - std::string name = "agent_" + std::to_string(agent.handle) + "_shader_engine_" + - std::to_string(se_id) + "_" + std::to_string(agent.handle) + ".att"; - -#ifdef ENABLE_ATT_FILES - { - std::ofstream file(tool.out_dir + name, std::ios::binary); - assert(file.is_open() && "Could not open ATT file for writing"); - file.write((char*) se_data, data_size); - } -#endif - - assert(se_data); - assert(data_size); - - auto _lk = std::unique_lock{tool.mut}; - tool.att_files.push_back(name); - - C_API_END + // const char* decoder_lib = std::getenv("ROCPROF_TRACE_DECODER_PATH"); + DECODER_CALL(rocprofiler_thread_trace_decoder_create(&decoder, "/opt/rocm/lib")); } void -finalize_json(void* userdata) +finalize(void* /* tool_data */) { - assert(userdata && "Dispatch callback passed null!"); + rocprofiler_thread_trace_decoder_destroy(decoder); - auto& tool = *reinterpret_cast(userdata); - auto _lk = std::unique_lock{tool.mut}; - assert(!tool.att_files.empty()); - -#ifdef ENABLE_ATT_FILES - nlohmann::json att_json; - for(auto& file : tool.att_files) - att_json.push_back(file); - - nlohmann::json codeobj_json; - nlohmann::json snapshot_json; - for(auto& file : tool.codeobjs) - { - nlohmann::json codeobj; - codeobj["code_object_id"] = file.id; - codeobj["load_delta"] = file.addr; - codeobj["load_size"] = file.size; - codeobj["uri"] = file.uri; - codeobj["filename"] = file.filename; - codeobj_json.push_back(codeobj); - - nlohmann::json pair_json; - pair_json["key"] = file.id; - pair_json["value"] = file.filename; - snapshot_json.push_back(pair_json); - } - nlohmann::json tool_json; - tool_json["strings"]["att_files"] = att_json; - tool_json["code_objects"] = codeobj_json; - tool_json["strings"]["code_object_snapshot_files"] = snapshot_json; - - nlohmann::json array; - array.push_back(tool_json); - - nlohmann::json sdk_json; - sdk_json["rocprofiler-sdk-tool"] = array; - - std::ofstream json_file(tool.out_dir + (std::to_string(getpid()) + "_results.json")); - assert(json_file.is_open() && "Could not open json file for writing!"); - json_file << sdk_json; -#endif + if(latency.load() == 0) std::cerr << "Error: No latency was assigned to the trace!"; } } // namespace Callbacks diff --git a/tests/thread-trace/trace_callbacks.hpp b/tests/thread-trace/trace_callbacks.hpp index 414d6c984a..9c175782b0 100644 --- a/tests/thread-trace/trace_callbacks.hpp +++ b/tests/thread-trace/trace_callbacks.hpp @@ -40,60 +40,21 @@ #include #include -#define ROCPROFILER_VAR_NAME_COMBINE(X, Y) X##Y -#define ROCPROFILER_VARIABLE(X, Y) ROCPROFILER_VAR_NAME_COMBINE(X, Y) - -#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; } - #define ROCPROFILER_CALL(result, msg) \ + if((result) != ROCPROFILER_STATUS_SUCCESS) \ { \ - rocprofiler_status_t CHECKSTATUS = result; \ - if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS) \ - { \ - std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS); \ - std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ - << " failed with error code " << CHECKSTATUS << ": " << status_msg \ - << std::endl; \ - std::stringstream errmsg{}; \ - errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ - << status_msg << ")"; \ - throw std::runtime_error(errmsg.str()); \ - } \ + std::cerr << "Error: " << msg << std::endl; \ + abort(); \ + } + +#define DECODER_CALL(result) \ + if((result) != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::cerr << "Error: Generic decoder error" << std::endl; \ } namespace Callbacks { -struct CodeobjInfo -{ - int64_t addr = 0; - size_t size = 0; - size_t id = 0; - std::string filename{}; - std::string uri{}; -}; - -struct ToolData -{ - ToolData(const char* out) - : out_dir(out){}; - - std::string out_dir{}; - std::mutex mut{}; - std::vector codeobjs{}; - std::vector att_files{}; - - std::unordered_map kernel_id_to_kernel_name = {}; -}; - void tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t*, @@ -107,6 +68,9 @@ shader_data_callback(rocprofiler_agent_id_t agent, rocprofiler_user_data_t userdata); void -finalize_json(void* userdata); +init(); + +void +finalize(void* /* tool_data */); } // namespace Callbacks From 2c7f260e623f5967ba5d5cbf6c5e3b497c6e361e Mon Sep 17 00:00:00 2001 From: "Kuricheti, Mythreya" Date: Wed, 23 Jul 2025 13:52:52 -0700 Subject: [PATCH 2/7] [SDK] Fix context tracing domain bitset overflow (#536) - Fix context tracing domain bitset overflow - Previous behavior would enable all flags above ROCPROFILER_BUFFER_TRACING_MARKER_CORE_RANGE_API when this domain was enabled. --- source/lib/rocprofiler-sdk/context/domain.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/lib/rocprofiler-sdk/context/domain.cpp b/source/lib/rocprofiler-sdk/context/domain.cpp index d0a97e45ec..6ca94bbec8 100644 --- a/source/lib/rocprofiler-sdk/context/domain.cpp +++ b/source/lib/rocprofiler-sdk/context/domain.cpp @@ -68,7 +68,7 @@ add_domain(domain_context& _cfg, DomainT _domain) if(_didx >= _cfg.array_size) return ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND; - _cfg.domains |= (1 << _didx); + _cfg.domains |= (1UL << _didx); return ROCPROFILER_STATUS_SUCCESS; } From 3a36fd13fed56f9b00ba452d23a27365e892e89e Mon Sep 17 00:00:00 2001 From: "U, Srihari" Date: Thu, 24 Jul 2025 12:11:36 +0530 Subject: [PATCH 3/7] [rocprofv3] rocpd doesn't generate output files for counter collection (#480) * Fix kernel dispatch for counter collection * Updated change log * Fix format * rename output csv file * Fix warnings * Address review comment * Address final review comment --- CHANGELOG.md | 1 + source/lib/output/generateRocpd.cpp | 242 ++++++++++++++----------- source/lib/python/rocpd/source/csv.cpp | 2 +- 3 files changed, 141 insertions(+), 104 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 5c992ae193..6550fc14d8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -216,6 +216,7 @@ Full documentation for ROCprofiler-SDK is available at [rocm.docs.amd.com/projec - Code object disassembly was missing function inlining information - Fixed queue preemption error and HSA_STATUS_ERROR_INVALID_PACKET_FORMAT error for stochastic PC-sampling for MI300X, leading to more stable runs. - Fixed the system hang issue for host-trap PC-sampling on MI300X. +- Fixed rocpd counter collection issue when counter collection alone is enabled, rocpd_kernel_dispatch table gets populated by counters data instead of kernel_dispatch data. ### Removed diff --git a/source/lib/output/generateRocpd.cpp b/source/lib/output/generateRocpd.cpp index 6899919ade..49970e4792 100644 --- a/source/lib/output/generateRocpd.cpp +++ b/source/lib/output/generateRocpd.cpp @@ -970,46 +970,45 @@ write_rocpd( } }; - auto insert_kernel_dispatch_data = [&conn, &tool_metadata, &string_entries, node_id, this_pid]( - const auto& _gen, auto& dispatch_evt_ids) { + auto insert_kernel_dispatch_data = [&, node_id, this_pid](auto& dispatch_evt_ids) { auto _sqlgenperf_rocpd = get_simple_timer("rocpd_kernel_dispatch"); - for(auto pitr : _gen) - { - auto _deferred = sql::deferred_transaction{conn}; - for(auto itr : _gen.get(pitr)) + auto process_dispatch = [&](uint64_t dispatch_id, + uint64_t kernel_id, + const auto& corr_id, + const auto& info, + const auto& kind, + uint32_t thread_id, + uint64_t queue_id, + uint64_t stream_id, + uint64_t start_timestamp, + uint64_t end_timestamp, + const auto& grid, + const auto& workgroup, + bool enable_duplicate_check) { + // Skip if we've already processed this dispatch_id + if(dispatch_evt_ids.size() > dispatch_id && dispatch_evt_ids[dispatch_id] != 0) return; + + auto kern_name = (kernel_id > 0) + ? tool_metadata.get_kernel_symbol(kernel_id)->formatted_kernel_name + : "unknown_kernel"; + + auto evt_id = create_event(conn, + { + insert_value("category_id", string_entries.at(kind)), + insert_value("stack_id", corr_id.internal), + insert_value("parent_stack_id", corr_id.internal), + insert_value("correlation_id", corr_id.external.value), + }); + + // Ensure dispatch_evt_ids is large enough + if(dispatch_evt_ids.size() < dispatch_id + 1) + common::container::resize(dispatch_evt_ids, dispatch_id + 1, 0UL); + + // Check for duplicates if requested + if(enable_duplicate_check && dispatch_evt_ids.at(dispatch_id) != 0) { - // insert thread info if it doesn't already exist - get_thread_id(itr.thread_id); - - auto kind = tool_metadata.buffer_names.at(itr.kind); - auto info = itr.dispatch_info; - auto kernel_id = info.kernel_id; - auto dispatch_id = info.dispatch_id; - auto corr_id = itr.correlation_id; - auto grid = info.grid_size; - auto workgroup = info.workgroup_size; - auto kern_name = tool_metadata.get_kernel_symbol(kernel_id)->formatted_kernel_name; - auto stream_id = get_stream_id(itr.stream_id); - auto queue_id = get_queue_id(info.queue_id); - auto region_name = - (corr_id.external.value > 0) - ? tool_metadata.get_kernel_name(kernel_id, corr_id.external.value) - : std::string_view{}; - - auto evt_id = - create_event(conn, - { - insert_value("category_id", string_entries.at(kind)), - insert_value("stack_id", corr_id.internal), - insert_value("parent_stack_id", corr_id.internal), - insert_value("correlation_id", corr_id.external.value), - }); - - if(dispatch_evt_ids.size() < dispatch_id + 1) - common::container::resize(dispatch_evt_ids, dispatch_id + 1, 0UL); - - ROCP_CI_LOG_IF(WARNING, dispatch_evt_ids.at(dispatch_id) != 0) + ROCP_CI_LOG(WARNING) << fmt::format("duplicate kernel dispatch id {} :: event_id={}, kernel_id={}, " "corr_id={}, name='{}'", dispatch_id, @@ -1017,45 +1016,115 @@ write_rocpd( kernel_id, corr_id.internal, kern_name); + } - dispatch_evt_ids.at(dispatch_id) = evt_id; + dispatch_evt_ids.at(dispatch_id) = evt_id; - auto stmt = get_insert_statement( - "rocpd_kernel_dispatch{{uuid}}", - { - insert_value("id", dispatch_id), - insert_value("nid", node_id), - insert_value("pid", this_pid), - insert_value("tid", itr.thread_id), - insert_value("agent_id", tool_metadata.get_agent(info.agent_id)->node_id), - insert_value("kernel_id", kernel_id), - insert_value("dispatch_id", dispatch_id), - insert_value("queue_id", queue_id), - insert_value("stream_id", stream_id), - insert_value("start", itr.start_timestamp), - insert_value("end", itr.end_timestamp), - insert_value("private_segment_size", info.private_segment_size), - insert_value("group_segment_size", info.group_segment_size), - insert_value("workgroup_size_x", workgroup.x), - insert_value("workgroup_size_y", workgroup.y), - insert_value("workgroup_size_z", workgroup.z), - insert_value("grid_size_x", grid.x), - insert_value("grid_size_y", grid.y), - insert_value("grid_size_z", grid.z), - insert_value("region_name_id", string_entries.at(region_name)), - insert_value("event_id", evt_id), - }); + auto region_name = + (corr_id.external.value > 0 && + (enable_duplicate_check || kernel_id > 0)) + ? tool_metadata.get_kernel_name(kernel_id, corr_id.external.value) + : std::string_view{}; - execute_raw_sql_statements(conn, stmt); + auto agent_node_id = tool_metadata.get_agent(info.agent_id)->node_id; + + // Insert into kernel dispatch table + auto stmt = get_insert_statement( + "rocpd_kernel_dispatch{{uuid}}", + { + insert_value("id", dispatch_id), + insert_value("nid", node_id), + insert_value("pid", this_pid), + insert_value("tid", thread_id), + insert_value("agent_id", agent_node_id), + insert_value("kernel_id", kernel_id), + insert_value("dispatch_id", dispatch_id), + insert_value("queue_id", queue_id), + insert_value("stream_id", stream_id), + insert_value("start", start_timestamp), + insert_value("end", end_timestamp), + insert_value("private_segment_size", info.private_segment_size), + insert_value("group_segment_size", info.group_segment_size), + insert_value("workgroup_size_x", workgroup.x), + insert_value("workgroup_size_y", workgroup.y), + insert_value("workgroup_size_z", workgroup.z), + insert_value("grid_size_x", grid.x), + insert_value("grid_size_y", grid.y), + insert_value("grid_size_z", grid.z), + insert_value("region_name_id", string_entries.at(region_name)), + insert_value("event_id", evt_id), + }); + + execute_raw_sql_statements(conn, stmt); + }; + + if(kernel_dispatch_gen.empty()) + { + for(auto pctr : counter_collection_gen) + { + auto _deferred = sql::deferred_transaction{conn}; + for(const auto& record : counter_collection_gen.get(pctr)) + { + const auto& dispatch_data = record.dispatch_data; + const auto& info = dispatch_data.dispatch_info; + + // Register thread ID + get_thread_id(record.thread_id); + + // Use buffer category for kernel dispatches + auto kind = + tool_metadata.buffer_names.at(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH); + + // Process this dispatch + process_dispatch(info.dispatch_id, // dispatch_id + info.kernel_id, // kernel_id + dispatch_data.correlation_id, // corr_id + info, // info + kind, // kind + record.thread_id, // thread_id + get_queue_id(info.queue_id), // queue_id + get_stream_id(record.stream_id), // stream_id + dispatch_data.start_timestamp, // start_timestamp + dispatch_data.end_timestamp, // end_timestamp + info.grid_size, // grid + info.workgroup_size, // workgroup + false // enable_duplicate_check + ); + } + } + } + else + { + for(auto pitr : kernel_dispatch_gen) + { + auto _deferred = sql::deferred_transaction{conn}; + for(auto itr : kernel_dispatch_gen.get(pitr)) + { + // Register thread ID + get_thread_id(itr.thread_id); + + // Process this dispatch + process_dispatch(itr.dispatch_info.dispatch_id, // dispatch_id + itr.dispatch_info.kernel_id, // kernel_id + itr.correlation_id, // corr_id + itr.dispatch_info, // info + tool_metadata.buffer_names.at(itr.kind), // kind + itr.thread_id, // thread_id + get_queue_id(itr.dispatch_info.queue_id), // queue_id + get_stream_id(itr.stream_id), // stream_id + itr.start_timestamp, // start_timestamp + itr.end_timestamp, // end_timestamp + itr.dispatch_info.grid_size, // grid + itr.dispatch_info.workgroup_size, // workgroup + true // enable_duplicate_check + ); + } } } }; - auto insert_pmc_event_data = [&conn, - &tool_metadata, - &string_entries, - &counter_collection_gen, - &kernel_dispatch_gen](auto& dispatch_evt_ids) { + auto insert_pmc_event_data = [&conn, &tool_metadata, &counter_collection_gen]( + auto& dispatch_evt_ids) { auto _sqlgenperf_rocpd = get_simple_timer("rocpd_pmc_event"); size_t idx = tool_metadata.pmc_event_offset; for(auto ditr : counter_collection_gen) @@ -1066,39 +1135,6 @@ write_rocpd( const auto& info = record.dispatch_data.dispatch_info; auto dispatch_id = info.dispatch_id; - if(kernel_dispatch_gen.empty()) - { - auto kind = - tool_metadata.buffer_names.at(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH); - auto kernel_id = info.kernel_id; - auto corr_id = record.dispatch_data.correlation_id; - auto kern_name = - tool_metadata.get_kernel_symbol(kernel_id)->formatted_kernel_name; - - auto evt_id = - create_event(conn, - { - insert_value("category_id", string_entries.at(kind)), - insert_value("stack_id", corr_id.internal), - insert_value("parent_stack_id", corr_id.internal), - insert_value("correlation_id", corr_id.external.value), - }); - - if(dispatch_evt_ids.size() < dispatch_id + 1) - common::container::resize(dispatch_evt_ids, dispatch_id + 1, 0UL); - - ROCP_CI_LOG_IF(WARNING, dispatch_evt_ids.at(dispatch_id) != 0) << fmt::format( - "duplicate kernel dispatch id {} :: event_id={}, kernel_id={}, " - "corr_id={}, name='{}'", - dispatch_id, - evt_id, - kernel_id, - corr_id.internal, - kern_name); - - dispatch_evt_ids.at(dispatch_id) = evt_id; - } - auto evt_id = dispatch_evt_ids.at(dispatch_id); for(const auto& count : record.read()) { @@ -1377,7 +1413,7 @@ write_rocpd( insert_api_data(rocdecode_api_gen); } - insert_kernel_dispatch_data(kernel_dispatch_gen, dispatch_to_evt_id); + insert_kernel_dispatch_data(dispatch_to_evt_id); insert_pmc_event_data(dispatch_to_evt_id); insert_memory_copy_data(memory_copy_gen); diff --git a/source/lib/python/rocpd/source/csv.cpp b/source/lib/python/rocpd/source/csv.cpp index 5be3f97977..0e011e0663 100644 --- a/source/lib/python/rocpd/source/csv.cpp +++ b/source/lib/python/rocpd/source/csv.cpp @@ -109,7 +109,7 @@ CsvManager::CsvManager(rocprofiler::tool::output_config output_cfg) {CsvType::ROCJPEG_API, {"rocjpeg_api_trace.csv", API_TRACE_HEADER}}, {CsvType::COUNTER, - {"counter_trace.csv", + {"counter_collection.csv", "\"Pid\",\"Correlation_Id\",\"Dispatch_Id\",\"Agent_Id\",\"Queue_Id\"," "\"Process_Id\"," "\"Thread_Id\"," From 735b5c3d4a8ace9ba41287a859486c842ef7686b Mon Sep 17 00:00:00 2001 From: "Madsen, Jonathan" Date: Thu, 24 Jul 2025 15:23:37 -0500 Subject: [PATCH 4/7] [CMake] Fix thread trace sample ENVIRONMENT test property (#544) Fix thread trace samples set tests properties --- samples/thread_trace/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/thread_trace/CMakeLists.txt b/samples/thread_trace/CMakeLists.txt index 0c9c805631..e3dfd1151f 100644 --- a/samples/thread_trace/CMakeLists.txt +++ b/samples/thread_trace/CMakeLists.txt @@ -75,5 +75,5 @@ rocprofiler_samples_get_preload_env(PRELOAD_ENV) list(APPEND PRELOAD_ENV "ROCPROFILER_TRACE_DECODER_LIB_PATH=${attdecoder_LIB_DIR}") set_tests_properties( - thread-trace-sample PROPERTIES TIMEOUT 60 ENVIRONMENT ${PRELOAD_ENV} LABELS + thread-trace-sample PROPERTIES TIMEOUT 60 ENVIRONMENT "${PRELOAD_ENV}" LABELS "samples;thread-trace" DISABLED ${IS_DISABLED}) From 3954cedd253a6b370cba9018edfe03291875d3aa Mon Sep 17 00:00:00 2001 From: "Hui, Young" Date: Thu, 24 Jul 2025 17:12:06 -0400 Subject: [PATCH 5/7] [rocpd] Adding summary module to generate summaries from rocpd database + query submodule + rocpd command-line tools (#488) * adding summary.py to generate tmp _summary views * migrating CSV summary to SDK method of writing CSVs - Add domain_view to summary.py - omit the C++ code of writing CSV because it gets revered later anyway * Add summary subparser and write_sql_view_to_csv function * adding all <>_summary views generation to summary.py * add summary_per_rank feature * add --summary-per-rank * reconstruct generate_summary_view and create_domain_view -introduce by_rank * remove sqr and variance in summary views * use RocpdImportData instead of connection * two fixes on summary.py --modify the generate_summary_view function to return a tuple with view name and sql code add if_not_exits parameter to generete_summary_view * Refactor summary.py to allow output path and filename args, and apply time_window - clean up summary table column headers - only generate by-rank views if that param is specified * Add ProcessID to Hostname output and csv, so users can identify the system in the by-rank summaries * Summary.py, just add hostname to by-rank summaries, instead of creating mapping table * Summary - migrate csv writer to pandas, for more future flexibility * Adding a few simple tests for summary.py * Linting fixes * add region_categories to summary options - Automatically retrieve region categories from the database if argument is None * add backticks for view_names * fix tests after rebase * Made code review changes - fixed whitespace in CMakelists.txt - adding query.py module & subparser in __main__.py - refactor summary function to return query - used query.py to output csv - used query.py to also output summary to console - provided new command line options to select summary output to csv or console * Made fix to jinja template in query.py, as suggested by copilot * Consolidated output calls to query in export_view function based on feedback - refactored: helpers, query functions, create view functions - extended formats to include what query supports (md, html, pdf, json) - added json format to query, and changed orient=records - adding jinja2 and reportlab to requirements.txt * Add version_info for rocpd and roctx * Add rocpd commandline tool * Add executable permissions to source/bin/rocpd.py * Removed rocpd2query, and cleaned up --help examples --------- Co-authored-by: acanadas Co-authored-by: Jin Tao Co-authored-by: a-canadasruiz Co-authored-by: Jonathan R. Madsen --- requirements.txt | 2 + source/bin/CMakeLists.txt | 36 ++ source/bin/rocpd.py | 74 +++ source/lib/python/rocpd/__init__.py | 16 + source/lib/python/rocpd/__main__.py | 235 ++++++-- source/lib/python/rocpd/query.py | 560 ++++++++++++++++++ source/lib/python/rocpd/summary.py | 532 +++++++++++++++++ source/lib/python/roctx/__init__.py | 16 + source/lib/python/utilities.cmake | 6 +- .../rocprofiler-sdk-rocpd/data_views.sql | 1 + .../rocprofiler-sdk-rocpd/summary_views.sql | 223 ------- tests/rocpd/CMakeLists.txt | 4 +- tests/rocprofv3/rocpd/CMakeLists.txt | 52 ++ 13 files changed, 1480 insertions(+), 277 deletions(-) create mode 100755 source/bin/rocpd.py create mode 100644 source/lib/python/rocpd/query.py create mode 100644 source/lib/python/rocpd/summary.py diff --git a/requirements.txt b/requirements.txt index 178bc1fa4b..6eb9da496e 100644 --- a/requirements.txt +++ b/requirements.txt @@ -5,6 +5,7 @@ cmake>=3.21.0 cmake-format dataclasses flake8 +jinja2 numpy otf2 pandas @@ -12,3 +13,4 @@ perfetto pycobertura pytest pyyaml +reportlab diff --git a/source/bin/CMakeLists.txt b/source/bin/CMakeLists.txt index cba31cf1cc..7f0bb05908 100644 --- a/source/bin/CMakeLists.txt +++ b/source/bin/CMakeLists.txt @@ -25,3 +25,39 @@ install( PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE COMPONENT tools) + +# for each entry here there must be a ROCPD_BIN_ list +set(ROCPD_EXECUTABLES "all" "csv" "otf2" "pftrace" "summary") + +# format is list: [, , ] +set(ROCPD_BIN_all "rocpd" "rocpd" "") +set(ROCPD_BIN_csv "rocpd2csv" "rocpd.csv" "") +set(ROCPD_BIN_otf2 "rocpd2otf2" "rocpd.otf2" "") +set(ROCPD_BIN_pftrace "rocpd2pftrace" "rocpd.pftrace" "") +set(ROCPD_BIN_summary "rocpd2summary" "rocpd.summary" "") + +foreach(_EXE IN LISTS ROCPD_EXECUTABLES) + list(GET ROCPD_BIN_${_EXE} 0 ROCPD_EXE_NAME) + list(GET ROCPD_BIN_${_EXE} 1 ROCPD_EXE_MODULE) + list(GET ROCPD_BIN_${_EXE} 2 ROCPD_EXE_MODULE_ARGS) + + if(NOT ROCPD_EXE_NAME + OR NOT ROCPD_EXE_MODULE + OR (NOT ROCPD_EXE_MODULE_ARGS AND NOT ROCPD_EXE_MODULE_ARGS STREQUAL "")) + message( + FATAL_ERROR "ROCPD_BIN_${_EXE} not properly defined: ${ROCPD_BIN_${_EXE}}") + endif() + + set(ROCPD_EXE_OUTPUT_FILE + ${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/${ROCPD_EXE_NAME}) + + # Adding main rocpd + configure_file(rocpd.py ${ROCPD_EXE_OUTPUT_FILE} @ONLY) + + install( + FILES ${ROCPD_EXE_OUTPUT_FILE} + DESTINATION ${CMAKE_INSTALL_BINDIR} + PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE + WORLD_READ WORLD_EXECUTE + COMPONENT rocpd) +endforeach() diff --git a/source/bin/rocpd.py b/source/bin/rocpd.py new file mode 100755 index 0000000000..ffe8ce4f9e --- /dev/null +++ b/source/bin/rocpd.py @@ -0,0 +1,74 @@ +#!/usr/bin/env python3 + +# 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. + +import os +import sys + +""" +Simple Python executable script for invoking `python3 -m @ROCPD_EXE_MODULE@` +""" + + +def main(argv=sys.argv[1:], environ=dict(os.environ)): + """ + Executes {sys.executable} -m @ROCPD_EXE_MODULE@ @ROCPD_EXE_MODULE_ARGS@ + """ + + ROCPD_SUPPORTED_PYTHON_VERSIONS = [ + ".".join(itr.split(".")[:2]) for itr in "@ROCPROFILER_PYTHON_VERSIONS@".split(";") + ] + ROCPD_MODULE_ARGS = [f"{itr}" for itr in "@ROCPD_EXE_MODULE_ARGS@".split(" ") if itr] + + this_dir = os.path.dirname(os.path.realpath(__file__)) + this_python_ver = f"{sys.version_info.major}.{sys.version_info.minor}" + if this_python_ver not in ROCPD_SUPPORTED_PYTHON_VERSIONS: + raise ImportError( + "@ROCPD_EXE_NAME@ not supported for Python version {} (sys.executable='{}').\n@ROCPD_EXE_NAME@ supported python versions: {}".format( + this_python_ver, + sys.executable, + ", ".join(ROCPD_SUPPORTED_PYTHON_VERSIONS), + ) + ) + + module_path = os.path.join( + this_dir, + "..", + "@CMAKE_INSTALL_LIBDIR@", + f"python{this_python_ver}", + "site-packages", + ) + + python_path = [module_path] + os.environ.get("PYTHONPATH", "").split(":") + + # update PYTHONPATH environment variable + environ["PYTHONPATH"] = ":".join(python_path) + + args = [f"{sys.executable}", "-m", "@ROCPD_EXE_MODULE@"] + ROCPD_MODULE_ARGS + argv + + # does not return + os.execvpe(args[0], args, env=environ) + + +if __name__ == "__main__": + main() diff --git a/source/lib/python/rocpd/__init__.py b/source/lib/python/rocpd/__init__.py index bd8b171833..dc9f798560 100644 --- a/source/lib/python/rocpd/__init__.py +++ b/source/lib/python/rocpd/__init__.py @@ -46,8 +46,24 @@ __all__ = [ "write_csv", "write_otf2", "RocpdImportData", + "version_info", ] +version_info = { + "version": "@PROJECT_VERSION@", + "major": int("@PROJECT_VERSION_MAJOR@"), + "minor": int("@PROJECT_VERSION_MINOR@"), + "patch": int("@PROJECT_VERSION_PATCH@"), + "git_revision": "@ROCPROFILER_SDK_GIT_REVISION@", + "library_arch": "@CMAKE_LIBRARY_ARCHITECTURE@", + "system_name": "@CMAKE_SYSTEM_NAME@", + "system_processor": "@CMAKE_SYSTEM_PROCESSOR@", + "system_version": "@CMAKE_SYSTEM_VERSION@", + "compiler_id": "@CMAKE_CXX_COMPILER_ID@", + "compiler_version": "@CMAKE_CXX_COMPILER_VERSION@", + "rocm_version": "@rocm_version_FULL_VERSION@", +} + def format_path(path, tag=os.path.basename(sys.executable)): return libpyrocpd.format_path(path, tag) diff --git a/source/lib/python/rocpd/__main__.py b/source/lib/python/rocpd/__main__.py index a056d497dd..92f5eaf4e9 100644 --- a/source/lib/python/rocpd/__main__.py +++ b/source/lib/python/rocpd/__main__.py @@ -38,11 +38,14 @@ def main(argv=None, config=None): """ import argparse - from . import time_window - from . import output_config - from . import pftrace from . import csv from . import otf2 + from . import output_config + from . import pftrace + from . import query + from . import summary + from . import time_window + from . import version_info from .importer import RocpdImportData convert_examples = """ @@ -50,25 +53,62 @@ def main(argv=None, config=None): Example usage: Convert 1 database, output perfetto trace - $ python3 -m rocpd convert -i db1.db --output-format pftrace + $ rocpd convert -i db1.db --output-format pftrace Convert 2 databases, output perfetto trace to path and filename, reduce time window to omit the first 30% - $ python3 -m rocpd convert -i db1.db db2.db --output-format pftrace -d "./output/" -o "twoFileTraces" --start 30% --end 100% + $ rocpd convert -i db1.db db2.db --output-format pftrace -d "./output/" -o "twoFileTraces" --start 30% --end 100% Convert 6 databases, output CSV and perfetto trace formats - $ python3 -m rocpd convert -i db{0..5}.db --output-format csv pftrace -d "~/output_folder/" -o "sixFileTraces" + $ rocpd convert -i db{0..5}.db --output-format csv pftrace -d "~/output_folder/" -o "sixFileTraces" Convert 2 databases, output CSV, OTF2, and perfetto trace formats - $ python3 -m rocpd convert -i db{3,4}.db --output-format csv otf2 pftrace + $ rocpd convert -i db{3,4}.db --output-format csv otf2 pftrace """ + query_examples = """ + +Example usage: + + Query the first 5 rows of the 'rocpd_info_agents' view and output to console + $ rocpd query -i db0.db --query "SELECT * FROM rocpd_info_agents LIMIT 5" + + Combine 4 databases and query the first 10 rows of the 'top_kernels' view and output to CSV file + $ rocpd query -i db{0..3}.db --query "SELECT * FROM top_kernels LIMIT 10" --format csv +""" + + summary_examples = """ + +Example usage: + + Output all summaries to console and include domain summary for 1 database + $ rocpd summary -i db1.db --domain-summary + + Aggregate 3 databases and output all summary files and include summary by rank/process ID, to csv file output + $ rocpd summary -i db{1..3}.db --summary-by-rank --format csv + + Output all summaries to console and exlude all regions to save processing time + $ rocpd summary -i db0.db --region-categories NONE + + Aggregate 2 databases and output all summary files to HTML, only include HIP and MARKER regions, include domain summary + $ rocpd summary -i db{0,1}.db --region-categories HIP MARKERS --domain-summary --format html + +""" + + # Add the subparsers parser = argparse.ArgumentParser( prog="rocpd", description="Aggregate and/or analyze ROCm Profiling Data (rocpd)", allow_abbrev=False, ) + parser.add_argument( + "-v", + "--version", + action="store_true", + help="Print the version information and exit", + ) + subparsers = parser.add_subparsers(dest="command") converter = subparsers.add_parser( "convert", @@ -78,20 +118,36 @@ Example usage: epilog=convert_examples, ) + query_reporter = subparsers.add_parser( + "query", + description="Generate output on a query", + allow_abbrev=False, + formatter_class=argparse.RawTextHelpFormatter, + epilog=query_examples, + ) + + generate_summary = subparsers.add_parser( + "summary", + description="Generate summary views from rocPD data", + allow_abbrev=False, + formatter_class=argparse.RawTextHelpFormatter, + epilog=summary_examples, + ) + def get_output_type(val): return val.lower().replace("perfetto", "pftrace") - required_params = converter.add_argument_group("Required arguments") - - required_params.add_argument( + # add required options for each subparser + converter_required_params = converter.add_argument_group("Required options") + converter_required_params.add_argument( "-i", "--input", required=True, type=output_config.check_file_exists, nargs="+", - help="Input path and filename to one or more database(s), separated by spaces", + help="Input path and filename to one or more database(s)", ) - required_params.add_argument( + converter_required_params.add_argument( "-f", "--output-format", help="For adding output format (supported formats: csv, pftrace, otf2)", @@ -102,7 +158,27 @@ Example usage: required=True, ) - # add args from any sub-modules + query_required_params = query_reporter.add_argument_group("Required options") + query_required_params.add_argument( + "-i", + "--input", + required=True, + type=output_config.check_file_exists, + nargs="+", + help="Input path and filename to one or more database(s)", + ) + + summary_required_params = generate_summary.add_argument_group("Required options") + summary_required_params.add_argument( + "-i", + "--input", + required=True, + type=output_config.check_file_exists, + nargs="+", + help="Input path and filename to one or more database(s)", + ) + + # converter: add args from any sub-modules valid_out_config_args = output_config.add_args(converter) valid_generic_args = output_config.add_generic_args(converter) valid_pftrace_args = pftrace.add_args(converter) @@ -110,51 +186,110 @@ Example usage: valid_otf2_args = otf2.add_args(converter) valid_time_window_args = time_window.add_args(converter) + # query: subparser args + valid_out_config_args = output_config.add_args(query_reporter) + valid_query_args = query.add_args(query_reporter) + valid_time_window_args = time_window.add_args(query_reporter) + + # summary: subparser args + valid_io_args = summary.add_io_args(generate_summary) + valid_summary_args = summary.add_args(generate_summary) + valid_time_window_args = time_window.add_args(generate_summary) + # parse the command line arguments args = parser.parse_args(argv) - # process the args - out_cfg_args = output_config.process_args(args, valid_out_config_args) - generic_out_cfg_args = output_config.process_generic_args(args, valid_generic_args) - pftrace_args = pftrace.process_args(args, valid_pftrace_args) - csv_args = csv.process_args(args, valid_csv_args) - otf2_args = otf2.process_args(args, valid_otf2_args) - window_args = time_window.process_args(args, valid_time_window_args) + if args.version: + for key, itr in version_info.items(): + if key in ["major", "minor", "patch"]: + continue + print(f" {key:>16}: {itr}") + return 0 - # now start processing the data. Import the data and merge the views - importData = RocpdImportData(args.input) + # error check the command line arguments, if no subparser command is given, print the help message + if args.command is None: + parser.print_help() + return - # adjust the time window view of the data - if window_args is not None: - time_window.apply_time_window(importData, **window_args) + # if the user requested converter, process the conversion + if args.command == "convert": + # process the args + out_cfg_args = output_config.process_args(args, valid_out_config_args) + generic_out_cfg_args = output_config.process_generic_args( + args, valid_generic_args + ) + pftrace_args = pftrace.process_args(args, valid_pftrace_args) + csv_args = csv.process_args(args, valid_csv_args) + otf2_args = otf2.process_args(args, valid_otf2_args) + window_args = time_window.process_args(args, valid_time_window_args) - all_args = { - **out_cfg_args, - **generic_out_cfg_args, - **pftrace_args, - **csv_args, - **otf2_args, - } - # setup the config args - config = ( - output_config.output_config(**all_args) - if config is None - else config.update(**all_args) - ) + # now start processing the data. Import the data and merge the views + importData = RocpdImportData(args.input) - # process each requested output format - format_handlers = { - "pftrace": pftrace.write_pftrace, - "csv": csv.write_csv, - "otf2": otf2.write_otf2, - } + # adjust the time window view of the data + if window_args is not None: + time_window.apply_time_window(importData, **window_args) - for out_format in args.output_format: - if out_format in format_handlers: - print(f"Converting database(s) to {out_format} format:") - format_handlers[out_format](importData, config) - else: - print(f"Warning: Unsupported output format '{out_format}'") + all_args = { + **out_cfg_args, + **generic_out_cfg_args, + **pftrace_args, + **csv_args, + **otf2_args, + } + # setup the config args + config = ( + output_config.output_config(**all_args) + if config is None + else config.update(**all_args) + ) + + # process each requested output format + format_handlers = { + "pftrace": pftrace.write_pftrace, + "csv": csv.write_csv, + "otf2": otf2.write_otf2, + } + + for out_format in args.output_format: + if out_format in format_handlers: + print(f"Converting database(s) to {out_format} format:") + format_handlers[out_format](importData, config) + else: + print(f"Warning: Unsupported output format '{out_format}'") + + # if the user requested query module, execute the query + elif args.command == "query": + # query subparser args + query_args = query.process_args(args, valid_query_args) + out_cfg_args = output_config.process_args(args, valid_out_config_args) + window_args = time_window.process_args(args, valid_time_window_args) + + all_args = {**query_args, **out_cfg_args} + + query.execute( + args.input, + args, + window_args=window_args, + **all_args, + ) + + # if the user requested a summary, generate the views + elif args.command == "summary": + # summary subparser args + summary_args = summary.process_args(args, valid_summary_args) + io_args = output_config.process_args(args, valid_io_args) + window_args = time_window.process_args(args, valid_time_window_args) + + # now start processing the data. Import the data and merge the views + importData = RocpdImportData(args.input) + + # adjust the time window view of the data + if window_args is not None: + time_window.apply_time_window(importData, **window_args) + + all_args = {**summary_args, **io_args} + summary.generate_all_summaries(importData, **all_args) print("Done. Exiting...") diff --git a/source/lib/python/rocpd/query.py b/source/lib/python/rocpd/query.py new file mode 100644 index 0000000000..ab369326db --- /dev/null +++ b/source/lib/python/rocpd/query.py @@ -0,0 +1,560 @@ +#!/usr/bin/env python3 +############################################################################### +# MIT License +# +# Copyright (c) 2025 Advanced Micro Devices, Inc. +# +# 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. +############################################################################### + +import os +import sys + +from typing import Union, Tuple, List, Optional +from datetime import datetime + +from . import output_config +from . import libpyrocpd +from .importer import RocpdImportData +from .time_window import apply_time_window + + +def export_sqlite_query( + conn: RocpdImportData, + query: str, + params: Union[Tuple, List] = (), + export_format: Optional[str] = None, + export_path: Optional[str] = None, + dashboard_template_path: Optional[str] = None, +) -> Optional[str]: + """ + Execute a SQLite query and print it to console. + Then, if export_format is specified, write the results to a file. + Returns the path to the exported file (or None if nothing was exported). + + Supported export_format values (case-insensitive): + - "csv" + - "html" + - "md" (markdown) + - "pdf" + - "dashboard" (templated HTML dashboard) + - "clipboard" + + If export_format == "dashboard", you may optionally pass a + dashboard_template_path (a Jinja2 template file). If omitted, + a built-in default template is used. + """ + + try: + import pandas as pd + + conn = conn.connection if isinstance(conn, RocpdImportData) else conn + + # 1) Run the query via pandas + df = pd.read_sql_query(query, conn, params=params) + + if df.empty: + sys.stderr.write(f"No results found for query: {query}\n") + sys.stderr.flush() + return None + + if export_format == "console" or export_format is None: + # 2) Print to console + print(df.to_string(index=False)) + return None + + elif export_format == "clipboard": + df.to_clipboard(excel=False) + return None + + export_format = export_format.lower() + ext = export_format + export_path = export_path or f"query_output.{ext}" + if not export_path.endswith(f".{ext}"): + export_path = f"{export_path}.{ext}" + export_path = os.path.abspath(libpyrocpd.format_path(export_path, "rocpd")) + + os.makedirs(os.path.dirname(export_path), exist_ok=True) + + def write_export(content): + with open(export_path, "w") as ofs: + ofs.write(f"{content}\n") + ofs.flush() + + # 3) Export based on format + if export_format == "csv": + df.to_csv(export_path, index=False) + + elif export_format == "html": + write_export(df.to_html(index=False)) + + elif export_format == "md": + # pandas 1.0+ has to_markdown + try: + write_export(df.to_markdown(index=False)) + except AttributeError: + # fallback: manually write markdown table + _df_to_markdown_fallback(df, export_path) + + elif export_format == "pdf": + _export_df_to_pdf(df, export_path) + + elif export_format == "dashboard": + _export_dashboard( + df, export_path=export_path, template_path=dashboard_template_path + ) + + elif export_format == "json": + df.to_json(export_path, index=False, indent=2, orient="records") + + else: + print(f"Unsupported export format: {export_format}") + return None + + print(f"Exported to: {export_path}\n") + return export_path + + except Exception as e: + print(f"Error: {e}") + return None + + +def _df_to_markdown_fallback(df, path: str): + """ + Simple fallback if pandas.DataFrame.to_markdown(...) is unavailable. + """ + headers = list(df.columns) + with open(path, "w", encoding="utf-8") as f: + # Header row + f.write("| " + " | ".join(headers) + " |\n") + # Separator + f.write("|" + "|".join("---" for _ in headers) + "|\n") + # Data rows + for row in df.itertuples(index=False): + line = "| " + " | ".join(str(v) for v in row) + " |\n" + f.write(line) + + +def _export_df_to_pdf(df, path: str): + """ + Render a DataFrame into a monospaced text table inside a PDF. + """ + from reportlab.lib.pagesizes import letter + from reportlab.pdfgen import canvas + from reportlab.lib.units import inch + + c = canvas.Canvas(path, pagesize=letter) + width, height = letter + x = 0.5 * inch + y = height - 1 * inch + row_height = 14 + + c.setFont("Courier", 9) + headers = list(df.columns) + header_line = " | ".join(headers) + c.drawString(x, y, header_line) + y -= row_height + c.drawString(x, y, "-" * len(header_line)) + y -= row_height + + for _, row in df.iterrows(): + row_line = " | ".join(str(v) for v in row) + # Clip at ~160 characters so it doesn’t overflow the page width + c.drawString(x, y, row_line[:160]) + y -= row_height + if y < 1 * inch: + c.showPage() + c.setFont("Courier", 9) + y = height - 1 * inch + + c.save() + + +def _export_dashboard(df, export_path: str, template_path: Optional[str] = None): + """ + Generate a templated HTML “dashboard” from df. If template_path is None, + use a built-in template. Otherwise, load the Jinja2 template from that path. + """ + from jinja2 import Environment, FileSystemLoader, select_autoescape + + # 1) Prepare Jinja2 environment + if template_path: + # User provided a .html (Jinja2) file + env = Environment( + loader=FileSystemLoader(os.path.dirname(template_path)), + autoescape=select_autoescape(["html", "xml"]), + ) + template = env.get_template(os.path.basename(template_path)) + else: + # Built-in default template + builtin_html = """ + + + + + Dashboard Report + + + +

{{ title }}

+

Generated on {{ timestamp }}

+
+ {{ table_html | safe }} +
+ + + """ + env = Environment(autoescape=select_autoescape(["html", "xml"])) + template = env.from_string(builtin_html) + + # 2) Render template with context + context = { + "title": "SQLite Query Dashboard", + "timestamp": datetime.now().strftime("%Y-%m-%d %H:%M:%S"), + "table_html": df.to_html(index=False, classes="dashboard-table"), + } + rendered = template.render(**context) + + # 3) Write to export_path + with open(export_path, "w", encoding="utf-8") as f: + f.write(rendered) + + +def zip_files(file_paths: List[str], zip_path: str) -> str: + """ + Zip up one or more files into zip_path. Overwrites existing zip if present. + Returns the path to the created zip. + """ + import zipfile + + with zipfile.ZipFile(zip_path, "w", zipfile.ZIP_DEFLATED) as zf: + for fp in file_paths: + if os.path.isfile(fp): + zf.write(fp, arcname=os.path.basename(fp)) + else: + raise FileNotFoundError(f"Cannot find file to zip: {fp}") + print(f"Created ZIP archive: {zip_path}") + return zip_path + + +def send_report_email( + file_paths: List[str], + to: Union[str, List[str]], + sender: str, + subject: str = "rocpd query Report", + inline_preview: bool = False, + smtp_server: str = "localhost", + smtp_port: int = 25, + smtp_user: Optional[str] = None, + smtp_password: Optional[str] = None, + zip_attachments: bool = False, +) -> None: + """ + Send an email with one or more attachments, optionally zipped, + and optionally with an inline preview (if the primary attachment is HTML). + + Args: + file_paths: List of file paths to attach (each must exist). + to: Recipient email address, or list of addresses. + sender: Sender email address. + subject: Subject line. + inline_preview: If True, and one of the attachments is HTML, use that + HTML as the email body (and still attach files). + smtp_server / smtp_port / smtp_user / smtp_password: SMTP credentials. + zip_attachments: If True, bundle all file_paths into a single ZIP named + "_attachments.zip" and attach that ZIP only. + """ + import smtplib + from email.message import EmailMessage + + # 1) Validate that files exist + for fp in file_paths: + if not os.path.isfile(fp): + raise FileNotFoundError(f"Attachment not found: {fp}") + + # 2) If zip_attachments is True, zip everything into one archive + actual_attachments: List[str] + if zip_attachments: + timestamp = datetime.now().strftime("%Y%m%d_%H%M%S") + zip_path = f"attachments_{timestamp}.zip" + zip_files(file_paths, zip_path) + actual_attachments = [zip_path] + else: + actual_attachments = file_paths.copy() + + # 3) Build the EmailMessage + msg = EmailMessage() + msg["Subject"] = subject + msg["From"] = sender + msg["To"] = ", ".join(to) if isinstance(to, list) else to + + # 4) If inline_preview is True, look for the first HTML attachment, + # read its content, and set as an HTML alternative in the email body. + if inline_preview: + html_body_found = False + for fp in actual_attachments: + if fp.lower().endswith(".html"): + with open(fp, "r", encoding="utf-8") as f: + html_content = f.read() + msg.set_content( + "This email contains an inline HTML preview. If your mail client " + "doesn’t display HTML, see the attachment." + ) + msg.add_alternative(html_content, subtype="html") + html_body_found = True + break + if not html_body_found: + # No HTML attachment found; create a simple text body + msg.set_content("Please see attached report file(s).") + + else: + # No inline preview desired; use a simple text body + msg.set_content("Please see attached report file(s).") + + # 5) Attach each file (or the single ZIP) + for fp in actual_attachments: + with open(fp, "rb") as f: + data = f.read() + ctype = "application" + subtype = "octet-stream" + filename = os.path.basename(fp) + msg.add_attachment(data, maintype=ctype, subtype=subtype, filename=filename) + + # 6) Connect to SMTP and send + with smtplib.SMTP(smtp_server, smtp_port) as server: + server.ehlo() + if smtp_user and smtp_password: + server.starttls() + server.login(smtp_user, smtp_password) + server.send_message(msg) + + print(f"Email sent to {msg['To']} with subject '{subject}'") + + +def add_args(parser): + """Add query arguments""" + + query_options = parser.add_argument_group("Query Options") + + # Common arguments + query_options.add_argument( + "--query", required=True, help="SQL SELECT query to execute (enclose in quotes)." + ) + + query_options.add_argument( + "--script", + required=False, + type=str, + help="Input SQL script which should be read before query (e.g. defines views)", + ) + + query_options.add_argument( + "--format", + help="Export format", + choices=("console", "csv", "html", "json", "md", "pdf", "dashboard", "clipboard"), + type=str.lower, + ) + + email_options = parser.add_argument_group("Query Email Options") + + # Email options (optional) + email_options.add_argument( + "--email-to", help="Recipient email address (or comma-separated list)." + ) + email_options.add_argument( + "--email-from", help="Sender email address (required if --email-to is used)." + ) + email_options.add_argument( + "--email-subject", + default="SQLite Query Report", + help="Subject line for the email (default: %(default)s).", + ) + email_options.add_argument( + "--smtp-server", + default="localhost", + help="SMTP server hostname (default: %(default)s).", + ) + email_options.add_argument( + "--smtp-port", + type=int, + default=25, + help="SMTP server port (default: %(default)d).", + ) + email_options.add_argument("--smtp-user", help="SMTP login username (if required).") + email_options.add_argument( + "--smtp-password", help="SMTP login password (if required)." + ) + email_options.add_argument( + "--zip-attachments", + action="store_true", + help="Zip all attachments into a single .zip file before sending.", + ) + email_options.add_argument( + "--inline-preview", + action="store_true", + help="Embed HTML report as inline body if an HTML attachment is present.", + ) + + dashboard_options = parser.add_argument_group("Query Dashboard Options") + + dashboard_options.add_argument( + "--template-path", help="Path to a Jinja2 HTML template for the dashboard" + ) + + return [ + "query", + "script", + "email_to", + "email_from", + "email_subject", + "smtp_server", + "smtp_port", + "smtp_user", + "smtp_password", + "inline_preview", + "zip_attachments", + "format", + "template_path", + ] + + +def process_args(args, valid_args): + # do not add any of the arguments to the output config dict + ret = {} + return ret + + +def execute(input, args, config=None, window_args=None, **kwargs): + + importData = RocpdImportData(input) + + apply_time_window(importData, **window_args) + + config = ( + output_config.output_config(**kwargs) + if config is None + else config.update(**kwargs) + ) + + if args.script: + # read script and execute statements + with open(args.script, "r") as ifs: + for itr in ifs.read().split(";"): + importData.execute(f"{itr}") + + # Prepare parameters for export + query = args.query + db = importData + export_format = args.format + export_path = os.path.join(config.output_path, config.output_file) + + # Dashboard-only extra + dashboard_template = kwargs.get("template_path", None) + + # 1) Run and export + exported_file = export_sqlite_query( + db, + query=query, + params=(), + export_format=export_format, + export_path=export_path, + dashboard_template_path=dashboard_template, + ) + + # 2) If --email-to was provided and we have a file, send it + if args.email_to: + if not args.email_from: + raise ValueError("--email-from is required when --email-to is used.") + if not exported_file: + print("No file was exported; skipping email.") + return + + recipients = [addr.strip() for addr in args.email_to.split(",")] + send_report_email( + file_paths=[exported_file], + to=recipients, + sender=args.email_from, + subject=args.email_subject, + inline_preview=args.inline_preview, + smtp_server=args.smtp_server, + smtp_port=args.smtp_port, + smtp_user=args.smtp_user, + smtp_password=args.smtp_password, + zip_attachments=args.zip_attachments, + ) + + +def main(argv=None): + import argparse + from .time_window import add_args as add_args_time_window + from .time_window import process_args as process_args_time_window + from .output_config import add_args as add_args_output_config + from .output_config import process_args as process_args_output_config + from .output_config import add_generic_args, process_generic_args + + parser = argparse.ArgumentParser( + description="Generate report for rocpd query", allow_abbrev=False + ) + + required_params = parser.add_argument_group("Required options") + + required_params.add_argument( + "-i", + "--input", + required=True, + type=output_config.check_file_exists, + nargs="+", + help="Input path and filename to one or more database(s), separated by spaces", + ) + + valid_out_config_args = add_args_output_config(parser) + valid_generic_args = add_generic_args(parser) + valid_time_window_args = add_args_time_window(parser) + valid_query_args = add_args(parser) + + args = parser.parse_args(argv) + + out_cfg_args = process_args_output_config(args, valid_out_config_args) + generic_out_cfg_args = process_generic_args(args, valid_generic_args) + window_args = process_args_time_window(args, valid_time_window_args) + query_args = process_args(args, valid_query_args) + + all_args = { + **query_args, + **out_cfg_args, + **generic_out_cfg_args, + } + + execute( + args.input, + args, + window_args=window_args, + **all_args, + ) + + +if __name__ == "__main__": + main() diff --git a/source/lib/python/rocpd/summary.py b/source/lib/python/rocpd/summary.py new file mode 100644 index 0000000000..e9295cb638 --- /dev/null +++ b/source/lib/python/rocpd/summary.py @@ -0,0 +1,532 @@ +#!/usr/bin/env python3 +############################################################################### +# MIT License +# +# Copyright (c) 2025 Advanced Micro Devices, Inc. +# +# 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. +############################################################################### + +import argparse +import os + +from typing import Any, List, Tuple +from .importer import RocpdImportData, execute_statement +from .query import export_sqlite_query +from . import output_config + + +def get_temp_view_names(connection: RocpdImportData) -> List[str]: + """Return the names of all temporary views in the SQLite connection.""" + return [ + v[0] + for v in execute_statement( + connection, "SELECT name FROM sqlite_temp_master WHERE type='view';" + ).fetchall() + ] + + +def get_temp_view_columns(connection: RocpdImportData, view_name: str) -> List[str]: + """Return the column names of a given temporary view.""" + cursor = connection.cursor() + cursor.execute(f"PRAGMA table_xinfo('{view_name}')") + return [row[1] for row in cursor.fetchall()] + + +def make_temp_view_query(view_name, query) -> str: + return "CREATE TEMPORARY VIEW IF NOT EXISTS `{}` AS {}".format(view_name, query) + + +def export_view( + connection: RocpdImportData, view_name, output_format, output_path, filename="" +) -> None: + """Write the contents of a SQL view to an output format.""" + + query = "SELECT * FROM `{}`".format(view_name) + query_one = "SELECT * FROM `{}` LIMIT 1".format(view_name) + + # just return if view is empty + if not connection.execute(query_one).fetchone(): + return + + # prepare the output filename + if not filename: + output_filename = view_name + else: + output_filename = f"{filename}_{view_name}" + + if output_format == "console": + print(f"\n{view_name.upper()}:") + + # call query module to export. query will append the extension + export_path = os.path.join(output_path, output_filename) + export_sqlite_query( + connection, query, export_format=output_format, export_path=export_path + ) + + +def generate_summary_query( + view_name: str, + name_column="name", + by_rank=False, +) -> Tuple[str, str]: + """Generate the SQL statement to create a summary view.""" + + if by_rank: + view_suffix = "_summary_by_rank" + group_by_columns = "guid, {name_column}".format(name_column=name_column) + aggregation_group_by = "T.guid, T.nid, T.{name_column}".format( + name_column=name_column + ) + total_duration_group_by = "guid" + additional_select_columns = "AD.pid AS ProcessID, P.hostname AS Hostname," + additional_aggregated_columns = """ + T.guid, + T.nid, + T.pid,""" + join_condition = "T.guid = A.guid AND T.{name_column} = A.name".format( + name_column=name_column + ) + total_duration_join = "JOIN total_duration TD ON AD.guid = TD.guid JOIN processes P ON AD.pid = P.pid" + else: + view_suffix = "_summary" + group_by_columns = name_column + aggregation_group_by = "T.{name_column}".format(name_column=name_column) + total_duration_group_by = "" + additional_select_columns = "" + additional_aggregated_columns = "" + join_condition = "T.{name_column} = A.name".format(name_column=name_column) + total_duration_join = "CROSS JOIN total_duration TD" + + full_view_name = f"{view_name}{view_suffix}" + + summary_query = f""" + WITH + avg_data AS ( + SELECT + {group_by_columns.replace(name_column, f"{name_column} AS name")}, + AVG(duration) AS avg_duration + FROM {view_name} + GROUP BY {group_by_columns} + ), + aggregated_data AS ( + SELECT{additional_aggregated_columns} + T.{name_column} as name, + COUNT(*) AS calls, + SUM(T.duration) AS total_duration, + A.avg_duration AS average_duration, + MIN(T.duration) AS min_duration, + MAX(T.duration) AS max_duration, + SQRT(SUM(CAST((T.duration - A.avg_duration) AS REAL) * CAST((T.duration - A.avg_duration) AS REAL)) / (COUNT(*) - 1)) AS std_dev_duration + FROM {view_name} T + JOIN avg_data A ON {join_condition} + GROUP BY {aggregation_group_by} + ), + total_duration AS ( + SELECT + {f"{total_duration_group_by}," if total_duration_group_by else ""} + SUM(total_duration) AS grand_total_duration + FROM + aggregated_data + {f"GROUP BY {total_duration_group_by}" if total_duration_group_by else ""} + ) + SELECT + {additional_select_columns} + AD.name AS Name, + AD.calls AS Calls, + AD.total_duration AS "DURATION (nsec)", + AD.average_duration AS "AVERAGE (nsec)", + (CAST(AD.total_duration AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", + AD.min_duration AS "MIN (nsec)", + AD.max_duration AS "MAX (nsec)", + AD.std_dev_duration AS "STD_DEV" + FROM + aggregated_data AD + {total_duration_join} + ORDER BY + {"AD.pid," if by_rank else ""} AD.total_duration DESC; + """ + + return (full_view_name, summary_query) + + +def generate_domain_query(connection: RocpdImportData, by_rank=False) -> Tuple[str, str]: + """Generate the SQL statement for domain summary by doing union over all summary views.""" + + if by_rank: + view_suffix = "_summary_by_rank" + view_name = "domain_summary_by_rank" + additional_group_columns = "ProcessID, Hostname," + additional_select_columns = "GD.ProcessID, GD.Hostname," + total_duration_group_by = "GROUP BY ProcessID" + join_condition = "JOIN total_duration TD ON GD.ProcessID = TD.ProcessID" + order_by = "ORDER BY GD.ProcessID" + else: + view_suffix = "_summary" + view_name = "domain_summary" + additional_group_columns = "" + additional_select_columns = "" + total_duration_group_by = "" + join_condition = "CROSS JOIN total_duration TD" + order_by = 'ORDER BY GD."DURATION (nsec)" DESC' + + summary_views = [ + itr for itr in get_temp_view_names(connection) if itr.endswith(view_suffix) + ] + + if len(summary_views) < 1: + return view_name + + union_selects = [ + f" SELECT '{s.replace(view_suffix, '').upper()}' as domain, * FROM {s} " + for s in summary_views + ] + + domain_select = f""" + WITH + all_domains AS ( + {f" UNION ALL ".join(union_selects)} + ), + grouped_domains AS ( + SELECT + domain, + {additional_group_columns} + SUM(calls) AS calls, + SUM("DURATION (nsec)") AS "DURATION (nsec)", + SUM("AVERAGE (nsec)") AS "AVERAGE (nsec)", + MIN("MIN (nsec)") AS "MIN (nsec)", + MAX("MAX (nsec)") AS "MAX (nsec)", + SUM("STD_DEV") AS "STD_DEV" + FROM all_domains + GROUP BY domain{", ProcessID" if by_rank else ""} + ), + total_duration AS ( + SELECT + {additional_group_columns} + SUM("DURATION (nsec)") AS grand_total_duration + FROM grouped_domains + {total_duration_group_by} + ) + SELECT + {additional_select_columns} + GD.domain AS Name, + GD.calls AS Calls, + GD."DURATION (nsec)", + GD."AVERAGE (nsec)", + (CAST(GD."DURATION (nsec)" AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", + GD."MIN (nsec)", + GD."MAX (nsec)", + GD."STD_DEV" + FROM + grouped_domains GD + {join_condition} + {order_by}; + """ + + return (view_name, domain_select) + + +def create_summary_views(connection: RocpdImportData, by_rank=False) -> None: + """Create summary views for eligible temporary views in the database.""" + + NAME_COLUMN_MAP = { + "memory_allocations": "type", + "scratch_memory": "operation", + } + + avoid_view_pattern = ("rocpd", "region", "counter", "pmc") + required_columns = {"duration"} + + views = get_temp_view_names(connection) + + for view_name in views: + if any(pattern in view_name for pattern in avoid_view_pattern): + continue + + columns = get_temp_view_columns(connection, view_name) + if not required_columns.issubset(columns): + continue + + # Create regular summary view + summary_view_name, summary_query = generate_summary_query( + view_name, name_column=NAME_COLUMN_MAP.get(view_name, "name") + ) + connection.execute(make_temp_view_query(summary_view_name, summary_query)) + + # Create per-rank summary + if by_rank: + per_rank_view_name, summary_by_rank_query = generate_summary_query( + view_name, + name_column=NAME_COLUMN_MAP.get(view_name, "name"), + by_rank=True, + ) + connection.execute( + make_temp_view_query(per_rank_view_name, summary_by_rank_query) + ) + + +def create_summary_region_views( + connection: RocpdImportData, by_rank=False, region_categories=None +) -> None: + """Create summary and region views""" + + query = "SELECT DISTINCT(category) FROM regions_and_samples;" + categories = execute_statement(connection, query).fetchall() + + if region_categories is None: + # Automatically retrieve region categories from the database + region_categories = set([cat[0].split("_")[0] for cat in categories]) + + category_map = { + cat.lower(): [c[0] for c in categories if c[0].startswith(cat + "_")] + for cat in region_categories + if "MARKER" not in cat.upper() + } + + for k, v in category_map.items(): + if len(v) > 0: + conditions = [f"category LIKE '{c}'" for c in v] + temp_region_view = f""" + CREATE TEMPORARY VIEW IF NOT EXISTS `{k}` AS + SELECT * + FROM regions_and_samples + WHERE {" OR ".join(conditions)}; + """ + + connection.execute(temp_region_view) + + # Create regular summary view + summary_view_name, summary_query = generate_summary_query(k) + connection.execute(make_temp_view_query(summary_view_name, summary_query)) + + # Create per-rank summary view + if by_rank: + per_rank_view_name, summary_by_rank_query = generate_summary_query( + k, by_rank=True + ) + connection.execute( + make_temp_view_query(per_rank_view_name, summary_by_rank_query) + ) + + # Markers + if "MARKER" not in region_categories: + return + + view_name = "markers" + markers_create = f""" + CREATE TEMPORARY VIEW IF NOT EXISTS `{view_name}` AS + SELECT JSON_EXTRACT(extdata, '$.message') AS marker_name, * + FROM regions_and_samples + WHERE category LIKE 'MARKER_%' + """ + connection.execute(markers_create) + + # Create regular summary view + summary_view_name, summary_query = generate_summary_query( + view_name, name_column="marker_name" + ) + connection.execute(make_temp_view_query(summary_view_name, summary_query)) + + # Create per-rank summary view + if by_rank: + per_rank_view_name, summary_by_rank_query = generate_summary_query( + view_name, name_column="marker_name", by_rank=True + ) + connection.execute( + make_temp_view_query(per_rank_view_name, summary_by_rank_query) + ) + + +def create_domain_view(connection: RocpdImportData, by_rank=False) -> str: + """Create a domain summary view by aggregating all summary views.""" + + view_name, domain_query = generate_domain_query(connection, by_rank=by_rank) + + # Create the domain summary view + connection.execute(make_temp_view_query(view_name, domain_query)) + + return view_name + + +def generate_all_summaries(connection: RocpdImportData, **kwargs: Any) -> None: + """Generate all summary views and write them to CSV files.""" + + domain_summary = kwargs.get("domain_summary", False) + by_rank = kwargs.get("summary_by_rank", False) + filename = kwargs.get("output_file", "") + output_path = kwargs.get("output_path", "./rocpd-output-data") + region_categories = kwargs.get("region_categories", None) + output_format = kwargs.get("format", "console") + + # create the temporary summary views + create_summary_views(connection, by_rank) + create_summary_region_views(connection, by_rank, region_categories=region_categories) + + if domain_summary: + create_domain_view(connection) + # Create domain summary per rank only if both domain_summary and summary_by_rank are enabled + if by_rank: + create_domain_view(connection, by_rank=True) + + # Write regular summary views + print("\nSummary files:") + summary_views = [ + itr for itr in get_temp_view_names(connection) if itr.endswith("_summary") + ] + for v in summary_views: + export_view(connection, v, output_format, output_path, filename) + + # Write per-rank summary views if flag is set + if by_rank: + print("\nSummary files by rank:") + summary_by_rank_views = [ + itr + for itr in get_temp_view_names(connection) + if itr.endswith("_summary_by_rank") + ] + for v in summary_by_rank_views: + export_view(connection, v, output_format, output_path, filename) + + +# +# Command-line interface functions +# + + +def add_io_args(parser): + """Add input/output arguments for summary.""" + io_options = parser.add_argument_group("I/O options") + + io_options.add_argument( + "-f", + "--format", + help="Sets the format the summaries are output to (default: console)", + choices=("console", "csv", "html", "json", "md", "pdf"), + default="console", + type=str, + required=False, + ) + io_options.add_argument( + "-o", + "--output-file", + help="Sets the base output file name", + default=os.environ.get("ROCPD_OUTPUT_NAME", ""), + type=str, + required=False, + ) + io_options.add_argument( + "-d", + "--output-path", + help="Sets the output path where the output files will be saved (default path: `./rocpd-output-data`)", + default=os.environ.get("ROCPD_OUTPUT_PATH", "./rocpd-output-data"), + type=str, + required=False, + ) + + return ["format", "output_file", "output_path"] + + +def add_args(parser): + """Add arguments for summary.""" + summary_options = parser.add_argument_group("Summary options") + summary_options.add_argument( + "--domain-summary", + action="store_true", + default=False, + help="Generate domain summary view", + ) + summary_options.add_argument( + "--summary-by-rank", + action="store_true", + default=False, + help="Generate summary views by-rank (or Process ID)", + ) + summary_options.add_argument( + "--region-categories", + nargs="+", + default=None, + help="Specify region categories to include in the summary (example: HIP, HSA, RCCL, ROCDECODE, ROCJPEG, MARKER). If not specified, categories will be automatically retrieved from the database.", + ) + + return ["domain_summary", "summary_by_rank", "region_categories"] + + +def process_args(args, valid_args): + + ret = {} + for itr in valid_args: + if hasattr(args, itr): + val = getattr(args, itr) + if val is not None: + ret[itr] = val + return ret + + +def execute(input, window_args=None, **kwargs: Any) -> RocpdImportData: + from .time_window import apply_time_window + + importData = RocpdImportData(input) + + apply_time_window(importData, **window_args) + + generate_all_summaries(importData, **kwargs) + + return importData + + +def main(argv=None) -> int: + """Main entry point for command line execution.""" + from .time_window import add_args as add_args_time_window + from .time_window import process_args as process_args_time_window + + parser = argparse.ArgumentParser( + description="Create ROCpd database summary region views" + ) + required_params = parser.add_argument_group("Required options") + + required_params.add_argument( + "-i", + "--input", + required=True, + type=output_config.check_file_exists, + nargs="+", + help="Input path and filename to one or more database(s), separated by spaces", + ) + + valid_io_args = add_io_args(parser) + valid_summary_args = add_args(parser) + valid_time_window_args = add_args_time_window(parser) + + args = parser.parse_args(argv) + + summary_args = process_args(args, valid_summary_args) + io_args = output_config.process_args(args, valid_io_args) + window_args = process_args_time_window(args, valid_time_window_args) + + all_args = {**summary_args, **io_args} + + execute( + args.input, + window_args=window_args, + **all_args, + ) + + +if __name__ == "__main__": + main() diff --git a/source/lib/python/roctx/__init__.py b/source/lib/python/roctx/__init__.py index 701121babb..64b874be83 100644 --- a/source/lib/python/roctx/__init__.py +++ b/source/lib/python/roctx/__init__.py @@ -38,8 +38,24 @@ __all__ = [ "nameOsThread", "nameHipDevice", "context_decorators", + "version_info", ] +version_info = { + "version": "@PROJECT_VERSION@", + "major": int("@PROJECT_VERSION_MAJOR@"), + "minor": int("@PROJECT_VERSION_MINOR@"), + "patch": int("@PROJECT_VERSION_PATCH@"), + "git_revision": "@ROCPROFILER_SDK_GIT_REVISION@", + "library_arch": "@CMAKE_LIBRARY_ARCHITECTURE@", + "system_name": "@CMAKE_SYSTEM_NAME@", + "system_processor": "@CMAKE_SYSTEM_PROCESSOR@", + "system_version": "@CMAKE_SYSTEM_VERSION@", + "compiler_id": "@CMAKE_CXX_COMPILER_ID@", + "compiler_version": "@CMAKE_CXX_COMPILER_VERSION@", + "rocm_version": "@rocm_version_FULL_VERSION@", +} + def mark(msg): return libpyroctx.roctxMark(msg) if msg is not None else None diff --git a/source/lib/python/utilities.cmake b/source/lib/python/utilities.cmake index 6bae034df2..9c8a987d76 100644 --- a/source/lib/python/utilities.cmake +++ b/source/lib/python/utilities.cmake @@ -97,7 +97,7 @@ function(rocprofiler_roctx_python_bindings _VERSION) foreach(_SOURCE ${roctx_PYTHON_SOURCES}) configure_file(${CMAKE_CURRENT_LIST_DIR}/${_SOURCE} - ${roctx_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} COPYONLY) + ${roctx_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} @ONLY) install( FILES ${roctx_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} DESTINATION ${roctx_PYTHON_INSTALL_DIRECTORY} @@ -154,12 +154,14 @@ function(rocprofiler_rocpd_python_bindings _VERSION) output_config.py otf2.py pftrace.py + query.py schema.py + summary.py time_window.py) foreach(_SOURCE ${rocpd_PYTHON_SOURCES}) configure_file(${CMAKE_CURRENT_LIST_DIR}/${_SOURCE} - ${rocpd_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} COPYONLY) + ${rocpd_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} @ONLY) install( FILES ${rocpd_PYTHON_OUTPUT_DIRECTORY}/${_SOURCE} DESTINATION ${rocpd_PYTHON_INSTALL_DIRECTORY} diff --git a/source/share/rocprofiler-sdk-rocpd/data_views.sql b/source/share/rocprofiler-sdk-rocpd/data_views.sql index 868542c99f..07095dc29d 100644 --- a/source/share/rocprofiler-sdk-rocpd/data_views.sql +++ b/source/share/rocprofiler-sdk-rocpd/data_views.sql @@ -596,6 +596,7 @@ SELECT JSON_EXTRACT(M.extdata, '$.flags') AS alloc_flags, M.start, M.end, + (M.end - M.start) AS duration, M.size, M.address, E.correlation_id, diff --git a/source/share/rocprofiler-sdk-rocpd/summary_views.sql b/source/share/rocprofiler-sdk-rocpd/summary_views.sql index 3575bf6f55..357eecec0e 100644 --- a/source/share/rocprofiler-sdk-rocpd/summary_views.sql +++ b/source/share/rocprofiler-sdk-rocpd/summary_views.sql @@ -151,226 +151,3 @@ GROUP BY name ORDER BY total_duration DESC; - --- Kernel summary by name -CREATE VIEW - `kernel_summary` AS -WITH - avg_data AS ( - SELECT - name, - AVG(duration) AS avg_duration - FROM - `kernels` - GROUP BY - name - ), - aggregated_data AS ( - SELECT - K.name, - COUNT(*) AS calls, - SUM(K.duration) AS total_duration, - SUM(CAST(K.duration AS REAL) * CAST(K.duration AS REAL)) AS sqr_duration, - A.avg_duration AS average_duration, - MIN(K.duration) AS min_duration, - MAX(K.duration) AS max_duration, - SUM(CAST((K.duration - A.avg_duration) AS REAL) * CAST((K.duration - A.avg_duration) AS REAL)) / (COUNT(*) - 1) AS variance_duration, - SQRT( - SUM(CAST((K.duration - A.avg_duration) AS REAL) * CAST((K.duration - A.avg_duration) AS REAL)) / (COUNT(*) - 1) - ) AS std_dev_duration - FROM - `kernels` K - JOIN avg_data A ON K.name = A.name - GROUP BY - K.name - ), - total_duration AS ( - SELECT - SUM(total_duration) AS grand_total_duration - FROM - aggregated_data - ) -SELECT - AD.name AS name, - AD.calls, - AD.total_duration AS "DURATION (nsec)", - AD.sqr_duration AS "SQR (nsec)", - AD.average_duration AS "AVERAGE (nsec)", - (CAST(AD.total_duration AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", - AD.min_duration AS "MIN (nsec)", - AD.max_duration AS "MAX (nsec)", - AD.variance_duration AS "VARIANCE", - AD.std_dev_duration AS "STD_DEV" -FROM - aggregated_data AD - CROSS JOIN total_duration TD; - --- --- Kernel summary by region name -CREATE VIEW - `kernel_summary_region` AS -WITH - avg_data AS ( - SELECT - region, - AVG(duration) AS avg_duration - FROM - `kernels` - GROUP BY - region - ), - aggregated_data AS ( - SELECT - K.region AS name, - COUNT(*) AS calls, - SUM(K.duration) AS total_duration, - SUM(CAST(K.duration AS REAL) * CAST(K.duration AS REAL)) AS sqr_duration, - A.avg_duration AS average_duration, - MIN(K.duration) AS min_duration, - MAX(K.duration) AS max_duration, - SUM(CAST((K.duration - A.avg_duration) AS REAL) * CAST((K.duration - A.avg_duration) AS REAL)) / (COUNT(*) - 1) AS variance_duration, - SQRT( - SUM(CAST((K.duration - A.avg_duration) AS REAL) * CAST((K.duration - A.avg_duration) AS REAL)) / (COUNT(*) - 1) - ) AS std_dev_duration - FROM - `kernels` K - JOIN avg_data A ON K.region = A.region - GROUP BY - K.region - ), - total_duration AS ( - SELECT - SUM(total_duration) AS grand_total_duration - FROM - aggregated_data - ) -SELECT - AD.name AS name, - AD.calls, - AD.total_duration AS "DURATION (nsec)", - AD.sqr_duration AS "SQR (nsec)", - AD.average_duration AS "AVERAGE (nsec)", - (CAST(AD.total_duration AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", - AD.min_duration AS "MIN (nsec)", - AD.max_duration AS "MAX (nsec)", - AD.variance_duration AS "VARIANCE", - AD.std_dev_duration AS "STD_DEV" -FROM - aggregated_data AD - CROSS JOIN total_duration TD; - --- --- Memory copy summary -CREATE VIEW - `memory_copy_summary` AS -WITH - avg_data AS ( - SELECT - name, - AVG(duration) AS avg_duration - FROM - `memory_copies` - GROUP BY - name - ), - aggregated_data AS ( - SELECT - MC.name, - COUNT(*) AS calls, - SUM(MC.duration) AS total_duration, - SUM(CAST(MC.duration AS REAL) * CAST(MC.duration AS REAL)) AS sqr_duration, - A.avg_duration AS average_duration, - MIN(MC.duration) AS min_duration, - MAX(MC.duration) AS max_duration, - SUM( - CAST((MC.duration - A.avg_duration) AS REAL) * CAST((MC.duration - A.avg_duration) AS REAL) - ) / (COUNT(*) - 1) AS variance_duration, - SQRT( - SUM( - CAST((MC.duration - A.avg_duration) AS REAL) * CAST((MC.duration - A.avg_duration) AS REAL) - ) / (COUNT(*) - 1) - ) AS std_dev_duration - FROM - `memory_copies` MC - JOIN avg_data A ON MC.name = A.name - GROUP BY - MC.name - ), - total_duration AS ( - SELECT - SUM(total_duration) AS grand_total_duration - FROM - aggregated_data - ) -SELECT - AD.name AS name, - AD.calls, - AD.total_duration AS "DURATION (nsec)", - AD.sqr_duration AS "SQR (nsec)", - AD.average_duration AS "AVERAGE (nsec)", - (CAST(AD.total_duration AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", - AD.min_duration AS "MIN (nsec)", - AD.max_duration AS "MAX (nsec)", - AD.variance_duration AS "VARIANCE", - AD.std_dev_duration AS "STD_DEV" -FROM - aggregated_data AD - CROSS JOIN total_duration TD; - --- --- Memory allocation summary -CREATE VIEW - `memory_allocation_summary` AS -WITH - avg_data AS ( - SELECT - type AS name, - AVG(duration) AS avg_duration - FROM - `memory_allocations` - GROUP BY - type - ), - aggregated_data AS ( - SELECT - MA.type AS name, - COUNT(*) AS calls, - SUM(MA.duration) AS total_duration, - SUM(CAST(MA.duration AS REAL) * CAST(MA.duration AS REAL)) AS sqr_duration, - A.avg_duration AS average_duration, - MIN(MA.duration) AS min_duration, - MAX(MA.duration) AS max_duration, - SUM( - CAST((MA.duration - A.avg_duration) AS REAL) * CAST((MA.duration - A.avg_duration) AS REAL) - ) / (COUNT(*) - 1) AS variance_duration, - SQRT( - SUM( - CAST((MA.duration - A.avg_duration) AS REAL) * CAST((MA.duration - A.avg_duration) AS REAL) - ) / (COUNT(*) - 1) - ) AS std_dev_duration - FROM - `memory_allocations` MA - JOIN avg_data A ON MA.type = A.name - GROUP BY - MA.type - ), - total_duration AS ( - SELECT - SUM(total_duration) AS grand_total_duration - FROM - aggregated_data - ) -SELECT - 'MEMORY_ALLOCATION_' || AD.name AS name, - AD.calls, - AD.total_duration AS "DURATION (nsec)", - AD.sqr_duration AS "SQR (nsec)", - AD.average_duration AS "AVERAGE (nsec)", - (CAST(AD.total_duration AS REAL) / TD.grand_total_duration) * 100 AS "PERCENT (INC)", - AD.min_duration AS "MIN (nsec)", - AD.max_duration AS "MAX (nsec)", - AD.variance_duration AS "VARIANCE", - AD.std_dev_duration AS "STD_DEV" -FROM - aggregated_data AD - CROSS JOIN total_duration TD; diff --git a/tests/rocpd/CMakeLists.txt b/tests/rocpd/CMakeLists.txt index 1d365006cf..8b55f1f145 100644 --- a/tests/rocpd/CMakeLists.txt +++ b/tests/rocpd/CMakeLists.txt @@ -22,7 +22,7 @@ set(rocpd-env # ######################################################################################### -foreach(_SUBPARSER "" "-convert") +foreach(_SUBPARSER "" "-convert" "-query" "-summary") string(REPLACE "-" "" _CMD "${_SUBPARSER}") add_test(NAME rocpd${_SUBPARSER}-help COMMAND ${Python3_EXECUTABLE} -m rocpd ${_CMD} --help) @@ -34,7 +34,7 @@ foreach(_SUBPARSER "" "-convert") "${ROCPROFILER_DEFAULT_FAIL_REGEX}") endforeach() -foreach(_MODULE "csv" "pftrace" "otf2") +foreach(_MODULE "csv" "pftrace" "otf2" "query" "summary") add_test(NAME rocpd-module-${_MODULE}-help COMMAND ${Python3_EXECUTABLE} -m rocpd.${_MODULE} --help) diff --git a/tests/rocprofv3/rocpd/CMakeLists.txt b/tests/rocprofv3/rocpd/CMakeLists.txt index 1bafc1730c..8d06c7599e 100644 --- a/tests/rocprofv3/rocpd/CMakeLists.txt +++ b/tests/rocprofv3/rocpd/CMakeLists.txt @@ -222,6 +222,58 @@ set_tests_properties( FIXTURES_REQUIRED rocprofv3-test-rocpd) +######################################################################################### +# +# Summary generate +# +######################################################################################### +add_test( + NAME rocprofv3-test-rocpd-summary-generation + COMMAND + ${Python3_EXECUTABLE} -m rocpd summary --domain-summary --summary-by-rank -f csv + -d ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/summary -i + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-input-data/out_results.db) + +set_tests_properties( + rocprofv3-test-rocpd-summary-generation + PROPERTIES TIMEOUT + 45 + LABELS + "integration-tests;rocpd" + ENVIRONMENT + "${rocprofv3-rocpd-env}" + DEPENDS + "rocprofv3-test-rocpd-execute" + FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}" + FIXTURES_SETUP + rocprofv3-test-rocpd-generation + FIXTURES_REQUIRED + rocprofv3-test-rocpd) + +add_test( + NAME rocprofv3-test-rocpd-summary-generation-multiproc + COMMAND + ${Python3_EXECUTABLE} -m rocpd summary --domain-summary --summary-by-rank -f csv + -d ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-test/summary -o out_mp -i + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-input-data-multiproc/out_mp_0_results.db + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-input-data-multiproc/out_mp_1_results.db) + +set_tests_properties( + rocprofv3-test-rocpd-summary-generation-multiproc + PROPERTIES TIMEOUT + 120 + LABELS + "integration-tests;rocpd" + ENVIRONMENT + "${rocprofv3-rocpd-env}" + FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}" + DISABLED + "${MULTIPROC_IS_DISABLED}" + FIXTURES_REQUIRED + rocprofv3-test-rocpd-multiproc) + ######################################################################################### # # Validation From 2d8936362e8a25d75c68c43367ef69be840b7efd Mon Sep 17 00:00:00 2001 From: "Indic, Vladimir" Date: Tue, 29 Jul 2025 21:11:15 +0200 Subject: [PATCH 6/7] PCS test: cast agent name to str (#546) * PCS test: cast agent name to str --- .../host-trap/transpose-multiple-agents/conftest.py | 9 ++++++++- .../stochastic/transpose-multiple-agents/conftest.py | 9 ++++++++- 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/tests/rocprofv3/pc-sampling/host-trap/transpose-multiple-agents/conftest.py b/tests/rocprofv3/pc-sampling/host-trap/transpose-multiple-agents/conftest.py index 9a19c5768f..d88b9bc4fc 100644 --- a/tests/rocprofv3/pc-sampling/host-trap/transpose-multiple-agents/conftest.py +++ b/tests/rocprofv3/pc-sampling/host-trap/transpose-multiple-agents/conftest.py @@ -84,4 +84,11 @@ def input_kernel_trace_csv(request): def input_agent_info_csv(request): filename = request.config.getoption("--input-agent-info-csv") with open(filename, "r") as inp: - return pd.read_csv(inp) + return pd.read_csv( + inp, + na_filter=False, # parse empty fields as "" + keep_default_na=False, # parse empty fields as "" + dtype={ + "Name": str, + }, + ) diff --git a/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/conftest.py b/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/conftest.py index 7472a41c13..56b39cca61 100644 --- a/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/conftest.py +++ b/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/conftest.py @@ -93,7 +93,14 @@ def input_kernel_trace_csv(request): def input_agent_info_csv(request): filename = request.config.getoption("--input-agent-info-csv") with open(filename, "r") as inp: - return pd.read_csv(inp) + return pd.read_csv( + inp, + na_filter=False, # parse empty fields as "" + keep_default_na=False, # parse empty fields as "" + dtype={ + "Name": str, + }, + ) @pytest.fixture From 1ba08cd4dfb7fe99a51765019210947dfcd199f7 Mon Sep 17 00:00:00 2001 From: "Baraldi, Giovanni" Date: Tue, 29 Jul 2025 22:47:40 +0200 Subject: [PATCH 7/7] Removing ATT buffer size limitation (#534) * Removing SQTT buffer size limitation * Update source/lib/rocprofiler-sdk/thread_trace/core.cpp * Added testing for buffer size. Formatting. * Add test as unstable * Increase default buffer size * Apply suggestions from code review Co-authored-by: Indic, Vladimir * Fix typo from code review * Update tests/thread-trace/agent.cpp --------- Co-authored-by: Giovanni Co-authored-by: Indic, Vladimir --- source/bin/rocprofv3.py | 2 +- source/lib/rocprofiler-sdk-tool/config.hpp | 3 +- .../lib/rocprofiler-sdk/aql/aql_profile_v2.h | 8 ++++++ .../rocprofiler-sdk/aql/packet_construct.cpp | 8 ++++-- .../lib/rocprofiler-sdk/thread_trace/core.cpp | 9 +++--- tests/thread-trace/CMakeLists.txt | 17 +++++++++++ tests/thread-trace/agent.cpp | 28 ++++++++++--------- tests/thread-trace/trace_callbacks.hpp | 1 - 8 files changed, 53 insertions(+), 23 deletions(-) diff --git a/source/bin/rocprofv3.py b/source/bin/rocprofv3.py index 571be911ba..1b2acd5dbc 100755 --- a/source/bin/rocprofv3.py +++ b/source/bin/rocprofv3.py @@ -769,7 +769,7 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins att_options.add_argument( "--att-buffer-size", - help="Thread trace buffer size. Default 96MB", + help="Thread trace buffer size. Default 256MB", default=None, type=str, ) diff --git a/source/lib/rocprofiler-sdk-tool/config.hpp b/source/lib/rocprofiler-sdk-tool/config.hpp index a253439542..7da04a9a9c 100644 --- a/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/source/lib/rocprofiler-sdk-tool/config.hpp @@ -139,7 +139,8 @@ struct config : output_config int mpi_rank = get_mpi_rank(); uint64_t att_param_shader_engine_mask = get_env("ROCPROF_ATT_PARAM_SHADER_ENGINE_MASK", 0x1); - uint64_t att_param_buffer_size = get_env("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x6000000); + // 256MB + uint64_t att_param_buffer_size = get_env("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x10000000); 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); diff --git a/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h b/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h index 3dcacbb3df..4771f1f204 100644 --- a/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h +++ b/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h @@ -184,6 +184,14 @@ aqlprofile_get_pmc_info(const aqlprofile_pmc_profile_t* profile, aqlprofile_pmc_info_type_t attribute, void* value); +typedef enum aqlprofile_att_parameter_name_ext_t +{ + /** + * HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE + 1 + */ + AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH = 11, +} aqlprofile_att_parameter_name_ext_t; + // Profile parameter object typedef struct { diff --git a/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index 4c739adb83..6425bd3884 100644 --- a/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -145,7 +145,8 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache& uint32_t cu = static_cast(params.target_cu); uint32_t shader_engine_mask = static_cast(params.shader_engine_mask); uint32_t simd = static_cast(params.simd_select); - uint32_t buffer_size = static_cast(params.buffer_size); + uint32_t buffer_size_lo = static_cast(params.buffer_size); + uint32_t buffer_size_hi = static_cast(params.buffer_size >> 32); uint32_t perf_ctrl = static_cast(params.perfcounter_ctrl); aql_params.clear(); @@ -153,7 +154,10 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache& aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET, {cu}}); aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SE_MASK, {shader_engine_mask}}); aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SIMD_SELECTION, {simd}}); - aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE, {buffer_size}}); + aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE, {buffer_size_lo}}); + + if(buffer_size_hi != 0) aql_params.push_back({static_cast( + AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH), {buffer_size_hi}}); if(perf_ctrl != 0 && !params.perfcounters.empty()) { diff --git a/source/lib/rocprofiler-sdk/thread_trace/core.cpp b/source/lib/rocprofiler-sdk/thread_trace/core.cpp index 16360ad5b2..8f50e08f00 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/core.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/core.cpp @@ -57,8 +57,7 @@ namespace rocprofiler namespace thread_trace { constexpr size_t QUEUE_SIZE = 128; -constexpr uint64_t MIN_BUFFER_SIZE = 1 << 18; // 2 pages per SE -constexpr uint64_t MAX_BUFFER_SIZE = std::numeric_limits::max(); // aqlprofile limit +constexpr uint64_t MIN_BUFFER_SIZE = 1 << 20; // 1MB struct cbdata_t { @@ -74,15 +73,15 @@ thread_trace_parameter_pack::are_params_valid() const { if(shader_cb_fn == nullptr) { - ROCP_WARNING << "Callback cannot be null!"; + ROCP_CI_LOG(WARNING) << "Callback cannot be null!"; return false; } if(shader_engine_mask == 0) return false; - if(buffer_size > MAX_BUFFER_SIZE || buffer_size < MIN_BUFFER_SIZE) + if(buffer_size < MIN_BUFFER_SIZE) { - ROCP_WARNING << "Invalid buffer size: " << buffer_size; + ROCP_CI_LOG(WARNING) << "Invalid buffer size: " << buffer_size; return false; } diff --git a/tests/thread-trace/CMakeLists.txt b/tests/thread-trace/CMakeLists.txt index 6bf784e5ba..4c43dc9f67 100644 --- a/tests/thread-trace/CMakeLists.txt +++ b/tests/thread-trace/CMakeLists.txt @@ -97,3 +97,20 @@ set_tests_properties( PROPERTIES TIMEOUT 10 LABELS "integration-tests" ENVIRONMENT "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +# Test large buffer sizes. 5120 == 5GB +add_test(NAME thread-trace-api-large-buffer-test + COMMAND $) + +set_tests_properties( + thread-trace-api-large-buffer-test + PROPERTIES TIMEOUT + 10 + LABELS + "integration-tests" + ENVIRONMENT + "${PRELOAD_ENV};ATT_BUFFER_SIZE_MB=5120" + FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}" + DISABLED + ${ROCPROFILER_DISABLE_UNSTABLE_CTESTS}) diff --git a/tests/thread-trace/agent.cpp b/tests/thread-trace/agent.cpp index 708ba44968..ccc2004995 100644 --- a/tests/thread-trace/agent.cpp +++ b/tests/thread-trace/agent.cpp @@ -47,19 +47,16 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record, assert(record.payload); auto* rdata = static_cast(record.payload); - int dispatch_id = (int) rdata->dispatch_info.dispatch_id; + auto dispatch_id = rdata->dispatch_info.dispatch_id; - auto get_int_var = [](const char* var_name, int def) { - const char* var = getenv(var_name); - if(var) return atoi(var); - return def; - }; - static int begin_dispatch = get_int_var("ROCPROFILER_THREAD_TRACE_BEGIN", 1); - static int end_dispatch = get_int_var("ROCPROFILER_THREAD_TRACE_END", 4); + // Choose two dispatches to begin(6) and end(10) the trace + constexpr uint64_t begin_dispatch = 6; + constexpr uint64_t end_dispatch = 10; static std::atomic isprofiling{false}; + static std::atomic stop_profiling{false}; - static std::mutex mut; - static std::set captured_ids; + static std::mutex mut{}; + static std::set captured_ids{}; if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) { @@ -73,6 +70,7 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record, std::unique_lock lk(mut); captured_ids.insert(dispatch_id); } + if(dispatch_id > end_dispatch) stop_profiling.store(true); return; } @@ -82,7 +80,7 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record, std::unique_lock lk(mut); captured_ids.erase(dispatch_id); - if(!captured_ids.empty()) return; + if(!captured_ids.empty() || stop_profiling == false) return; bool _exp = true; if(!isprofiling.compare_exchange_strong(_exp, false, std::memory_order_relaxed)) return; @@ -104,11 +102,15 @@ query_available_agents(rocprofiler_agent_version_t /* version */, const auto* agent = static_cast(agents[idx]); if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; + // Check if we are testing for large buffers + static const char* var = getenv("ATT_BUFFER_SIZE_MB"); + static uint64_t buffer_size_mb = (var ? atoi(var) : 96) * 1024ul * 1024ul; + std::vector parameters; parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU, 1}); parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, 0xF}); - parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, 0x6000000}); - parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, 0x11}); + parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, buffer_size_mb}); + parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, 0x1}); parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL, 0}); ROCPROFILER_CALL( diff --git a/tests/thread-trace/trace_callbacks.hpp b/tests/thread-trace/trace_callbacks.hpp index 9c175782b0..32afefe053 100644 --- a/tests/thread-trace/trace_callbacks.hpp +++ b/tests/thread-trace/trace_callbacks.hpp @@ -30,7 +30,6 @@ #include #include -#include #include #include #include