From 1525b34107dffc04c1c8de3702ca1933cc2b2f24 Mon Sep 17 00:00:00 2001 From: Robert Esclapez Garcia Date: Thu, 2 Jul 2026 01:28:46 -0700 Subject: [PATCH] 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=; 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 so the report covers only the measured run, not warmup. The declaration and call are guarded by #ifdef GGML_HIP_ROOFLINE (the macro is added to the llama-bench target only when the option is on), so they compile out entirely in non-roofline builds. 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. When the backend fuses several ops into one kernel (matmul+bias, ffn up+gate+GLU, rms_norm+mul, rope+view+set_rows, ...), the row is labelled by the head op and adds a fused_ops array with each fused node's geometry; bytes/dst_bytes are corrected to the fused group's real traffic (intermediates excluded), and MoE expert weights are scaled to the routed experts (top_k taken from the ids tensor). Co-Authored-By: Claude Opus 4 (1M context) --- .github/workflows/build-gfx11-rocm.yml | 1 + ggml/CMakeLists.txt | 1 + ggml/src/ggml-cuda/ggml-cuda-roofline.cpp | 573 ++++++++++++++++++++++ ggml/src/ggml-cuda/ggml-cuda-roofline.h | 37 ++ ggml/src/ggml-cuda/ggml-cuda.cu | 19 + ggml/src/ggml-hip/CMakeLists.txt | 16 + tools/llama-bench/CMakeLists.txt | 7 + tools/llama-bench/llama-bench.cpp | 13 + 8 files changed, 667 insertions(+) create mode 100644 ggml/src/ggml-cuda/ggml-cuda-roofline.cpp create mode 100644 ggml/src/ggml-cuda/ggml-cuda-roofline.h diff --git a/.github/workflows/build-gfx11-rocm.yml b/.github/workflows/build-gfx11-rocm.yml index 6dabeabb5cd0..71d7d94302ce 100644 --- a/.github/workflows/build-gfx11-rocm.yml +++ b/.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 \ diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index a0cd4e7158f1..d823bb775b11 100644 --- a/ggml/CMakeLists.txt +++ b/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) diff --git a/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp b/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp new file mode 100644 index 000000000000..91a1479d991a --- /dev/null +++ b/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp @@ -0,0 +1,573 @@ +// Per-op GPU roofline profiling for the HIP/ROCm backend via rocprofiler-sdk. +// +// Attributes each GPU kernel's device-measured duration to the ggml op that launched +// it, using a unique external correlation id per op invocation, and writes a JSON +// report on exit with one row per invocation listing the kernels it dispatched. +// Aggregation (grouping identical ops, averaging, counting) is left to the consumer. +// Activated at runtime by the environment variable GGML_ROOFLINE_OUT=; +// every entry point is a no-op unless it is set. +// +// rocprofiler-sdk is loaded with dlopen and configured with rocprofiler_force_configure +// at runtime, and its entry points are resolved with dlsym. It is intentionally not +// linked: its global constructors abort when the library is loaded early as part of +// another shared object. + +#include "ggml-cuda-roofline.h" +#include "ggml-impl.h" + +#include +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr int n_op_params = GGML_MAX_OP_PARAMS / (int) sizeof(int32_t); // 16 + +// Geometry and memory traffic captured for one distinct ggml op shape. +struct op_record { + const char * op = ""; // op / unary / glu name (ggml_op_desc) + const char * dtype = ""; // destination type + const char * quant = ""; // src0 type (weight quantization for matmul) + int64_t ne[4] = {0, 0, 0, 0}; // destination shape + int64_t src_ne[GGML_MAX_SRC][4] = {}; // source shapes + const char * src_types[GGML_MAX_SRC] = {}; // source types (ggml_type_name) + int n_src = 0; + int32_t op_params[n_op_params] = {}; // op parameters (conv / pool / rope / ... geometry) + int64_t bytes = 0; // total HBM traffic: destination + all sources + int64_t dst_bytes = 0; // ggml_nbytes(destination) + int64_t src_bytes[GGML_MAX_SRC] = {}; // ggml_nbytes(source) + int64_t M = 0, N = 0, K = 0, n_experts = 0, top_k = 0; // matmul dimensions + std::vector fused_nodes; // per-node geometry when this row is a fused group (else empty) +}; + +// One kernel dispatch: which kernel ran, when it started, and for how long. The start +// time is kept only to order kernels within an op (buffer records arrive unordered). +struct dispatch { + uint64_t start_ns = 0; + uint64_t kernel_id = 0; + uint64_t duration_ns = 0; +}; + +std::mutex g_mutex; +std::unordered_map g_records; // geometry id -> geometry +std::unordered_map g_invocations; // invocation id -> geometry id +std::unordered_map> g_dispatches; // invocation id -> dispatches +std::unordered_map g_kernel_names; // kernel id -> demangled symbol +std::atomic g_next_invocation{1}; +thread_local uint64_t g_current_invocation = 0; // id pushed by the last begin_op on this thread + +bool g_active = false; +std::string g_out_path; +std::string g_device; // GPU architecture, e.g. "gfx1151" +rocprofiler_context_id_t g_context{0}; +rocprofiler_buffer_id_t g_buffer{0}; + +// rocprofiler-sdk entry points, resolved at runtime via dlsym. +using fn_create_context = rocprofiler_status_t (*)(rocprofiler_context_id_t *); +using fn_create_buffer = rocprofiler_status_t (*)(rocprofiler_context_id_t, size_t, size_t, + rocprofiler_buffer_policy_t, + rocprofiler_buffer_tracing_cb_t, void *, + rocprofiler_buffer_id_t *); +using fn_config_buffer = rocprofiler_status_t (*)(rocprofiler_context_id_t, + rocprofiler_buffer_tracing_kind_t, + const rocprofiler_tracing_operation_t *, + size_t, rocprofiler_buffer_id_t); +using fn_config_callback = rocprofiler_status_t (*)(rocprofiler_context_id_t, + rocprofiler_callback_tracing_kind_t, + const rocprofiler_tracing_operation_t *, + size_t, rocprofiler_callback_tracing_cb_t, void *); +using fn_start = rocprofiler_status_t (*)(rocprofiler_context_id_t); +using fn_flush = rocprofiler_status_t (*)(rocprofiler_buffer_id_t); +using fn_get_thread_id = rocprofiler_status_t (*)(rocprofiler_thread_id_t *); +using fn_push_id = rocprofiler_status_t (*)(rocprofiler_context_id_t, rocprofiler_thread_id_t, + rocprofiler_user_data_t); +using fn_pop_id = rocprofiler_status_t (*)(rocprofiler_context_id_t, rocprofiler_thread_id_t, + rocprofiler_user_data_t *); + +fn_create_context p_create_context = nullptr; +fn_create_buffer p_create_buffer = nullptr; +fn_config_buffer p_config_buffer = nullptr; +fn_config_callback p_config_callback = nullptr; +fn_start p_start = nullptr; +fn_flush p_flush = nullptr; +fn_get_thread_id p_get_thread_id = nullptr; +fn_push_id p_push_id = nullptr; +fn_pop_id p_pop_id = nullptr; + +// Resolve the required entry points; returns false if any is missing. The code-object +// callback service is optional (kernel-symbol names) and resolved separately. +bool resolve_symbols() { + p_create_context = (fn_create_context) dlsym(RTLD_DEFAULT, "rocprofiler_create_context"); + p_create_buffer = (fn_create_buffer) dlsym(RTLD_DEFAULT, "rocprofiler_create_buffer"); + p_config_buffer = (fn_config_buffer) dlsym(RTLD_DEFAULT, "rocprofiler_configure_buffer_tracing_service"); + p_start = (fn_start) dlsym(RTLD_DEFAULT, "rocprofiler_start_context"); + p_flush = (fn_flush) dlsym(RTLD_DEFAULT, "rocprofiler_flush_buffer"); + p_get_thread_id = (fn_get_thread_id) dlsym(RTLD_DEFAULT, "rocprofiler_get_thread_id"); + p_push_id = (fn_push_id) dlsym(RTLD_DEFAULT, "rocprofiler_push_external_correlation_id"); + p_pop_id = (fn_pop_id) dlsym(RTLD_DEFAULT, "rocprofiler_pop_external_correlation_id"); + p_config_callback = (fn_config_callback) dlsym(RTLD_DEFAULT, "rocprofiler_configure_callback_tracing_service"); + return p_create_context && p_create_buffer && p_config_buffer && p_start && p_flush && + p_get_thread_id && p_push_id && p_pop_id; +} + +std::string demangle(const char * name) { + if (!name) return ""; + int status = 0; + char * demangled = abi::__cxa_demangle(name, nullptr, nullptr, &status); + std::string result = (status == 0 && demangled) ? demangled : name; + free(demangled); + return result; +} + +void json_escape(std::ostringstream & out, const std::string & str) { + for (char c : str) { + if (c == '"' || c == '\\') out << '\\' << c; + else out << c; + } +} + +uint64_t hash_mix(uint64_t hash, uint64_t value) { + hash ^= value + 0x9e3779b97f4a7c15ULL + (hash << 6) + (hash >> 2); + return hash; +} + +// Fill a record's geometry and single-node HBM byte fields from one ggml node. +void fill_head_record(op_record & rec, const ggml_tensor * node) { + const ggml_tensor * src0 = node->src[0]; + const ggml_tensor * src1 = node->src[1]; + + rec.op = ggml_op_desc(node); + rec.dtype = ggml_type_name(node->type); + rec.quant = src0 ? ggml_type_name(src0->type) : ""; + for (int d = 0; d < 4; d++) rec.ne[d] = node->ne[d]; + for (int j = 0; j < GGML_MAX_SRC; j++) { + if (node->src[j]) { + for (int d = 0; d < 4; d++) rec.src_ne[j][d] = node->src[j]->ne[d]; + rec.src_types[j] = ggml_type_name(node->src[j]->type); + rec.n_src = j + 1; + } + } + memcpy(rec.op_params, node->op_params, sizeof(rec.op_params)); + + // total HBM traffic for this op: write the destination and read every source. + rec.dst_bytes = (int64_t) ggml_nbytes(node); + rec.bytes = rec.dst_bytes; + for (int j = 0; j < GGML_MAX_SRC; j++) { + if (node->src[j]) { + rec.src_bytes[j] = (int64_t) ggml_nbytes(node->src[j]); + rec.bytes += rec.src_bytes[j]; + } + } + + if ((node->op == GGML_OP_MUL_MAT || node->op == GGML_OP_MUL_MAT_ID) && src0 && src1) { + rec.K = src0->ne[0]; + rec.N = src0->ne[1]; + if (node->op == GGML_OP_MUL_MAT) { + rec.M = src1->ne[1]; + } else { + // MUL_MAT_ID (MoE): src0 = experts [K,N,E], src1 = input (often broadcast + // [K,1,n_tokens]), src2 = ids [n_expert_used, n_tokens]. top_k is the number of + // experts routed per token = ids->ne[0]; src1->ne[1] is 1 when the input is + // broadcast, so it must not be used. Fall back to src1->ne[1] only if ids absent. + rec.M = src1->ne[2]; + rec.n_experts = src0->ne[2]; + rec.top_k = node->src[2] ? node->src[2]->ne[0] : src1->ne[1]; + } + } +} + +// Dedup hash of one node's geometry (destination, all sources, op params, types); distinct +// shapes get distinct ids so the report can be deduplicated. +uint64_t head_geometry_id(const op_record & rec, const ggml_tensor * node) { + const ggml_tensor * src0 = node->src[0]; + uint64_t geometry_id = hash_mix(0, (uint64_t) node->op); + for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) rec.ne[d]); + for (int j = 0; j < GGML_MAX_SRC; j++) { + for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) rec.src_ne[j][d]); + } + for (int p = 0; p < n_op_params; p++) geometry_id = hash_mix(geometry_id, (uint64_t) (uint32_t) rec.op_params[p]); + geometry_id = hash_mix(geometry_id, (uint64_t) (src0 ? src0->type : 0)); + geometry_id = hash_mix(geometry_id, (uint64_t) node->type); + return geometry_id; +} + +// Buffer callback: records each kernel dispatch (kernel id + device time) under the +// invocation that launched it. Invoked asynchronously as the tracing buffer fills. +void buffer_callback(rocprofiler_context_id_t, rocprofiler_buffer_id_t, + rocprofiler_record_header_t ** headers, size_t n_headers, void *, uint64_t) { + std::lock_guard lock(g_mutex); + for (size_t i = 0; i < n_headers; i++) { + rocprofiler_record_header_t * header = headers[i]; + if (header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING && + header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) { + auto * record = static_cast(header->payload); + const uint64_t invocation = record->correlation_id.external.value; + g_dispatches[invocation].push_back({record->start_timestamp, + record->dispatch_info.kernel_id, + record->end_timestamp - record->start_timestamp}); + } + } +} + +// Code-object callback: records kernel id -> demangled symbol as kernels load. The SDK +// frees the name string on unload, so the demangled copy is kept. +void code_object_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t *, void *) { + if (record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT || + record.operation != ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER || + record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) { + return; + } + auto * symbol = static_cast< + rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t *>(record.payload); + std::lock_guard lock(g_mutex); + g_kernel_names[symbol->kernel_id] = demangle(symbol->kernel_name); +} + +void write_shape(std::ostringstream & out, const char * key, const int64_t ne[4]) { + out << "\"" << key << "\": [" << ne[0] << ", " << ne[1] << ", " << ne[2] << ", " << ne[3] << "]"; +} + +// One fused node's geometry: op name, types, shapes, params, and matmul dims. Enough for +// the consumer to sum exact FLOPs and render each fused op's config. Byte fields are +// omitted: the fused row's group-level bytes are authoritative. +void write_fused_node(std::ostringstream & out, const op_record & rec) { + out << "{\"ggml_op\": \"" << rec.op << "\", " + << "\"dtype\": \"" << rec.dtype << "\", \"quant\": \"" << rec.quant << "\", "; + write_shape(out, "ne", rec.ne); out << ", "; + out << "\"src_ne\": ["; + for (int j = 0; j < rec.n_src; j++) { + if (j) out << ", "; + out << "[" << rec.src_ne[j][0] << ", " << rec.src_ne[j][1] << ", " + << rec.src_ne[j][2] << ", " << rec.src_ne[j][3] << "]"; + } + out << "], \"src_types\": ["; + for (int j = 0; j < rec.n_src; j++) { + if (j) out << ", "; + out << "\"" << (rec.src_types[j] ? rec.src_types[j] : "") << "\""; + } + out << "], \"op_params\": ["; + for (int p = 0; p < n_op_params; p++) { if (p) out << ", "; out << rec.op_params[p]; } + out << "], \"M\": " << rec.M << ", \"N\": " << rec.N << ", \"K\": " << rec.K + << ", \"top_k\": " << rec.top_k << "}"; +} + +// Write the per-invocation JSON report. Registered with atexit while profiling is active. +void write_report() { + if (!g_active || !p_flush) return; + p_flush(g_buffer); + + std::lock_guard lock(g_mutex); + double total_us = 0.0; + for (auto & [invocation, dispatches] : g_dispatches) { + for (const auto & d : dispatches) total_us += d.duration_ns / 1e3; + // order kernels within an op causally (buffer records arrive unordered) + std::sort(dispatches.begin(), dispatches.end(), + [](const dispatch & a, const dispatch & b) { return a.start_ns < b.start_ns; }); + } + + std::ostringstream out; + out << "{\n \"device\": \"" << g_device << "\",\n"; + out << " \"total_gpu_time_us\": " << total_us << ",\n \"rows\": [\n"; + bool first = true; + for (const auto & [invocation, geometry_id] : g_invocations) { + auto dispatch_it = g_dispatches.find(invocation); + if (dispatch_it == g_dispatches.end()) continue; // no kernels recorded for this op + auto record_it = g_records.find(geometry_id); + if (record_it == g_records.end()) continue; + const op_record & rec = record_it->second; + + if (!first) out << ",\n"; + first = false; + + out << " {\"ggml_op\": \"" << rec.op << "\", "; + if (!rec.fused_nodes.empty()) { + out << "\"fused_ops\": ["; + for (size_t k = 0; k < rec.fused_nodes.size(); k++) { + if (k) out << ", "; + write_fused_node(out, rec.fused_nodes[k]); + } + out << "], "; + } + out << "\"dtype\": \"" << rec.dtype << "\", \"quant\": \"" << rec.quant << "\", "; + write_shape(out, "ne", rec.ne); out << ", "; + out << "\"src_ne\": ["; + for (int j = 0; j < rec.n_src; j++) { + if (j) out << ", "; + out << "[" << rec.src_ne[j][0] << ", " << rec.src_ne[j][1] << ", " + << rec.src_ne[j][2] << ", " << rec.src_ne[j][3] << "]"; + } + out << "], "; + out << "\"src_types\": ["; + for (int j = 0; j < rec.n_src; j++) { + if (j) out << ", "; + out << "\"" << (rec.src_types[j] ? rec.src_types[j] : "") << "\""; + } + out << "], "; + out << "\"op_params\": ["; + for (int p = 0; p < n_op_params; p++) { if (p) out << ", "; out << rec.op_params[p]; } + out << "], "; + out << "\"bytes\": " << rec.bytes << ", "; + out << "\"dst_bytes\": " << rec.dst_bytes << ", "; + out << "\"src_bytes\": ["; + for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_bytes[j]; } + out << "], "; + out << "\"M\": " << rec.M << ", \"N\": " << rec.N << ", \"K\": " << rec.K + << ", \"n_experts\": " << rec.n_experts << ", \"top_k\": " << rec.top_k << ", "; + out << "\"kernels\": ["; + bool kernel_first = true; + for (const auto & d : dispatch_it->second) { + if (!kernel_first) out << ", "; + kernel_first = false; + auto name_it = g_kernel_names.find(d.kernel_id); + out << "{\"name\": \""; + json_escape(out, name_it != g_kernel_names.end() ? name_it->second : ""); + out << "\", \"gpu_time_us\": " << d.duration_ns / 1e3 << "}"; + } + out << "]}"; + } + out << "\n ]\n}\n"; + + std::ofstream file(g_out_path); + file << out.str(); + // runs from atexit, where the ggml logger is already torn down, so write directly + fprintf(stderr, "ggml-roofline: wrote %s (%zu invocations, %zu unique ops, total GPU %.1f us)\n", + g_out_path.c_str(), g_invocations.size(), g_records.size(), total_us); +} + +int tool_init(rocprofiler_client_finalize_t, void *) { + if (!resolve_symbols()) { + GGML_LOG_ERROR("ggml-roofline: could not resolve rocprofiler-sdk symbols; disabled\n"); + g_active = false; + return -1; + } + if (p_create_context(&g_context) != ROCPROFILER_STATUS_SUCCESS) return -1; + // 256 MiB buffer, 4 MiB watermark: batch dispatch records so the callback fires in + // large chunks instead of per record. + if (p_create_buffer(g_context, 256 * 1024 * 1024, 4 * 1024 * 1024, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, buffer_callback, nullptr, &g_buffer) + != ROCPROFILER_STATUS_SUCCESS) return -1; + if (p_config_buffer(g_context, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, g_buffer) + != ROCPROFILER_STATUS_SUCCESS) return -1; + // Kernel-symbol names are optional: timing works without them, kernels just omit "name". + if (p_config_callback) { + p_config_callback(g_context, ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, nullptr, 0, + code_object_callback, nullptr); + } + if (p_start(g_context) != ROCPROFILER_STATUS_SUCCESS) return -1; + return 0; +} + +void tool_fini(void *) {} + +rocprofiler_tool_configure_result_t g_configure_result{ + sizeof(rocprofiler_tool_configure_result_t), &tool_init, &tool_fini, nullptr}; + +rocprofiler_tool_configure_result_t * +configure(uint32_t, const char *, uint32_t, rocprofiler_client_id_t * client_id) { + client_id->name = "ggml-roofline"; + g_active = true; + return &g_configure_result; +} + +using fn_force_configure = rocprofiler_status_t (*)( + rocprofiler_tool_configure_result_t * (*)(uint32_t, const char *, uint32_t, rocprofiler_client_id_t *)); + +} // namespace + +void ggml_cuda_roofline_init(void) { + static std::once_flag once; + std::call_once(once, [] { + const char * out_path = getenv("GGML_ROOFLINE_OUT"); + if (!out_path || !out_path[0]) return; + g_out_path = out_path; + + // Per-op attribution requires the eager node loop, so graphs must be disabled. + // The backend reads this env var lazily and only checks for its presence; set it + // here before that happens, without overriding an explicit user setting. + if (setenv("GGML_CUDA_DISABLE_GRAPHS", "1", 0) == 0) { + GGML_LOG_INFO("ggml-roofline: graphs disabled (required for per-op attribution)\n"); + } + + void * handle = dlopen("librocprofiler-sdk.so.1", RTLD_NOW | RTLD_GLOBAL); + if (!handle) handle = dlopen("librocprofiler-sdk.so", RTLD_NOW | RTLD_GLOBAL); + if (!handle) { + GGML_LOG_ERROR("ggml-roofline: cannot dlopen librocprofiler-sdk: %s\n", dlerror()); + return; + } + auto force_configure = (fn_force_configure) dlsym(handle, "rocprofiler_force_configure"); + if (!force_configure || force_configure(&configure) != ROCPROFILER_STATUS_SUCCESS) { + GGML_LOG_ERROR("ggml-roofline: rocprofiler_force_configure failed\n"); + return; + } + std::atexit(write_report); + GGML_LOG_INFO("ggml-roofline: enabled -> %s\n", g_out_path.c_str()); + }); +} + +void ggml_cuda_roofline_set_device(const char * arch) { + if (g_active && arch) g_device = arch; +} + +void ggml_cuda_roofline_reset(void) { + if (!g_active) return; + // Drain any pending records first so warmup dispatches are accounted, then dropped. + if (p_flush) p_flush(g_buffer); + std::lock_guard lock(g_mutex); + g_records.clear(); + g_invocations.clear(); + g_dispatches.clear(); + // g_next_invocation stays monotonic so a late warmup record cannot collide with a + // post-reset invocation id; g_kernel_names is kept (code objects do not reload). +} + +void ggml_cuda_roofline_begin_op(const struct ggml_tensor * node) { + if (!g_active || node == nullptr || !p_push_id) return; + + op_record rec; + fill_head_record(rec, node); + const uint64_t geometry_id = head_geometry_id(rec, node); + + // Each op invocation gets a unique correlation id so its kernels stay separate; the + // shared geometry is stored once per shape. + const uint64_t invocation = g_next_invocation.fetch_add(1, std::memory_order_relaxed); + { + std::lock_guard lock(g_mutex); + g_invocations.emplace(invocation, geometry_id); + if (g_records.find(geometry_id) == g_records.end()) g_records.emplace(geometry_id, rec); + } + + // Tag the kernels launched until the next op with this invocation id. rocprofiler + // keeps a per-thread stack, so pop the previous id before pushing the new one. + static thread_local rocprofiler_thread_id_t thread_id = [] { + rocprofiler_thread_id_t t = 0; if (p_get_thread_id) p_get_thread_id(&t); return t; + }(); + static thread_local bool pushed = false; + if (pushed) { + rocprofiler_user_data_t previous; + p_pop_id(g_context, thread_id, &previous); + } + rocprofiler_user_data_t current; + current.value = invocation; + p_push_id(g_context, thread_id, current); + pushed = true; + + // Remember which invocation this thread just tagged so a following fusion can override + // its record to cover the whole fused span (see ggml_cuda_roofline_fuse_ops). + g_current_invocation = invocation; +} + +void ggml_cuda_roofline_fuse_ops(const struct ggml_cgraph * cgraph, int node_idx, int node_count) { + if (!g_active || cgraph == nullptr || node_count < 2 || g_current_invocation == 0) return; + + const ggml_tensor * head = cgraph->nodes[node_idx]; + + // Keep the head op's geometry (op name, shapes, M/N/K) so the row is still labelled by + // the head op; the byte fields below are replaced with the fused-group traffic. + op_record rec; + fill_head_record(rec, head); + + // Record every fused node's geometry (op name + shapes + params + matmul dims) so the + // consumer can sum exact FLOPs across the whole fused group. FLOPs are additive under + // fusion (only memory traffic is saved), so per-node geometry is all it needs; byte + // fields on the sub-records are left unused (the group total below is authoritative). + rec.fused_nodes.reserve(node_count); + for (int j = node_idx; j < node_idx + node_count; ++j) { + op_record sub; + fill_head_record(sub, cgraph->nodes[j]); + rec.fused_nodes.push_back(std::move(sub)); + } + + // Correct HBM traffic for the fused kernel: intermediate tensors produced and consumed + // inside the span never reach global memory, so count only external inputs and outputs. + // A tensor is internal iff it is the output of one of the fused nodes (in ggml the node + // tensor is its own output), matching ggml_cuda_check_fusion_memory_ranges. + std::unordered_set produced; + for (int j = node_idx; j < node_idx + node_count; ++j) { + produced.insert(cgraph->nodes[j]); + } + + // External inputs (weights + activations), deduplicated; intermediates skipped. Unlike + // the fusion overlap check, leaf tensors (op == GGML_OP_NONE, e.g. weights) ARE counted. + // MoE (MUL_MAT_ID) expert weights are stored for all experts but only the routed ones + // (min(M*top_k, n_experts)) are read, so src0 of such a node is scaled by that fraction + // to avoid over-counting; every other source is read in full. + const auto input_bytes = [](const ggml_tensor * n, int s, const ggml_tensor * src) -> int64_t { + const int64_t full = (int64_t) ggml_nbytes(src); + if (n->op == GGML_OP_MUL_MAT_ID && s == 0 && n->src[0] && n->src[1]) { + const int64_t n_experts = n->src[0]->ne[2]; + // top_k = experts routed per token = ids (src2) ->ne[0]; src1->ne[1] is 1 when + // the input is broadcast, so use ids and fall back only if it is absent. + const int64_t top_k = n->src[2] ? n->src[2]->ne[0] : n->src[1]->ne[1]; + const int64_t m = n->src[1]->ne[2]; + if (n_experts > 0) { + const int64_t used = m > 0 ? std::min(m * top_k, n_experts) : n_experts; + return full * used / n_experts; + } + } + return full; + }; + int64_t bytes = 0; + std::unordered_set counted_inputs; + for (int j = node_idx; j < node_idx + node_count; ++j) { + const ggml_tensor * n = cgraph->nodes[j]; + for (int s = 0; s < GGML_MAX_SRC; ++s) { + const ggml_tensor * src = n->src[s]; + if (src && !produced.count(src) && counted_inputs.insert(src).second) { + bytes += input_bytes(n, s, src); + } + } + } + + // External outputs: fused nodes not consumed as a source by any other fused node. + int64_t dst_bytes = 0; + for (int j = node_idx; j < node_idx + node_count; ++j) { + const ggml_tensor * n = cgraph->nodes[j]; + bool consumed = false; + for (int k = node_idx; k < node_idx + node_count && !consumed; ++k) { + for (int s = 0; s < GGML_MAX_SRC; ++s) { + if (cgraph->nodes[k]->src[s] == n) { consumed = true; break; } + } + } + if (!consumed) { + dst_bytes += (int64_t) ggml_nbytes(n); + } + } + bytes += dst_bytes; + + rec.bytes = bytes; + rec.dst_bytes = dst_bytes; + + // Fused geometry id: head geometry plus each fused node's op and destination shape, so + // identical fusions deduplicate and distinct ones stay separate. + uint64_t geometry_id = head_geometry_id(rec, head); + for (int j = node_idx; j < node_idx + node_count; ++j) { + const ggml_tensor * n = cgraph->nodes[j]; + geometry_id = hash_mix(geometry_id, (uint64_t) n->op); + for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) n->ne[d]); + } + + // Re-point the current invocation at the fused record. The provisional head-only record + // from begin_op stays in g_records; if no non-fused invocation references it, it is + // simply never emitted (the report iterates g_invocations). + std::lock_guard lock(g_mutex); + g_invocations[g_current_invocation] = geometry_id; + if (g_records.find(geometry_id) == g_records.end()) g_records.emplace(geometry_id, std::move(rec)); +} diff --git a/ggml/src/ggml-cuda/ggml-cuda-roofline.h b/ggml/src/ggml-cuda/ggml-cuda-roofline.h new file mode 100644 index 000000000000..2cf5a0b385a7 --- /dev/null +++ b/ggml/src/ggml-cuda/ggml-cuda-roofline.h @@ -0,0 +1,37 @@ +#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=. 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); + +// Override the record of the op tagged by the last begin_op so it covers a fused span of +// node_count nodes (cgraph->nodes[node_idx .. node_idx+node_count-1]): lists every fused op +// and reports the fused group's HBM traffic with intermediates discarded. Call right after a +// successful fusion, before advancing the loop. No-op unless active. +void ggml_cuda_roofline_fuse_ops(const struct ggml_cgraph * cgraph, int node_idx, int node_count); + +#ifdef __cplusplus +} +#endif diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index cca70592f807..ab7590434de6 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -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,9 +4384,16 @@ 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) { +#ifdef GGML_HIP_ROOFLINE + ggml_cuda_roofline_fuse_ops(cgraph, i, nodes_to_skip + 1); +#endif i += nodes_to_skip; continue; } @@ -5648,6 +5659,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; diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index a7d4e0ea2b53..dcfdb2bdb599 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -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() diff --git a/tools/llama-bench/CMakeLists.txt b/tools/llama-bench/CMakeLists.txt index b1c35ee88a5f..e5119c24380d 100644 --- a/tools/llama-bench/CMakeLists.txt +++ b/tools/llama-bench/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() diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp index 2695f58785e2..e651e904d287 100644 --- a/tools/llama-bench/llama-bench.cpp +++ b/tools/llama-bench/llama-bench.cpp @@ -35,6 +35,13 @@ # include #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);