feat(amd): add device-wide PM sampling

CMakeLists.txt

@@ -708,6 +708,7 @@ if(GPUFL_ENABLE_AMD)
             target_sources(gpufl PRIVATE
                 include/gpufl/backends/amd/monitor_adapter_amd.cpp
                 include/gpufl/backends/amd/rocprofiler_backend.cpp
+                include/gpufl/backends/amd/engine/device_counter_engine.cpp
                 include/gpufl/backends/amd/engine/dispatch_counter_engine.cpp
             )
             message(STATUS "Found ROCprofiler-SDK support")

example/amd/CMakeLists.txt

@@ -48,6 +48,7 @@ list(APPEND HIP_HIPCC_FLAGS "-g")
 set(AMD_EXAMPLE_SOURCES
   check_device.cpp
   gpufl_scope_demo.cpp
+  pm_sampling_sample_rows.cpp
   vector_add_benchmark.cpp
 )
 
@@ -58,14 +59,17 @@ set_source_files_properties(
 
 hip_add_executable(amd_check_device check_device.cpp)
 hip_add_executable(amd_gpufl_scope_demo gpufl_scope_demo.cpp)
+hip_add_executable(amd_pm_sampling_sample_rows pm_sampling_sample_rows.cpp)
 hip_add_executable(amd_vector_add_benchmark vector_add_benchmark.cpp)
 
 if(TARGET hip::host)
   target_link_libraries(amd_check_device PRIVATE hip::host)
   target_link_libraries(amd_gpufl_scope_demo PRIVATE hip::host)
+  target_link_libraries(amd_pm_sampling_sample_rows PRIVATE hip::host)
   target_link_libraries(amd_vector_add_benchmark PRIVATE hip::host)
 endif()
 
 target_link_libraries(amd_check_device PRIVATE gpufl::gpufl)
 target_link_libraries(amd_gpufl_scope_demo PRIVATE gpufl::gpufl)
+target_link_libraries(amd_pm_sampling_sample_rows PRIVATE gpufl::gpufl)
 target_link_libraries(amd_vector_add_benchmark PRIVATE gpufl::gpufl)

example/amd/README.md

@@ -8,24 +8,28 @@ This folder mirrors the CUDA example area with runnable HIP examples for AMD GPU
 - AMD static device inventory via HIP
 - AMD kernel dispatch tracing via `rocprofiler-sdk`
 - AMD memcpy tracing via `rocprofiler-sdk`
+- Per-dispatch AMD hardware counters via ROCprofiler dispatch counting
+- Device-wide `PmSampling` timelines via ROCprofiler device counting
 - `gpufl` initialization with `backend = gpufl::BackendKind::Amd`
 - User-defined scope logging via `GFL_SCOPE(...)`
 - HIP example programs that run on ROCm hardware
 
 ## What Does Not Work Yet
 
-- AMD profiling engines equivalent to CUPTI PC Sampling / SASS Metrics / Range Profiler
+- AMD PC sampling
+- Instruction-level SASS metrics and NVIDIA-compatible Range Profiler metrics
 
 Today, the AMD backend is useful for:
 
 - system metric logging
 - device inventory
 - automatic HIP kernel and memcpy tracing
+- per-dispatch and device-wide hardware-counter profiling
 - scope-level application instrumentation
 
 It is not yet useful for:
 
-- instruction-level or hardware-counter profiling
+- PC sampling or instruction-level profiling
 
 ## Targets
 
@@ -35,6 +39,8 @@ It is not yet useful for:
   - HIP vector add benchmark with result verification
 - `amd_gpufl_scope_demo`
   - Initializes `gpufl` with the AMD backend, runs HIP work inside scopes, and writes logs
+- `amd_pm_sampling_sample_rows`
+  - Selects AMD device counting and exits successfully only when each of two named scopes emits PM sample rows
 
 ## Build
 
@@ -51,6 +57,7 @@ cmake -S . -B build-rocm-examples \
 cmake --build build-rocm-examples --target amd_check_device
 cmake --build build-rocm-examples --target amd_vector_add_benchmark
 cmake --build build-rocm-examples --target amd_gpufl_scope_demo
+cmake --build build-rocm-examples --target amd_pm_sampling_sample_rows
 ```
 
 The AMD example targets are only added when CMake detects HIP successfully.
@@ -106,8 +113,27 @@ subproject and disables the parent example/test targets to avoid recursion.
 ./build-rocm-examples/example/amd/amd_check_device
 ./build-rocm-examples/example/amd/amd_vector_add_benchmark
 ./build-rocm-examples/example/amd/amd_gpufl_scope_demo
+./build-rocm-examples/example/amd/amd_pm_sampling_sample_rows
 ```
 
+`amd_pm_sampling_sample_rows` requests the portable `GPUBusy` counter, runs
+GPU work in `pm_rows_phase_a` and `pm_rows_phase_b`, and checks that the PM row
+count increases after each scope. It returns exit code 2 when AMD device
+counting is unavailable or either scope does not produce a row. Its generated
+report shows the same rows grouped by scope for manual inspection.
+
+The scope demo selects per-dispatch counters by default. To exercise the
+device-wide PM timeline instead:
+
+```bash
+GPUFL_PROFILING_ENGINE=PmSampling \
+  ./build-rocm-examples/example/amd/amd_gpufl_scope_demo
+```
+
+`PmSampling` uses the portable `GPUBusy` derived counter by default. Set
+`pm_sampling_metrics` programmatically to request other native ROCprofiler
+counter names.
+
 On a working ROCm system, `amd_check_device` should print output similar to:
 
 ```text
@@ -123,17 +149,25 @@ Success! Device 0: AMD Radeon RX 9070 XT (arch gfx1201, capability 12.0)
 gfl_amd_scope
 ```
 
+`amd_pm_sampling_sample_rows` writes logs with prefix:
+
+```bash
+gfl_amd_pm_rows
+```
+
 With `rocprofiler-sdk` available, expect:
 
 - `job_start` inventory
 - kernel dictionaries
 - `kernel_event_batch`
 - `kernel_detail`
 - `memcpy_event_batch`
+- `profile_sample_batch` for dispatch-counting requests
+- `pm_sampling_config` and `pm_sample_batch` for `PmSampling`
 - system metric samples
 - scope events
 
 Without `rocprofiler-sdk`, expect only telemetry, static inventory, and scope
 events.
 
-Do not expect AMD profiling samples such as PC sampling or hardware counters.
+PC samples and instruction-level SASS samples are not available on AMD yet.

example/amd/pm_sampling_sample_rows.cpp

@@ -0,0 +1,145 @@
+#include <hip/hip_runtime.h>
+
+#include <cstdint>
+#include <iostream>
+
+#include "gpufl/core/monitor.hpp"
+#include "gpufl/gpufl.hpp"
+
+namespace {
+
+constexpr int kElementCount = 1 << 20;
+constexpr int kBlockSize = 256;
+constexpr int kLaunchesPerPhase = 8;
+constexpr int kIterationsPerLaunch = 1024;
+
+bool CheckHip(const hipError_t status, const char* what) {
+    if (status == hipSuccess) return true;
+    std::cerr << what << " failed: " << hipGetErrorString(status) << "\n";
+    return false;
+}
+
+__global__ void sampleRowsWorkload(float* values, const int count,
+                                   const int iterations) {
+    const int index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (index >= count) return;
+
+    float x = values[index] +
+              static_cast<float>((index & 1023) + 1) * 0.0001f;
+    float y = static_cast<float>((threadIdx.x & 31) + 1) * 0.00001f;
+    for (int iteration = 0; iteration < iterations; ++iteration) {
+        x = x * 1.0000001f + y;
+        y = y * 0.9999999f + x * 0.000001f;
+        if (x > 4096.0f) x -= 4096.0f;
+    }
+    values[index] = x + y;
+}
+
+bool RunPhase(float* values, const int launches, const int iterations) {
+    const dim3 block(kBlockSize);
+    const dim3 grid((kElementCount + block.x - 1) / block.x);
+    for (int launch = 0; launch < launches; ++launch) {
+        hipLaunchKernelGGL(sampleRowsWorkload, grid, block, 0, 0, values,
+                           kElementCount, iterations);
+    }
+    return CheckHip(hipGetLastError(), "sampleRowsWorkload launch") &&
+           CheckHip(hipDeviceSynchronize(), "sampleRowsWorkload sync");
+}
+
+}  // namespace
+
+int main() {
+    gpufl::InitOptions opts;
+    opts.app_name = "amd_pm_sampling_sample_rows";
+    opts.log_path = "gfl_amd_pm_rows";
+    opts.backend = gpufl::BackendKind::Amd;
+    opts.profiling_engine = gpufl::ProfilingEngine::PmSampling;
+    opts.pm_sampling_interval_us = 1000;
+    opts.pm_sampling_max_samples = 4096;
+    opts.pm_sampling_preset = "overview";
+    opts.pm_sampling_metrics = {"GPUBusy"};
+    opts.pm_sampling_scope_only = true;
+    opts.continuous_system_sampling = false;
+    opts.enable_debug_output = true;
+    opts.enable_stack_trace = false;
+
+    if (!gpufl::init(opts)) {
+        std::cerr << "Failed to initialize gpufl for AMD PM sampling\n";
+        return 1;
+    }
+
+    const std::string engine =
+        gpufl::Monitor::ResolvedProfilingEngineWireName();
+    const bool engine_ok = engine == "amd.device_counting";
+    std::cout << "=== GPUFL AMD PM Sample Rows ===\n"
+              << "Resolved engine: " << engine << "\n";
+    if (!engine_ok) {
+        std::cerr << "Expected amd.device_counting; ROCprofiler device "
+                     "counting may be unavailable\n";
+    }
+
+    float* device_values = nullptr;
+    bool workload_ok =
+        CheckHip(hipMalloc(&device_values,
+                           static_cast<size_t>(kElementCount) * sizeof(float)),
+                 "hipMalloc(device_values)");
+    if (workload_ok) {
+        workload_ok = CheckHip(
+            hipMemset(device_values, 0,
+                      static_cast<size_t>(kElementCount) * sizeof(float)),
+            "hipMemset(device_values)");
+    }
+
+    // Warm up HIP and load the kernel before opening a measured scope. This
+    // also gives ROCprofiler's deferred device-counting callback time to
+    // accept the configured profile.
+    if (workload_ok) workload_ok = RunPhase(device_values, 1, 64);
+
+    const uint64_t rows_before = gpufl::Monitor::PmSampleRowsSeen();
+    bool phase_a_ok = false;
+    if (workload_ok) {
+        GFL_SCOPE("pm_rows_phase_a") {
+            phase_a_ok = RunPhase(device_values, kLaunchesPerPhase,
+                                  kIterationsPerLaunch);
+        }
+    }
+    const uint64_t rows_after_a = gpufl::Monitor::PmSampleRowsSeen();
+
+    bool phase_b_ok = false;
+    if (workload_ok && phase_a_ok) {
+        GFL_SCOPE("pm_rows_phase_b") {
+            phase_b_ok = RunPhase(device_values, kLaunchesPerPhase,
+                                  kIterationsPerLaunch);
+        }
+    }
+    const uint64_t rows_after_b = gpufl::Monitor::PmSampleRowsSeen();
+
+    std::cout << "PM rows: before=" << rows_before
+              << ", after phase A=" << rows_after_a
+              << ", after phase B=" << rows_after_b << "\n";
+
+    const bool phase_a_rows = rows_after_a > rows_before;
+    const bool phase_b_rows = rows_after_b > rows_after_a;
+    if (!phase_a_rows) {
+        std::cerr << "Phase A did not emit a PM sample row\n";
+    }
+    if (!phase_b_rows) {
+        std::cerr << "Phase B did not emit a PM sample row\n";
+    }
+
+    if (device_values != nullptr) {
+        (void) hipFree(device_values);
+    }
+
+    gpufl::shutdown();
+    gpufl::generateReport();
+
+    const bool passed = engine_ok && workload_ok && phase_a_ok && phase_b_ok &&
+                        phase_a_rows && phase_b_rows;
+    if (!passed) return 2;
+
+    std::cout << "\nPASS: both named scopes emitted GPUBusy sample rows.\n"
+              << "Inspect logs with prefix " << opts.log_path
+              << " for pm_sampling_config and pm_sample_batch events.\n";
+    return 0;
+}

include/gpufl/backends/amd/amd_capture_capabilities.cpp

@@ -129,17 +129,20 @@ CaptureCapabilitiesEvent BuildAmdCaptureCapabilitiesEvent(
         !device_counting_requested
             ? "not_requested"
             : (device_counting_selected
-                   ? (input.profiling_sample_rows > 0 ? "collected"
-                                                      : "enabled_no_data")
+                   ? (input.pm_sample_rows > 0 ? "collected"
+                                               : "enabled_no_data")
                    : "skipped"),
         device_counting_selected ? "rocprofiler_device_counting_service"
                                  : "disabled",
         device_counting_selected
-            ? (input.profiling_sample_rows > 0 ? "" : "enabled_but_no_records")
+            ? (input.pm_sample_rows > 0 ? "" : "enabled_but_no_records")
             : (device_counting_requested ? input.plan.reason_code
                                          : "not_selected"),
         device_counting_selected
-            ? "AMD device counting was selected."
+            ? (input.pm_sample_rows > 0
+                   ? "Device-wide AMD hardware-counter samples were collected."
+                   : "AMD device counting was selected but emitted no samples "
+                     "this session.")
             : "AMD device counting is not available in the current "
               "implementation.");
 

include/gpufl/backends/amd/amd_capture_capabilities.hpp

@@ -15,6 +15,7 @@ struct AmdCaptureCapabilityInput {
     uint64_t kernel_rows = 0;
     uint64_t memcpy_rows = 0;
     uint64_t profiling_sample_rows = 0;
+    uint64_t pm_sample_rows = 0;
     uint64_t dropped_trace_records = 0;
     uint64_t dropped_client_records = 0;
     uint64_t trace_buffer_flush_failures = 0;

include/gpufl/backends/amd/amd_profiling_policy.cpp

@@ -1,5 +1,7 @@
 #include "gpufl/backends/amd/amd_profiling_policy.hpp"
 
+#include <unordered_set>
+
 namespace gpufl::amd {
 namespace {
 
@@ -60,6 +62,22 @@ bool AmdRequestNeedsDeviceCounting(const ProfilingEngine engine) {
            engine == ProfilingEngine::Deep;
 }
 
+std::vector<std::string> ResolveAmdDeviceCountingMetrics(
+    const std::vector<std::string>& requested_metrics) {
+    if (requested_metrics.empty()) return {"GPUBusy"};
+
+    std::vector<std::string> resolved;
+    resolved.reserve(requested_metrics.size());
+    std::unordered_set<std::string> seen;
+    for (const auto& metric : requested_metrics) {
+        if (!metric.empty() && seen.emplace(metric).second) {
+            resolved.push_back(metric);
+        }
+    }
+    return resolved.empty() ? std::vector<std::string>{"GPUBusy"}
+                            : resolved;
+}
+
 std::optional<uint32_t> ResolveAmdDispatchDeviceId(
     const uint64_t configured_agent_handle,
     const uint32_t configured_device_id,

include/gpufl/backends/amd/amd_profiling_policy.hpp

@@ -3,6 +3,7 @@
 #include <cstdint>
 #include <optional>
 #include <string>
+#include <vector>
 
 #include "gpufl/core/monitor.hpp"
 
@@ -43,6 +44,12 @@ bool AmdRequestNeedsDispatchCounting(ProfilingEngine engine);
 bool AmdRequestNeedsPcSampling(ProfilingEngine engine);
 bool AmdRequestNeedsDeviceCounting(ProfilingEngine engine);
 
+// AMD's device-counting service consumes native ROCprofiler counter names.
+// Explicit user metrics win; the default is a portable device-utilization
+// signal that is available as a derived counter across supported AMD GPUs.
+std::vector<std::string> ResolveAmdDeviceCountingMetrics(
+    const std::vector<std::string>& requested_metrics);
+
 // Dispatch counting is currently configured for one GPU agent. Resolve only
 // records from that agent so a secondary GPU can never be mislabeled as device
 // zero (or as the configured primary device).

include/gpufl/backends/amd/engine/amd_profiling_engine.hpp

@@ -32,6 +32,10 @@ class AmdProfilingEngine {
     /// Periodically drain buffered profiling data into the monitor ring buffer.
     virtual void drain() = 0;
 
+    /// Cheap collector-loop service tick. Pull-based engines use this to
+    /// honor their sampling cadence; callback/buffer engines leave it idle.
+    virtual void service() {}
+
     /// Release resources.
     virtual void shutdown() = 0;