ggml-cuda: add per-op GPU roofline profiling for the HIP backend

[roberteg16]

Summary

This PR adds an optional, off-by-default profiler for the HIP/ROCm backend that attributes each GPU kernel's device-measured time to the ggml op that launched it and writes a per-op JSON report on exit. It is intended to feed roofline analysis (arithmetic intensity, % of peak) with quant-aware, per-op numbers.

It is enabled at build time with GGML_HIP_ROOFLINE and activated at runtime with the GGML_ROOFLINE_OUT=<path.json> environment variable. When the variable is not set, every entry point is a no-op and the build behaves exactly like a normal one.

Accuracy

Kernel durations are read from device dispatch timestamps, the same source rocprofv3 uses, and match rocprofv3's measurements. Verified apples-to-apples with graphs disabled on both a small dense model and a large MoE model: token-generation throughput was within run-to-run noise of rocprofv3 and prefill was equal or slightly faster.

An earlier approach based on host CUDA/HIP events (cudaEventElapsedTime around each op) was giving wrong results: it captured kernel-launch bubbles and overstated GPU time by roughly 1.7-3x. That approach was discarded in favor of device timestamps, which are bubble-free.

How it works

The profiler uses rocprofiler-sdk, loaded with dlopen and configured at runtime via rocprofiler_force_configure, with its entry points resolved through dlsym. It is deliberately not linked: rocprofiler-sdk's global constructors abort when the library is linked into an early-loaded shared object, so linking it into libggml-hip is not viable.

Each op invocation is tagged with a unique external correlation id, so its kernels stay separate from every other invocation; the shared op geometry (destination and source shapes, types, and op params) is stored once per distinct shape. The kernel-dispatch records produced by rocprofiler carry the invocation id back, and each invocation's kernels are recorded with their individual device time. Kernel symbol names are resolved through the code-object callback and demangled. Grouping identical ops, counting, and averaging are left to the consumer.

While the profiler is active it disables GPU graphs (it sets GGML_CUDA_DISABLE_GRAPHS if the user has not set it). Per-op attribution requires the eager node loop, because a replayed graph dispatches all of its kernels in a single host call and the correlation id cannot be updated per kernel. This has negligible cost: disabling graphs was within noise on the dense model and slightly faster on the MoE model in testing.

To keep the report to the measured run only, the profiler exposes ggml_cuda_roofline_reset(), which drains any pending records and discards everything captured so far (kernel-name cache and the monotonic invocation counter are kept). llama-bench calls it after its warmup block and before the timed reps, so the artifact covers exactly one profiled run rather than warmup plus reps. Both the declaration and the call are guarded by #ifdef GGML_HIP_ROOFLINE (the same macro is added to the llama-bench target only when the option is enabled), so the code is compiled out entirely in non-roofline builds.

Report contents

The JSON report has one row per op invocation. Each row records the ggml op name, destination and source shapes and types, HBM byte traffic (total and per tensor), matmul dimensions, and the list of GPU kernels the op dispatched, each with its demangled symbol name and device time.

When the CUDA backend fuses several ops into one kernel (e.g. MUL_MAT+bias, ffn up+gate+SwiGLU, RMS_NORM+MUL, ROPE+VIEW+SET_ROWS), the row is labelled by the head op and adds a fused_ops array holding each fused node's geometry (op name, types, shapes, params, matmul dims). The row's bytes/dst_bytes are corrected to the fused group's real HBM traffic — intermediates produced and consumed inside the fused span never reach global memory and are excluded — while the per-node geometry lets a consumer sum exact FLOPs across the group (fusion changes memory traffic, not arithmetic).

Example

A single output-projection MUL_MAT invocation (the lm_head, q8_0 weight), showing the two kernels it dispatched (activation quantization + the matmul):

{
  "device": "gfx1151",
  "total_gpu_time_us": 63981.4,
  "rows": [
    {
      "ggml_op": "MUL_MAT",
      "dtype": "f32",
      "quant": "q8_0",
      "ne": [151936, 1, 1, 1],
      "src_ne": [[896, 151936, 1, 1], [896, 1, 1, 1]],
      "src_types": ["q8_0", "f32"],
      "op_params": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
      "bytes": 145254400,
      "dst_bytes": 607744,
      "src_bytes": [144643072, 3584],
      "M": 1, "N": 151936, "K": 896, "n_experts": 0, "top_k": 0,
      "kernels": [
        { "name": "quantize_q8_1(float const*, void*, ...) [clone .kd]", "gpu_time_us": 1.04 },
        { "name": "void mul_mat_vec_q<(ggml_type)8, 1, false, false>(void const*, ...) [clone .kd]", "gpu_time_us": 608.7 }
      ]
    }
  ]
}

ne is the destination shape and src_ne/src_types the sources (here src0 is the q8_0 weight [896, 151936], src1 the f32 activation). bytes is the total HBM traffic (dst_bytes + src_bytes). M/N/K are the matmul dims (n_experts/top_k are non-zero only for MUL_MAT_ID). op_params carries op geometry for conv/pool/rope/etc. Kernel symbol names are truncated above for readability; the artifact stores the full demangled signatures.

A fused row adds the fused_ops array. Here a down-projection MUL_MAT is fused with its residual ADD into one kernel; bytes/dst_bytes are the fused group's external traffic (the matmul output is an intermediate consumed by the ADD, so it is not counted), and each fused node keeps its own geometry so a consumer can sum exact FLOPs (2·M·N·K + numel here):

{
  "ggml_op": "MUL_MAT",
  "fused_ops": [
    {"ggml_op": "MUL_MAT", "dtype": "f32", "quant": "q4_K", "ne": [896, 1, 1, 1],
     "src_ne": [[4864, 896, 1, 1], [4864, 1, 1, 1]], "src_types": ["q4_K", "f32"],
     "op_params": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
     "M": 1, "N": 896, "K": 4864, "top_k": 0},
    {"ggml_op": "ADD", "dtype": "f32", "quant": "f32", "ne": [896, 1, 1, 1],
     "src_ne": [[896, 1, 1, 1], [896, 1, 1, 1]], "src_types": ["f32", "f32"],
     "op_params": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
     "M": 0, "N": 0, "K": 0, "top_k": 0}
  ],
  "dtype": "f32", "quant": "q4_K", "ne": [896, 1, 1, 1],
  "src_ne": [[4864, 896, 1, 1], [4864, 1, 1, 1]], "src_types": ["q4_K", "f32"],
  "op_params": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
  "bytes": 2478080, "dst_bytes": 3584, "src_bytes": [2451456, 19456],
  "M": 1, "N": 896, "K": 4864, "n_experts": 0, "top_k": 0,
  "kernels": [
    { "name": "quantize_q8_1(float const*, void*, ...) [clone .kd]", "gpu_time_us": 1.16 },
    { "name": "void mul_mat_vec_q<(ggml_type)12, 1, true, false>(void const*, ...) [clone .kd]", "gpu_time_us": 15.36 }
  ]
}

For a fused MUL_MAT_ID group (MoE), the expert-weight bytes in the group total are scaled to the experts actually routed (min(M·top_k, n_experts)), and top_k is taken from the ids tensor (the number of experts selected per token), so decode traffic is not over-counted.

Fusion is genuine, not just adjacent nodes

The fused_ops labels reflect real kernel-level fusion, verified two ways. First, structurally: every fused row dispatches exactly one compute kernel (matmul rows additionally carry a separate activation-quantize prep kernel, which is reported but not counted as the fused work) — no fused claim ever maps to multiple kernels. Second, by reading each kernel's implementation to confirm it performs the whole group in a single pass. The template/boolean parameters baked into the demangled kernel names even line up with the fusion flags:

• RMS_NORM+MUL → rms_norm_f32<block, do_multiply, do_add>; the recorded <…, true, false> sets do_multiply=true, so the kernel multiplies by the second operand in-kernel after normalizing (the <…, true, true> variant additionally fuses an ADD).
• SILU+MUL, SIGMOID+MUL → unary_gated_op_kernel<op, T>, which computes dst = op(x) * gate in one pass — the activation and the multiply are the same kernel.
• ROPE+VIEW+SET_ROWS → rope_neox<…, __half>, which takes row_indices/set_rows_stride and scatters the roped result straight into the KV cache row (the __half output type is the cache), so the rope compute and the set_rows write are fused.
• MUL_MAT+ADD and MUL_MAT+MUL_MAT+GLU → mul_mat_vec_q<type, ncols, has_fusion, …>; the recorded has_fusion=true path reads fusion.gate (the second projection), fusion.x_bias/fusion.gate_bias (the bias ADD), and active_glu, so the matmul(s) + bias + GLU run together.
• ADD+ADD+… → k_bin_bcast<op_add, n_fuse>, a kernel variadic over its source pointers and launched with make_index_sequence<n_fuse>, so an n-way residual add is one kernel.
• topk-moe routing (SOFT_MAX+RESHAPE+ARGSORT+VIEW+GET_ROWS+RESHAPE+SUM_ROWS+CLAMP+DIV+RESHAPE) → topk_moe_cuda, documented as the "fusion of softmax → top-k → get_rows pipeline", doing softmax + top-k + normalize/clamp in-kernel.

Usage

Build with the flag enabled: cmake -S . -B build -DGGML_HIP=ON -DGGML_HIP_ROOFLINE=ON && cmake --build build -j. The prebuilt ROCm multiarch release binaries (llama-b*-ubuntu-rocm-multiarch-x64.tar.gz) are built with GGML_HIP_ROOFLINE=ON, so they are profiling-capable out of the box — no rebuild needed, just set GGML_ROOFLINE_OUT at runtime.

Run any binary with the output path set: GGML_ROOFLINE_OUT=/tmp/roofline.json ./build/bin/llama-bench -m model.gguf -ngl 999 -p 256 -n 128. The report is written to that path on exit. Under llama-bench the artifact covers a single profiled run (warmup is reset out), so -r 1 gives a clean one-pass report.

The profiler loads librocprofiler-sdk.so at runtime; it ships in the same ROCm library directory as hipBLAS/rocBLAS, so no setup beyond what already lets the HIP binary find those libraries is needed. If the ROCm runtime is not on the default loader path (e.g. a pip/venv install), put that directory on LD_LIBRARY_PATH — the same as for any normal run of the binary.

Scope and impact

The change is contained to the HIP backend and a single new source file plus a one-line hook in the CUDA graph-compute loop, guarded by GGML_HIP_ROOFLINE. Builds without the flag are unaffected, and builds with the flag are unaffected unless GGML_ROOFLINE_OUT is set at runtime.

By design it touches as little existing code surface as possible so that pulling in upstream changes stays easy: all of the logic lives in the new ggml-cuda-roofline.{cpp,h}, and the edits to existing files are only the guarded one-line hook and init call in ggml-cuda.cu, the GGML_HIP_ROOFLINE option in ggml/CMakeLists.txt, the roofline build wiring in ggml-hip/CMakeLists.txt, and, for llama-bench, a guarded post-warmup reset call (behind #ifdef GGML_HIP_ROOFLINE) plus a two-line target_compile_definitions in its CMakeLists.txt that defines that macro for the llama-bench translation unit only when the option is on. Everything is compiled out unless the option is enabled, which keeps merge conflicts against upstream minimal and confined to a few well-isolated points.

The only non-source change is a one-line -DGGML_HIP_ROOFLINE=ON added to the ROCm release build in .github/workflows/build-gfx11-rocm.yml, so the published multiarch binaries ship profiling-capable. The rocprofiler-sdk headers this requires are already present in TheRock's multiarch ROCm tarball the workflow extracts to /opt/rocm, so no extra dependency install is needed.

[roberteg16]

Roofline counter part PR: https://gitenterprise.xilinx.com/FaaSApps/rocm-scripts/pull/597

[roberteg16]

@mgehre-amd this PR is hidden behind a cmake flag, it is safe to merge. It would simplify my workflow having it merge so I don't have to cherry-pick the commits on every branch