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);