rocprofv3 --pmc crashes with unordered_map::at on gfx1201 (RDNA4, Navi 48)

[srgtuszy]

Problem Description

Any PMC counter collection with rocprofv3 --pmc crashes immediately on gfx1201 (RDNA4) with terminate called after throwing an instance of 'std::out_of_range' / what(): unordered_map::at.

This affects every counter (tested SQ_WAVES, VALUBusy, OccupancyPercent, MemUnitBusy, LDSBankConflict) and every workload (including a trivial 1-kernel vector_add program). Kernel tracing (--kernel-trace, -r) works correctly on the same hardware.

The crash occurs inside the HSA queue callback during counter collection setup, before the profiled kernel completes.

rocprofv2 explicitly rejects gfx1201 with "Unsupported hardware. Use rocprofv3 for navi4x," making rocprofv3 --pmc the only counter collection path for RDNA4 — and it is non-functional.

Operating System

Ubuntu 24.04.4 LTS, kernel 7.0.0-28-generic

CPU

AMD Ryzen 7 7700 (8C/16T)

GPU

AMD Radeon AI PRO R9700 (Navi 48, RDNA4, gfx1201)

ROCm Version

7.2.4

ROCm Component

rocprofiler-sdk 1.1.0 (git revision 97f5574fe2fdc7bef44fb01545347912ee9f1779)

Steps to Reproduce

1. Minimal reproducer

// repro_pmc.cpp
#include <hip/hip_runtime.h>
#include <cstdio>

__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    constexpr int N = 1024;
    float *d_a, *d_b, *d_c;
    (void)hipMalloc(&d_a, N * sizeof(float));
    (void)hipMalloc(&d_b, N * sizeof(float));
    (void)hipMalloc(&d_c, N * sizeof(float));
    vector_add<<<(N+255)/256, 256>>>(d_a, d_b, d_c, N);
    (void)hipDeviceSynchronize();
    (void)hipFree(d_a);
    (void)hipFree(d_b);
    (void)hipFree(d_c);
    printf("Done.\n");
    return 0;
}

2. Build

hipcc -o repro_pmc repro_pmc.cpp --offload-arch=gfx1201

3. Control test — tracing works ✅

$ HIP_VISIBLE_DEVICES=0 rocprofv3 --kernel-trace -S -- ./repro_pmc

Done.

ROCPROFV3 SUMMARY:

    |                        NAME                         |     DOMAIN      |      CALLS      | ...
    | vector_add(float const*, float const*, float*, int) | KERNEL_DISPATCH |               1 | ...

# Exit code: 0

4. Crash — PMC counter collection ❌

$ HIP_VISIBLE_DEVICES=0 rocprofv3 --pmc SQ_WAVES -f csv -d /tmp/pmc_out -- ./repro_pmc

Crash output

W... metadata.cpp:330] rocprofiler_iterate_agent_supported_counters failed for agent 2 (gfx1036) :: Agent HW architecture is not supported, no counter metrics found.
W... simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.073674 sec
W... simple_timer.cpp:55] [rocprofv3] './repro_pmc' ::     0.000000 sec
W... tool.cpp:2422] HSA version 8.20.4 initialized (instance=0)
terminate called after throwing an instance of 'std::out_of_range'
  what():  unordered_map::at
W... tool.cpp:3104] [rocprofv3_error_signal_handler] rocprofv3 caught signal 6...
W... tool.cpp:3142] [rocprofv3_error_signal_handler] rocprofv3 finalizing after signal 6...
W... queue.cpp:938] Timeout while waiting for queue sync: 1 kernels still active
W... pool.hpp:201] Pool object at index 0 is still in use during pool clear
W... correlation_id.cpp:231] retiring dangling correlation ID 1 ... remaining reference count: 1
W... queue.cpp:938] Timeout while waiting for queue sync: 1 kernels still active

No output files are generated. Exit code: non-zero (SIGABRT).

Additional Information

• The gfx1036 warning about unsupported counters is expected — that's the Ryzen iGPU. Using HIP_VISIBLE_DEVICES=0 isolates to the discrete gfx1201 GPU but the crash persists.
• rocprofv3-avail list --pmc does list counters for gfx1201 (SQ_WAVES, VALUBusy, etc.), so the tool reports the counters as available but crashes when attempting to collect them.
• Tested all combinations: single counter, multiple counters, with/without --kernel-include-regex, with/without HIP graph capture disabled. All crash identically.
• rocprofv2 rejects the hardware: "ROCProfiler: fatal error: Unsupported hardware. Use rocprofv3 tool for navi4x, mi35x and later."
• This means there is no working path for hardware counter collection on gfx1201 with ROCm 7.2.4.
• Possibly related to [Issue]: Issue when profiling llama.cpp #86 (similar unordered_map::at crash, closed/migrated) and GL2C (L2 cache) performance counters read 0 on gfx1201 (RDNA4 / Navi 48) — every path, every privilege; compute counters work #155 (GL2C counters read 0 on gfx1201).

Impact

This blocks GPU performance analysis (occupancy, VALU utilization, memory bandwidth, LDS pressure) on RDNA4 hardware, which is needed for kernel tuning work such as flash attention optimization in projects like llama.cpp.