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

Optional profiler (GGML_HIP_ROOFLINE, off by default) that attributes each GPU
kernel's device-measured time to the ggml op that launched it and writes a JSON
report, one row per op invocation listing every kernel that op dispatched.
Aggregation (grouping identical ops, averaging, counting) is left to the consumer.
Activated at runtime by GGML_ROOFLINE_OUT=<path>; otherwise every entry point is a
no-op.

Kernel durations are read from device dispatch timestamps (the same source as
rocprofv3) and match rocprofv3's measurements. An earlier approach based on host
CUDA/HIP events (cudaEventElapsedTime around each op) was giving wrong results: it
included kernel-launch bubbles and overstated GPU time by roughly 1.7-3x.

Uses rocprofiler-sdk, loaded with dlopen and configured at runtime with symbols
resolved via dlsym (the library is not linked, since its constructors abort when
linked into an early-loaded shared object). Disables GPU graphs while active so each
op's kernels are dispatched individually, which is required for per-op attribution.

llama-bench calls ggml_cuda_roofline_reset() after its warmup run (via a weak symbol
that is a no-op in non-roofline builds), so the report covers only the measured run,
not warmup.

Each row carries the op's destination/source shapes and types, HBM byte traffic,
matmul dimensions, and the name and device time of every GPU kernel it dispatched.

Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>

Author: roberteg16

.github/workflows/build-gfx11-rocm.yml

@@ -220,6 +220,7 @@ jobs:
           -DGGML_CUDA_FORCE_CUBLAS=OFF \
           -DGGML_RPC=ON \
           -DGGML_HIP_ROCWMMA_FATTN=OFF \
+          -DGGML_HIP_ROOFLINE=ON \
           -DLLAMA_BUILD_BORINGSSL=ON \
           -DGGML_NATIVE=OFF \
           -DGGML_STATIC=OFF \

ggml/CMakeLists.txt

@@ -219,6 +219,7 @@ option(GGML_HIP_NO_VMM                      "ggml: do not try to use HIP VMM"
 option(GGML_HIP_ROCWMMA_FATTN               "ggml: enable rocWMMA for FlashAttention"         OFF)
 option(GGML_HIP_MMQ_MFMA                    "ggml: enable MFMA MMA for CDNA in MMQ"           ON)
 option(GGML_HIP_EXPORT_METRICS              "ggml: enable kernel perf metrics output"         OFF)
+option(GGML_HIP_ROOFLINE                    "ggml: enable per-op roofline profiling"          OFF)
 option(GGML_MUSA_GRAPHS                     "ggml: use MUSA graph, experimental, unstable"    OFF)
 option(GGML_MUSA_MUDNN_COPY                 "ggml: enable muDNN for accelerated copy"         OFF)
 option(GGML_VULKAN                          "ggml: use Vulkan"                                OFF)

ggml/src/ggml-cuda/ggml-cuda-roofline.cpp

@@ -0,0 +1,31 @@
+#pragma once
+
+// Per-op GPU roofline profiling for the HIP/ROCm backend (see ggml-cuda-roofline.cpp).
+// Compiled only when GGML_HIP_ROOFLINE is defined and activated at runtime by the
+// environment variable GGML_ROOFLINE_OUT=<path.json>. Every entry point is a no-op
+// unless that variable is set.
+
+#include "ggml.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+// Load and configure rocprofiler-sdk when GGML_ROOFLINE_OUT is set. Call once during
+// backend registration, before any GPU stream is created. Idempotent; no-op otherwise.
+void ggml_cuda_roofline_init(void);
+
+// Record the GPU architecture (e.g. "gfx1151") stored in the report. No-op unless active.
+void ggml_cuda_roofline_set_device(const char * arch);
+
+// Discard everything captured so far (drains pending records first). Call after a
+// warmup run so the report covers only the measured run. No-op unless active.
+void ggml_cuda_roofline_reset(void);
+
+// Tag the GPU kernels launched for this op so their device time is attributed to it.
+// Call once per op, before its kernel(s) are dispatched. No-op unless active.
+void ggml_cuda_roofline_begin_op(const struct ggml_tensor * node);
+
+#ifdef __cplusplus
+}
+#endif

ggml/src/ggml-cuda/ggml-cuda-roofline.h

@@ -2,6 +2,10 @@
 #include "ggml-impl.h"
 #include "ggml-backend-impl.h"
 
+#ifdef GGML_HIP_ROOFLINE
+#include "ggml-cuda/ggml-cuda-roofline.h"
+#endif
+
 #include "ggml-cuda/allreduce.cuh"
 #include "ggml-cuda/common.cuh"
 #include "ggml-cuda/acc.cuh"
@@ -4380,6 +4384,10 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
                     continue;
                 }
 
+#ifdef GGML_HIP_ROOFLINE
+                ggml_cuda_roofline_begin_op(node);
+#endif
+
                 int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i);
 
                 if (nodes_to_skip != 0) {
@@ -5648,6 +5656,14 @@ static const ggml_backend_reg_i ggml_backend_cuda_reg_interface = {
 
 // backend registry
 ggml_backend_reg_t ggml_backend_cuda_reg() {
+#ifdef GGML_HIP_ROOFLINE
+    ggml_cuda_roofline_init();  // configure rocprofiler tracing before any stream is created
+    if (ggml_cuda_info().device_count > 0) {
+        char arch[32];
+        snprintf(arch, sizeof(arch), "gfx%x", ggml_cuda_info().devices[0].cc & 0xffff);
+        ggml_cuda_roofline_set_device(arch);
+    }
+#endif
     static ggml_backend_reg reg;
     static bool initialized = false;
 

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -155,3 +155,19 @@ if (GGML_HIP_RCCL)
 endif()
 
 target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)
+
+# Per-op roofline profiling via rocprofiler-sdk. The library is loaded at runtime with
+# dlopen and its symbols resolved with dlsym, so only its headers and libdl are needed
+# to build; it is not linked (its constructors abort when linked into an early-loaded
+# shared object).
+if (GGML_HIP_ROOFLINE)
+    find_path(ROCPROFILER_SDK_INCLUDE_DIR rocprofiler-sdk/rocprofiler.h)
+    if (NOT ROCPROFILER_SDK_INCLUDE_DIR)
+        message(FATAL_ERROR "GGML_HIP_ROOFLINE requires rocprofiler-sdk headers (set CMAKE_PREFIX_PATH to the ROCm root)")
+    endif()
+    target_sources(ggml-hip PRIVATE ../ggml-cuda/ggml-cuda-roofline.cpp)
+    set_source_files_properties(../ggml-cuda/ggml-cuda-roofline.cpp PROPERTIES LANGUAGE CXX)
+    target_include_directories(ggml-hip PRIVATE ${ROCPROFILER_SDK_INCLUDE_DIR})
+    target_compile_definitions(ggml-hip PRIVATE GGML_HIP_ROOFLINE)
+    target_link_libraries(ggml-hip PRIVATE ${CMAKE_DL_LIBS})
+endif()

ggml/src/ggml-hip/CMakeLists.txt

@@ -8,6 +8,13 @@ set_target_properties(${TARGET} PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS ON)
 target_include_directories(${TARGET} PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
 target_link_libraries(${TARGET} PUBLIC llama-common llama ${CMAKE_THREAD_LIBS_INIT})
 
+# Compile in the post-warmup roofline reset only when the HIP profiler is built
+# (ggml-cuda-roofline.cpp, linked via ggml-hip). No-op in every other build.
+if (GGML_HIP_ROOFLINE)
+    target_compile_definitions(${TARGET} PRIVATE GGML_HIP_ROOFLINE)
+    target_include_directories(${TARGET} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../ggml/src/ggml-cuda)
+endif()
+
 if(LLAMA_TOOLS_INSTALL)
     install(TARGETS ${TARGET} LIBRARY)
 endif()

tools/llama-bench/CMakeLists.txt

@@ -35,6 +35,13 @@
 #    include <windows.h>
 #endif
 
+// Optional per-op roofline profiler (HIP build with GGML_HIP_ROOFLINE). Used to drop
+// warmup from the report. Compiled out entirely otherwise. See
+// ggml/src/ggml-cuda/ggml-cuda-roofline.cpp.
+#ifdef GGML_HIP_ROOFLINE
+#include "ggml-cuda-roofline.h"
+#endif
+
 // utils
 static uint64_t get_time_ns() {
     using clock = std::chrono::high_resolution_clock;
@@ -2355,6 +2362,12 @@ int llama_bench(int argc, char ** argv) {
             }
         }
 
+        // Discard warmup from the per-op roofline profiler so the report covers only
+        // the measured runs below. Compiled out unless GGML_HIP_ROOFLINE is set.
+#ifdef GGML_HIP_ROOFLINE
+        ggml_cuda_roofline_reset();
+#endif
+
         for (int i = 0; i < params.reps; i++) {
             llama_memory_clear(llama_get_memory(ctx), false);