fix(rocprofiler-sdk): serializer race fixes and hook tag range for counters (#8891)

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp

@@ -537,10 +537,9 @@ QueueController::serializer(const Queue* queue)
                 return false;
             },
             [&](auto& m) {
-                ret =
-                    m.emplace(agent_id,
-                              std::make_shared<common::Synchronized<hsa::profiler_serializer>>())
-                        .first->second.get();
+                ret = m.emplace(agent_id,
+                                std::make_shared<common::Synchronized<hsa::profiler_serializer>>())
+                          .first->second.get();
                 if(should_serialize)
                 {
                     ret->wlock([&](auto& serializer) { serializer.enable({}); });
@@ -610,16 +609,19 @@ QueueController::update_serialization(const agent_handle_set_t& agents, bool ena
                 {
                     auto itr = was_enabled.find(agent_id);
                     if(itr == was_enabled.end()) continue;
-                    if(itr->second == state.enabled(agent_id)) continue;
+                    const bool now_enabled = state.enabled(agent_id);
+                    if(itr->second == now_enabled) continue;
 
                     auto queues = hsa_barrier::queue_map_ptr_t{};
-                    if(auto it = pd_map.find(agent_id); it != pd_map.end())
-                        queues = it->second;
+                    if(auto it = pd_map.find(agent_id); it != pd_map.end()) queues = it->second;
 
                     auto ser_itr = serializers.find(agent_id);
                     if(ser_itr == serializers.end() || !ser_itr->second) continue;
+                    // Derived from the post-update refcount rather than the caller's `enable`
+                    // flag, so the applied state cannot disagree with the count that produced
+                    // the transition.
                     ser_itr->second->wlock([&](auto& serializer) {
-                        if(enable)
+                        if(now_enabled)
                             serializer.enable(queues);
                         else
                             serializer.disable(queues);

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp

@@ -30,30 +30,35 @@ namespace hsa
 {
 namespace queue_hooks
 {
-// Tags identifying the producer subsystem of each inst_pkt_t entry.
-// Only distinctness matters; values are fixed for test-order stability.
-// Services are migrated off the per-queue callback registry one at a time; the
-// unused ids are reserved so the tag values stay stable across those PRs.
+// Tags identifying the producer subsystem of each inst_pkt_t entry, stored in the
+// inst_pkt_t pair as an hsa::ClientID.
 //
-// The underlying type is fixed to int64_t to match hsa::ClientID, which is how the
-// tag is stored in inst_pkt_t. The enum is deliberately unscoped rather than an
-// enum class: unscoped keeps the implicit conversion to ClientID, so the emplace
-// sites need no cast, while still grouping the values under one type.
+// The underlying type is fixed to int64_t to match hsa::ClientID. The enum is
+// deliberately unscoped rather than an enum class: unscoped keeps the implicit
+// conversion to ClientID, so the emplace sites need no cast, while still grouping
+// the values under one type.
 //
-// These values numerically overlap the per-queue registry's auto-incrementing
-// ClientID, which also starts at 1 and is still in use by the services that have
-// not been migrated yet. That is inert only because no consumer routes on the tag:
-// counters' completed_cb and thread trace's post_kernel_call both identify their
-// own packets by pointer lookup or dynamic_cast. Before adding a consumer that
-// dispatches on these values, either finish migrating the remaining services off
-// the registry or move this enum to a range the registry cannot produce.
+// The values are negative because they share the inst_pkt_t tag space with the
+// per-queue callback registry, whose QueueController::add_callback hands out
+// strictly positive ids counting up from 1. Migrated services (counters, thread
+// trace, SPM) select their own packets by comparing against these tags, so a tag
+// drawn from the registry's range could match a packet the service does not own.
+// Services that have not migrated yet still receive registry ids, so the two
+// ranges must stay disjoint until the registry is gone. Any tag added here must
+// likewise be negative; -1 is left unused because QueueController::add_callback
+// uses it as a local "unassigned" initializer.
 enum client_id : int64_t
 {
-    COUNTERS_CLIENT_ID     = 1,
-    THREAD_TRACE_CLIENT_ID = 2,
-    PC_SAMPLING_CLIENT_ID  = 3,
-    SPM_CLIENT_ID          = 4,
+    COUNTERS_CLIENT_ID     = -1001,
+    THREAD_TRACE_CLIENT_ID = -1002,
+    PC_SAMPLING_CLIENT_ID  = -1003,
+    SPM_CLIENT_ID          = -1004,
 };
+
+static_assert(COUNTERS_CLIENT_ID < 0 && THREAD_TRACE_CLIENT_ID < 0 && PC_SAMPLING_CLIENT_ID < 0 &&
+                  SPM_CLIENT_ID < 0,
+              "queue hook tags must stay disjoint from the positive ids handed out by "
+              "QueueController::add_callback");
 }  // namespace queue_hooks
 }  // namespace hsa
 }  // namespace rocprofiler

projects/rocprofiler-sdk/tests/CMakeLists.txt

@@ -85,6 +85,7 @@ add_subdirectory(pc_sampling)
 add_subdirectory(hip-graph-tracing)
 add_subdirectory(counter-collection)
 add_subdirectory(counter-collection-per-agent)
+add_subdirectory(queue-hooks-perf)
 add_subdirectory(openmp-tools)
 add_subdirectory(rocdecode)
 add_subdirectory(rocjpeg)

projects/rocprofiler-sdk/tests/queue-hooks-perf/CMakeLists.txt

@@ -0,0 +1,84 @@
+#
+# Queue interposition perf tests — dispatch counter path (callback registry vs direct hooks).
+#
+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-tests-queue-hooks-perf LANGUAGES CXX HIP)
+
+foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
+    if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
+        set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
+    endif()
+endforeach()
+
+find_package(rocprofiler-sdk REQUIRED)
+find_package(Python3 REQUIRED)
+
+set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP)
+add_executable(queue-hooks-perf-testapp)
+target_sources(queue-hooks-perf-testapp PRIVATE main.cpp)
+target_link_libraries(queue-hooks-perf-testapp PRIVATE Threads::Threads)
+
+add_test(
+    NAME test-queue-hooks-perf-scaling
+    COMMAND
+        ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/run_and_validate.py --testapp
+        $<TARGET_FILE:queue-hooks-perf-testapp> --preload
+        $<TARGET_FILE:rocprofiler-sdk-json-tool> --ballast-mb 8 --launches-low 8
+        --launches-high 32 --max-scaling-ratio 6.0)
+
+set_tests_properties(
+    test-queue-hooks-perf-scaling
+    PROPERTIES TIMEOUT
+               180
+               LABELS
+               "integration-tests;queue-hooks;queue-hooks-perf"
+               RESOURCE_LOCK
+               GPU
+               DEPENDS
+               "queue-hooks-perf-testapp;rocprofiler-sdk-json-tool"
+               ENVIRONMENT
+               "QH_PERF_RESULTS_JSON=${CMAKE_CURRENT_BINARY_DIR}/qh_perf_scaling.json"
+               PASS_REGULAR_EXPRESSION
+               "\\[qh-perf-run\\] PASS"
+               FAIL_REGULAR_EXPRESSION
+               "${ROCPROFILER_DEFAULT_FAIL_REGEX}")
+
+add_test(
+    NAME test-queue-hooks-perf-launch-sweep
+    COMMAND
+        ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/run_launch_sweep.py --testapp
+        $<TARGET_FILE:queue-hooks-perf-testapp> --preload
+        $<TARGET_FILE:rocprofiler-sdk-json-tool> --ballast-mb 8 --launches 4 8 16 32
+        --slack 2.5)
+
+set_tests_properties(
+    test-queue-hooks-perf-launch-sweep
+    PROPERTIES TIMEOUT
+               240
+               LABELS
+               "integration-tests;queue-hooks;queue-hooks-perf"
+               RESOURCE_LOCK
+               GPU
+               DEPENDS
+               "queue-hooks-perf-testapp;rocprofiler-sdk-json-tool"
+               ENVIRONMENT
+               "QH_PERF_SWEEP_JSON=${CMAKE_CURRENT_BINARY_DIR}/qh_perf_launch_sweep.json"
+               PASS_REGULAR_EXPRESSION
+               "\\[qh-perf-sweep\\] PASS overall"
+               FAIL_REGULAR_EXPRESSION
+               "${ROCPROFILER_DEFAULT_FAIL_REGEX}")

projects/rocprofiler-sdk/tests/queue-hooks-perf/main.cpp

@@ -0,0 +1,83 @@
+// MIT License
+//
+// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+#include <hip/hip_runtime.h>
+
+#include <chrono>
+#include <cstdio>
+#include <cstdlib>
+#include <vector>
+
+#define HIP_CHECK(call)                                                                            \
+    do                                                                                             \
+    {                                                                                              \
+        hipError_t _err = (call);                                                                  \
+        if(_err != hipSuccess)                                                                     \
+        {                                                                                          \
+            fprintf(                                                                               \
+                stderr, "HIP error '%s' at %s:%d\n", hipGetErrorString(_err), __FILE__, __LINE__); \
+            return EXIT_FAILURE;                                                                   \
+        }                                                                                          \
+    } while(0)
+
+__global__ void
+touch(float* buf, int n, int* counter)
+{
+    const int stride = blockDim.x * gridDim.x;
+    for(int i = blockDim.x * blockIdx.x + threadIdx.x; i < n; i += stride)
+        buf[i] = buf[i] * 1.0001f + 0.001f;
+    if(blockIdx.x == 0 && threadIdx.x == 0) atomicAdd(counter, 1);
+}
+
+int
+main(int argc, char** argv)
+{
+    int ballast_mb = 8;
+    int launches   = 32;
+    if(argc > 1) ballast_mb = std::atoi(argv[1]);
+    if(argc > 2) launches = std::atoi(argv[2]);
+    if(ballast_mb < 1) ballast_mb = 1;
+    if(launches < 1) launches = 1;
+
+    const size_t ballast_bytes = static_cast<size_t>(ballast_mb) * 1024U * 1024U;
+    const int    ballast_elems = static_cast<int>(ballast_bytes / sizeof(float));
+
+    float* ballast = nullptr;
+    int*   counter = nullptr;
+    HIP_CHECK(hipMalloc(&ballast, ballast_bytes));
+    HIP_CHECK(hipMalloc(&counter, sizeof(int)));
+    HIP_CHECK(hipMemset(ballast, 0, ballast_bytes));
+    HIP_CHECK(hipMemset(counter, 0, sizeof(int)));
+
+    using clock      = std::chrono::steady_clock;
+    const auto start = clock::now();
+    for(int i = 0; i < launches; ++i)
+    {
+        touch<<<256, 64>>>(ballast, ballast_elems, counter);
+        HIP_CHECK(hipGetLastError());
+        HIP_CHECK(hipDeviceSynchronize());
+    }
+    const auto   end     = clock::now();
+    const double wall_ms = std::chrono::duration<double, std::milli>(end - start).count();
+
+    int counter_h = 0;
+    HIP_CHECK(hipMemcpy(&counter_h, counter, sizeof(int), hipMemcpyDeviceToHost));
+    HIP_CHECK(hipFree(ballast));
+    HIP_CHECK(hipFree(counter));
+
+    std::printf("[qh-perf] ballast_mb=%d launches=%d wall_ms=%.3f counter=%d\n",
+                ballast_mb,
+                launches,
+                wall_ms,
+                counter_h);
+
+    if(counter_h != launches)
+    {
+        std::fprintf(stderr, "[qh-perf] FAIL counter=%d expected %d\n", counter_h, launches);
+        return EXIT_FAILURE;
+    }
+
+    std::printf("[qh-perf] PASS\n");
+    return EXIT_SUCCESS;
+}

projects/rocprofiler-sdk/tests/queue-hooks-perf/perf_cost_model.py

@@ -0,0 +1,35 @@
+# MIT License
+#
+# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Cost model for dispatch-counter queue interposition perf tests.
+# Architecture-agnostic ceilings for CI regression detection.
+
+import os
+
+# Per-dispatch counter instrumentation allowance (ms).
+PER_DISPATCH_MS = float(os.environ.get("ROCPROFILER_QH_PER_DISPATCH_MS", "3.0"))
+
+# Fixed attach / context setup (ms).
+FIXED_OVERHEAD_MS = float(os.environ.get("ROCPROFILER_QH_FIXED_OVERHEAD_MS", "80.0"))
+
+# Slack over linear launch scaling.
+OVERHEAD_MARGIN = float(os.environ.get("ROCPROFILER_QH_OVERHEAD_MARGIN", "8.0"))
+
+
+def model_max_ms(
+    launches: int,
+    per_dispatch_ms: float | None = None,
+    fixed_ms: float | None = None,
+    margin: float | None = None,
+) -> float:
+    per = PER_DISPATCH_MS if per_dispatch_ms is None else per_dispatch_ms
+    fixed = FIXED_OVERHEAD_MS if fixed_ms is None else fixed_ms
+    m = OVERHEAD_MARGIN if margin is None else margin
+    return fixed + launches * per * m
+
+
+def max_launch_scaling_ratio(l_low: int, l_high: int, slack: float = 2.0) -> float:
+    if l_low <= 0:
+        return slack * float(l_high)
+    return (float(l_high) / float(l_low)) * slack