Merge commit '1ba08cd4dfb7fe99a51765019210947dfcd199f7' into develop

projects/rocprofiler-sdk/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
 

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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)

projects/rocprofiler-sdk/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 $<TARGET_FILE:thread-trace-sample>)
+
+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})

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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 <rocprofiler-sdk/cxx/codeobj/code_printing.hpp>
+
+#include <rocprofiler-sdk/buffer.h>
+#include <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/experimental/thread_trace.h>
+#include <rocprofiler-sdk/fwd.h>
+#include <rocprofiler-sdk/registration.h>
+#include <rocprofiler-sdk/rocprofiler.h>
+
+#include <algorithm>
+#include <atomic>
+#include <cstdint>
+#include <cstdlib>
+#include <iostream>
+#include <mutex>
+#include <string>
+#include <vector>
+
+#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<rocprofiler_thread_trace_decoder_pc_t, Latency, address_sort_t>;
+// 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<pcinfo_t, Latency>;
+
+    std::vector<Element> 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<rocprofiler_callback_tracing_code_object_load_data_t*>(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<const void*>(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<rocprofiler_thread_trace_decoder_occupancy_t*>(events)[i];
+
+                if(event.start)
+                {
+                    Results::wave_lifetime -= static_cast<int64_t>(event.time);
+                    Results::waves_started++;
+                }
+                else
+                {
+                    Results::wave_lifetime += static_cast<int64_t>(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<rocprofiler_thread_trace_decoder_wave_t*>(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<const rocprofiler_agent_v0_t*>(agents[idx]);
+        if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue;
+
+        auto parameters = std::vector<rocprofiler_thread_trace_parameter_t>{};
+        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;
+}

projects/rocprofiler-sdk/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 <cstdint>
+#include <cstdlib>
+#include <iostream>
+#include <vector>
+
+#include <rocprofiler-sdk-roctx/roctx.h>
+
+#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<HipStream, 3>              streams{};
+    std::vector<decltype(divide_kernel)*> 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;
+}

projects/rocprofiler-sdk/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_<entry> list
+set(ROCPD_EXECUTABLES "all" "csv" "otf2" "pftrace" "summary")
+
+# format is list: [<exe-name>, <python-module>, <extra-args-to-python-module>]
+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()

projects/rocprofiler-sdk/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()

projects/rocprofiler-sdk/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,
     )

projects/rocprofiler-sdk/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(

projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/hash.hpp

@@ -23,6 +23,7 @@
 
 #pragma once
 
+#include <rocprofiler-sdk/experimental/thread-trace/trace_decoder.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/hsa.h>
 #include <rocprofiler-sdk/internal_threading.h>
@@ -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

projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/operators.hpp

@@ -26,6 +26,7 @@
 #include <rocprofiler-sdk/agent.h>
 #include <rocprofiler-sdk/counters.h>
 #include <rocprofiler-sdk/defines.h>
+#include <rocprofiler-sdk/experimental/thread-trace/trace_decoder.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/hsa.h>
 #include <rocprofiler-sdk/internal_threading.h>
@@ -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

projects/rocprofiler-sdk/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;
 
 /** @} */

projects/rocprofiler-sdk/source/lib/att-tool/att_lib_wrapper.cpp

@@ -41,9 +41,9 @@ namespace rocprofiler
 {
 namespace att_wrapper
 {
-ATTFileMgr::ATTFileMgr(Fspath                                    _dir,
-                       std::vector<std::string>                  _counters,
-                       rocprofiler_thread_trace_decoder_handle_t _decoder)
+ATTFileMgr::ATTFileMgr(Fspath                                _dir,
+                       std::vector<std::string>              _counters,
+                       rocprofiler_thread_trace_decoder_id_t _decoder)
 : dir(std::move(_dir))
 , decoder(_decoder)
 {

projects/rocprofiler-sdk/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<std::string>                  _counters,
-               rocprofiler_thread_trace_decoder_handle_t _decoder);
+    ATTFileMgr(Fspath                                _dir,
+               std::vector<std::string>              _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<AddressTable>              table{nullptr};
     std::map<size_t, std::vector<occupancy_t>> occupancy{};
     std::vector<uint64_t>                      codeobjs_to_delete{};
-    rocprofiler_thread_trace_decoder_handle_t  decoder{};
+    rocprofiler_thread_trace_decoder_id_t      decoder{};
 
     std::array<std::shared_ptr<class WstatesFile>, ROCPROFILER_THREAD_TRACE_DECODER_WSTATE_LAST>
         wstates;

projects/rocprofiler-sdk/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<char>&                        _data,
-                   WaveConfig&                               _config,
-                   rocprofiler_thread_trace_decoder_handle_t _decoder)
+ToolData::ToolData(std::vector<char>&                    _data,
+                   WaveConfig&                           _config,
+                   rocprofiler_thread_trace_decoder_id_t _decoder)
 : cfile(_config.code)
 , config(_config)
 , decoder(_decoder)

projects/rocprofiler-sdk/source/lib/att-tool/profile_interface.hpp

@@ -60,9 +60,9 @@ using SymbolInfo  = rocprofiler::sdk::codeobj::disassembly::SymbolInfo;
 
 struct ToolData
 {
-    ToolData(std::vector<char>&                        data,
-             WaveConfig&                               config,
-             rocprofiler_thread_trace_decoder_handle_t decoder);
+    ToolData(std::vector<char>&                    data,
+             WaveConfig&                           config,
+             rocprofiler_thread_trace_decoder_id_t decoder);
     ~ToolData();
 
     CodeLine& get(pcinfo_t pc);
@@ -73,7 +73,7 @@ struct ToolData
     std::vector<char> shader_data{};
     size_t            num_waves = 0;
 
-    rocprofiler_thread_trace_decoder_handle_t decoder{};
+    rocprofiler_thread_trace_decoder_id_t decoder{};
 };
 
 }  // namespace att_wrapper

projects/rocprofiler-sdk/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);
 

projects/rocprofiler-sdk/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)

projects/rocprofiler-sdk/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...")
 

projects/rocprofiler-sdk/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 = """
+        <!DOCTYPE html>
+        <html lang="en">
+        <head>
+            <meta charset="UTF-8" />
+            <title>Dashboard Report</title>
+            <style>
+                body { font-family: Arial, sans-serif; margin: 20px; }
+                h1 { color: #333; }
+                table { border-collapse: collapse; width: 100%; }
+                th, td { border: 1px solid #aaa; padding: 8px; text-align: left; }
+                th { background-color: #f0f0f0; }
+                tr:nth-child(even) { background-color: #fafafa; }
+            </style>
+        </head>
+        <body>
+            <h1>{{ title }}</h1>
+            <p><em>Generated on {{ timestamp }}</em></p>
+            <div>
+                {{ table_html | safe }}
+            </div>
+        </body>
+        </html>
+        """
+        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
+                         "<timestamp>_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()

projects/rocprofiler-sdk/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\","

projects/rocprofiler-sdk/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()

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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}

projects/rocprofiler-sdk/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<uint64_t>("ROCPROF_ATT_PARAM_SHADER_ENGINE_MASK", 0x1);
-    uint64_t att_param_buffer_size = get_env<uint64_t>("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x6000000);
+    // 256MB
+    uint64_t att_param_buffer_size = get_env<uint64_t>("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x10000000);
     uint64_t att_param_simd_select = get_env<uint64_t>("ROCPROF_ATT_PARAM_SIMD_SELECT", 0xF);
     uint64_t att_param_target_cu   = get_env<uint64_t>("ROCPROF_ATT_PARAM_TARGET_CU", 1);
     uint64_t att_param_perf_ctrl   = get_env<uint64_t>("ROCPROF_ATT_PARAM_PERFCOUNTER_CTRL", 0);

projects/rocprofiler-sdk/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
 {

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp

@@ -145,7 +145,8 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache&
     uint32_t cu                 = static_cast<uint32_t>(params.target_cu);
     uint32_t shader_engine_mask = static_cast<uint32_t>(params.shader_engine_mask);
     uint32_t simd               = static_cast<uint32_t>(params.simd_select);
-    uint32_t buffer_size        = static_cast<uint32_t>(params.buffer_size);
+    uint32_t buffer_size_lo     = static_cast<uint32_t>(params.buffer_size);
+    uint32_t buffer_size_hi     = static_cast<uint32_t>(params.buffer_size >> 32);
     uint32_t perf_ctrl          = static_cast<uint32_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(
+        AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH), {buffer_size_hi}});
 
     if(perf_ctrl != 0 && !params.perfcounters.empty())
     {

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/domain.cpp

@@ -68,7 +68,7 @@ add_domain(domain_context<DomainT>& _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;
 }
 

projects/rocprofiler-sdk/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<int32_t>::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;
     }
 

projects/rocprofiler-sdk/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 <rocprofiler-sdk/experimental/thread-trace/trace_decoder.h>
 #include <rocprofiler-sdk/experimental/thread_trace.h>
+#include <rocprofiler-sdk/cxx/hash.hpp>
+#include <rocprofiler-sdk/cxx/operators.hpp>
 
 #include <glog/logging.h>
 
+#include <atomic>
 #include <cstdint>
 
 namespace
 {
 using DL           = rocprofiler::thread_trace::DL;
 using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
+using LockedTable  = rocprofiler::common::Synchronized<AddressTable>;
 
 class DecoderInstance
 {
@@ -46,72 +51,72 @@ public:
     : dl(std::move(_dl))
     {}
 
-    std::unique_ptr<DL> dl{nullptr};
-    AddressTable        table{};
+    const std::unique_ptr<const DL> dl{nullptr};
+
+    LockedTable table{};
 };
 
-std::mutex map_mut;
+using DecoderMap =
+    std::unordered_map<rocprofiler_thread_trace_decoder_id_t, std::shared_ptr<DecoderInstance>>;
+using LockedMap = rocprofiler::common::Synchronized<DecoderMap>;
 
 auto&
-get_dlopens()
+get_dlmap()
 {
-    static auto*& _v = rocprofiler::common::static_object<
-        std::unordered_map<uint64_t, std::shared_ptr<DecoderInstance>>>::construct();
+    static auto*& _v = rocprofiler::common::static_object<LockedMap>::construct();
     return *CHECK_NOTNULL(_v);
 }
 
 std::shared_ptr<DecoderInstance>
-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<DecoderInstance> {
+        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<DL>(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<uint64_t> count{1};
+    handle->handle = count.fetch_add(1);
 
-    auto instance = std::make_shared<DecoderInstance>(std::move(dl));
-
-    handle->handle                = count++;
-    get_dlopens()[handle->handle] = std::move(instance);
+    get_dlmap().wlock(
+        [&](DecoderMap& map) { map[*handle] = std::make_shared<DecoderInstance>(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<trace_data_t*>(userdata)->decoder->table;
-
-    std::unique_ptr<Instruction> instruction{nullptr};
+    auto decoder = static_cast<trace_data_t*>(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;

projects/rocprofiler-sdk/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,

projects/rocprofiler-sdk/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;

projects/rocprofiler-sdk/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)
 

projects/rocprofiler-sdk/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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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 --
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${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 --
         $<TARGET_FILE:hsa_code_object_testapp>)
 
 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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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} --
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS}/json_input -i
+        ${CMAKE_CURRENT_BINARY_DIR}/att_input.json --
         $<TARGET_FILE:hsa_code_object_testapp>)
 
 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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
-        ${CMAKE_CURRENT_BINARY_DIR}/att_input.yml --log-level env --echo --
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS_NO_LIB}/yaml -i
+        ${CMAKE_CURRENT_BINARY_DIR}/att_input.yml --echo --
         $<TARGET_FILE:hsa_code_object_testapp>)
 
 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 $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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 $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --log-level env --echo
-                 -- $<TARGET_FILE:hsa_code_object_testapp>)
+add_test(
+    NAME rocprofv3-test-att-library-path-env-var-will-fail
+    COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS_NO_LIB}/envfail
+            --att-library-path . --echo -- $<TARGET_FILE:hsa_code_object_testapp>)
 
 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 $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS_NO_LIB}/envvar
+                 -- $<TARGET_FILE:hsa_code_object_testapp>)
+
+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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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} -- $<TARGET_FILE:vector-ops>)
+add_test(NAME rocprofv3-test-att-hsa-multiqueue-plus-pmc-execute
+         COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS}/cmd_input
+                 --pmc SQ_WAVES -o out -- $<TARGET_FILE:vector-ops>)
 
 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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --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} -- $<TARGET_FILE:vector-ops>)
+add_test(NAME rocprofv3-test-att-hsa-multiqueue-activity-pmc-will-fail
+         COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS}/cmd_input
+                 --pmc SQ_WAVES -o out --att-activity 8 -- $<TARGET_FILE:vector-ops>)
 
 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
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --log-level env --att
-        --att-perfcounter-ctrl 8 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o
-        out --att-activity 8 ${PRELOAD_ARGS} -- $<TARGET_FILE:vector-ops>)
+    NAME rocprofv3-test-att-hsa-multiqueue-activity-perf-will-fail
+    COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS}/cmd_input
+            --att-perfcounter-ctrl 8 -o out --att-activity 8 -- $<TARGET_FILE:vector-ops>)
 
 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 $<TARGET_FILE:rocprofiler-sdk::rocprofv3> ${COMMON_PARAMS}/cmd_input
+                 ${PCS_ARGS} -o out -- $<TARGET_FILE:vector-ops>)
+
+set_tests_properties(
+    rocprofv3-test-att-plus-pc-sampling
+    PROPERTIES TIMEOUT 90 LABELS "integration-tests;thread-trace;pc-sampling" DISABLED
+               ${ATT_PLUS_PCS_DISABLE})

projects/rocprofiler-sdk/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@

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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,
+            },
+        )

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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

projects/rocprofiler-sdk/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 $<TARGET_FILE:thread-trace-api-agent-test>)
+
+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})

projects/rocprofiler-sdk/tests/thread-trace/agent.cpp

@@ -29,17 +29,6 @@
 
 #include <set>
 
-#define C_API_BEGIN                                                                                \
-    try                                                                                            \
-    {
-#define C_API_END                                                                                  \
-    }                                                                                              \
-    catch(std::exception & e)                                                                      \
-    {                                                                                              \
-        std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl;   \
-    }                                                                                              \
-    catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; }
-
 namespace ATTTest
 {
 namespace Agent
@@ -58,19 +47,16 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record,
 
     assert(record.payload);
     auto* rdata = static_cast<rocprofiler_callback_tracing_kernel_dispatch_data_t*>(record.payload);
-    int   dispatch_id = (int) rdata->dispatch_info.dispatch_id;
+    auto  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<bool> isprofiling{false};
+    static std::atomic<bool> stop_profiling{false};
 
-    static std::mutex    mut;
-    static std::set<int> captured_ids;
+    static std::mutex    mut{};
+    static std::set<int> captured_ids{};
 
     if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER)
     {
@@ -84,6 +70,7 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record,
             std::unique_lock<std::mutex> lk(mut);
             captured_ids.insert(dispatch_id);
         }
+        if(dispatch_id > end_dispatch) stop_profiling.store(true);
         return;
     }
 
@@ -93,7 +80,7 @@ dispatch_tracing_callback(rocprofiler_callback_tracing_record_t record,
 
     std::unique_lock<std::mutex> 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;
@@ -115,11 +102,15 @@ query_available_agents(rocprofiler_agent_version_t /* version */,
         const auto* agent = static_cast<const rocprofiler_agent_v0_t*>(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<rocprofiler_thread_trace_parameter_t> 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(
@@ -135,9 +126,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 +139,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 +148,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 +169,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<Callbacks::ToolData*>(tool_data);
-}
-
 }  // namespace Agent
 }  // namespace ATTTest
 
@@ -206,8 +191,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;

projects/rocprofiler-sdk/tests/thread-trace/multi_dispatch.cpp

@@ -48,15 +48,17 @@ dispatch_callback(rocprofiler_agent_id_t /* agent */,
     static std::atomic<size_t> 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<rocprofiler_thread_trace_parameter_t> 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<Callbacks::ToolData*>(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<void*>(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;

projects/rocprofiler-sdk/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<Callbacks::ToolData*>(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<rocprofiler_agent_id_t> 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<Callbacks::ToolData*>(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<void*>(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;

projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp

@@ -26,11 +26,6 @@
 #endif
 
 #include "trace_callbacks.hpp"
-#include <rocprofiler-sdk/cxx/codeobj/code_printing.hpp>
-
-#ifdef ENABLE_ATT_FILES
-#    include <nlohmann/json.hpp>
-#endif
 
 #include <unistd.h>
 #include <cassert>
@@ -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<size_t>                   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<Callbacks::ToolData*>(userdata);
-
-    if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
-    {
-        auto* data = static_cast<kernel_symbol_data_t*>(record.payload);
-        tool.kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name);
-    }
-
     if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return;
 
-    auto* data = static_cast<code_obj_load_data_t*>(record.payload);
+    auto* data = static_cast<rocprofiler_callback_tracing_code_object_load_data_t*>(record.payload);
+    if(data->storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) return;
 
-    static std::atomic<int> filecnt{0};
-    std::string             name = "codeobj_" + std::to_string(filecnt.fetch_add(1)) + ".out";
-
-#ifdef ENABLE_ATT_FILES
-    if(std::string_view(data->uri).find("file:///") == 0)
+    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<const void*>(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<rocprofiler_thread_trace_decoder_wave_t*>(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<Callbacks::ToolData*>(userdata.ptr);
-
-    std::string name = "agent_" + std::to_string(agent.handle) + "_shader_engine_" +
-                       std::to_string(se_id) + "_" + std::to_string(agent.handle) + ".att";
-
-#ifdef ENABLE_ATT_FILES
-    {
-        std::ofstream file(tool.out_dir + name, std::ios::binary);
-        assert(file.is_open() && "Could not open ATT file for writing");
-        file.write((char*) se_data, data_size);
-    }
-#endif
-
-    assert(se_data);
-    assert(data_size);
-
-    auto _lk = std::unique_lock{tool.mut};
-    tool.att_files.push_back(name);
-
-    C_API_END
+    // 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<Callbacks::ToolData*>(userdata);
-    auto  _lk  = std::unique_lock{tool.mut};
-    assert(!tool.att_files.empty());
-
-#ifdef ENABLE_ATT_FILES
-    nlohmann::json att_json;
-    for(auto& file : tool.att_files)
-        att_json.push_back(file);
-
-    nlohmann::json codeobj_json;
-    nlohmann::json snapshot_json;
-    for(auto& file : tool.codeobjs)
-    {
-        nlohmann::json codeobj;
-        codeobj["code_object_id"] = file.id;
-        codeobj["load_delta"]     = file.addr;
-        codeobj["load_size"]      = file.size;
-        codeobj["uri"]            = file.uri;
-        codeobj["filename"]       = file.filename;
-        codeobj_json.push_back(codeobj);
-
-        nlohmann::json pair_json;
-        pair_json["key"]   = file.id;
-        pair_json["value"] = file.filename;
-        snapshot_json.push_back(pair_json);
-    }
-    nlohmann::json tool_json;
-    tool_json["strings"]["att_files"]                  = att_json;
-    tool_json["code_objects"]                          = codeobj_json;
-    tool_json["strings"]["code_object_snapshot_files"] = snapshot_json;
-
-    nlohmann::json array;
-    array.push_back(tool_json);
-
-    nlohmann::json sdk_json;
-    sdk_json["rocprofiler-sdk-tool"] = array;
-
-    std::ofstream json_file(tool.out_dir + (std::to_string(getpid()) + "_results.json"));
-    assert(json_file.is_open() && "Could not open json file for writing!");
-    json_file << sdk_json;
-#endif
+    if(latency.load() == 0) std::cerr << "Error: No latency was assigned to the trace!";
 }
 
 }  // namespace Callbacks

projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.hpp

@@ -30,7 +30,6 @@
 #include <rocprofiler-sdk/rocprofiler.h>
 
 #include <atomic>
-#include <cassert>
 #include <cstdint>
 #include <cstdlib>
 #include <iostream>
@@ -40,60 +39,21 @@
 #include <unordered_map>
 #include <vector>
 
-#define ROCPROFILER_VAR_NAME_COMBINE(X, Y) X##Y
-#define ROCPROFILER_VARIABLE(X, Y)         ROCPROFILER_VAR_NAME_COMBINE(X, Y)
-
-#define 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<CodeobjInfo> codeobjs{};
-    std::vector<std::string> att_files{};
-
-    std::unordered_map<uint64_t, std::string> kernel_id_to_kernel_name = {};
-};
-
 void
 tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
                               rocprofiler_user_data_t*,
@@ -107,6 +67,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