diff --git a/common/arg.cpp b/common/arg.cpp index 015196ca14..c217002f40 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1483,6 +1483,21 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.server_base = value; } ).set_examples({LLAMA_EXAMPLE_CLI})); + add_opt(common_arg( + {"--profile"}, + "enable cross-backend profiling (CPU, BLAS, CUDA)", + [](common_params & params) { + params.profiling = true; + } + ).set_examples({LLAMA_EXAMPLE_CLI, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_DEBUG})); + add_opt(common_arg( + {"--profile-output"}, "FNAME", + "write profiling JSON output to FNAME (default: stdout)", + [](common_params & params, const std::string & value) { + params.profiling = true; + params.profiling_output = value; + } + ).set_examples({LLAMA_EXAMPLE_CLI, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_DEBUG})); add_opt(common_arg( {"--verbose-prompt"}, string_format("print a verbose prompt before generation (default: %s)", params.verbose_prompt ? "true" : "false"), @@ -3177,6 +3192,13 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } ).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_FINETUNE, LLAMA_EXAMPLE_RESULTS, LLAMA_EXAMPLE_EXPORT_GRAPH_OPS, LLAMA_EXAMPLE_CLI})); + add_opt(common_arg( + {"--with-backends"}, + "export graph ops with backend assignments (default: CPU only)", + [](common_params & params) { + params.with_backends = true; + } + ).set_examples({LLAMA_EXAMPLE_EXPORT_GRAPH_OPS})); add_opt(common_arg( {"-ofreq", "--output-frequency"}, "N", string_format("output the imatrix every N iterations (default: %d)", params.n_out_freq), diff --git a/common/common.cpp b/common/common.cpp index d162a38800..72f39561e8 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -3,6 +3,7 @@ #include "build-info.h" #include "common.h" +#include "ggml-profiler.h" #include "fit.h" #include "log.h" #include "llama.h" @@ -1401,6 +1402,14 @@ common_init_result::common_init_result(common_params & params, bool model_only) return; } + if (params.profiling) { + ggml_backend_sched_t sched = llama_context_get_sched(lctx); + if (sched != nullptr) { + ggml_backend_sched_set_profiling(sched, true); + LOG_INF("%s: profiling enabled\n", __func__); + } + } + pimpl->context.reset(lctx); set_process_priority(params.cpuparams.priority); diff --git a/common/common.h b/common/common.h index 63d0badd0f..e3589fb4f7 100644 --- a/common/common.h +++ b/common/common.h @@ -5,6 +5,7 @@ #include "llama-cpp.h" #include "ggml-opt.h" +#include "ggml-profiler.h" #include "ggml.h" #include "llama.h" @@ -476,6 +477,7 @@ struct common_params { bool fit_params = true; // whether to fit unset model/context parameters to free device memory bool fit_params_print = false; // print the estimated required memory to run the model int32_t fit_params_min_ctx = 4096; // minimum context size to set when trying to reduce memory use + bool with_backends = false; // export graph ops with backend assignments // margin per device in bytes for fitting parameters to free memory: std::vector fit_params_target = std::vector(llama_max_devices(), 1024 * 1024*1024); @@ -733,6 +735,10 @@ struct common_params { bool spm_infill = false; // suffix/prefix/middle pattern for infill + // profiling + bool profiling = false; // enable cross-backend profiling + std::string profiling_output; // path to write profiling JSON output (empty = stdout) + // batched-bench params bool batched_bench_output_jsonl = false; diff --git a/docs/cross-profiler.md b/docs/cross-profiler.md new file mode 100644 index 0000000000..ad90b9e20a --- /dev/null +++ b/docs/cross-profiler.md @@ -0,0 +1,225 @@ +# Cross-Backend Profiler + +llama.cpp includes a built-in cross-backend profiler that captures per-operation timing, data transfer costs, and tensor shapes across all compute backends. It works with any application built on the ggml scheduler — no source changes needed. + +## Supported Backends + +| Backend | Status | Timing method | +|---------|--------|---------------| +| CPU | Supported | Wall-clock (`CLOCK_MONOTONIC_RAW`) | +| CUDA | Supported | `cudaEvent` GPU timestamps | +| Vulkan | Supported | GPU timestamp queries | +| BLAS | Supported | Wall-clock | +| Metal | Not yet supported | — | +| OpenCL | Not yet supported | — | + +The scheduler also profiles **data copies** (H2D, D2H, D2D) between backends regardless of which backends have native profiler support. + +## Enabling the Profiler + +There are two independent ways to enable profiling. They can be used separately or together. + +### CLI flags (`--profile`, `--profile-output`) + +Available in `llama-cli`, `llama-completion`, `llama-server`, and `debug`: + +```bash +# Print summary to stdout +llama-completion -m model.gguf --profile -p "Hello world" + +# Export to JSON +llama-completion -m model.gguf --profile --profile-output profile.json -p "Hello world" + +# Export to plain text +llama-completion -m model.gguf --profile --profile-output profile.txt -p "Hello world" +``` + +The output format is chosen by file extension: `.json` for JSON, `.txt` for plain text. Any other extension defaults to JSON. + +### Environment variable (`GGML_PROFILE`) + +The `GGML_PROFILE` environment variable enables profiling at the ggml scheduler level. This works with **any** application that uses the scheduler — including third-party tools like `sd.cpp` — without CLI flag support. + +```bash +# Print summary to stdout +GGML_PROFILE=1 llama-completion -m model.gguf -p "Hello world" + +# Export JSON +GGML_PROFILE=profile.json llama-completion -m model.gguf -p "Hello world" + +# Export plain text +GGML_PROFILE=profile.txt llama-completion -m model.gguf -p "Hello world" + +# Works with any ggml-based application +GGML_PROFILE=1 sd -m model.gguf -p "a cat" +``` + +| Value | Behavior | +|-------|----------| +| `1`, `stdout`, or empty | Print summary to stdout | +| `path.json` | Export JSON to file | +| `path.txt` | Export plain text to file | +| Any other path | Export JSON to file | + +The export happens automatically when the scheduler is freed (typically at program exit). + +## Output Formats + +### Console summary (stdout) + +The default when `--profile` is used without `--profile-output`, or `GGML_PROFILE=1`: + +``` +=== Profiling Summary === + [OP ] backend 0 MUL_MAT 45.2% count=1200 total= 120.50 ms avg= 100.42 us ... 12.30 GB/s [4096 x 4096] + [OP ] backend 1 MUL_MAT_ID 30.1% count= 600 total= 80.20 ms avg= 133.67 us ... 0.08 GB/s [2688 x 1856 x 128] + [COPY] backend 0 copy_H2D 5.3% count= 200 total= 14.10 ms avg= 70.50 us ... 2.50 GB/s + ... +``` + +Each line shows: event type (OP or COPY), backend index, operation name, percentage of total time, call count, timing stats, bandwidth, and representative tensor shape. + +### Plain text (`.txt`) + +A more detailed report with three sections: + +1. **Profiling Summary** — total time, record count, unique ops +2. **Per-Backend Summary** — ops and copies per backend with aggregate bandwidth +3. **Operations table** — full breakdown with bandwidth and tensor shapes for all source tensors + +### JSON (`.json`) + +Machine-readable format suitable for the Python analysis tool. Contains: + +- `version`: Format version (currently `2`) +- `backends[]`: Backend metadata (name, device, device type) +- `records[]`: Every profiling event with: + - `type`: `0` = OP, `1` = COPY + - `name`: Operation name (e.g. `"MUL_MAT"`, `"copy_H2D"`) + - `backend_id`, `split_id`: Scheduler indices + - `start_ns`, `duration_ns`: Timing in nanoseconds + - `bytes`: Output tensor size (OPs) or transfer size (COPYs) + - `extra`: Fusion name for fused ops, or `null` + - `ne_src0`, `ne_src1`, `ne_src2`: Source tensor dimensions (4-element arrays) + +`ne_src2` is populated only for `MUL_MAT_ID` (expert selection indices); it is `[0,0,0,0]` for all other ops. + +## Python Analysis Tool + +The `tools/profiler/profiler.py` script reads JSON exports and produces analysis reports and visualizations. + +### Basic usage + +```bash +# Print summary +python -m tools.profiler.profiler profile.json + +# Show top 10 operations by time +python -m tools.profiler.profiler profile.json --top-ops 10 + +# Show top 10 longest individual kernels +python -m tools.profiler.profiler profile.json --top-kernels 10 + +# Show inefficiency ranking (highest time-per-byte) +python -m tools.profiler.profiler profile.json --inefficiency +``` + +### Export visualizations + +```bash +# Interactive HTML timeline (self-contained, no dependencies) +python -m tools.profiler.profiler profile.json --html-viewer timeline.html + +# Chrome Trace format (open in chrome://tracing or Perfetto) +python -m tools.profiler.profiler profile.json --chrome-trace trace.json + +# Downsample large traces for the HTML viewer +python -m tools.profiler.profiler profile.json --html-viewer timeline.html --html-max-records 50000 +``` + +Multiple exports can be combined in a single invocation: + +```bash +python -m tools.profiler.profiler profile.json --html-viewer timeline.html --chrome-trace trace.json --top-ops 20 +``` + +### CLI reference + +| Argument | Description | +|----------|-------------| +| `profile` (positional) | Path to profiler JSON file | +| `--chrome-trace FILE` | Export Chrome Trace Event format | +| `--html-viewer FILE` | Export interactive HTML timeline | +| `--html-max-records N` | Limit records in HTML output (0 = unlimited) | +| `--top-ops N` | Show top N operations by total time | +| `--top-kernels N` | Show top N longest individual kernels | +| `--inefficiency` | Rank operations by time per byte (higher = worse) | + +### HTML viewer features + +The HTML viewer is a self-contained file with no external dependencies: + +- **Canvas timeline** with per-backend lanes and color-coded operations +- **Zoom controls** (1s / 100ms / 1ms / 100us) and mouse drag navigation +- **Minimap** showing the full trace with a viewport indicator +- **Hover tooltips** with operation name, duration, shape, and bytes +- **Stats table** with collapsible tree: Operation → Backend → Tensor shape, showing % time, count, avg/min/max, and bandwidth +- **Legend** showing the most frequent operation types + +## What Gets Measured + +### OP events + +Every tensor operation (MUL_MAT, ADD, UNARY, FLASH_ATTN_EXT, etc.) is recorded with: + +- **Timing**: Start/end timestamps (nanosecond precision) +- **Bytes**: Output tensor size (`ggml_nbytes(node)`) +- **Tensor shapes**: Dimensions of `src[0]`, `src[1]`, and `src[2]` (when applicable) +- **Bandwidth**: Computed as `bytes / duration` — useful for identifying memory-bound vs compute-bound operations + +### COPY events + +Data transfers between backends: + +- **Direction**: `copy_H2D` (host→device), `copy_D2H` (device→host), `copy_D2D` (device→device) +- **Bytes**: Exact transfer size +- **Bandwidth**: Transfer throughput + +### MoE weight copies + +When `--cpu-moe` is used, the scheduler selectively copies only the active experts. These partial copies are recorded as individual COPY events with the actual bytes transferred. + +## Programmatic API + +For custom applications, the profiler can be controlled through the C API defined in `ggml/include/ggml-profiler.h`: + +```c +// Enable profiling on a scheduler +ggml_backend_sched_set_profiling(sched, true); + +// ... run inference ... + +// Get raw records +const ggml_profile_record * records; +int n = ggml_backend_sched_get_profiling_records(sched, &records); + +// Or export directly +ggml_backend_sched_print_profiling(sched); // stdout +ggml_backend_sched_export_profiling_json(sched, "profile.json"); // JSON file +ggml_backend_sched_export_profiling_text(sched, "profile.txt"); // text file +ggml_backend_sched_write_profiling_json(sched, fp); // JSON to FILE* +ggml_backend_sched_write_profiling_text(sched, fp); // text to FILE* + +// Reset for next measurement window +ggml_backend_sched_reset_profiling(sched); +``` + +Records accumulate across multiple `graph_compute` calls until explicitly reset or the scheduler is freed. + +## Tips + +- **Prompt eval vs generation**: The profiler captures all graph computes. During prompt evaluation you'll see larger batch sizes in tensor shapes; during generation, batch size is typically 1-2. +- **Vulkan concurrent mode**: When Vulkan dispatches multiple operations concurrently, they are reported as a single combined record spanning the full GPU time interval. +- **Bandwidth interpretation**: For compute ops, bandwidth = `output_bytes / duration`. This is not memory bandwidth — it's a proxy for throughput. MUL_MAT with low bandwidth typically indicates compute-bound behavior; high bandwidth indicates memory-bound. +- **Large traces**: For long inference runs, the JSON can be large. Use `--html-max-records` to downsample the HTML viewer, or use Chrome Trace format which handles large files well. +- **Multiple backends**: Backend IDs in the output correspond to the scheduler's priority order (0 = highest priority, typically GPU; last = CPU). diff --git a/examples/debug/debug.cpp b/examples/debug/debug.cpp index 761e7a2db5..7f8a46fc30 100644 --- a/examples/debug/debug.cpp +++ b/examples/debug/debug.cpp @@ -252,6 +252,29 @@ int main(int argc, char ** argv) { return 1; } + // Export profiling data if profiling was enabled + if (params.profiling) { + ggml_backend_sched_t sched = llama_context_get_sched(ctx); + if (sched != nullptr) { + if (params.profiling_output.empty()) { + ggml_backend_sched_print_profiling(sched); + } else { + const std::string & path = params.profiling_output; + int ret; + if (path.size() >= 4 && path.compare(path.size() - 4, 4, ".txt") == 0) { + ret = ggml_backend_sched_export_profiling_text(sched, path.c_str()); + } else { + ret = ggml_backend_sched_export_profiling_json(sched, path.c_str()); + } + if (ret == 0) { + LOG("\nProfiling data exported to: %s\n", path.c_str()); + } else { + LOG_ERR("\nFailed to export profiling data to: %s\n", path.c_str()); + } + } + } + } + LOG("\n"); llama_perf_context_print(ctx); diff --git a/ggml/include/ggml-cpu.h b/ggml/include/ggml-cpu.h index dc6453c6ea..4f1e652939 100644 --- a/ggml/include/ggml-cpu.h +++ b/ggml/include/ggml-cpu.h @@ -22,6 +22,24 @@ extern "C" { // use only reference implementations bool use_ref; + + // profiler context (set by backend when profiling is enabled, NULL otherwise) + // when non-NULL, the compute loop will record per-node timing + void * profiling_context; + + // callback for recording a profile record from C code (set by backend when profiling) + // The callback receives the full tensor node so it can extract all sources, types, + // op_params, and sub-op information directly. + // params: context, type, name, split_id, start_ns, end_ns, bytes, extra, node + void (*profiling_record_fn)(void * context, + int type, + const char * name, + int split_id, + uint64_t start_ns, + uint64_t end_ns, + uint64_t bytes, + const char * extra, + const struct ggml_tensor * node); }; // numa strategies diff --git a/ggml/include/ggml-profiler.h b/ggml/include/ggml-profiler.h new file mode 100644 index 0000000000..d802dba498 --- /dev/null +++ b/ggml/include/ggml-profiler.h @@ -0,0 +1,134 @@ +#pragma once + +#include "ggml-backend.h" +#include "ggml.h" + +#include + +#ifdef __cplusplus +extern "C" { +#endif + +// +// Profiler +// + +// Profile event types +enum ggml_profile_event_type { + GGML_PROFILE_EVENT_OP, // single operation execution (computation kernel) + GGML_PROFILE_EVENT_COPY, // data transfer between devices +}; + +// A single profiling record representing a timed interval +typedef struct ggml_profile_record { + enum ggml_profile_event_type type; + const char * name; // operation name (e.g., "mul_mat", "copy_H2D") + int backend_id; // scheduler's backend index (0 = highest priority) + int split_id; // which graph split (0..n_splits-1) + uint64_t start_ns; // start timestamp in nanoseconds + uint64_t end_ns; // end timestamp in nanoseconds + uint64_t bytes; // bytes transferred (for copy) or tensor size (for ops) + const char * extra; // fusion name for fused ops, or NULL + + // Output tensor info + char tensor_name[GGML_MAX_NAME]; // output tensor name (e.g. "ffn_out-0"), "" if unnamed + int64_t ne[4]; // output tensor dimensions + int out_type; // output tensor type (ggml_type), -1 if N/A + + // Source tensors (up to GGML_MAX_SRC). n_src is the actual number populated. + int n_src; + int64_t ne_src[GGML_MAX_SRC][4]; // per-source dimensions + int64_t nb_src[GGML_MAX_SRC][4]; // per-source strides (bytes) + int type_src[GGML_MAX_SRC]; // per-source ggml_type, -1 if not present + + // Operation parameters (raw bytes copied from ggml_tensor::op_params) + int32_t op_params[GGML_MAX_OP_PARAMS / sizeof(int32_t)]; + + int sub_op; // sub-operation (ggml_unary_op or ggml_glu_op), -1 if N/A +} ggml_profile_record; + +// Backend profiler interface - each backend optionally implements this +// to provide fine-grained operation timing +struct ggml_backend_profiler { + void * context; // backend-specific profiler context + + // Enable or disable profiling on this backend + void (*enable)(void * context, bool enable); + + // Clear all recorded data + void (*reset)(void * context); + + // Set the current split ID (called by scheduler before graph_compute) + void (*set_split_id)(void * context, int split_id); + + // Get recorded profiling data + // Returns the number of records; sets *out to point to internal storage + // The returned pointer remains valid until the next reset or disable call + int (*get_records)(void * context, const ggml_profile_record ** out); + + // Free the profiler context + void (*free_context)(void * context); +}; + +typedef struct ggml_backend_profiler * ggml_backend_profiler_t; + +// Populate the per-node fields of a ggml_profile_record from a ggml_tensor node: +// ne, out_type, n_src, ne_src, nb_src, type_src, op_params, sub_op. +// All other fields (type/name/backend_id/split_id/timestamps/bytes/extra) must +// be filled in separately by the backend that records the event. +GGML_API void ggml_profile_record_from_tensor(struct ggml_profile_record * rec, + const struct ggml_tensor * node); + +// Register a profiler on a backend (called by backend during init) +// The profiler is owned by the backend and will be freed when the backend is freed +GGML_API void ggml_backend_set_profiler(ggml_backend_t backend, ggml_backend_profiler_t profiler); + +// Get the profiler associated with a backend (returns NULL if none) +GGML_API ggml_backend_profiler_t ggml_backend_get_profiler(ggml_backend_t backend); + +// +// Scheduler profiling API +// + +// Enable or disable profiling on a scheduler +// When enabled, the scheduler will: +// - Time data copy operations between backends +// - Enable profiling on all backends that support it +// - Collect profiling records from all backends after each graph compute +GGML_API void ggml_backend_sched_set_profiling(ggml_backend_sched_t sched, bool enable); + +// Check if profiling is enabled on a scheduler +GGML_API bool ggml_backend_sched_get_profiling(ggml_backend_sched_t sched); + +// Get profiling data from the last graph compute +// Records are owned by the scheduler; valid until the next compute or reset +// Returns the number of records +GGML_API int ggml_backend_sched_get_profiling_records(ggml_backend_sched_t sched, const ggml_profile_record ** records); + +// Print a human-readable summary of the last profiling run to stdout +// Groups records by operation name and shows total/count/min/max/avg time +GGML_API void ggml_backend_sched_print_profiling(ggml_backend_sched_t sched); + +// Reset profiling data (clear all recorded data) +GGML_API void ggml_backend_sched_reset_profiling(ggml_backend_sched_t sched); + +// Get current time in nanoseconds (for manual profiling if needed) +GGML_API uint64_t ggml_profiler_time_ns(void); + +// Export profiling data as JSON to a file +// Returns 0 on success, -1 on error +GGML_API int ggml_backend_sched_export_profiling_json(ggml_backend_sched_t sched, const char * filepath); + +// Export profiling data as JSON to a FILE pointer +GGML_API int ggml_backend_sched_write_profiling_json(ggml_backend_sched_t sched, FILE * fp); + +// Export profiling data as plain text statistics to a file +// Returns 0 on success, -1 on error +GGML_API int ggml_backend_sched_export_profiling_text(ggml_backend_sched_t sched, const char * filepath); + +// Export profiling data as plain text statistics to a FILE pointer +GGML_API int ggml_backend_sched_write_profiling_text(ggml_backend_sched_t sched, FILE * fp); + +#ifdef __cplusplus +} +#endif diff --git a/ggml/src/CMakeLists.txt b/ggml/src/CMakeLists.txt index 9477320002..7cf26645eb 100644 --- a/ggml/src/CMakeLists.txt +++ b/ggml/src/CMakeLists.txt @@ -195,6 +195,7 @@ add_library(ggml-base ../include/ggml-backend.h ../include/ggml-cpp.h ../include/ggml-opt.h + ../include/ggml-profiler.h ../include/gguf.h ggml.c ggml.cpp @@ -202,6 +203,7 @@ add_library(ggml-base ggml-backend.cpp ggml-backend-meta.cpp ggml-opt.cpp + ggml-profiler.cpp ggml-threading.cpp ggml-threading.h ggml-quants.c diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index ef05905cf9..dd9ce5a13c 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -3,6 +3,7 @@ // ggml-backend internal header #include "ggml-backend.h" +#include "ggml-profiler.h" #ifdef __cplusplus extern "C" { @@ -160,6 +161,9 @@ extern "C" { struct ggml_backend_i iface; ggml_backend_dev_t device; void * context; + + // Optional profiler (set by backend during init, NULL if not profiling) + ggml_backend_profiler_t profiler; }; struct ggml_backend_event { diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 40e50c5c9d..71d825a5ad 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -12,6 +12,7 @@ #include "ggml-backend-impl.h" #include "ggml-alloc.h" #include "ggml-impl.h" +#include "ggml-profiler.h" #include #include @@ -20,6 +21,7 @@ #include #include #include +#include #include #include @@ -242,6 +244,15 @@ void ggml_backend_free(ggml_backend_t backend) { return; } + // Clean up profiler if present (before backend frees its context) + if (backend->profiler != NULL) { + if (backend->profiler->free_context != NULL) { + backend->profiler->free_context(backend->profiler->context); + } + delete backend->profiler; + backend->profiler = NULL; + } + backend->iface.free(backend); } @@ -838,6 +849,20 @@ struct ggml_backend_sched { int debug_realloc; int debug_graph_size; int debug_prev_graph_size; + + // profiling + bool profiling_enabled; + std::string profiling_env_path; // GGML_PROFILE env var value (for auto-export on free) + std::vector copy_records; // copy events recorded by the scheduler + std::vector profiling_records; // merged records from all sources + + // Cached backend metadata for safe access during auto-export (backends may be freed first) + struct backend_meta { + std::string name; + std::string device; + int device_type; + }; + std::vector profiling_backend_meta; }; #define hash_id(tensor) ggml_hash_find_or_insert(&sched->hash_set, tensor) @@ -1640,6 +1665,39 @@ static bool ggml_backend_sched_alloc_splits(ggml_backend_sched_t sched) { return true; } +// Build a COPY profiling record. Copies have no real ggml_tensor "node" backing +// them, so we synthesize one source describing the input tensor that was moved. +static ggml_profile_record make_copy_record(const char * copy_dir, int backend_id, int split_id, + uint64_t start_ns, uint64_t end_ns, uint64_t bytes, + const struct ggml_tensor * input) { + ggml_profile_record rec = {}; + rec.type = GGML_PROFILE_EVENT_COPY; + rec.name = copy_dir; + rec.backend_id = backend_id; + rec.split_id = split_id; + rec.start_ns = start_ns; + rec.end_ns = end_ns; + rec.bytes = bytes; + rec.extra = input ? input->name : NULL; + snprintf(rec.tensor_name, sizeof(rec.tensor_name), "%s", input ? input->name : ""); + rec.out_type = -1; + rec.sub_op = -1; + rec.n_src = 0; + if (input != NULL) { + // Describe the input tensor as src[0] so consumers can inspect its shape. + rec.n_src = 1; + memcpy(rec.ne_src[0], input->ne, sizeof(rec.ne_src[0])); + for (int d = 0; d < 4; d++) { + rec.nb_src[0][d] = (int64_t) input->nb[d]; + } + rec.type_src[0] = (int) input->type; + } + for (int i = rec.n_src; i < GGML_MAX_SRC; i++) { + rec.type_src[i] = -1; + } + return rec; +} + static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t sched) { GGML_ASSERT(sched); struct ggml_backend_sched_split * splits = sched->splits; @@ -1650,6 +1708,11 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s int prev_backend_id = -1; + // Profiling: reset copy records for this compute pass + if (sched->profiling_enabled) { + sched->copy_records.clear(); + } + for (int split_id = 0; split_id < sched->n_splits; split_id++) { struct ggml_backend_sched_split * split = &splits[split_id]; int split_backend_id = split->backend_id; @@ -1665,6 +1728,18 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } + // Profiling: set split ID and enable backend profiling + if (sched->profiling_enabled) { + if (split_backend->profiler != NULL) { + if (split_backend->profiler->enable != NULL) { + split_backend->profiler->enable(split_backend->profiler->context, true); + } + if (split_backend->profiler->set_split_id != NULL) { + split_backend->profiler->set_split_id(split_backend->profiler->context, split_id); + } + } + } + // copy the input tensors to the split backend for (int input_id = 0; input_id < split->n_inputs; input_id++) { ggml_backend_t input_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[input_id]); @@ -1678,7 +1753,25 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } else { ggml_backend_synchronize(split_backend); } - ggml_backend_tensor_copy(input, input_cpy); + if (sched->profiling_enabled) { + uint64_t copy_start = ggml_profiler_time_ns(); + ggml_backend_tensor_copy(input, input_cpy); + uint64_t copy_end = ggml_profiler_time_ns(); + + enum ggml_backend_dev_type src_type = ggml_backend_dev_type(input_backend->device); + enum ggml_backend_dev_type dst_type = ggml_backend_dev_type(split_backend->device); + const char * copy_dir = "copy_D2D"; + if (src_type == GGML_BACKEND_DEVICE_TYPE_CPU && dst_type != GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_H2D"; + } else if (src_type != GGML_BACKEND_DEVICE_TYPE_CPU && dst_type == GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_D2H"; + } + + sched->copy_records.push_back(make_copy_record(copy_dir, split_backend_id, split_id, + copy_start, copy_end, ggml_nbytes(input), input)); + } else { + ggml_backend_tensor_copy(input, input_cpy); + } } else { // wait for the split backend to finish using the input before overwriting it if (sched->events[split_backend_id][sched->cur_copy] != NULL) { @@ -1735,12 +1828,14 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } // group consecutive experts and copy them together + size_t total_copied_bytes = 0; auto copy_experts = [&](int32_t first_id, int32_t last_id) { const size_t expert_offset = first_id * expert_size; const size_t expert_size_copy = (last_id - first_id + 1) * expert_size; const size_t padding = std::min(expert_size, 512); const size_t padding_end = last_id < n_expert - 1 ? padding : 0; + total_copied_bytes += expert_size_copy + padding_end; ggml_backend_tensor_set_async(split_backend, input_cpy, (const uint8_t *)input->data + expert_offset, expert_offset, @@ -1749,6 +1844,11 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s expert_size_copy + padding_end); }; + uint64_t moe_copy_start = 0; + if (sched->profiling_enabled) { + moe_copy_start = ggml_profiler_time_ns(); + } + int id = 0; while (!ggml_bitset_get(used_ids.data(), id)) { id++; @@ -1772,9 +1872,34 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s last_id = id; } copy_experts(first_id, last_id); + + if (sched->profiling_enabled) { + uint64_t moe_copy_end = ggml_profiler_time_ns(); + + enum ggml_backend_dev_type src_type = ggml_backend_dev_type(input_backend->device); + enum ggml_backend_dev_type dst_type = ggml_backend_dev_type(split_backend->device); + const char * copy_dir = "copy_D2D"; + if (src_type == GGML_BACKEND_DEVICE_TYPE_CPU && dst_type != GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_H2D"; + } else if (src_type != GGML_BACKEND_DEVICE_TYPE_CPU && + dst_type == GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_D2H"; + } + + sched->copy_records.push_back(make_copy_record(copy_dir, split_backend_id, split_id, + moe_copy_start, moe_copy_end, + (uint64_t) total_copied_bytes, input)); + } } else { // try async copy, but if not possible, we can still use a sync copy without synchronizing the dst backend, since we handle the synchronization here with multiple copies and events // TODO: add public function to facilitate this, since applications do not have direct access to the backend interface + + // Capture timestamp before async attempt so we can record launch time + uint64_t copy_start = 0; + if (sched->profiling_enabled) { + copy_start = ggml_profiler_time_ns(); + } + if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) { ggml_backend_synchronize(input_backend); if (sched->events[split_backend_id][sched->cur_copy] != NULL) { @@ -1782,7 +1907,45 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } else { ggml_backend_synchronize(split_backend); } - ggml_backend_tensor_copy(input, input_cpy); + if (sched->profiling_enabled) { + // Re-take start after sync for accurate sync copy measurement + copy_start = ggml_profiler_time_ns(); + ggml_backend_tensor_copy(input, input_cpy); + uint64_t copy_end = ggml_profiler_time_ns(); + + enum ggml_backend_dev_type src_type = ggml_backend_dev_type(input_backend->device); + enum ggml_backend_dev_type dst_type = ggml_backend_dev_type(split_backend->device); + const char * copy_dir = "copy_D2D"; + if (src_type == GGML_BACKEND_DEVICE_TYPE_CPU && dst_type != GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_H2D"; + } else if (src_type != GGML_BACKEND_DEVICE_TYPE_CPU && + dst_type == GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_D2H"; + } + + sched->copy_records.push_back(make_copy_record(copy_dir, split_backend_id, split_id, + copy_start, copy_end, ggml_nbytes(input), input)); + } else { + ggml_backend_tensor_copy(input, input_cpy); + } + } else { + // async copy was launched — record the time spanning the async call + if (sched->profiling_enabled) { + uint64_t copy_end = ggml_profiler_time_ns(); + + enum ggml_backend_dev_type src_type = ggml_backend_dev_type(input_backend->device); + enum ggml_backend_dev_type dst_type = ggml_backend_dev_type(split_backend->device); + const char * copy_dir = "copy_D2D"; + if (src_type == GGML_BACKEND_DEVICE_TYPE_CPU && dst_type != GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_H2D"; + } else if (src_type != GGML_BACKEND_DEVICE_TYPE_CPU && + dst_type == GGML_BACKEND_DEVICE_TYPE_CPU) { + copy_dir = "copy_D2H"; + } + + sched->copy_records.push_back(make_copy_record(copy_dir, split_backend_id, split_id, + copy_start, copy_end, ggml_nbytes(input), input)); + } } } } @@ -1835,6 +1998,32 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s prev_backend_id = split_backend_id; } + // Profiling: collect records from all backends and append to accumulated records + if (sched->profiling_enabled) { + // Collect backend operation records + for (int b = 0; b < sched->n_backends; b++) { + ggml_backend_t backend = sched->backends[b]; + if (backend->profiler != NULL && backend->profiler->get_records != NULL) { + const ggml_profile_record * backend_recs = NULL; + int count = backend->profiler->get_records(backend->profiler->context, &backend_recs); + for (int r = 0; r < count; r++) { + ggml_profile_record rec = backend_recs[r]; + rec.backend_id = b; // stamp correct scheduler backend index + sched->profiling_records.push_back(rec); + } + // Reset backend records (but keep profiling enabled for next compute) + if (backend->profiler->reset != NULL) { + backend->profiler->reset(backend->profiler->context); + } + } + } + + // Append copy records + for (const auto & rec : sched->copy_records) { + sched->profiling_records.push_back(rec); + } + } + return GGML_STATUS_SUCCESS; } @@ -1904,6 +2093,24 @@ ggml_backend_sched_t ggml_backend_sched_new( sched->galloc = ggml_gallocr_new_n(sched->bufts, n_backends); sched->op_offload = op_offload; + const char * profile_env = getenv("GGML_PROFILE"); + if (profile_env != NULL) { + sched->profiling_enabled = true; + sched->profiling_env_path = profile_env; + } + + // Cache backend metadata for safe access during auto-export + for (int b = 0; b < n_backends; b++) { + ggml_backend_sched::backend_meta meta; + meta.name = ggml_backend_name(backends[b]); + meta.device = "unknown"; + meta.device_type = 0; + if (backends[b]->device != NULL) { + meta.device = ggml_backend_dev_name(backends[b]->device); + meta.device_type = (int) ggml_backend_dev_type(backends[b]->device); + } + sched->profiling_backend_meta.push_back(std::move(meta)); + } ggml_backend_sched_reset(sched); @@ -1914,6 +2121,33 @@ void ggml_backend_sched_free(ggml_backend_sched_t sched) { if (sched == NULL) { return; } + + // Auto-export profiling data if enabled via GGML_PROFILE env var + // GGML_PROFILE=1 or GGML_PROFILE="" → print to stdout + // GGML_PROFILE=file.json → export JSON + // GGML_PROFILE=file.txt → export text + if (!sched->profiling_records.empty() && getenv("GGML_PROFILE") != NULL) { + const std::string & path = sched->profiling_env_path; + if (path.empty() || path == "1" || path == "stdout") { + ggml_backend_sched_print_profiling(sched); + } else if (path.size() >= 4 && path.compare(path.size() - 4, 4, ".txt") == 0) { + int ret = ggml_backend_sched_export_profiling_text(sched, path.c_str()); + if (ret == 0) { + GGML_LOG_INFO("[profiler] Data exported to: %s\n", path.c_str()); + } else { + GGML_LOG_ERROR("[profiler] Failed to export data to: %s\n", path.c_str()); + } + } else { + // Default to JSON for any other path (including .json) + int ret = ggml_backend_sched_export_profiling_json(sched, path.c_str()); + if (ret == 0) { + GGML_LOG_INFO("[profiler] Data exported to: %s\n", path.c_str()); + } else { + GGML_LOG_ERROR("[profiler] Failed to export data to: %s\n", path.c_str()); + } + } + } + for (int b = 0; b < sched->n_backends; b++) { for (int c = 0; c < sched->n_copies; c++) { ggml_backend_event_free(sched->events[b][c]); @@ -2504,3 +2738,509 @@ ggml_backend_buffer_t ggml_backend_cpu_buffer_from_ptr(void * ptr, size_t size) GGML_ASSERT((uintptr_t)ptr % TENSOR_ALIGNMENT == 0 && "buffer pointer must be aligned"); return ggml_backend_buffer_init(ggml_backend_cpu_buffer_from_ptr_type(), ggml_backend_cpu_buffer_from_ptr_i, ptr, size); } + +// +// Scheduler profiling +// + +void ggml_backend_sched_set_profiling(ggml_backend_sched_t sched, bool enable) { + GGML_ASSERT(sched); + sched->profiling_enabled = enable; + + if (!enable) { + ggml_backend_sched_reset_profiling(sched); + } +} + +bool ggml_backend_sched_get_profiling(ggml_backend_sched_t sched) { + GGML_ASSERT(sched); + return sched->profiling_enabled; +} + +int ggml_backend_sched_get_profiling_records(ggml_backend_sched_t sched, const ggml_profile_record ** records) { + GGML_ASSERT(sched); + GGML_ASSERT(records != NULL); + + *records = sched->profiling_records.data(); + return (int) sched->profiling_records.size(); +} + +void ggml_backend_sched_reset_profiling(ggml_backend_sched_t sched) { + GGML_ASSERT(sched); + sched->profiling_records.clear(); + sched->copy_records.clear(); +} + +void ggml_backend_sched_print_profiling(ggml_backend_sched_t sched) { + GGML_ASSERT(sched); + + if (sched->profiling_records.empty()) { + GGML_LOG_INFO("[profiler] No profiling data available\n"); + return; + } + + GGML_LOG_INFO("\n=== Profiling Summary ===\n"); + + // Aggregate by (name, type, backend_id) + struct op_stats { + const char * name; + enum ggml_profile_event_type type; + int backend_id; + uint64_t total_ns; + uint64_t min_ns; + uint64_t max_ns; + int count; + uint64_t total_bytes; + int64_t representative_ne[4]; + }; + + std::vector stats; + for (const auto & rec : sched->profiling_records) { + bool found = false; + for (auto & s : stats) { + if (s.type == rec.type && s.backend_id == rec.backend_id && strcmp(s.name, rec.name) == 0) { + uint64_t dur = (rec.end_ns > rec.start_ns) ? (rec.end_ns - rec.start_ns) : 0; + s.total_ns += dur; + s.min_ns = std::min(s.min_ns, dur); + s.max_ns = std::max(s.max_ns, dur); + s.count++; + s.total_bytes += rec.bytes; + found = true; + break; + } + } + if (!found) { + uint64_t dur = (rec.end_ns > rec.start_ns) ? (rec.end_ns - rec.start_ns) : 0; + op_stats s; + s.name = rec.name; + s.type = rec.type; + s.backend_id = rec.backend_id; + s.total_ns = dur; + s.min_ns = dur; + s.max_ns = dur; + s.count = 1; + s.total_bytes = rec.bytes; + memcpy(s.representative_ne, rec.ne_src[0], sizeof(s.representative_ne)); + stats.push_back(s); + } + } + + // Sort by total time descending + std::sort(stats.begin(), stats.end(), + [](const op_stats & a, const op_stats & b) { return a.total_ns > b.total_ns; }); + + uint64_t grand_total = 0; + for (const auto & s : stats) { + grand_total += s.total_ns; + } + + const char * type_str[] = { "OP ", "COPY" }; + for (const auto & s : stats) { + double pct = 100.0 * (double) s.total_ns / (double) grand_total; + double avg_us = (double) s.total_ns / (double) s.count / 1000.0; + double min_us = (double) s.min_ns / 1000.0; + double max_us = (double) s.max_ns / 1000.0; + + GGML_LOG_INFO( + " [%s] backend %d %-28s %7.1f%% count=%-6d total=%8.2f ms avg=%8.2f us min=%8.2f us max=%8.2f us", + type_str[s.type], s.backend_id, s.name, pct, s.count, (double) s.total_ns / 1e6, avg_us, min_us, max_us); + if (s.total_bytes > 0 && s.total_ns > 0) { + double bw_gbps = (double) s.total_bytes / (double) s.total_ns; + if (bw_gbps >= 1000.0) { + GGML_LOG_INFO(" %6.2f TB/s", bw_gbps / 1000.0); + } else { + GGML_LOG_INFO(" %6.2f GB/s", bw_gbps); + } + } + // Print representative tensor shape (first record's ne) + if (s.representative_ne[0] > 0 || s.representative_ne[1] > 0) { + GGML_LOG_INFO(" [%lld x %lld", (long long) s.representative_ne[0], (long long) s.representative_ne[1]); + if (s.representative_ne[2] > 1) { + GGML_LOG_INFO(" x %lld", (long long) s.representative_ne[2]); + } + if (s.representative_ne[3] > 1) { + GGML_LOG_INFO(" x %lld", (long long) s.representative_ne[3]); + } + GGML_LOG_INFO("]"); + } + GGML_LOG_INFO("\n"); + } + + GGML_LOG_INFO(" ---\n"); + GGML_LOG_INFO(" Total: %.2f ms (%d records, %d unique ops)\n\n", (double) grand_total / 1e6, + (int) sched->profiling_records.size(), (int) stats.size()); +} + +int ggml_backend_sched_write_profiling_json(ggml_backend_sched_t sched, FILE * fp) { + GGML_ASSERT(sched); + GGML_ASSERT(fp != NULL); + + uint64_t total_ns = 0; + for (const auto & rec : sched->profiling_records) { + total_ns += (rec.end_ns > rec.start_ns) ? (rec.end_ns - rec.start_ns) : 0; + } + + fprintf(fp, "{\n"); + fprintf(fp, " \"version\": 3,\n"); + fprintf(fp, " \"profiler\": \"ggml\",\n"); + fprintf(fp, " \"total_records\": %d,\n", (int) sched->profiling_records.size()); + fprintf(fp, " \"total_ns\": %llu,\n", (unsigned long long) total_ns); + + // Backend metadata (use cached data if available, fall back to live pointers) + fprintf(fp, " \"backends\": [\n"); + for (int b = 0; b < sched->n_backends; b++) { + const char * name = "unknown"; + const char * dev_name = "unknown"; + int dev_type = 0; + if (b < (int) sched->profiling_backend_meta.size()) { + name = sched->profiling_backend_meta[b].name.c_str(); + dev_name = sched->profiling_backend_meta[b].device.c_str(); + dev_type = sched->profiling_backend_meta[b].device_type; + } else if (sched->backends[b] != NULL) { + name = ggml_backend_name(sched->backends[b]); + if (sched->backends[b]->device != NULL) { + dev_name = ggml_backend_dev_name(sched->backends[b]->device); + dev_type = (int) ggml_backend_dev_type(sched->backends[b]->device); + } + } + fprintf(fp, " {\"id\": %d, \"name\": \"%s\", \"device\": \"%s\", \"device_type\": %d}%s\n", b, name, + dev_name, dev_type, (b < sched->n_backends - 1) ? "," : ""); + } + fprintf(fp, " ],\n"); + + // Records + fprintf(fp, " \"records\": [\n"); + + for (int i = 0; i < (int) sched->profiling_records.size(); i++) { + const auto & rec = sched->profiling_records[i]; + uint64_t duration_ns = (rec.end_ns > rec.start_ns) ? (rec.end_ns - rec.start_ns) : 0; + + fprintf(fp, + " {\"type\": %d, \"name\": \"%s\", \"backend_id\": %d, \"split_id\": %d, " + "\"start_ns\": %llu, \"duration_ns\": %llu, \"bytes\": %llu, \"extra\": ", + (int) rec.type, rec.name ? rec.name : "unknown", rec.backend_id, rec.split_id, + (unsigned long long) rec.start_ns, (unsigned long long) duration_ns, (unsigned long long) rec.bytes); + + if (rec.extra != NULL) { + fprintf(fp, "\"%s\"", rec.extra); + } else { + fprintf(fp, "null"); + } + + // Output tensor info + if (rec.tensor_name[0] != '\0') { + fprintf(fp, ", \"tensor_name\": \"%s\"", rec.tensor_name); + } else { + fprintf(fp, ", \"tensor_name\": null"); + } + fprintf(fp, ", \"ne\": [%lld, %lld, %lld, %lld]", (long long) rec.ne[0], (long long) rec.ne[1], + (long long) rec.ne[2], (long long) rec.ne[3]); + fprintf(fp, ", \"out_type\": %d", rec.out_type); + + // Source tensors + fprintf(fp, ", \"n_src\": %d", rec.n_src); + fprintf(fp, ", \"ne_src\": ["); + for (int s = 0; s < rec.n_src; s++) { + fprintf(fp, "%s[%lld, %lld, %lld, %lld]", s == 0 ? "" : ", ", + (long long) rec.ne_src[s][0], (long long) rec.ne_src[s][1], + (long long) rec.ne_src[s][2], (long long) rec.ne_src[s][3]); + } + fprintf(fp, "]"); + fprintf(fp, ", \"nb_src\": ["); + for (int s = 0; s < rec.n_src; s++) { + fprintf(fp, "%s[%lld, %lld, %lld, %lld]", s == 0 ? "" : ", ", + (long long) rec.nb_src[s][0], (long long) rec.nb_src[s][1], + (long long) rec.nb_src[s][2], (long long) rec.nb_src[s][3]); + } + fprintf(fp, "]"); + fprintf(fp, ", \"type_src\": ["); + for (int s = 0; s < rec.n_src; s++) { + fprintf(fp, "%s%d", s == 0 ? "" : ", ", rec.type_src[s]); + } + fprintf(fp, "]"); + + // op_params (full 16-int32 block, matching export-graph-ops format) + fprintf(fp, ", \"op_params\": ["); + const int n_op_params = (int) (sizeof(rec.op_params) / sizeof(rec.op_params[0])); + for (int p = 0; p < n_op_params; p++) { + fprintf(fp, "%s%d", p == 0 ? "" : ", ", rec.op_params[p]); + } + fprintf(fp, "]"); + + fprintf(fp, ", \"sub_op\": %d", rec.sub_op); + + fprintf(fp, "}%s\n", (i < (int) sched->profiling_records.size() - 1) ? "," : ""); + } + + fprintf(fp, " ]\n"); + fprintf(fp, "}\n"); + + return 0; +} + +int ggml_backend_sched_export_profiling_json(ggml_backend_sched_t sched, const char * filepath) { + GGML_ASSERT(sched); + GGML_ASSERT(filepath != NULL); + + FILE * fp = fopen(filepath, "w"); + if (fp == NULL) { + GGML_LOG_ERROR("%s: failed to open %s for writing\n", __func__, filepath); + return -1; + } + + int ret = ggml_backend_sched_write_profiling_json(sched, fp); + fclose(fp); + + return ret; +} + +// Helper: format ne dimensions as string, e.g. "[4096, 4096, 1]" +static void fmt_ne(char * buf, size_t bufsize, const int64_t ne[4]) { + if (ne[0] == 0 && ne[1] == 0 && ne[2] == 0 && ne[3] == 0) { + buf[0] = '\0'; + return; + } + int ndims = 4; + while (ndims > 1 && ne[ndims - 1] <= 1) { + ndims--; + } + int pos = snprintf(buf, bufsize, "["); + for (int i = 0; i < ndims && pos < (int) bufsize - 1; i++) { + pos += snprintf(buf + pos, bufsize - pos, "%s%lld", i > 0 ? ", " : "", (long long) ne[i]); + } + snprintf(buf + pos, bufsize - pos, "]"); +} + +// Helper: format bandwidth as string +static void fmt_bandwidth(char * buf, size_t bufsize, uint64_t bytes, uint64_t ns) { + if (ns == 0 || bytes == 0) { + buf[0] = '\0'; + return; + } + double bw_gbps = (double) bytes / (double) ns; + if (bw_gbps >= 1000.0) { + snprintf(buf, bufsize, "%.2f TB/s", bw_gbps / 1000.0); + } else { + snprintf(buf, bufsize, "%.2f GB/s", bw_gbps); + } +} + +int ggml_backend_sched_write_profiling_text(ggml_backend_sched_t sched, FILE * fp) { + GGML_ASSERT(sched); + GGML_ASSERT(fp != NULL); + + if (sched->profiling_records.empty()) { + fprintf(fp, "No profiling data available.\n"); + return 0; + } + + // Aggregate by (name, type, backend_id) + struct op_stats { + const char * name; + enum ggml_profile_event_type type; + int backend_id; + uint64_t total_ns; + uint64_t min_ns; + uint64_t max_ns; + int count; + uint64_t total_bytes; + int64_t representative_ne_src0[4]; + int64_t representative_ne_src1[4]; + int64_t representative_ne_src2[4]; + }; + + std::vector stats; + for (const auto & rec : sched->profiling_records) { + uint64_t dur = (rec.end_ns > rec.start_ns) ? (rec.end_ns - rec.start_ns) : 0; + bool found = false; + for (auto & s : stats) { + if (s.type == rec.type && s.backend_id == rec.backend_id && strcmp(s.name, rec.name) == 0) { + s.total_ns += dur; + s.min_ns = std::min(s.min_ns, dur); + s.max_ns = std::max(s.max_ns, dur); + s.count++; + s.total_bytes += rec.bytes; + found = true; + break; + } + } + if (!found) { + op_stats s = {}; + s.name = rec.name; + s.type = rec.type; + s.backend_id = rec.backend_id; + s.total_ns = dur; + s.min_ns = dur; + s.max_ns = dur; + s.count = 1; + s.total_bytes = rec.bytes; + memcpy(s.representative_ne_src0, rec.ne_src[0], sizeof(s.representative_ne_src0)); + memcpy(s.representative_ne_src1, rec.ne_src[1], sizeof(s.representative_ne_src1)); + memcpy(s.representative_ne_src2, rec.ne_src[2], sizeof(s.representative_ne_src2)); + stats.push_back(s); + } + } + + std::sort(stats.begin(), stats.end(), + [](const op_stats & a, const op_stats & b) { return a.total_ns > b.total_ns; }); + + uint64_t grand_total = 0; + for (const auto & s : stats) { + grand_total += s.total_ns; + } + + // --- Section 1: Overall summary --- + fprintf(fp, "=== Profiling Summary ===\n"); + fprintf(fp, "Total time: %.2f ms\n", (double) grand_total / 1e6); + fprintf(fp, "Total records: %d\n", (int) sched->profiling_records.size()); + fprintf(fp, "Unique ops: %d\n\n", (int) stats.size()); + + // --- Section 2: Per-backend breakdown --- + fprintf(fp, "=== Per-Backend Summary ===\n"); + { + struct backend_stats { + int backend_id; + int op_count; + int copy_count; + uint64_t op_ns; + uint64_t copy_ns; + uint64_t op_bytes; + uint64_t copy_bytes; + }; + std::vector bstats; + for (const auto & s : stats) { + bool found = false; + for (auto & bs : bstats) { + if (bs.backend_id == s.backend_id) { + if (s.type == GGML_PROFILE_EVENT_OP) { + bs.op_count += s.count; + bs.op_ns += s.total_ns; + bs.op_bytes += s.total_bytes; + } else { + bs.copy_count += s.count; + bs.copy_ns += s.total_ns; + bs.copy_bytes += s.total_bytes; + } + found = true; + break; + } + } + if (!found) { + backend_stats bs = {}; + bs.backend_id = s.backend_id; + if (s.type == GGML_PROFILE_EVENT_OP) { + bs.op_count = s.count; + bs.op_ns = s.total_ns; + bs.op_bytes = s.total_bytes; + } else { + bs.copy_count = s.count; + bs.copy_ns = s.total_ns; + bs.copy_bytes = s.total_bytes; + } + bstats.push_back(bs); + } + } + std::sort(bstats.begin(), bstats.end(), + [](const backend_stats & a, const backend_stats & b) { + return (a.op_ns + a.copy_ns) > (b.op_ns + b.copy_ns); + }); + + for (const auto & bs : bstats) { + uint64_t total = bs.op_ns + bs.copy_ns; + double pct = grand_total > 0 ? 100.0 * (double) total / (double) grand_total : 0; + + const char * bname = "unknown"; + if (bs.backend_id >= 0 && bs.backend_id < (int) sched->profiling_backend_meta.size()) { + bname = sched->profiling_backend_meta[bs.backend_id].name.c_str(); + } else if (bs.backend_id >= 0 && bs.backend_id < sched->n_backends && sched->backends[bs.backend_id] != NULL) { + bname = ggml_backend_name(sched->backends[bs.backend_id]); + } + + fprintf(fp, " Backend %d (%s): %.2f ms (%.1f%%)\n", bs.backend_id, bname, (double) total / 1e6, pct); + if (bs.op_count > 0) { + char bw_buf[32]; + fmt_bandwidth(bw_buf, sizeof(bw_buf), bs.op_bytes, bs.op_ns); + fprintf(fp, " OPs: %d calls, %.2f ms", bs.op_count, (double) bs.op_ns / 1e6); + if (bw_buf[0]) { + fprintf(fp, ", %s", bw_buf); + } + fprintf(fp, "\n"); + } + if (bs.copy_count > 0) { + char bw_buf[32]; + fmt_bandwidth(bw_buf, sizeof(bw_buf), bs.copy_bytes, bs.copy_ns); + fprintf(fp, " COPYs: %d calls, %.2f ms", bs.copy_count, (double) bs.copy_ns / 1e6); + if (bw_buf[0]) { + fprintf(fp, ", %s", bw_buf); + } + fprintf(fp, "\n"); + } + } + } + fprintf(fp, "\n"); + + // --- Section 3: Detailed operation table --- + fprintf(fp, "=== Operations (sorted by total time) ===\n"); + fprintf(fp, "%-5s %4s %-28s %7s %6s %10s %10s %10s %10s %12s %s\n", + "TYPE", "BKND", "Operation", "%Time", "Count", "Total(ms)", "Avg(us)", "Min(us)", "Max(us)", "Bandwidth", "Tensors"); + fprintf(fp, "%-5s %4s %-28s %7s %6s %10s %10s %10s %10s %12s %s\n", + "-----", "----", "----------------------------", "-------", "------", + "----------", "----------", "----------", "----------", "------------", "-------"); + + const char * type_str[] = { "OP", "COPY" }; + for (const auto & s : stats) { + double pct = grand_total > 0 ? 100.0 * (double) s.total_ns / (double) grand_total : 0; + double avg_us = (double) s.total_ns / (double) s.count / 1000.0; + double min_us = (double) s.min_ns / 1000.0; + double max_us = (double) s.max_ns / 1000.0; + + char bw_buf[32] = ""; + fmt_bandwidth(bw_buf, sizeof(bw_buf), s.total_bytes, s.total_ns); + + char ne0_buf[64]; + char ne1_buf[64]; + char ne2_buf[64]; + fmt_ne(ne0_buf, sizeof(ne0_buf), s.representative_ne_src0); + fmt_ne(ne1_buf, sizeof(ne1_buf), s.representative_ne_src1); + fmt_ne(ne2_buf, sizeof(ne2_buf), s.representative_ne_src2); + + // Build tensor shapes string + char tensors_buf[256] = ""; + int tpos = 0; + if (ne0_buf[0]) { + tpos += snprintf(tensors_buf + tpos, sizeof(tensors_buf) - tpos, "%s", ne0_buf); + } + if (ne1_buf[0]) { + tpos += snprintf(tensors_buf + tpos, sizeof(tensors_buf) - tpos, " x %s", ne1_buf); + } + if (ne2_buf[0]) { + tpos += snprintf(tensors_buf + tpos, sizeof(tensors_buf) - tpos, " x %s", ne2_buf); + } + + fprintf(fp, "%-5s %4d %-28s %6.1f%% %6d %10.2f %10.2f %10.2f %10.2f %12s %s\n", + type_str[s.type], s.backend_id, s.name, pct, s.count, + (double) s.total_ns / 1e6, avg_us, min_us, max_us, + bw_buf, tensors_buf); + } + + fprintf(fp, "\nTotal: %.2f ms (%d records, %d unique ops)\n", (double) grand_total / 1e6, + (int) sched->profiling_records.size(), (int) stats.size()); + + return 0; +} + +int ggml_backend_sched_export_profiling_text(ggml_backend_sched_t sched, const char * filepath) { + GGML_ASSERT(sched); + GGML_ASSERT(filepath != NULL); + + FILE * fp = fopen(filepath, "w"); + if (fp == NULL) { + GGML_LOG_ERROR("%s: failed to open %s for writing\n", __func__, filepath); + return -1; + } + + int ret = ggml_backend_sched_write_profiling_text(sched, fp); + fclose(fp); + + return ret; +} diff --git a/ggml/src/ggml-blas/ggml-blas.cpp b/ggml/src/ggml-blas/ggml-blas.cpp index e4b5bd2547..c8b6b37990 100644 --- a/ggml/src/ggml-blas/ggml-blas.cpp +++ b/ggml/src/ggml-blas/ggml-blas.cpp @@ -2,6 +2,7 @@ #include "ggml-impl.h" #include "ggml-blas.h" #include "ggml-backend-impl.h" +#include "ggml-profiler.h" #include #include @@ -26,6 +27,11 @@ struct ggml_backend_blas_context { #ifndef GGML_USE_OPENMP std::vector> tasks; #endif + + // Profiling state + bool profiling_enabled = false; + int profiling_split_id = -1; + std::vector profiling_records; }; static void ggml_backend_blas_mul_mat(ggml_backend_blas_context * ctx, struct ggml_tensor * dst) { @@ -233,6 +239,18 @@ static enum ggml_status ggml_backend_blas_graph_compute(ggml_backend_t backend, continue; } + // Skip view/identity ops + if (node->op == GGML_OP_NONE || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_VIEW || + node->op == GGML_OP_PERMUTE || node->op == GGML_OP_TRANSPOSE) { + continue; + } + + // Profiling: time this operation + uint64_t t_start = 0; + if (ctx->profiling_enabled) { + t_start = ggml_profiler_time_ns(); + } + switch (node->op) { case GGML_OP_MUL_MAT: ggml_backend_blas_mul_mat(ctx, node); @@ -242,16 +260,24 @@ static enum ggml_status ggml_backend_blas_graph_compute(ggml_backend_t backend, ggml_backend_blas_out_prod(ctx, node); break; - case GGML_OP_NONE: - case GGML_OP_RESHAPE: - case GGML_OP_VIEW: - case GGML_OP_PERMUTE: - case GGML_OP_TRANSPOSE: - break; - default: GGML_ABORT("%s: unsupported op %s\n", __func__, ggml_op_desc(node)); } + + if (ctx->profiling_enabled) { + uint64_t t_end = ggml_profiler_time_ns(); + ggml_profile_record rec; + rec.type = GGML_PROFILE_EVENT_OP; + rec.name = ggml_op_name(node->op); + rec.backend_id = 0; + rec.split_id = ctx->profiling_split_id; + rec.start_ns = t_start; + rec.end_ns = t_end; + rec.bytes = ggml_nbytes(node); + rec.extra = NULL; + ggml_profile_record_from_tensor(&rec, node); + ctx->profiling_records.push_back(rec); + } } return GGML_STATUS_SUCCESS; @@ -287,10 +313,11 @@ ggml_backend_t ggml_backend_blas_init(void) { ggml_backend_blas_context * ctx = new ggml_backend_blas_context; ggml_backend_t backend = new ggml_backend { - /* .guid = */ ggml_backend_blas_guid(), - /* .iface = */ blas_backend_i, - /* .device = */ ggml_backend_reg_dev_get(ggml_backend_blas_reg(), 0), - /* .context = */ ctx, + /* .guid = */ ggml_backend_blas_guid(), + /* .iface = */ blas_backend_i, + /* .device = */ ggml_backend_reg_dev_get(ggml_backend_blas_reg(), 0), + /* .context = */ ctx, + /* .profiler = */ nullptr, }; #if defined(GGML_BLAS_USE_OPENBLAS) && defined(GGML_USE_OPENMP) @@ -303,6 +330,44 @@ ggml_backend_t ggml_backend_blas_init(void) { GGML_LOG_DEBUG("%s: warning: ggml is using OpenMP, but BLIS was compiled without OpenMP support\n", __func__); #endif + // Register profiler + ggml_backend_blas_context * blas_ctx = ctx; // ctx is already defined above + + static auto blas_prof_enable = [](void * ctx, bool enable) { + auto * bctx = (ggml_backend_blas_context *) ctx; + bctx->profiling_enabled = enable; + if (!enable) { + bctx->profiling_records.clear(); + } + }; + static auto blas_prof_reset = [](void * ctx) { + auto * bctx = (ggml_backend_blas_context *) ctx; + bctx->profiling_records.clear(); + bctx->profiling_split_id = -1; + }; + static auto blas_prof_set_split_id = [](void * ctx, int split_id) { + auto * bctx = (ggml_backend_blas_context *) ctx; + bctx->profiling_split_id = split_id; + }; + static auto blas_prof_get_records = [](void * ctx, const ggml_profile_record ** out) -> int { + auto * bctx = (ggml_backend_blas_context *) ctx; + *out = bctx->profiling_records.data(); + return (int) bctx->profiling_records.size(); + }; + static auto blas_prof_free = [](void * ctx) { + (void) ctx; + }; + + auto * profiler = new ggml_backend_profiler{ + /* .context = */ blas_ctx, + /* .enable = */ blas_prof_enable, + /* .reset = */ blas_prof_reset, + /* .set_split_id = */ blas_prof_set_split_id, + /* .get_records = */ blas_prof_get_records, + /* .free_context = */ blas_prof_free, + }; + ggml_backend_set_profiler(backend, profiler); + return backend; } diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp index c2745014a1..f6b2fc0772 100644 --- a/ggml/src/ggml-cann/ggml-cann.cpp +++ b/ggml/src/ggml-cann/ggml-cann.cpp @@ -3043,7 +3043,8 @@ ggml_backend_t ggml_backend_cann_init(int32_t device) { new ggml_backend{ /* .guid = */ ggml_backend_cann_guid(), /* .interface = */ ggml_backend_cann_interface, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cann_reg(), device), - /* .context = */ ctx }; + /* .context = */ ctx, + /* .profiler = */ nullptr }; return cann_backend; } diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 87a329f269..df677d5a5b 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -7,6 +7,7 @@ #include "iqp.h" #include "ggml-cpu-impl.h" #include "ggml-impl.h" +#include "ggml-profiler.h" #include "quants.h" #include "ggml-threading.h" #include "unary-ops.h" @@ -1179,8 +1180,8 @@ static void ggml_compute_forward_mul_mat_one_chunk( const bool src1_cont = ggml_is_contiguous(src1); - ggml_vec_dot_t const vec_dot = type_traits_cpu[type].vec_dot; - enum ggml_type const vec_dot_type = type_traits_cpu[type].vec_dot_type; + const ggml_vec_dot_t vec_dot = type_traits_cpu[type].vec_dot; + const enum ggml_type vec_dot_type = type_traits_cpu[type].vec_dot_type; // broadcast factors const int64_t r2 = ne12 / ne02; @@ -1270,9 +1271,9 @@ void ggml_compute_forward_mul_mat( const int ith = params->ith; const int nth = params->nth; - enum ggml_type const vec_dot_type = type_traits_cpu[src0->type].vec_dot_type; - ggml_from_float_t const from_float = type_traits_cpu[vec_dot_type].from_float; - int64_t const vec_dot_num_rows = type_traits_cpu[src0->type].nrows; + const enum ggml_type vec_dot_type = type_traits_cpu[src0->type].vec_dot_type; + const ggml_from_float_t from_float = type_traits_cpu[vec_dot_type].from_float; + const int64_t vec_dot_num_rows = type_traits_cpu[src0->type].nrows; GGML_ASSERT(ne0 == ne01); GGML_ASSERT(ne1 == ne11); @@ -1488,8 +1489,8 @@ static void ggml_compute_forward_mul_mat_id_one_chunk( const enum ggml_type type = src0->type; - ggml_vec_dot_t const vec_dot = type_traits_cpu[type].vec_dot; - enum ggml_type const vec_dot_type = type_traits_cpu[type].vec_dot_type; + const ggml_vec_dot_t vec_dot = type_traits_cpu[type].vec_dot; + const enum ggml_type vec_dot_type = type_traits_cpu[type].vec_dot_type; const int64_t blck_0 = 16; const int64_t blck_1 = 16; @@ -1556,8 +1557,8 @@ static void ggml_compute_forward_mul_mat_id( const bool src1_cont = ggml_is_contiguous(src1); - enum ggml_type const vec_dot_type = type_traits_cpu[type].vec_dot_type; - ggml_from_float_t const from_float = type_traits_cpu[vec_dot_type].from_float; + const enum ggml_type vec_dot_type = type_traits_cpu[type].vec_dot_type; + const ggml_from_float_t from_float = type_traits_cpu[vec_dot_type].from_float; // we don't support permuted src0 or src1 GGML_ASSERT(nb00 == ggml_type_size(type)); @@ -3126,17 +3127,55 @@ static thread_ret_t ggml_graph_compute_thread(void * data) { GGML_PRINT_DEBUG("thread #%d compute-start cplan %p last-graph %d\n", state->ith, (const void *)cplan, state->last_graph); #endif - for (int node_n = 0; node_n < cgraph->n_nodes && atomic_load_explicit(&tp->abort, memory_order_relaxed) != node_n; node_n++) { - struct ggml_tensor * node = cgraph->nodes[node_n]; + // Profiling state + if (cplan->profiling_context != NULL && cplan->profiling_record_fn != NULL) { + for (int node_n = 0; node_n < cgraph->n_nodes && atomic_load_explicit(&tp->abort, memory_order_relaxed) != node_n; node_n++) { + struct ggml_tensor * node = cgraph->nodes[node_n]; - if (ggml_op_is_empty(node->op)) { - // skip NOPs - continue; - } + if (ggml_op_is_empty(node->op)) { + continue; + } - if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { - continue; + if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { + continue; + } + + // Only thread 0 records timing (after barrier = total node time) + uint64_t t_start = 0; + if (state->ith == 0) { + t_start = ggml_profiler_time_ns(); + } + + ggml_compute_forward(¶ms, node); + + if (node_n + 1 < cgraph->n_nodes) { + ggml_barrier(state->threadpool); + } + + if (state->ith == 0) { + uint64_t t_end = ggml_profiler_time_ns(); + cplan->profiling_record_fn(cplan->profiling_context, 0 /* GGML_PROFILE_EVENT_OP */, + ggml_op_name(node->op), -1, t_start, t_end, ggml_nbytes(node), NULL, + node); + } + + if (state->ith == 0 && cplan->abort_callback && cplan->abort_callback(cplan->abort_callback_data)) { + atomic_store_explicit(&tp->abort, node_n + 1, memory_order_relaxed); + tp->ec = GGML_STATUS_ABORTED; + } } + } else { + for (int node_n = 0; node_n < cgraph->n_nodes && atomic_load_explicit(&tp->abort, memory_order_relaxed) != node_n; node_n++) { + struct ggml_tensor * node = cgraph->nodes[node_n]; + + if (ggml_op_is_empty(node->op)) { + // skip NOPs + continue; + } + + if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) { + continue; + } // TODO: move fused-op detection into ggml_graph_plan so fusion decisions are made once at planning time // Try fused ops, fall back to normal compute @@ -3147,14 +3186,15 @@ static thread_ret_t ggml_graph_compute_thread(void * data) { ggml_compute_forward(¶ms, node); } - if (state->ith == 0 && cplan->abort_callback && - cplan->abort_callback(cplan->abort_callback_data)) { - atomic_store_explicit(&tp->abort, node_n + 1, memory_order_relaxed); - tp->ec = GGML_STATUS_ABORTED; - } + if (state->ith == 0 && cplan->abort_callback && + cplan->abort_callback(cplan->abort_callback_data)) { + atomic_store_explicit(&tp->abort, node_n + 1, memory_order_relaxed); + tp->ec = GGML_STATUS_ABORTED; + } - if (node_n + 1 < cgraph->n_nodes) { - ggml_barrier(state->threadpool); + if (node_n + 1 < cgraph->n_nodes) { + ggml_barrier(state->threadpool); + } } } diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 8cece71f18..3abab16091 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -1,6 +1,7 @@ #include "ggml-backend.h" #include "ggml-backend-impl.h" #include "ggml-cpu.h" +#include "ggml-profiler.h" #include "repack.h" #include "traits.h" #include "ggml-impl.h" @@ -107,6 +108,11 @@ struct ggml_backend_cpu_context { void * abort_callback_data; bool use_ref; // use reference implementation + + // Profiling state + bool profiling_enabled; + int profiling_split_id; + std::vector profiling_records; }; static const char * ggml_backend_cpu_get_name(ggml_backend_t backend) { @@ -167,6 +173,30 @@ static enum ggml_status ggml_backend_cpu_graph_plan_compute(ggml_backend_t backe GGML_UNUSED(backend); } +// Callback function for recording CPU profiling events from C code (ggml-cpu.c) +static void ggml_cpu_profiler_record_callback(void * context, + int type, + const char * name, + int split_id, + uint64_t start_ns, + uint64_t end_ns, + uint64_t bytes, + const char * extra, + const struct ggml_tensor * node) { + auto * cpu_ctx = (ggml_backend_cpu_context *) context; + ggml_profile_record rec; + rec.type = (enum ggml_profile_event_type) type; + rec.name = name; + rec.backend_id = 0; // will be overwritten by scheduler + rec.split_id = split_id != -1 ? split_id : cpu_ctx->profiling_split_id; + rec.start_ns = start_ns; + rec.end_ns = end_ns; + rec.bytes = bytes; + rec.extra = extra; + ggml_profile_record_from_tensor(&rec, node); + cpu_ctx->profiling_records.push_back(rec); +} + static enum ggml_status ggml_backend_cpu_graph_compute(ggml_backend_t backend, struct ggml_cgraph * cgraph) { struct ggml_backend_cpu_context * cpu_ctx = (struct ggml_backend_cpu_context *)backend->context; @@ -187,6 +217,9 @@ static enum ggml_status ggml_backend_cpu_graph_compute(ggml_backend_t backend, s cplan.abort_callback_data = cpu_ctx->abort_callback_data; cplan.use_ref = cpu_ctx->use_ref; + cplan.profiling_context = cpu_ctx->profiling_enabled ? cpu_ctx : NULL; + cplan.profiling_record_fn = cpu_ctx->profiling_enabled ? ggml_cpu_profiler_record_callback : NULL; + return ggml_graph_compute(cgraph, &cplan); } @@ -230,12 +263,15 @@ ggml_backend_t ggml_backend_cpu_init(void) { ctx->abort_callback = NULL; ctx->abort_callback_data = NULL; ctx->use_ref = false; + ctx->profiling_enabled = false; + ctx->profiling_split_id = -1; ggml_backend_t cpu_backend = new ggml_backend { - /* .guid = */ ggml_backend_cpu_guid(), - /* .iface = */ ggml_backend_cpu_i, - /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cpu_reg(), 0), - /* .context = */ ctx, + /* .guid = */ ggml_backend_cpu_guid(), + /* .iface = */ ggml_backend_cpu_i, + /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cpu_reg(), 0), + /* .context = */ ctx, + /* .profiler = */ nullptr, }; if (cpu_backend == NULL) { @@ -243,6 +279,43 @@ ggml_backend_t ggml_backend_cpu_init(void) { return NULL; } + // Register profiler + static auto cpu_prof_enable = [](void * ctx, bool enable) { + auto * cpu_ctx = (ggml_backend_cpu_context *) ctx; + cpu_ctx->profiling_enabled = enable; + if (!enable) { + cpu_ctx->profiling_records.clear(); + } + }; + static auto cpu_prof_reset = [](void * ctx) { + auto * cpu_ctx = (ggml_backend_cpu_context *) ctx; + cpu_ctx->profiling_records.clear(); + cpu_ctx->profiling_split_id = -1; + }; + static auto cpu_prof_set_split_id = [](void * ctx, int split_id) { + auto * cpu_ctx = (ggml_backend_cpu_context *) ctx; + cpu_ctx->profiling_split_id = split_id; + }; + static auto cpu_prof_get_records = [](void * ctx, const ggml_profile_record ** out) -> int { + auto * cpu_ctx = (ggml_backend_cpu_context *) ctx; + *out = cpu_ctx->profiling_records.data(); + return (int) cpu_ctx->profiling_records.size(); + }; + static auto cpu_prof_free = [](void * ctx) { + // Nothing to free - records are in the CPU context's vector + (void) ctx; + }; + + auto * profiler = new ggml_backend_profiler{ + /* .context = */ ctx, + /* .enable = */ cpu_prof_enable, + /* .reset = */ cpu_prof_reset, + /* .set_split_id = */ cpu_prof_set_split_id, + /* .get_records = */ cpu_prof_get_records, + /* .free_context = */ cpu_prof_free, + }; + ggml_backend_set_profiler(cpu_backend, profiler); + return cpu_backend; } diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index ed0ea60bdd..5640bbc3fc 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1419,6 +1419,9 @@ struct ggml_cuda_stream_context { } }; +// Forward declaration for profiler state (defined in ggml-cuda.cu) +struct ggml_cuda_profiler_state; + struct ggml_backend_cuda_context { int device; std::string name; @@ -1537,6 +1540,9 @@ struct ggml_backend_cuda_context { ggml_cuda_pool & pool() { return pool(device); } + + // Profiling + ggml_cuda_profiler_state * profiler_state = nullptr; }; struct ggml_cuda_mm_fusion_args_host { diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu old mode 100644 new mode 100755 index 45e9537f0e..95dca37dca --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1,6 +1,7 @@ #include "ggml-cuda.h" #include "ggml-impl.h" #include "ggml-backend-impl.h" +#include "ggml-profiler.h" #include "ggml-cuda/allreduce.cuh" #include "ggml-cuda/common.cuh" @@ -93,6 +94,92 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size"); +// CUDA profiler state +struct ggml_cuda_profiler_state { + bool enabled = false; + int split_id = -1; + cudaStream_t stream = nullptr; + + static constexpr int MAX_PENDING_EVENTS = 4096; + std::vector start_events; + std::vector end_events; + std::vector cpu_timestamps; // CPU-side timestamps for global ordering + int event_count = 0; + + std::vector records; + std::vector record_event_indices; + + void init(cudaStream_t stream) { + this->stream = stream; + start_events.reserve(MAX_PENDING_EVENTS); + end_events.reserve(MAX_PENDING_EVENTS); + cpu_timestamps.reserve(MAX_PENDING_EVENTS); + } + + void reset() { + for (auto & ev : start_events) { + (void) cudaEventDestroy(ev); + } + for (auto & ev : end_events) { + (void) cudaEventDestroy(ev); + } + start_events.clear(); + end_events.clear(); + cpu_timestamps.clear(); + event_count = 0; + records.clear(); + record_event_indices.clear(); + } + + ~ggml_cuda_profiler_state() { + reset(); + } + + void record_start() { + cudaEvent_t ev; + (void) cudaEventCreate(&ev); + (void) cudaEventRecord(ev, stream); + start_events.push_back(ev); + cpu_timestamps.push_back(ggml_profiler_time_ns()); + event_count++; + } + + void record_end(const char * name, int backend_id, int split_id, uint64_t bytes, const char * extra, + const ggml_tensor * node) { + cudaEvent_t ev; + (void) cudaEventCreate(&ev); + (void) cudaEventRecord(ev, stream); + end_events.push_back(ev); + record_event_indices.push_back(records.size()); + + ggml_profile_record rec; + rec.type = GGML_PROFILE_EVENT_OP; + rec.name = name; + rec.backend_id = backend_id; + rec.split_id = split_id; + rec.start_ns = 0; + rec.end_ns = 0; + rec.bytes = bytes; + rec.extra = extra; + ggml_profile_record_from_tensor(&rec, node); + records.push_back(rec); + } + + void finalize() { + (void) cudaStreamSynchronize(stream); + + for (int i = 0; i < (int)record_event_indices.size(); i++) { + float ms = 0.0f; + (void) cudaEventElapsedTime(&ms, start_events[i], end_events[i]); + uint64_t duration_ns = (uint64_t)(ms * 1e6f); + int rec_idx = record_event_indices[i]; + // Use CPU-side timestamp for global ordering, GPU-measured duration for accuracy + records[rec_idx].start_ns = cpu_timestamps[i]; + records[rec_idx].end_ns = cpu_timestamps[i] + duration_ns; + } + } +}; + #define GGML_LOG_WARN_ONCE(str) \ { static std::once_flag warn_flag; std::call_once(warn_flag, []() { GGML_LOG_WARN(str); }); } @@ -4345,8 +4432,23 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud #else GGML_UNUSED(integrated); #endif // NDEBUG + if (cuda_ctx->profiler_state != nullptr && cuda_ctx->profiler_state->enabled) { + cuda_ctx->profiler_state->record_start(); + } bool ok = ggml_cuda_compute_forward(*cuda_ctx, node); + + if (cuda_ctx->profiler_state != nullptr && cuda_ctx->profiler_state->enabled) { + cuda_ctx->profiler_state->record_end( + ggml_op_name(node->op), + -1, + cuda_ctx->profiler_state->split_id, + ggml_nbytes(node), + nullptr, + node + ); + } + if (!ok) { GGML_LOG_ERROR("%s: op not supported %s (%s)\n", __func__, node->name, ggml_op_name(node->op)); } @@ -4417,6 +4519,19 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cuda_set_device(cuda_ctx->device); + // Disable CUDA graphs when profiling (we need per-node timing) + bool was_graph_enabled = false; + if (cuda_ctx->profiler_state != nullptr && cuda_ctx->profiler_state->enabled) { +#ifdef USE_CUDA_GRAPH + const void * graph_key = ggml_cuda_graph_get_key(cgraph); + ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key); + was_graph_enabled = graph->is_enabled(); + if (was_graph_enabled) { + graph->disable_due_to_gpu_arch = true; + } +#endif + } + bool use_cuda_graph = false; bool cuda_graph_update_required = false; const void * graph_key = nullptr; @@ -4468,6 +4583,15 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cuda_graph_evaluate_and_capture(cuda_ctx, cgraph, use_cuda_graph, cuda_graph_update_required, graph_key); + // Restore CUDA graph enabled state after profiling + if (was_graph_enabled) { +#ifdef USE_CUDA_GRAPH + const void * graph_key_prof = ggml_cuda_graph_get_key(cgraph); + ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key_prof); + graph->disable_due_to_gpu_arch = false; +#endif + } + return GGML_STATUS_SUCCESS; } @@ -5014,6 +5138,22 @@ static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend : GGML_BACKEND_DEVICE_TYPE_GPU; } +static bool ggml_backend_cuda_host_buffer_supported() { + return getenv("GGML_CUDA_NO_PINNED") == nullptr; +} + +static bool ggml_backend_cuda_device_supports_cuda_host_buft(int device) { +#if defined(GGML_USE_HIP) + if (ggml_cuda_info().devices[device].integrated) { + return false; + } +#else + GGML_UNUSED(device); +#endif + + return ggml_backend_cuda_host_buffer_supported(); +} + static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_backend_dev_props * props) { ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context; @@ -5023,7 +5163,7 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back props->device_id = ctx->pci_bus_id.empty() ? nullptr : ctx->pci_bus_id.c_str(); ggml_backend_cuda_device_get_memory(dev, &props->memory_free, &props->memory_total); - bool host_buffer = getenv("GGML_CUDA_NO_PINNED") == nullptr; + bool host_buffer = ggml_backend_cuda_host_buffer_supported(); #ifdef GGML_CUDA_NO_PEER_COPY bool events = false; #else @@ -5052,6 +5192,10 @@ static ggml_backend_buffer_type_t ggml_backend_cuda_device_get_buffer_type(ggml_ static ggml_backend_buffer_type_t ggml_backend_cuda_device_get_host_buffer_type(ggml_backend_dev_t dev) { GGML_UNUSED(dev); + if (!ggml_backend_cuda_host_buffer_supported()) { + return nullptr; + } + return ggml_backend_cuda_host_buffer_type(); } @@ -5518,7 +5662,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g static bool ggml_backend_cuda_device_supports_buft(ggml_backend_dev_t dev, ggml_backend_buffer_type_t buft) { ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *) dev->context; const bool integrated = ggml_cuda_info().devices[dev_ctx->device].integrated; - return (ggml_backend_buft_is_cuda(buft) && buft->device == dev) || (integrated && ggml_backend_buft_is_cuda_host(buft)); + return (ggml_backend_buft_is_cuda(buft) && buft->device == dev) || + (integrated && + ggml_backend_buft_is_cuda_host(buft) && + ggml_backend_cuda_device_supports_cuda_host_buft(dev_ctx->device)); } static int64_t get_op_batch_size(const ggml_tensor * op) { @@ -5769,12 +5916,68 @@ ggml_backend_t ggml_backend_cuda_init(int device) { } ggml_backend_t cuda_backend = new ggml_backend { - /* .guid = */ ggml_backend_cuda_guid(), - /* .iface = */ ggml_backend_cuda_interface, - /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cuda_reg(), device), - /* .context = */ ctx, + /* .guid = */ ggml_backend_cuda_guid(), + /* .iface = */ ggml_backend_cuda_interface, + /* .device = */ ggml_backend_reg_dev_get(ggml_backend_cuda_reg(), device), + /* .context = */ ctx, + /* .profiler = */ nullptr, }; + // Register profiler + auto * prof_state = new ggml_cuda_profiler_state(); + prof_state->init(ctx->stream()); + ctx->profiler_state = prof_state; + + static auto cuda_prof_enable = [](void * ctx, bool enable) { + auto * cuda_ctx = (ggml_backend_cuda_context *)ctx; + if (cuda_ctx->profiler_state != nullptr) { + cuda_ctx->profiler_state->enabled = enable; + if (!enable) { + cuda_ctx->profiler_state->reset(); + } + } + }; + static auto cuda_prof_reset = [](void * ctx) { + auto * cuda_ctx = (ggml_backend_cuda_context *)ctx; + if (cuda_ctx->profiler_state != nullptr) { + cuda_ctx->profiler_state->reset(); + cuda_ctx->profiler_state->split_id = -1; + } + }; + static auto cuda_prof_set_split_id = [](void * ctx, int split_id) { + auto * cuda_ctx = (ggml_backend_cuda_context *)ctx; + if (cuda_ctx->profiler_state != nullptr) { + cuda_ctx->profiler_state->split_id = split_id; + } + }; + static auto cuda_prof_get_records = [](void * ctx, const ggml_profile_record ** out) -> int { + auto * cuda_ctx = (ggml_backend_cuda_context *)ctx; + if (cuda_ctx->profiler_state != nullptr) { + cuda_ctx->profiler_state->finalize(); + *out = cuda_ctx->profiler_state->records.data(); + return (int)cuda_ctx->profiler_state->records.size(); + } + *out = nullptr; + return 0; + }; + static auto cuda_prof_free = [](void * ctx) { + auto * cuda_ctx = (ggml_backend_cuda_context *)ctx; + if (cuda_ctx->profiler_state != nullptr) { + delete cuda_ctx->profiler_state; + cuda_ctx->profiler_state = nullptr; + } + }; + + auto * profiler = new ggml_backend_profiler { + /* .context = */ ctx, + /* .enable = */ cuda_prof_enable, + /* .reset = */ cuda_prof_reset, + /* .set_split_id = */ cuda_prof_set_split_id, + /* .get_records = */ cuda_prof_get_records, + /* .free_context = */ cuda_prof_free, + }; + ggml_backend_set_profiler(cuda_backend, profiler); + return cuda_backend; } diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index 9aa558f3f4..b6386c6bb0 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -60,8 +60,10 @@ #define cudaErrorMemoryAllocation hipErrorOutOfMemory #define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled #define cudaErrorPeerAccessNotEnabled hipErrorPeerAccessNotEnabled +#define cudaEventCreate hipEventCreate #define cudaEventCreateWithFlags hipEventCreateWithFlags #define cudaEventDisableTiming hipEventDisableTiming +#define cudaEventElapsedTime hipEventElapsedTime #define cudaEventRecord hipEventRecord #define cudaEventSynchronize hipEventSynchronize #define cudaEvent_t hipEvent_t diff --git a/ggml/src/ggml-cuda/vendors/musa.h b/ggml/src/ggml-cuda/vendors/musa.h index 6d725c7ec1..ec4115aba2 100644 --- a/ggml/src/ggml-cuda/vendors/musa.h +++ b/ggml/src/ggml-cuda/vendors/musa.h @@ -48,8 +48,10 @@ #define cudaErrorMemoryAllocation musaErrorMemoryAllocation #define cudaErrorPeerAccessAlreadyEnabled musaErrorPeerAccessAlreadyEnabled #define cudaErrorPeerAccessNotEnabled musaErrorPeerAccessNotEnabled +#define cudaEventCreate musaEventCreate #define cudaEventCreateWithFlags musaEventCreateWithFlags #define cudaEventDisableTiming musaEventDisableTiming +#define cudaEventElapsedTime musaEventElapsedTime #define cudaEventRecord musaEventRecord #define cudaEventSynchronize musaEventSynchronize #define cudaEvent_t musaEvent_t diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index 104201daff..5f5c28cced 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -5684,6 +5684,7 @@ static ggml_backend_t ggml_backend_hexagon_device_init(ggml_backend_dev_t dev, c /* .interface = */ hexagon_backend_i, /* .device = */ dev, /* .context = */ sess, + /* .profiler = */ nullptr, }; GGML_UNUSED(params); diff --git a/ggml/src/ggml-metal/ggml-metal-context.h b/ggml/src/ggml-metal/ggml-metal-context.h index abf4b06ed2..8f61a826be 100644 --- a/ggml/src/ggml-metal/ggml-metal-context.h +++ b/ggml/src/ggml-metal/ggml-metal-context.h @@ -36,6 +36,32 @@ void ggml_metal_set_abort_callback (ggml_metal_t ctx, ggml_abort_callback abort bool ggml_metal_supports_family (ggml_metal_t ctx, int family); void ggml_metal_capture_next_compute(ggml_metal_t ctx); +// +// Profiling +// + +// Opaque profiler state, owned by the C++ backend layer (ggml-metal.cpp). Holds the std::vector of +// records returned via the ggml_backend_profiler interface. ggml_metal_t keeps a borrowed pointer +// so that graph_compute can push records when sampling is active. +struct ggml_metal_profiler_state; + +// Inject (or clear, with NULL) the profiler state pointer. Called once at backend init. +void ggml_metal_set_profiler_state(ggml_metal_t ctx, struct ggml_metal_profiler_state * state); + +// Bridge function implemented in ggml-metal.cpp. Used by graph_compute (in .m) to push records. +void ggml_metal_profiler_push_record( + struct ggml_metal_profiler_state * state, + const struct ggml_tensor * node, + uint64_t start_ns, + uint64_t end_ns); + +// Query whether the injected profiler state is currently enabled. +// (Avoids exposing the C++ struct layout to the .m file.) +bool ggml_metal_profiler_is_enabled(struct ggml_metal_profiler_state * state); + +// Query the split-id currently set on the profiler state. +int ggml_metal_profiler_get_split_id(struct ggml_metal_profiler_state * state); + #ifdef __cplusplus } #endif diff --git a/ggml/src/ggml-metal/ggml-metal-context.m b/ggml/src/ggml-metal/ggml-metal-context.m index 6cdc4006bc..991abdb8ab 100644 --- a/ggml/src/ggml-metal/ggml-metal-context.m +++ b/ggml/src/ggml-metal/ggml-metal-context.m @@ -6,6 +6,7 @@ #import "ggml-metal-impl.h" #import "ggml-metal-common.h" #import "ggml-metal-ops.h" +#import "ggml-profiler.h" #import @@ -83,6 +84,19 @@ struct ggml_metal { // error state - set when a command buffer fails during synchronize // once set, graph_compute will return GGML_STATUS_FAILED until the backend is recreated bool has_error; + + // Profiling + // Borrowed; owned by the C++ backend layer (ggml-metal.cpp). NULL when profiling is unavailable. + struct ggml_metal_profiler_state * profiler_state; + + // Per-graph-compute scratch state populated when profiling is active for the current invocation. + // Lifetime: from start of ggml_metal_graph_compute until its end. + bool profiler_active; + ggml_metal_sample_buf_t profiler_sample_buf; + size_t profiler_total_slots; + uint64_t profiler_cpu_anchor_ns; + uint64_t profiler_gpu_anchor_ns; + struct ggml_metal_op_sample_slot * profiler_slot_map; // size = gf->n_nodes }; ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) { @@ -471,6 +485,34 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph * // keep the memory wired ggml_metal_device_rsets_keep_alive(ctx->dev); + // Decide whether profiling is active for this invocation. Activation is sticky for the whole + // call: allocate sample buffer + slot map up front so the encode_async block can use them. + ctx->profiler_active = false; + if (ctx->profiler_state != NULL && ggml_metal_profiler_is_enabled(ctx->profiler_state)) { + if (ggml_metal_device_supports_profiling(ctx->dev) && gf->n_nodes > 0) { + const size_t total_slots = 2 * (size_t) gf->n_nodes; + ctx->profiler_sample_buf = ggml_metal_device_create_sample_buf(ctx->dev, total_slots); + if (ctx->profiler_sample_buf != NULL) { + ctx->profiler_total_slots = total_slots; + ctx->profiler_slot_map = (struct ggml_metal_op_sample_slot *) calloc( + (size_t) gf->n_nodes, sizeof(struct ggml_metal_op_sample_slot)); + if (ctx->profiler_slot_map != NULL) { + // Mark all entries unused (node_idx < 0). + for (int i = 0; i < gf->n_nodes; ++i) { + ctx->profiler_slot_map[i].node_idx = -1; + } + ggml_metal_device_sample_timestamps(ctx->dev, + &ctx->profiler_cpu_anchor_ns, + &ctx->profiler_gpu_anchor_ns); + ctx->profiler_active = true; + } else { + ggml_metal_sample_buf_free(ctx->profiler_sample_buf); + ctx->profiler_sample_buf = NULL; + } + } + } + } + // submit the ggml compute graph to the GPU by creating command buffers and encoding the ops in them // the first n_nodes_0 are encoded and submitted for processing directly by the calling thread // while these nodes are processing, we start n_cb threads to enqueue the rest of the nodes @@ -630,6 +672,66 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph * ctx->capture_started = false; } + + // Profiling drain: wait for all command buffers to complete, resolve timestamps, push records. + // This forces a synchronous wait — Vulkan does the same when its profiler is active. + if (ctx->profiler_active) { + { + id cmd_buf = ctx->cmd_bufs[n_cb].obj; + if (cmd_buf) { + [cmd_buf waitUntilCompleted]; + } + } + for (int i = 0; i < n_cb; ++i) { + id cmd_buf = ctx->cmd_bufs[i].obj; + if (cmd_buf) { + // Ensure cmd_bufs that were not auto-enqueued get committed. + if ([cmd_buf status] == MTLCommandBufferStatusNotEnqueued) { + [cmd_buf commit]; + } + [cmd_buf waitUntilCompleted]; + } + } + + uint64_t * ns = (uint64_t *) calloc(ctx->profiler_total_slots, sizeof(uint64_t)); + if (ns != NULL) { + ggml_metal_sample_buf_resolve(ctx->profiler_sample_buf, + /*base=*/0, + ctx->profiler_total_slots, + ctx->profiler_cpu_anchor_ns, + ctx->profiler_gpu_anchor_ns, + ns); + + for (int i = 0; i < gf->n_nodes; ++i) { + const struct ggml_metal_op_sample_slot * slot = &ctx->profiler_slot_map[i]; + if (slot->node_idx < 0) { + continue; + } + if (slot->slot_start >= ctx->profiler_total_slots || + slot->slot_end >= ctx->profiler_total_slots) { + continue; + } + const uint64_t t0 = ns[slot->slot_start]; + const uint64_t t1 = ns[slot->slot_end]; + if (t0 == 0 || t1 == 0 || t1 < t0) { + continue; + } + ggml_metal_profiler_push_record(ctx->profiler_state, + ggml_graph_node(gf, slot->node_idx), + t0, + t1); + } + + free(ns); + } + + free(ctx->profiler_slot_map); + ctx->profiler_slot_map = NULL; + ggml_metal_sample_buf_free(ctx->profiler_sample_buf); + ctx->profiler_sample_buf = NULL; + ctx->profiler_total_slots = 0; + ctx->profiler_active = false; + } } return GGML_STATUS_SUCCESS; @@ -681,6 +783,10 @@ ggml_metal_event_t ggml_metal_get_ev_cpy(ggml_metal_t ctx) { return ctx->ev_cpy; } +void ggml_metal_set_profiler_state(ggml_metal_t ctx, struct ggml_metal_profiler_state * state) { + ctx->profiler_state = state; +} + void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) { if (ctx->n_cb != n_cb) { ctx->n_cb = MIN(n_cb, GGML_METAL_MAX_COMMAND_BUFFERS); @@ -725,6 +831,17 @@ void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) { ctx->debug_graph, ctx->debug_fusion); + if (ctx->profiler_active) { + // Base slot for this command buffer is 2 * (first graph-node index it processes). + // The encoder uses one (start, end) pair per encoded op group; we over-allocate so + // that empty/filtered nodes simply leave gaps in the sample buffer. + const size_t base_slot = 2 * (size_t) idx_start; + ggml_metal_op_enable_profiling(ctx_op, + ctx->profiler_sample_buf, + base_slot, + ctx->profiler_slot_map); + } + for (int idx = 0; idx < ggml_metal_op_n_nodes(ctx_op); ++idx) { const int res = ggml_metal_op_encode(ctx_op, idx); if (res == 0) { diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 31fc07d44d..a36fd0cd01 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -91,6 +91,17 @@ void ggml_metal_encoder_memory_barrier(ggml_metal_encoder_t encoder); void ggml_metal_encoder_end_encoding(ggml_metal_encoder_t encoder); +// +// MTLCounterSampleBuffer wrapper (used by the profiler) +// + +typedef struct ggml_metal_sample_buf * ggml_metal_sample_buf_t; + +// Insert a GPU timestamp sample on the encoder at the given slot. +// Caller must ensure (a) the encoder belongs to a command buffer using a counter sample buffer that +// supports MTLCounterSamplingPointAtDispatchBoundary, and (b) `index` is unique within the buffer. +void ggml_metal_encoder_sample_timestamp(ggml_metal_encoder_t encoder, ggml_metal_sample_buf_t buf, size_t index); + // // MTLLibrary wrapper // @@ -325,6 +336,35 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te const struct ggml_metal_device_props * ggml_metal_device_get_props(ggml_metal_device_t dev); +// +// Profiling helpers +// + +// Returns true if the device supports MTLCounterSamplingPointAtDispatchBoundary on compute encoders +// AND exposes the MTLCommonCounterSetTimestamp counter set. +bool ggml_metal_device_supports_profiling(ggml_metal_device_t dev); + +// Allocate a counter sample buffer with `sample_count` slots backed by the timestamp counter set. +// Returns NULL on failure (e.g. unsupported device). +ggml_metal_sample_buf_t ggml_metal_device_create_sample_buf(ggml_metal_device_t dev, size_t sample_count); + +void ggml_metal_sample_buf_free(ggml_metal_sample_buf_t buf); + +// Capture correlated CPU/GPU timestamps (both in Mach-absolute units, i.e. nanoseconds on +// current Apple Silicon; the conversion via mach_timebase_info is applied internally). +// Used to anchor the GPU timestamp domain to the CPU clock returned by ggml_profiler_time_ns(). +void ggml_metal_device_sample_timestamps(ggml_metal_device_t dev, uint64_t * cpu_ns, uint64_t * gpu_ns); + +// Resolve `count` consecutive samples starting at `base` into nanosecond timestamps anchored against +// `cpu_anchor_ns` / `gpu_anchor_ns` (obtained from ggml_metal_device_sample_timestamps). +// `out_ns` must hold at least `count` uint64_t entries. +void ggml_metal_sample_buf_resolve(ggml_metal_sample_buf_t buf, + size_t base, + size_t count, + uint64_t cpu_anchor_ns, + uint64_t gpu_anchor_ns, + uint64_t * out_ns); + // // device buffers // diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index afd6f52101..73a6fd7180 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -10,6 +10,7 @@ #include #include +#include #ifndef TARGET_OS_VISION #define TARGET_OS_VISION 0 @@ -882,6 +883,21 @@ void ggml_metal_encoder_end_encoding(ggml_metal_encoder_t encoder) { [encoder->obj endEncoding]; } +// +// MTLCounterSampleBuffer wrapper +// + +struct ggml_metal_sample_buf { + id obj; + size_t sample_count; +}; + +void ggml_metal_encoder_sample_timestamp(ggml_metal_encoder_t encoder, ggml_metal_sample_buf_t buf, size_t index) { + if (@available(macOS 11.0, iOS 14.0, *)) { + [encoder->obj sampleCountersInBuffer:buf->obj atSampleIndex:index withBarrier:YES]; + } +} + struct ggml_metal_device { id mtl_device; @@ -1935,6 +1951,139 @@ static void ggml_metal_device_disable_tensor(ggml_metal_device_t dev) { dev->props.has_tensor = false; } +// +// Profiling helpers +// + +// Look up the MTLCommonCounterSetTimestamp counter set, or nil if the device doesn't expose it. +static id ggml_metal_device_get_timestamp_counter_set(ggml_metal_device_t dev) { + if (@available(macOS 10.15, iOS 14.0, *)) { + NSArray> * sets = [dev->mtl_device counterSets]; + for (id cs in sets) { + if ([[cs name] isEqualToString:MTLCommonCounterSetTimestamp]) { + return cs; + } + } + } + return nil; +} + +bool ggml_metal_device_supports_profiling(ggml_metal_device_t dev) { + if (@available(macOS 11.0, iOS 14.0, *)) { + if (![dev->mtl_device supportsCounterSampling:MTLCounterSamplingPointAtDispatchBoundary]) { + return false; + } + return ggml_metal_device_get_timestamp_counter_set(dev) != nil; + } + return false; +} + +ggml_metal_sample_buf_t ggml_metal_device_create_sample_buf(ggml_metal_device_t dev, size_t sample_count) { + if (sample_count == 0) { + return NULL; + } + if (@available(macOS 10.15, iOS 14.0, *)) { + id cs = ggml_metal_device_get_timestamp_counter_set(dev); + if (cs == nil) { + return NULL; + } + + MTLCounterSampleBufferDescriptor * desc = [[MTLCounterSampleBufferDescriptor alloc] init]; + desc.counterSet = cs; + desc.label = @"ggml-metal-profiler"; + desc.storageMode = MTLStorageModeShared; + desc.sampleCount = sample_count; + + NSError * error = nil; + id sb = [dev->mtl_device newCounterSampleBufferWithDescriptor:desc error:&error]; + [desc release]; + + if (sb == nil) { + GGML_LOG_WARN("%s: failed to create counter sample buffer: %s\n", __func__, + error ? [[error localizedDescription] UTF8String] : "unknown"); + return NULL; + } + + ggml_metal_sample_buf_t res = calloc(1, sizeof(struct ggml_metal_sample_buf)); + res->obj = sb; + res->sample_count = sample_count; + return res; + } + return NULL; +} + +void ggml_metal_sample_buf_free(ggml_metal_sample_buf_t buf) { + if (buf == NULL) { + return; + } + [buf->obj release]; + free(buf); +} + +void ggml_metal_device_sample_timestamps(ggml_metal_device_t dev, uint64_t * cpu_ns, uint64_t * gpu_ns) { + if (@available(macOS 10.15, iOS 14.0, *)) { + MTLTimestamp cpu_t = 0; + MTLTimestamp gpu_t = 0; + [dev->mtl_device sampleTimestamps:&cpu_t gpuTimestamp:&gpu_t]; + + // Apple docs: both timestamps are in Mach-absolute time units. Convert to nanoseconds. + static mach_timebase_info_data_t tb = {0, 0}; + if (tb.denom == 0) { + mach_timebase_info(&tb); + } + if (cpu_ns) *cpu_ns = (uint64_t)((__uint128_t)cpu_t * tb.numer / tb.denom); + if (gpu_ns) *gpu_ns = (uint64_t)((__uint128_t)gpu_t * tb.numer / tb.denom); + } else { + if (cpu_ns) *cpu_ns = 0; + if (gpu_ns) *gpu_ns = 0; + } +} + +void ggml_metal_sample_buf_resolve(ggml_metal_sample_buf_t buf, + size_t base, + size_t count, + uint64_t cpu_anchor_ns, + uint64_t gpu_anchor_ns, + uint64_t * out_ns) { + if (buf == NULL || count == 0 || out_ns == NULL) { + return; + } + if (@available(macOS 10.15, iOS 14.0, *)) { + NSRange range = NSMakeRange(base, count); + NSData * data = [buf->obj resolveCounterRange:range]; + if (data == nil || [data length] < count * sizeof(MTLCounterResultTimestamp)) { + // Failed resolve: fill with zeros so the caller can detect dropped samples. + for (size_t i = 0; i < count; ++i) { + out_ns[i] = 0; + } + return; + } + + static mach_timebase_info_data_t tb = {0, 0}; + if (tb.denom == 0) { + mach_timebase_info(&tb); + } + + const MTLCounterResultTimestamp * src = (const MTLCounterResultTimestamp *) [data bytes]; + for (size_t i = 0; i < count; ++i) { + // MTLCounterErrorValue (~0ULL) marks an unrecorded sample slot; map to 0. + if (src[i].timestamp == MTLCounterErrorValue) { + out_ns[i] = 0; + continue; + } + const uint64_t gpu_ns = (uint64_t)((__uint128_t) src[i].timestamp * tb.numer / tb.denom); + // Anchor to CPU clock returned by ggml_profiler_time_ns(). + out_ns[i] = (gpu_ns >= gpu_anchor_ns) + ? cpu_anchor_ns + (gpu_ns - gpu_anchor_ns) + : cpu_anchor_ns - (gpu_anchor_ns - gpu_ns); + } + } else { + for (size_t i = 0; i < count; ++i) { + out_ns[i] = 0; + } + } +} + // // device buffers // diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 3db8bca437..41d4c4eab5 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -78,6 +78,11 @@ struct ggml_metal_op { return ggml_graph_node(gf, idxs[i]); } + int node_global_idx(int i) const { + assert(i >= 0 && i < (int) idxs.size()); + return idxs[i]; + } + bool can_fuse(int i0, const ggml_op * ops, int n_ops) const { assert(use_fusion); assert(i0 >= 0 && i0 < n_nodes()); @@ -101,6 +106,12 @@ struct ggml_metal_op { int debug_graph; int debug_fusion; + // Profiling: when sample_buf is non-null, ggml_metal_op_encode brackets each impl call with + // two timestamp samples. The (node_idx, slot_start, slot_end) tuple is recorded in slot_map. + ggml_metal_sample_buf_t sample_buf = nullptr; + size_t next_slot = 0; + ggml_metal_op_sample_slot * slot_map = nullptr; + private: ggml_cgraph * gf; @@ -145,6 +156,16 @@ int ggml_metal_op_n_nodes(ggml_metal_op_t ctx) { return ctx->n_nodes(); } +void ggml_metal_op_enable_profiling( + ggml_metal_op_t ctx, + ggml_metal_sample_buf_t sample_buf, + size_t slot_base, + ggml_metal_op_sample_slot * slot_map) { + ctx->sample_buf = sample_buf; + ctx->next_slot = slot_base; + ctx->slot_map = slot_map; +} + static bool ggml_metal_op_concurrency_reset(ggml_metal_op_t ctx) { if (!ctx->mem_ranges) { return true; @@ -524,12 +545,31 @@ int ggml_metal_op_encode(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_debug_group_push(ctx->enc, ggml_op_desc(ctx->node(idx))); } + const size_t slot_start = ctx->next_slot; + const bool profiling = (ctx->sample_buf != nullptr && ctx->slot_map != nullptr); + if (profiling) { + ggml_metal_encoder_sample_timestamp(ctx->enc, ctx->sample_buf, slot_start); + ctx->next_slot++; + } + int res = ggml_metal_op_encode_impl(ctx, idx); if (idx + res > ctx->n_nodes()) { GGML_ABORT("fusion error: nodes spanning multiple encoders have been fused. this indicates a bug in the fusion logic %s", "https://github.com/ggml-org/llama.cpp/pull/14849"); } + if (profiling) { + const size_t slot_end = ctx->next_slot; + ggml_metal_encoder_sample_timestamp(ctx->enc, ctx->sample_buf, slot_end); + ctx->next_slot++; + + const int gidx = ctx->node_global_idx(idx); + ctx->slot_map[gidx].node_idx = gidx; + ctx->slot_map[gidx].n_fused = res; + ctx->slot_map[gidx].slot_start = slot_start; + ctx->slot_map[gidx].slot_end = slot_end; + } + if (ctx->use_capture) { ggml_metal_encoder_debug_group_pop(ctx->enc); } diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index f8fe50b468..3b17512b68 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -26,6 +26,29 @@ int ggml_metal_op_n_nodes(ggml_metal_op_t ctx); int ggml_metal_op_encode(ggml_metal_op_t ctx, int idx); +// +// Profiling: per-op timestamp sampling +// + +// (node_idx, slot_start, slot_end) triple recorded for each encoded op group. +// node_idx < 0 marks an unused slot. +struct ggml_metal_op_sample_slot { + int node_idx; + int n_fused; + size_t slot_start; + size_t slot_end; +}; + +// Enable profiling on this op encoder. `slot_base` is the first index in `sample_buf` available +// to this encoder; `slot_map` is an array (size = number of graph nodes) into which +// the encoder will record one entry per non-empty encoded op group. +// Must be called before any ggml_metal_op_encode(). +void ggml_metal_op_enable_profiling( + ggml_metal_op_t ctx, + ggml_metal_sample_buf_t sample_buf, + size_t slot_base, + struct ggml_metal_op_sample_slot * slot_map); + // // available ops: // diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp index 3bd6abd06f..ffbc8fbe58 100644 --- a/ggml/src/ggml-metal/ggml-metal.cpp +++ b/ggml/src/ggml-metal/ggml-metal.cpp @@ -2,6 +2,7 @@ #include "ggml-impl.h" #include "ggml-backend-impl.h" +#include "ggml-profiler.h" #include "ggml-metal-device.h" #include "ggml-metal-context.h" @@ -10,6 +11,7 @@ #include #include +#include #define GGML_METAL_NAME "MTL" #define GGML_METAL_MAX_DEVICES 16 @@ -21,6 +23,59 @@ static int g_devices = 1; // forward declaration static bool ggml_backend_buffer_is_metal(ggml_backend_buffer_t buffer); +// +// Profiler state (mirrors the per-backend profiler state from the Vulkan integration) +// +// Owned by ggml_backend_metal_init / device_init_backend; lifetime is the backend's lifetime. +// The Metal context holds a borrowed pointer and calls into the bridge functions below to push +// records during graph_compute. +// + +struct ggml_metal_profiler_state { + bool enabled = false; + int split_id = -1; + std::vector records; + + void reset() { + records.clear(); + split_id = -1; + } +}; + +extern "C" { + +void ggml_metal_profiler_push_record( + ggml_metal_profiler_state * state, + const ggml_tensor * node, + uint64_t start_ns, + uint64_t end_ns) { + if (state == nullptr || node == nullptr) { + return; + } + + ggml_profile_record rec; + rec.type = GGML_PROFILE_EVENT_OP; + rec.name = ggml_op_name(node->op); + rec.backend_id = -1; + rec.split_id = state->split_id; + rec.start_ns = start_ns; + rec.end_ns = end_ns; + rec.bytes = ggml_nbytes(node); + rec.extra = nullptr; + ggml_profile_record_from_tensor(&rec, node); + state->records.push_back(rec); +} + +bool ggml_metal_profiler_is_enabled(ggml_metal_profiler_state * state) { + return state != nullptr && state->enabled; +} + +int ggml_metal_profiler_get_split_id(ggml_metal_profiler_state * state) { + return state != nullptr ? state->split_id : -1; +} + +} // extern "C" + //////////////////////////////////////////////////////////////////////////////// // backend interface //////////////////////////////////////////////////////////////////////////////// @@ -599,6 +654,48 @@ static ggml_guid_t ggml_backend_metal_guid(void) { return &guid; } +// Register the per-backend ggml_backend_profiler interface on `backend` and inject the borrowed +// state pointer into the Metal context so graph_compute can push records. +static void ggml_backend_metal_register_profiler(ggml_backend_t backend, ggml_metal_t ctx) { + auto * state = new ggml_metal_profiler_state(); + ggml_metal_set_profiler_state(ctx, state); + + static auto metal_prof_enable = [](void * context, bool enable) { + auto * state = (ggml_metal_profiler_state *) context; + state->enabled = enable; + if (!enable) { + state->reset(); + } + }; + static auto metal_prof_reset = [](void * context) { + auto * state = (ggml_metal_profiler_state *) context; + state->reset(); + }; + static auto metal_prof_set_split_id = [](void * context, int split_id) { + auto * state = (ggml_metal_profiler_state *) context; + state->split_id = split_id; + }; + static auto metal_prof_get_records = [](void * context, const ggml_profile_record ** out) -> int { + auto * state = (ggml_metal_profiler_state *) context; + *out = state->records.data(); + return (int) state->records.size(); + }; + static auto metal_prof_free = [](void * context) { + auto * state = (ggml_metal_profiler_state *) context; + delete state; + }; + + auto * profiler = new ggml_backend_profiler { + /* .context = */ state, + /* .enable = */ metal_prof_enable, + /* .reset = */ metal_prof_reset, + /* .set_split_id = */ metal_prof_set_split_id, + /* .get_records = */ metal_prof_get_records, + /* .free_context = */ metal_prof_free, + }; + ggml_backend_set_profiler(backend, profiler); +} + ggml_backend_t ggml_backend_metal_init(void) { ggml_backend_dev_t dev = ggml_backend_reg_dev_get(ggml_backend_metal_reg(), 0); ggml_metal_device_t ctx_dev = (ggml_metal_device_t)dev->context; @@ -616,10 +713,13 @@ ggml_backend_t ggml_backend_metal_init(void) { /* .interface = */ ggml_backend_metal_i, /* .device = */ dev, /* .context = */ ctx, + /* .profiler = */ NULL, }; ggml_backend_metal_set_n_cb(backend, 1); + ggml_backend_metal_register_profiler(backend, ctx); + return backend; } @@ -711,10 +811,13 @@ static ggml_backend_t ggml_backend_metal_device_init_backend(ggml_backend_dev_t /* .interface = */ ggml_backend_metal_i, /* .device = */ dev, /* .context = */ ctx, + /* .profiler = */ NULL, }; ggml_backend_metal_set_n_cb(backend, 1); + ggml_backend_metal_register_profiler(backend, ctx); + return backend; GGML_UNUSED(params); diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index d737aea122..2c3845da77 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -8888,7 +8888,8 @@ ggml_backend_t ggml_backend_opencl_init(void) { /* .guid = */ ggml_backend_opencl_guid(), /* .iface = */ ggml_backend_opencl_i, /* .device = */ dev, - /* .context = */ backend_ctx + /* .context = */ backend_ctx, + /* .profiler = */ nullptr, }; return backend; @@ -12345,6 +12346,7 @@ static ggml_backend_t ggml_backend_opencl_device_init(ggml_backend_dev_t dev, co /* .interface = */ ggml_backend_opencl_i, /* .device = */ dev, /* .context = */ backend_ctx, + /* .profiler = */ nullptr, }; ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *) dev->context; diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp index 4b1789713d..194c434f28 100644 --- a/ggml/src/ggml-openvino/ggml-openvino.cpp +++ b/ggml/src/ggml-openvino/ggml-openvino.cpp @@ -803,6 +803,7 @@ GGML_BACKEND_API ggml_backend_t ggml_backend_openvino_init(int device) { /* .interface = */ ggml_backend_openvino_interface, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_openvino_reg(), device), /* .context = */ ctx, + /* .profiler = */ nullptr, }; return openvino_backend; diff --git a/ggml/src/ggml-profiler.cpp b/ggml/src/ggml-profiler.cpp new file mode 100644 index 0000000000..71024dbf64 --- /dev/null +++ b/ggml/src/ggml-profiler.cpp @@ -0,0 +1,121 @@ +#include "ggml-profiler.h" + +#include "ggml-backend-impl.h" +#include "ggml-impl.h" +#include "ggml.h" + +#include +#include + +#ifdef _WIN32 +# define WIN32_LEAN_AND_MEAN +# ifndef NOMINMAX +# define NOMINMAX +# endif +# include +#else +# include +# include +#endif + +// +// Time utilities +// + +uint64_t ggml_profiler_time_ns(void) { +#ifdef _WIN32 + LARGE_INTEGER freq, count; + QueryPerformanceFrequency(&freq); + QueryPerformanceCounter(&count); + return (uint64_t) (count.QuadPart * 1000000000ULL / freq.QuadPart); +#elif defined(CLOCK_MONOTONIC_RAW) + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC_RAW, &ts); + return (uint64_t) ts.tv_sec * 1000000000ULL + (uint64_t) ts.tv_nsec; +#else + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC, &ts); + return (uint64_t) ts.tv_sec * 1000000000ULL + (uint64_t) ts.tv_nsec; +#endif +} + +// +// Record helpers +// + +void ggml_profile_record_from_tensor(ggml_profile_record * rec, const struct ggml_tensor * node) { + if (rec == NULL) { + return; + } + + // Output tensor info + if (node != NULL) { + memcpy(rec->ne, node->ne, sizeof(rec->ne)); + rec->out_type = (int) node->type; + memcpy(rec->op_params, node->op_params, sizeof(rec->op_params)); + // Copy the tensor name (rather than aliasing node->name) so the record stays + // valid after the graph's meta context is reused on the next build. + snprintf(rec->tensor_name, sizeof(rec->tensor_name), "%s", node->name); + } else { + memset(rec->ne, 0, sizeof(rec->ne)); + rec->out_type = -1; + memset(rec->op_params, 0, sizeof(rec->op_params)); + rec->tensor_name[0] = '\0'; + } + + // Sub-op (UNARY/GLU) + rec->sub_op = -1; + if (node != NULL) { + if (node->op == GGML_OP_UNARY) { + rec->sub_op = (int) ggml_get_unary_op(node); + } else if (node->op == GGML_OP_GLU) { + rec->sub_op = (int) ggml_get_glu_op(node); + } + } + + // Source tensors + rec->n_src = 0; + for (int i = 0; i < GGML_MAX_SRC; i++) { + const struct ggml_tensor * src = (node != NULL) ? node->src[i] : NULL; + if (src == NULL) { + memset(rec->ne_src[i], 0, sizeof(rec->ne_src[i])); + memset(rec->nb_src[i], 0, sizeof(rec->nb_src[i])); + rec->type_src[i] = -1; + } else { + memcpy(rec->ne_src[i], src->ne, sizeof(rec->ne_src[i])); + for (int d = 0; d < 4; d++) { + rec->nb_src[i][d] = (int64_t) src->nb[d]; + } + rec->type_src[i] = (int) src->type; + rec->n_src = i + 1; + } + } +} + +// +// Backend profiler registration +// + +void ggml_backend_set_profiler(ggml_backend_t backend, ggml_backend_profiler_t profiler) { + if (backend == NULL) { + return; + } + + // Free any existing profiler + if (backend->profiler != NULL) { + if (backend->profiler->free_context != NULL) { + backend->profiler->free_context(backend->profiler->context); + } + delete backend->profiler; + backend->profiler = NULL; + } + + backend->profiler = profiler; +} + +ggml_backend_profiler_t ggml_backend_get_profiler(ggml_backend_t backend) { + if (backend == NULL) { + return NULL; + } + return backend->profiler; +} diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp index cc7d720693..e38097ff83 100644 --- a/ggml/src/ggml-rpc/ggml-rpc.cpp +++ b/ggml/src/ggml-rpc/ggml-rpc.cpp @@ -1106,7 +1106,8 @@ ggml_backend_t ggml_backend_rpc_init(const char * endpoint, uint32_t device) { /* .guid = */ ggml_backend_rpc_guid(), /* .iface = */ ggml_backend_rpc_interface, /* .device = */ ggml_backend_reg_dev_get(reg, device), - /* .context = */ ctx + /* .context = */ ctx, + /* .profiler = */ nullptr, }; return backend; } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index bfe6f1016b..9ab815cab1 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -7038,7 +7038,8 @@ ggml_backend_t ggml_backend_sycl_init(int device) { /* .guid = */ ggml_backend_sycl_guid(), /* .iface = */ ggml_backend_sycl_interface, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_sycl_reg(), device), - /* .context = */ ctx + /* .context = */ ctx, + /* .profiler = */ nullptr, }; return sycl_backend; diff --git a/ggml/src/ggml-virtgpu/ggml-backend.cpp b/ggml/src/ggml-virtgpu/ggml-backend.cpp index 996c57e358..f77754a3d8 100644 --- a/ggml/src/ggml-virtgpu/ggml-backend.cpp +++ b/ggml/src/ggml-virtgpu/ggml-backend.cpp @@ -66,6 +66,7 @@ ggml_backend_t ggml_backend_remoting_device_init(ggml_backend_dev_t dev, const c /* .interface = */ ggml_backend_remoting_interface, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_virtgpu_reg(), ctx->device), /* .context = */ ctx, + /* .profiler = */ nullptr, }; return remoting_backend; diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 9b47c6c958..8080c3978d 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1,4 +1,5 @@ #include "ggml-vulkan.h" +#include "ggml-profiler.h" #include #if defined(GGML_VULKAN_RUN_TESTS) || defined(GGML_VULKAN_CHECK_RESULTS) #include @@ -2212,8 +2213,8 @@ private: std::mutex vk_memory_logger::log_mutex; -static bool vk_perf_logger_enabled = false; -static bool vk_perf_logger_concurrent = false; +static bool vk_perf_logger_enabled = false; // deprecated: use --profile instead +static bool vk_perf_logger_concurrent = false; // GGML_VK_PERF_LOGGER_CONCURRENT: use concurrent timestamp mode static bool vk_enable_sync_logger = false; // number of calls between perf logger prints static uint32_t vk_perf_logger_frequency = 1; @@ -2433,6 +2434,21 @@ class vk_perf_logger { uint32_t print_count {}; }; +// Profiler state for the new ggml_backend_profiler interface +struct ggml_vk_profiler_state { + bool enabled = false; + int split_id = -1; + + std::vector records; + std::vector cpu_timestamps; // CPU-side timestamps for global ordering + + void reset() { + records.clear(); + cpu_timestamps.clear(); + split_id = -1; + } +}; + struct ggml_backend_vk_context { std::string name; @@ -2495,8 +2511,9 @@ struct ggml_backend_vk_context { bool fused_topk_qsa {}; rms_norm_mode fused_rms_norm_mode {RMS_NORM_COUNT}; - // for GGML_VK_PERF_LOGGER - std::unique_ptr perf_logger; + // Profiling + std::unique_ptr perf_logger; // legacy env-var profiler + ggml_vk_profiler_state * profiler_state = nullptr; vk::QueryPool query_pool; std::vector query_fusion_names; std::vector query_fusion_node_count; @@ -15897,10 +15914,14 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr ctx->unsynced_nodes_read.clear(); ggml_vk_sync_buffers(ctx, compute_ctx); - if (vk_perf_logger_enabled && vk_perf_logger_concurrent) { + if ((vk_perf_logger_enabled || (ctx->profiler_state != nullptr && ctx->profiler_state->enabled)) + && vk_perf_logger_concurrent) { ctx->query_node_idx[ctx->query_idx] = node_idx; compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); ggml_vk_sync_buffers(ctx, compute_ctx); + if (ctx->profiler_state != nullptr && ctx->profiler_state->enabled) { + ctx->profiler_state->cpu_timestamps.push_back(ggml_profiler_time_ns()); + } } } // Add all fused nodes to the unsynchronized lists. @@ -16480,7 +16501,7 @@ static void ggml_vk_cleanup(ggml_backend_vk_context * ctx) { ctx->transfer_cmd_pool.destroy(ctx->device->device); } - if (vk_perf_logger_enabled) { + if (ctx->perf_logger) { ctx->perf_logger->print_timings(true); } } @@ -17773,7 +17794,9 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ggml_vk_submit_transfer_ctx(ctx); vk_context compute_ctx; - if (vk_perf_logger_enabled) { + bool profiling = vk_perf_logger_enabled || + (ctx->profiler_state != nullptr && ctx->profiler_state->enabled); + if (profiling) { // allocate/resize the query pool if (ctx->num_queries < cgraph->n_nodes + 1) { if (ctx->query_pool) { @@ -17801,6 +17824,10 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->query_idx = 0; compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); ggml_vk_sync_buffers(ctx, compute_ctx); + if (ctx->profiler_state != nullptr && ctx->profiler_state->enabled) { + ctx->profiler_state->cpu_timestamps.clear(); + ctx->profiler_state->cpu_timestamps.push_back(ggml_profiler_time_ns()); + } } ctx->prealloc_y_last_pipeline_used = nullptr; @@ -18142,7 +18169,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg bool enqueued = ggml_vk_build_graph(ctx, cgraph, i, cgraph->nodes[submit_node_idx], submit_node_idx, i + ctx->num_additional_fused_ops >= last_node, almost_ready, submit); - if (vk_perf_logger_enabled && enqueued) { + if (profiling && enqueued) { compute_ctx = ggml_vk_get_compute_ctx(ctx); if (!vk_perf_logger_concurrent) { // track a single node/fusion for the current query @@ -18150,6 +18177,9 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->query_fusion_names[ctx->query_idx] = fusion_string; compute_ctx->s->buffer->buf.writeTimestamp(vk::PipelineStageFlagBits::eAllCommands, ctx->query_pool, ctx->query_idx++); ggml_vk_sync_buffers(ctx, compute_ctx); + if (ctx->profiler_state != nullptr && ctx->profiler_state->enabled) { + ctx->profiler_state->cpu_timestamps.push_back(ggml_profiler_time_ns()); + } } else { // track a fusion string and number of fused ops for the current node_idx ctx->query_fusion_names[i] = fusion_string; @@ -18177,7 +18207,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->last_total_flops = total_flops; - if (vk_perf_logger_enabled) { + if (profiling) { // End the command buffer and submit/wait GGML_ASSERT(!ctx->compute_ctx.expired()); compute_ctx = ctx->compute_ctx.lock(); @@ -18191,15 +18221,40 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg // Get the results and pass them to the logger std::vector timestamps(cgraph->n_nodes + 1); VK_CHECK(ctx->device->device.getQueryPoolResults(ctx->query_pool, 0, ctx->query_idx, (cgraph->n_nodes + 1)*sizeof(uint64_t), timestamps.data(), sizeof(uint64_t), vk::QueryResultFlagBits::e64 | vk::QueryResultFlagBits::eWait), "get timestamp results", ctx->device); + + const double ts_period = ctx->device->properties.limits.timestampPeriod; + const bool has_profiler = ctx->profiler_state != nullptr && ctx->profiler_state->enabled; + if (!vk_perf_logger_concurrent) { // Log each op separately for (int i = 1; i < ctx->query_idx; i++) { auto node = ctx->query_nodes[i]; auto name = ctx->query_fusion_names[i]; - ctx->perf_logger->log_timing(node, name, uint64_t((timestamps[i] - timestamps[i-1]) * ctx->device->properties.limits.timestampPeriod)); + uint64_t duration_ns = uint64_t((timestamps[i] - timestamps[i-1]) * ts_period); + + if (ctx->perf_logger) { + ctx->perf_logger->log_timing(node, name, duration_ns); + } + + if (has_profiler && node != nullptr) { + uint64_t cpu_ts = (i < (int)ctx->profiler_state->cpu_timestamps.size()) + ? ctx->profiler_state->cpu_timestamps[i] : 0; + + ggml_profile_record rec; + rec.type = GGML_PROFILE_EVENT_OP; + rec.name = ggml_op_name(node->op); + rec.backend_id = -1; + rec.split_id = ctx->profiler_state->split_id; + rec.start_ns = cpu_ts; + rec.end_ns = cpu_ts + duration_ns; + rec.bytes = ggml_nbytes(node); + rec.extra = name; // fusion name or NULL + ggml_profile_record_from_tensor(&rec, node); + ctx->profiler_state->records.push_back(rec); + } } } else { - // Log each group of nodes + // Log each group of nodes (concurrent mode) int prev_node_idx = 0; for (int i = 1; i < ctx->query_idx; i++) { auto cur_node_idx = ctx->query_node_idx[i]; @@ -18214,10 +18269,40 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg node_idx += ctx->query_fusion_node_count[node_idx]; } prev_node_idx = cur_node_idx; - ctx->perf_logger->log_timing(nodes, names, uint64_t((timestamps[i] - timestamps[i-1]) * ctx->device->properties.limits.timestampPeriod)); + uint64_t duration_ns = uint64_t((timestamps[i] - timestamps[i-1]) * ts_period); + + if (ctx->perf_logger) { + ctx->perf_logger->log_timing(nodes, names, duration_ns); + } + + if (has_profiler && !nodes.empty()) { + uint64_t cpu_ts = (i < (int)ctx->profiler_state->cpu_timestamps.size()) + ? ctx->profiler_state->cpu_timestamps[i] : 0; + // In concurrent mode, report the group as a single combined operation + auto * node = nodes[0]; + + uint64_t total_bytes = 0; + for (size_t j = 0; j < nodes.size(); j++) { + total_bytes += ggml_nbytes(nodes[j]); + } + + ggml_profile_record rec; + rec.type = GGML_PROFILE_EVENT_OP; + rec.name = ggml_op_name(node->op); + rec.backend_id = -1; + rec.split_id = ctx->profiler_state->split_id; + rec.start_ns = cpu_ts; + rec.end_ns = cpu_ts + duration_ns; + rec.bytes = total_bytes; + rec.extra = names[0]; // fusion name of first op, or NULL + ggml_profile_record_from_tensor(&rec, node); + ctx->profiler_state->records.push_back(rec); + } } } - ctx->perf_logger->print_timings(); + if (ctx->perf_logger) { + ctx->perf_logger->print_timings(); + } } if (!ctx->device->support_async) { @@ -18640,13 +18725,66 @@ ggml_backend_t ggml_backend_vk_init(size_t dev_num) { /* .guid = */ ggml_backend_vk_guid(), /* .iface = */ ggml_backend_vk_interface, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_vk_reg(), dev_num), - /* .context = */ ctx, + /* .context = */ ctx, + /* .profiler = */ nullptr, }; if (!ctx->device->support_async) { vk_backend->iface.get_tensor_async = nullptr; } + // Register profiler + auto * prof_state = new ggml_vk_profiler_state(); + ctx->profiler_state = prof_state; + + static auto vk_prof_enable = [](void * context, bool enable) { + auto * vk_ctx = (ggml_backend_vk_context *)context; + if (vk_ctx->profiler_state != nullptr) { + vk_ctx->profiler_state->enabled = enable; + if (!enable) { + vk_ctx->profiler_state->reset(); + } + } + }; + static auto vk_prof_reset = [](void * context) { + auto * vk_ctx = (ggml_backend_vk_context *)context; + if (vk_ctx->profiler_state != nullptr) { + vk_ctx->profiler_state->reset(); + } + }; + static auto vk_prof_set_split_id = [](void * context, int split_id) { + auto * vk_ctx = (ggml_backend_vk_context *)context; + if (vk_ctx->profiler_state != nullptr) { + vk_ctx->profiler_state->split_id = split_id; + } + }; + static auto vk_prof_get_records = [](void * context, const ggml_profile_record ** out) -> int { + auto * vk_ctx = (ggml_backend_vk_context *)context; + if (vk_ctx->profiler_state != nullptr) { + *out = vk_ctx->profiler_state->records.data(); + return (int)vk_ctx->profiler_state->records.size(); + } + *out = nullptr; + return 0; + }; + static auto vk_prof_free = [](void * context) { + auto * vk_ctx = (ggml_backend_vk_context *)context; + if (vk_ctx->profiler_state != nullptr) { + delete vk_ctx->profiler_state; + vk_ctx->profiler_state = nullptr; + } + }; + + auto * profiler = new ggml_backend_profiler { + /* .context = */ ctx, + /* .enable = */ vk_prof_enable, + /* .reset = */ vk_prof_reset, + /* .set_split_id = */ vk_prof_set_split_id, + /* .get_records = */ vk_prof_get_records, + /* .free_context = */ vk_prof_free, + }; + ggml_backend_set_profiler(vk_backend, profiler); + return vk_backend; } diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl index a19c7f2f4e..8fd8b895eb 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl @@ -6,6 +6,7 @@ #extension GL_EXT_shader_explicit_arithmetic_types_int16 : require #extension GL_EXT_shader_explicit_arithmetic_types_int8 : require #extension GL_EXT_shader_16bit_storage : require +#extension GL_EXT_control_flow_attributes : require #ifdef USE_OCP_FP4 #extension GL_EXT_float_e2m1 : require @@ -1851,7 +1852,9 @@ shared FLOAT_TYPE kvalues_iq4nl[16]; void init_iq_shmem(uvec3 wgsize) { // copy the table into shared memory and sync - for (uint i = gl_LocalInvocationIndex.x; i < kvalues_iq4nl.length(); i += wgsize.x) { + // Use gl_LocalInvocationIndex (uint scalar) directly - each thread loads its portion + // with stride equal to workgroup size for coalesced shared memory writes + [[unroll]] for (uint i = gl_LocalInvocationIndex; i < kvalues_iq4nl.length(); i += wgsize.x) { kvalues_iq4nl[i] = FLOAT_TYPE(kvalues_iq4nl_const[i]); } barrier(); @@ -1891,11 +1894,12 @@ float ue4m3_to_fp32_build(uint u) { void init_iq_shmem(uvec3 wgsize) { // copy the table into shared memory and sync - for (uint i = gl_LocalInvocationIndex.x; i < kvalues_mxfp4.length(); i += wgsize.x) { + // Use gl_LocalInvocationIndex (uint scalar) directly for thread-strided initialization + [[unroll]] for (uint i = gl_LocalInvocationIndex; i < kvalues_mxfp4.length(); i += wgsize.x) { kvalues_mxfp4[i] = kvalues_mxfp4_const[i]; } #if defined(DATA_A_NVFP4) - for (uint i = gl_LocalInvocationIndex.x; i < 128u; i += wgsize.x) { + [[unroll]] for (uint i = gl_LocalInvocationIndex; i < 128u; i += wgsize.x) { ue4m3_fp32_lut[i] = ue4m3_to_fp32_build(i); } #endif diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 1a43c72733..55377c996b 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -4208,6 +4208,7 @@ static ggml_backend_t ggml_backend_webgpu_backend_init(ggml_backend_dev_t dev, c /* .interface = */ ggml_backend_webgpu_i, /* .device = */ dev, /* .context = */ backend_ctx, + /* .profiler = */ nullptr, }; return backend; } diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp index 4007ac9dfc..cb7d4f6062 100644 --- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp +++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp @@ -504,7 +504,8 @@ static ggml_backend_t ggml_backend_zdnn_device_init(ggml_backend_dev_t dev, cons /* .guid = */ ggml_backend_zdnn_guid(), /* .iface = */ ggml_backend_zdnn_i, /* .device = */ dev, - /* .context = */ ctx + /* .context = */ ctx, + /* .profiler = */ NULL, }; return backend; diff --git a/ggml/src/ggml-zendnn/ggml-zendnn.cpp b/ggml/src/ggml-zendnn/ggml-zendnn.cpp index ec7ce23314..26ff2498ee 100644 --- a/ggml/src/ggml-zendnn/ggml-zendnn.cpp +++ b/ggml/src/ggml-zendnn/ggml-zendnn.cpp @@ -597,7 +597,8 @@ ggml_backend_t ggml_backend_zendnn_init(void) { /* .guid = */ ggml_backend_zendnn_guid(), /* .iface = */ ggml_backend_zendnn_i, /* .device = */ ggml_backend_reg_dev_get(ggml_backend_zendnn_reg(), 0), - /* .context = */ ctx, + /* .context = */ ctx, + /* .profiler = */ nullptr, }; return backend; diff --git a/include/llama.h b/include/llama.h index ef7a012c43..e958a68378 100644 --- a/include/llama.h +++ b/include/llama.h @@ -582,6 +582,8 @@ extern "C" { LLAMA_API llama_memory_t llama_get_memory (const struct llama_context * ctx); LLAMA_API enum llama_pooling_type llama_pooling_type(const struct llama_context * ctx); // TODO: rename to llama_get_pooling_type + LLAMA_API struct ggml_backend_sched * llama_context_get_sched(const struct llama_context * ctx); + LLAMA_API const struct llama_vocab * llama_model_get_vocab(const struct llama_model * model); LLAMA_API enum llama_rope_type llama_model_rope_type(const struct llama_model * model); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index c1ef12f56b..ee7e8cc391 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -1,6 +1,7 @@ #include "llama-context.h" #include "ggml.h" +#include "ggml-profiler.h" #include "llama-arch.h" #include "llama-graph.h" #include "llama-impl.h" @@ -2513,6 +2514,11 @@ ggml_status llama_context::graph_compute( LLAMA_LOG_ERROR("%s: ggml_backend_sched_graph_compute_async failed with error %d\n", __func__, status); } + // If profiling is enabled, synchronize to ensure records are complete + if (ggml_backend_sched_get_profiling(sched.get())) { + ggml_backend_sched_synchronize(sched.get()); + } + // fprintf(stderr, "splits: %d\n", ggml_backend_sched_get_n_splits(sched)); return status; @@ -3804,6 +3810,10 @@ enum llama_pooling_type llama_pooling_type(const llama_context * ctx) { return ctx->pooling_type(); } +ggml_backend_sched_t llama_context_get_sched(const llama_context * ctx) { + return ctx->get_sched(); +} + void llama_attach_threadpool( llama_context * ctx, ggml_threadpool_t threadpool, diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 6a8cfd46d0..af17c05986 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -11160,7 +11160,7 @@ static std::vector> make_test_cases_perf() { return test_cases; } -static std::vector> make_test_cases_from_file(const char * path) { +static std::vector> make_test_cases_from_file(const char * path, const char * backend_name = nullptr) { std::ifstream f(path); if (!f.is_open()) { @@ -11218,6 +11218,17 @@ static std::vector> make_test_cases_from_file(const c name = ""; } + std::string file_backend; + if (iss >> file_backend) { + if (file_backend.length() == 1 && file_backend[0] == '-') { + file_backend = ""; + } + } + + if (backend_name != nullptr && !file_backend.empty() && file_backend != backend_name) { + continue; + } + test_cases.emplace_back(new test_generic_op(op, type, ne, op_params, sources, std::move(name))); } @@ -11352,7 +11363,7 @@ static bool test_backend(ggml_backend_t backend, ggml_backend_dev_t dev, test_mo break; } } else { - test_cases = make_test_cases_from_file(test_file_path); + test_cases = make_test_cases_from_file(test_file_path, ggml_backend_name(backend)); } filter_test_cases(test_cases, params_filter); @@ -11617,12 +11628,13 @@ static void usage(char ** argv) { printf(" --output specifies output format (default: console, options: console, sql, csv)\n"); printf(" --list-ops lists all available GGML operations\n"); printf(" --show-coverage shows test coverage\n"); - printf(" --test-file reads test operators from a test file generated by test-export-graph-ops\n"); + printf(" --test-file reads test operators from a test file generated by llama-export-graph-ops or the profiler\n"); printf(" -j runs tests using parallel worker threads (default: 1, test mode only)\n"); } int main(int argc, char ** argv) { test_mode mode = MODE_TEST; + bool mode_explicit = false; output_formats output_format = CONSOLE; const char * op_names_filter = nullptr; const char * backend_filter = nullptr; @@ -11633,12 +11645,16 @@ int main(int argc, char ** argv) { for (int i = 1; i < argc; i++) { if (strcmp(argv[i], "test") == 0) { mode = MODE_TEST; + mode_explicit = true; } else if (strcmp(argv[i], "perf") == 0) { mode = MODE_PERF; + mode_explicit = true; } else if (strcmp(argv[i], "grad") == 0) { mode = MODE_GRAD; + mode_explicit = true; } else if (strcmp(argv[i], "support") == 0) { mode = MODE_SUPPORT; + mode_explicit = true; } else if (strcmp(argv[i], "-o") == 0) { if (i + 1 < argc) { op_names_filter = argv[++i]; @@ -11700,6 +11716,10 @@ int main(int argc, char ** argv) { } } + if (test_file_path != nullptr && !mode_explicit) { + mode = MODE_PERF; + } + // load and enumerate backends ggml_backend_load_all(); @@ -11723,7 +11743,8 @@ int main(int argc, char ** argv) { continue; } - if (backend_filter == NULL && ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU && mode != MODE_GRAD) { + if (backend_filter == NULL && test_file_path == NULL && + ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU && mode != MODE_GRAD) { output_printer->print_backend_init(backend_init_info( i, ggml_backend_dev_count(), ggml_backend_dev_name(dev), true, "Skipping CPU backend")); n_ok++; @@ -11740,7 +11761,7 @@ int main(int argc, char ** argv) { ggml_backend_set_n_threads_fn(backend.get(), N_THREADS); } - size_t free, total; // NOLINT + size_t free, total; ggml_backend_dev_memory(dev, &free, &total); output_printer->print_backend_init(backend_init_info(i, ggml_backend_dev_count(), ggml_backend_dev_name(dev), false, "", ggml_backend_dev_description(dev), diff --git a/tests/test-export-graph-ops.cpp b/tests/test-export-graph-ops.cpp index 46ded13985..3820810c3d 100644 --- a/tests/test-export-graph-ops.cpp +++ b/tests/test-export-graph-ops.cpp @@ -5,7 +5,6 @@ #include "../src/llama-ext.h" #include "ggml.h" #include "gguf-model-data.h" -#include "gguf.h" #include "ggml-backend.h" #include "download.h" @@ -14,7 +13,6 @@ #include #include #include -#include // Noop because weights are not needed static void set_tensor_data(struct ggml_tensor * tensor, void * userdata) { @@ -55,6 +53,7 @@ struct test_object { std::vector op_params; std::vector sources; std::string name; + std::string backend_name; void serialize(std::ostream& out) const { out << op << ' ' << type << ' '; @@ -78,16 +77,21 @@ struct test_object { out << '-'; } + if (!backend_name.empty()) { + out << ' ' << backend_name; + } + out << '\n'; } bool operator<(const test_object &b) const { - return std::tie(op, type, ne, op_params, sources) < - std::tie(b.op, b.type, b.ne, b.op_params, b.sources); + return std::tie(op, type, ne, op_params, sources, backend_name) < + std::tie(b.op, b.type, b.ne, b.op_params, b.sources, b.backend_name); } }; -static void extract_graph_ops(ggml_cgraph * cgraph, const char * label, std::set & tests) { +static void extract_graph_ops(ggml_cgraph * cgraph, const char * label, std::set & tests, + ggml_backend_sched_t sched = nullptr) { int n_nodes = ggml_graph_n_nodes(cgraph); int n_skipped = 0; int n_before = (int) tests.size(); @@ -117,6 +121,14 @@ static void extract_graph_ops(ggml_cgraph * cgraph, const char * label, std::set } test.name = node->name; + + if (sched) { + ggml_backend_t backend = ggml_backend_sched_get_tensor_backend(sched, node); + if (backend) { + test.backend_name = ggml_backend_name(backend); + } + } + tests.insert(test); } @@ -135,11 +147,12 @@ int main(int argc, char ** argv) { return 1; } - // Load CPU-only - ggml_backend_dev_t cpu_device = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); - params.devices = { cpu_device, nullptr }; - params.fit_params = false; - params.n_gpu_layers = 0; + if (!params.with_backends) { + ggml_backend_dev_t cpu_device = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); + params.devices = { cpu_device, nullptr }; + params.fit_params = false; + params.n_gpu_layers = 0; + } params.warmup = false; @@ -199,19 +212,21 @@ int main(int argc, char ** argv) { std::set tests; + ggml_backend_sched_t sched = params.with_backends ? llama_context_get_sched(ctx) : nullptr; + auto * gf_pp = llama_graph_reserve(ctx, n_tokens, n_seqs, n_tokens); if (!gf_pp) { LOG_ERR("failed to reserve prompt processing graph\n"); return 1; } - extract_graph_ops(gf_pp, "pp", tests); + extract_graph_ops(gf_pp, "pp", tests, sched); auto * gf_tg = llama_graph_reserve(ctx, n_seqs, n_seqs, n_seqs); if (!gf_tg) { LOG_ERR("failed to reserve token generation graph\n"); return 1; } - extract_graph_ops(gf_tg, "tg", tests); + extract_graph_ops(gf_tg, "tg", tests, sched); LOG_INF("%d unique ops total\n", (int) tests.size()); diff --git a/tools/completion/completion.cpp b/tools/completion/completion.cpp index 941b7399b2..4517d75b30 100644 --- a/tools/completion/completion.cpp +++ b/tools/completion/completion.cpp @@ -947,6 +947,29 @@ int llama_completion(int argc, char ** argv) { } + // Export profiling data if profiling was enabled + if (params.profiling) { + ggml_backend_sched_t sched = llama_context_get_sched(ctx); + if (sched != nullptr) { + if (params.profiling_output.empty()) { + ggml_backend_sched_print_profiling(sched); + } else { + const std::string & path = params.profiling_output; + int ret; + if (path.size() >= 4 && path.compare(path.size() - 4, 4, ".txt") == 0) { + ret = ggml_backend_sched_export_profiling_text(sched, path.c_str()); + } else { + ret = ggml_backend_sched_export_profiling_json(sched, path.c_str()); + } + if (ret == 0) { + LOG("\nProfiling data exported to: %s\n", path.c_str()); + } else { + LOG_ERR("\nFailed to export profiling data to: %s\n", path.c_str()); + } + } + } + } + LOG("\n\n"); common_perf_print(ctx, smpl); diff --git a/tools/profiler/__init__.py b/tools/profiler/__init__.py new file mode 100644 index 0000000000..cdd8dd543e --- /dev/null +++ b/tools/profiler/__init__.py @@ -0,0 +1 @@ +# llama.cpp profiler analysis tools diff --git a/tools/profiler/profiler.py b/tools/profiler/profiler.py new file mode 100644 index 0000000000..7bb7d73722 --- /dev/null +++ b/tools/profiler/profiler.py @@ -0,0 +1,1362 @@ +#!/usr/bin/env python3 +"""llama.cpp cross-backend profiler analysis tool. + +Usage: + python -m tools.profiler.profiler profile.json + python -m tools.profiler.profiler profile.json --chrome-trace trace.json +""" + +from __future__ import annotations + +import json +import sys +from dataclasses import dataclass, field +from pathlib import Path +from typing import Optional + + +OP_EVENT = 0 +COPY_EVENT = 1 + +TYPE_NAMES = {0: "OP", 1: "COPY"} + +GGML_TYPE_NAMES = { + 0: "F32", 1: "F16", 2: "Q4_0", 3: "Q4_1", + # 4, 5 removed + 6: "Q5_0", 7: "Q5_1", 8: "Q8_0", 9: "Q8_1", + 10: "Q2_K", 11: "Q3_K", 12: "Q4_K", 13: "Q5_K", + 14: "Q6_K", 15: "Q8_K", 16: "IQ2_XXS", 17: "IQ2_XS", + 18: "IQ3_XXS", 19: "IQ1_S", 20: "IQ4_NL", 21: "IQ3_S", + 22: "IQ2_S", 23: "IQ4_XS", 24: "I8", 25: "I16", 26: "I32", + 27: "I64", 28: "F64", 29: "IQ1_M", 30: "BF16", + # 31-33 removed + 34: "TQ1_0", 35: "TQ2_0", + # 36-38 removed + 39: "MXFP4", 40: "NVFP4", +} + +GGML_UNARY_OP_NAMES = { + 0: "ABS", 1: "SGN", 2: "NEG", 3: "STEP", + 4: "TANH", 5: "ELU", 6: "RELU", 7: "SIGMOID", + 8: "GELU", 9: "GELU_QUICK", 10: "SILU", 11: "HARDSWISH", + 12: "HARDSIGMOID", 13: "EXP", 14: "EXPM1", 15: "SOFTPLUS", + 16: "GELU_ERF", 17: "XIELU", 18: "FLOOR", 19: "CEIL", + 20: "ROUND", 21: "TRUNC", +} + +GGML_GLU_OP_NAMES = { + 0: "REGLU", 1: "GEGLU", 2: "SWIGLU", 3: "GEGLU_ERF", + 4: "GEGLU_QUICK", 5: "SWIGLU_OAI", +} + +GGML_OP_NAMES = { + 0: "NONE", 1: "DUP", 2: "ADD", 3: "ADD_ID", 4: "ADD1", + 5: "ACC", 6: "SUB", 7: "MUL", 8: "DIV", 9: "SQR", + 10: "SQRT", 11: "LOG", 12: "SIN", 13: "COS", 14: "SUM", + 15: "SUM_ROWS", 16: "CUMSUM", 17: "MEAN", 18: "ARGMAX", + 19: "COUNT_EQUAL", 20: "REPEAT", 21: "REPEAT_BACK", 22: "CONCAT", + 23: "SILU_BACK", 24: "NORM", 25: "RMS_NORM", 26: "RMS_NORM_BACK", + 27: "GROUP_NORM", 28: "L2_NORM", 29: "MUL_MAT", 30: "MUL_MAT_ID", + 31: "OUT_PROD", 32: "SCALE", 33: "SET", 34: "CPY", 35: "CONT", + 36: "RESHAPE", 37: "VIEW", 38: "PERMUTE", 39: "TRANSPOSE", + 40: "GET_ROWS", 41: "GET_ROWS_BACK", 42: "SET_ROWS", 43: "DIAG", + 44: "DIAG_MASK_INF", 45: "DIAG_MASK_ZERO", 46: "SOFT_MAX", + 47: "SOFT_MAX_BACK", 48: "ROPE", 49: "ROPE_BACK", 50: "CLAMP", + 51: "CONV_TRANSPOSE_1D", 52: "IM2COL", 53: "IM2COL_BACK", 54: "IM2COL_3D", + 55: "CONV_2D", 56: "CONV_3D", 57: "CONV_2D_DW", 58: "CONV_TRANSPOSE_2D", + 59: "POOL_1D", 60: "POOL_2D", 61: "POOL_2D_BACK", 62: "UPSCALE", + 63: "PAD", 64: "PAD_REFLECT_1D", 65: "ROLL", 66: "ARANGE", + 67: "TIMESTEP_EMBEDDING", 68: "ARGSORT", 69: "TOP_K", 70: "LEAKY_RELU", + 71: "TRI", 72: "FILL", 73: "FLASH_ATTN_EXT", 74: "FLASH_ATTN_BACK", + 75: "SSM_CONV", 76: "SSM_SCAN", 77: "WIN_PART", 78: "WIN_UNPART", + 79: "GET_REL_POS", 80: "ADD_REL_POS", 81: "RWKV_WKV6", + 82: "GATED_LINEAR_ATTN", 83: "RWKV_WKV7", 84: "SOLVE_TRI", + 85: "GATED_DELTA_NET", 86: "UNARY", 87: "MAP_CUSTOM1", + 88: "MAP_CUSTOM2", 89: "MAP_CUSTOM3", 90: "CUSTOM", + 91: "CROSS_ENTROPY_LOSS", 92: "CROSS_ENTROPY_LOSS_BACK", + 93: "OPT_STEP_ADAMW", 94: "OPT_STEP_SGD", 95: "GLU", + 96: "COUNT", +} + +GGML_TYPE_NAMES_TO_ID = {v: k for k, v in GGML_TYPE_NAMES.items()} + +GGML_OP_NAMES_TO_ID = {v: k for k, v in GGML_OP_NAMES.items()} + + +_EXPORT_SKIP_OPS = frozenset({ + 33, # SET + 34, # CPY + 35, # CONT + 36, # RESHAPE + 37, # VIEW + 38, # PERMUTE + 39, # TRANSPOSE + 41, # GET_ROWS_BACK + 42, # SET_ROWS + 43, # DIAG + 44, # DIAG_MASK_INF + 45, # DIAG_MASK_ZERO + 47, # SOFT_MAX_BACK + 49, # ROPE_BACK + 51, # CONV_TRANSPOSE_1D + 52, # IM2COL + 53, # IM2COL_BACK + 54, # IM2COL_3D + 58, # CONV_TRANSPOSE_2D + 61, # POOL_2D_BACK + 63, # PAD + 64, # PAD_REFLECT_1D + 65, # ROLL + 66, # ARANGE + 70, # LEAKY_RELU (covered by UNARY) + 71, # TRI + 72, # FILL + 77, # WIN_PART + 78, # WIN_UNPART + 92, # CROSS_ENTROPY_LOSS_BACK + 93, # OPT_STEP_ADAMW + 94, # OPT_STEP_SGD + 96, # COUNT +}) + + +def _compute_output_ne(op_id: int, ne0: list, ne1: list, ne2: list) -> list | None: + if op_id == 29: # MUL_MAT + return [ne0[1], ne1[1], max(ne0[2], ne1[2]), max(ne0[3], ne1[3])] + if op_id == 30: # MUL_MAT_ID + return [ne0[1], ne2[0], ne1[2], 1] + if op_id in (2, 7, 8, 32): # ADD, MUL, DIV, SCALE + return [max(ne0[i], ne1[i]) for i in range(4)] + if op_id == 4: # ADD1 + return list(ne0) + if op_id in (9, 86): # SQR, UNARY + return list(ne0) + if op_id in (46, 25): # SOFT_MAX, RMS_NORM + return list(ne0) + if op_id == 73: # FLASH_ATTN_EXT + # Per ggml_flash_attn_ext: result.ne = { v->ne[0], q->ne[2], q->ne[1], q->ne[3] } + # When V was not captured (legacy records), fall back to hsk == hsv (q->ne[0]). + hsv = ne2[0] if (ne2 and ne2[0] > 0) else ne0[0] + return [hsv, ne0[2], ne0[1], ne0[3]] + if op_id == 40: # GET_ROWS + return [ne0[0], ne1[1], ne1[2], ne1[3]] + if op_id == 41: # GET_ROWS_BACK + return [ne0[0], ne1[1], ne1[2], ne1[3]] + if op_id == 42: # SET_ROWS + return list(ne0) + if op_id == 31: # OUT_PROD + return [ne0[0], ne1[0], max(ne0[2], ne1[2]), max(ne0[3], ne1[3])] + if op_id == 22: # CONCAT + return [ne0[0] + ne1[0], max(ne0[1], ne1[1]), + max(ne0[2], ne1[2]), max(ne0[3], ne1[3])] + if op_id in (34, 35, 50, 60, 53, 68): # CPY, CONT, CLAMP, POOL_2D, IM2COL_BACK, ARGSORT + return list(ne0) + return None + + +GGML_MAX_SRC = 10 +GGML_MAX_OP_PARAMS_I32 = 16 # 64 bytes / sizeof(int32) + + +@dataclass +class ProfileRecord: + type: int + name: str + backend_id: int + split_id: int + start_ns: int + duration_ns: int + bytes: int + extra: Optional[str] + # Output tensor info + tensor_name: Optional[str] = None + ne: list[int] = field(default_factory=lambda: [0, 0, 0, 0]) + out_type: int = -1 + # Source tensors (variable length, up to GGML_MAX_SRC) + ne_src: list[list[int]] = field(default_factory=list) + nb_src: list[list[int]] = field(default_factory=list) + type_src: list[int] = field(default_factory=list) + # Operation parameters (16 int32, raw from ggml_tensor::op_params) + op_params: list[int] = field(default_factory=lambda: [0] * GGML_MAX_OP_PARAMS_I32) + sub_op: int = -1 + + # --- Convenience accessors mirroring the old API --- + @property + def ne_src0(self) -> list[int]: + return self.ne_src[0] if len(self.ne_src) > 0 else [0, 0, 0, 0] + + @property + def ne_src1(self) -> list[int]: + return self.ne_src[1] if len(self.ne_src) > 1 else [0, 0, 0, 0] + + @property + def ne_src2(self) -> list[int]: + return self.ne_src[2] if len(self.ne_src) > 2 else [0, 0, 0, 0] + + @property + def type_src0(self) -> int: + return self.type_src[0] if len(self.type_src) > 0 else -1 + + @property + def type_src1(self) -> int: + return self.type_src[1] if len(self.type_src) > 1 else -1 + + @property + def type_src2(self) -> int: + return self.type_src[2] if len(self.type_src) > 2 else -1 + + @property + def sub_op_name(self) -> str: + if self.sub_op < 0: + return "" + if self.name == "UNARY": + return GGML_UNARY_OP_NAMES.get(self.sub_op, f"UNARY_OP({self.sub_op})") + if self.name == "GLU": + return GGML_GLU_OP_NAMES.get(self.sub_op, f"GLU_OP({self.sub_op})") + return str(self.sub_op) + + @property + def type_name(self) -> str: + return TYPE_NAMES.get(self.type, f"UNKNOWN({self.type})") + + @property + def duration_us(self) -> float: + return self.duration_ns / 1000.0 + + @property + def duration_ms(self) -> float: + return self.duration_ns / 1_000_000.0 + + @property + def bandwidth_gbps(self) -> float: + """Bandwidth in GB/s.""" + if self.duration_ns == 0 or self.bytes == 0: + return 0.0 + return self.bytes / self.duration_ns + + @staticmethod + def _fmt_ne(ne: list[int]) -> str: + dims = [n for n in ne if n > 0] + if not dims: + return "" + return "[" + ", ".join(str(d) for d in dims) + "]" + + @property + def shape_str(self) -> str: + """Human-readable tensor shapes, e.g. '[4096, 4096] x [4096, 1] x [8, 1]'.""" + parts = [] + for ne, gt in zip(self.ne_src, self.type_src): + s = self._fmt_ne(ne) + if s: + type_name = GGML_TYPE_NAMES.get(gt, None) + if type_name: + s = f"{s} ({type_name})" + parts.append(s) + result = " x ".join(parts) + if self.sub_op_name: + result = f"[{self.sub_op_name}] {result}" + return result + + def to_dict(self) -> dict: + return { + "type": self.type, + "name": self.name, + "backend_id": self.backend_id, + "split_id": self.split_id, + "start_ns": self.start_ns, + "duration_ns": self.duration_ns, + "bytes": self.bytes, + "extra": self.extra, + "tensor_name": self.tensor_name, + "ne": self.ne, + "out_type": self.out_type, + "n_src": len(self.ne_src), + "ne_src": self.ne_src, + "nb_src": self.nb_src, + "type_src": self.type_src, + "op_params": self.op_params, + "sub_op": self.sub_op, + } + + +@dataclass +class OpStats: + name: str + event_type: int + backend_id: int + count: int = 0 + total_ns: int = 0 + min_ns: int = 0 + max_ns: int = 0 + total_bytes: int = 0 + representative_ne: list[int] = field(default_factory=lambda: [0, 0, 0, 0]) + + @property + def avg_ns(self) -> float: + return self.total_ns / self.count if self.count > 0 else 0 + + @property + def avg_us(self) -> float: + return self.avg_ns / 1000.0 + + @property + def total_ms(self) -> float: + return self.total_ns / 1_000_000.0 + + @property + def min_us(self) -> float: + return self.min_ns / 1000.0 + + @property + def max_us(self) -> float: + return self.max_ns / 1000.0 + + @property + def bandwidth_gbps(self) -> float: + if self.total_ns == 0 or self.total_bytes == 0: + return 0.0 + return self.total_bytes / self.total_ns + + @property + def time_per_byte_ns(self) -> float: + """Time per byte (lower = more efficient).""" + if self.total_bytes == 0: + return float("inf") + return self.total_ns / self.total_bytes + + @property + def type_name(self) -> str: + return TYPE_NAMES.get(self.event_type, f"UNKNOWN({self.event_type})") + + +class ProfileData: + def __init__(self, records: list[ProfileRecord], metadata: dict): + self.records = records + self.metadata = metadata + + @classmethod + def load(cls, filepath: str | Path) -> ProfileData: + """Load a profiler JSON file.""" + with open(filepath, "r") as f: + data = json.load(f) + + if data.get("profiler") != "ggml": + print(f"Warning: file may not be a ggml profiler output (profiler={data.get('profiler')})") + + records = [] + def _pad_ne(v): + if isinstance(v, list) and len(v) < 4: + return list(v) + [0] * (4 - len(v)) + if not isinstance(v, list): + return [0, 0, 0, 0] + return list(v) + + def _load_sources(r) -> tuple[list[list[int]], list[list[int]], list[int]]: + """Read source tensor arrays, supporting both v3+ (arrays) and v2 (ne_src0/1/2) JSON.""" + ne_list_raw = r.get("ne_src") + if isinstance(ne_list_raw, list): + # v3+ format + ne_src = [_pad_ne(x) for x in ne_list_raw] + nb_raw = r.get("nb_src", []) + if isinstance(nb_raw, list): + nb_src = [_pad_ne(x) for x in nb_raw] + else: + nb_src = [] + while len(nb_src) < len(ne_src): + nb_src.append([0, 0, 0, 0]) + type_raw = r.get("type_src", []) + if isinstance(type_raw, list): + type_src = [int(t) for t in type_raw] + else: + type_src = [] + while len(type_src) < len(ne_src): + type_src.append(-1) + return ne_src, nb_src[:len(ne_src)], type_src[:len(ne_src)] + + # Legacy v2 fallback + ne_src: list[list[int]] = [] + type_src: list[int] = [] + for i in range(3): + key_ne = f"ne_src{i}" + key_type = f"type_src{i}" + ne_v = r.get(key_ne) + if ne_v is None and i == 0: + ne_v = r.get("ne") + if ne_v is None: + break + ne_padded = _pad_ne(ne_v) + if all(v == 0 for v in ne_padded) and i > 0: + break + ne_src.append(ne_padded) + type_src.append(int(r.get(key_type, -1))) + nb_src = [[0, 0, 0, 0] for _ in ne_src] + return ne_src, nb_src, type_src + + def _load_op_params(r) -> list[int]: + raw = r.get("op_params") + if isinstance(raw, list): + ops = [int(x) for x in raw[:GGML_MAX_OP_PARAMS_I32]] + while len(ops) < GGML_MAX_OP_PARAMS_I32: + ops.append(0) + return ops + return [0] * GGML_MAX_OP_PARAMS_I32 + + for r in data.get("records", []): + ne_src, nb_src, type_src = _load_sources(r) + + # v3+ records have a real "ne" (output shape). Legacy v2 records did not — + # leave zero so export_graph_ops falls back to op-specific shape inference. + ne_raw = r.get("ne") if "ne_src" in r else None + ne_out = _pad_ne(ne_raw) if isinstance(ne_raw, list) and ne_raw else [0, 0, 0, 0] + + records.append(ProfileRecord( + type=r.get("type", 0), + name=r.get("name", "unknown"), + backend_id=r.get("backend_id", 0), + split_id=r.get("split_id", 0), + start_ns=r.get("start_ns", 0), + duration_ns=r.get("duration_ns", 0), + bytes=r.get("bytes", 0), + extra=r.get("extra"), + tensor_name=r.get("tensor_name"), + ne=ne_out, + out_type=int(r.get("out_type", -1)), + ne_src=ne_src, + nb_src=nb_src, + type_src=type_src, + op_params=_load_op_params(r), + sub_op=int(r.get("sub_op", -1)), + )) + + backends_raw = data.get("backends", []) + backends = [] + for b in backends_raw: + backends.append({ + "id": b.get("id", 0), + "name": b.get("name", "unknown"), + "device": b.get("device", "unknown"), + "device_type": b.get("device_type", 0), + }) + + metadata = { + "version": data.get("version", 0), + "total_records": data.get("total_records", len(records)), + "total_ns": data.get("total_ns", sum(r.duration_ns for r in records)), + "backends": backends, + } + + return cls(records, metadata) + + @property + def total_ns(self) -> int: + return sum(r.duration_ns for r in self.records) + + @property + def total_ms(self) -> float: + return self.total_ns / 1_000_000.0 + + def stats(self) -> list[OpStats]: + """Aggregate stats grouped by (name, type, backend_id).""" + groups: dict[tuple, OpStats] = {} + for rec in self.records: + key = (rec.name, rec.type, rec.backend_id) + if key not in groups: + groups[key] = OpStats( + name=rec.name, + event_type=rec.type, + backend_id=rec.backend_id, + min_ns=rec.duration_ns, + max_ns=rec.duration_ns, + representative_ne=list(rec.ne_src0), + ) + s = groups[key] + s.count += 1 + s.total_ns += rec.duration_ns + s.min_ns = min(s.min_ns, rec.duration_ns) + s.max_ns = max(s.max_ns, rec.duration_ns) + s.total_bytes += rec.bytes + + # Track the ne from the longest individual call + if rec.duration_ns >= s.max_ns: + s.representative_ne = list(rec.ne_src0) + + return sorted(groups.values(), key=lambda s: s.total_ns, reverse=True) + + def top_operations(self, n: int = 10) -> list[OpStats]: + """Return the N most time-consuming operations (aggregated).""" + return self.stats()[:n] + + def top_kernels(self, n: int = 10) -> list[ProfileRecord]: + """Return the N longest individual kernel executions.""" + return sorted(self.records, key=lambda r: r.duration_ns, reverse=True)[:n] + + def by_backend(self) -> dict[int, list[ProfileRecord]]: + """Group records by backend ID.""" + groups: dict[int, list[ProfileRecord]] = {} + for rec in self.records: + groups.setdefault(rec.backend_id, []).append(rec) + return dict(sorted(groups.items())) + + def timeline(self) -> list[ProfileRecord]: + """Return records sorted by start_ns for timeline visualization.""" + return sorted(self.records, key=lambda r: r.start_ns) + + def inefficiency_ranking(self, n: int = 10) -> list[OpStats]: + """Rank operations by time per byte (inefficiency). Lower is better.""" + all_stats = [s for s in self.stats() if s.total_bytes > 0 and s.event_type == OP_EVENT] + return sorted(all_stats, key=lambda s: s.time_per_byte_ns, reverse=True)[:n] + + def summary(self) -> None: + """Print a formatted summary table to stdout.""" + print(f"\n{'='*80}") + print(f" ggml Profiler Summary") + print(f"{'='*80}") + print(f" Total records: {len(self.records)}") + print(f" Total time: {self.total_ms:.2f} ms") + print(f" Unique ops: {len(set((r.name, r.type, r.backend_id) for r in self.records))}") + print(f"{'='*80}\n") + + stats = self.stats() + if not stats: + print(" No profiling data.\n") + return + + print(f" {'TYPE':<5} {'BKND':>4} {'Operation':<28} {'%Time':>7} {'Count':>6} " + f"{'Total':>10} {'Avg':>10} {'Min':>10} {'Max':>10} {'Bandwidth':>12}") + print(f" {'':->5} {'':->4} {'':->28} {'':->7} {'':->6} " + f"{'(ms)':>10} {'(us)':>10} {'(us)':>10} {'(us)':>10} {'':->12}") + + for s in stats: + pct = 100.0 * s.total_ns / self.total_ns if self.total_ns > 0 else 0 + + line = (f" {s.type_name:<5} {s.backend_id:>4} {s.name:<28} {pct:>6.1f}% " + f"{s.count:>6} {s.total_ms:>10.2f} {s.avg_us:>10.2f} " + f"{s.min_us:>10.2f} {s.max_us:>10.2f}") + + if s.total_bytes > 0 and s.total_ns > 0: + bw = s.bandwidth_gbps + if bw >= 1000.0: + line += f" {bw / 1000.0:>9.2f} TB/s" + else: + line += f" {bw:>9.2f} GB/s" + else: + line += f" {'':>12}" + + # Tensor shape from longest call + shape_dims = [n for n in s.representative_ne if n > 0] + if shape_dims: + line += f" [{', '.join(str(d) for d in shape_dims)}]" + + print(line) + + backend_groups = self.by_backend() + if len(backend_groups) > 1: + print(f"\n --- By Backend ---") + for bid, recs in sorted(backend_groups.items()): + bk_total = sum(r.duration_ns for r in recs) + bk_pct = 100.0 * bk_total / self.total_ns if self.total_ns > 0 else 0 + print(f" Backend {bid}: {bk_total / 1e6:.2f} ms ({bk_pct:.1f}%) — {len(recs)} records") + + inef = self.inefficiency_ranking(5) + if inef: + print(f"\n --- Top 5 Inefficient Operations (time/byte) ---") + for s in inef: + print(f" {s.name:<28} {s.time_per_byte_ns / 1000:.2f} us/byte " + f"({s.count} calls, {s.total_bytes / 1e6:.1f} MB)") + + top_k = self.top_kernels(5) + print(f"\n --- Top 5 Longest Kernels ---") + for rec in top_k: + shape = f" {rec.shape_str}" if rec.shape_str else "" + tname = f" '{rec.tensor_name}'" if rec.tensor_name else "" + print(f" {rec.type_name:<5} {rec.name:<28}{tname} {rec.duration_us:>10.2f} us{shape} " + f"(split={rec.split_id}, backend={rec.backend_id})") + + print() + + def export_chrome_trace(self, filepath: str | Path) -> None: + """Export as Chrome Trace Event format for chrome://tracing.""" + events = [] + + # Build backend name mapping and remap to non-negative PIDs + # (Chrome cannot handle negative PIDs) + backend_ids = sorted(set(rec.backend_id for rec in self.records)) + backend_names: dict[int, str] = {} + pid_map: dict[int, int] = {} + + # Use metadata from JSON if available + metadata_backends = self.metadata.get("backends", []) + backend_by_id: dict[int, dict] = {b["id"]: b for b in metadata_backends} + + device_type_names = {0: "CPU", 1: "GPU", 2: "ACCEL"} + for idx, bid in enumerate(backend_ids): + pid_map[bid] = idx + if bid in backend_by_id: + binfo = backend_by_id[bid] + dev_type = binfo.get("device_type", 0) + dev_name = binfo.get("device", "") + type_name = device_type_names.get(dev_type, "Device") + if dev_name and dev_name != "unknown": + backend_names[bid] = f"{type_name}: {dev_name}" + else: + backend_names[bid] = f"{type_name}: {binfo.get('name', f'Backend {bid}')}" + else: + backend_names[bid] = f"Backend {bid}" + + # Process metadata events + for bid in backend_ids: + pid = pid_map[bid] + events.append({ + "ph": "M", # metadata + "pid": pid, + "name": "process_name", + "args": {"name": backend_names[bid]}, + }) + + # Use real timestamps, but prevent overlaps within each track. + # GPU kernels are launched rapidly (small start_ns gaps) but have long + # durations, so naive real timestamps overlap. Sweep-line per track: + # sort by start_ns, then place each event at max(start, prev_end). + from collections import defaultdict + tracks: dict[tuple, list[ProfileRecord]] = defaultdict(list) + for rec in self.records: + tracks[(rec.backend_id, rec.split_id)].append(rec) + + for key in tracks: + tracks[key].sort(key=lambda r: r.start_ns) + + for key, recs in tracks.items(): + pid = pid_map[key[0]] + tid = f"split_{key[1]}" + cursor = 0.0 + for rec in recs: + ts = max(rec.start_ns / 1000.0, cursor) + dur = rec.duration_ns / 1000.0 + cat = "copy" if rec.type == COPY_EVENT else "compute" + events.append({ + "ph": "X", # complete event + "pid": pid, + "tid": tid, + "name": rec.name, + "ts": ts, + "dur": dur, + "cat": cat, + "args": { + "bytes": rec.bytes, + "duration_us": dur, + "shape": rec.shape_str, + }, + }) + cursor = ts + dur + + trace = {"traceEvents": events} + with open(filepath, "w") as f: + json.dump(trace, f, indent=2) + + print(f"Chrome trace exported to: {filepath}") + print(f"Open chrome://tracing in Chrome/Edge and load this file.") + + def export_graph_ops(self, filepath: str | Path) -> None: + """Export operations in export-graph-ops format for test-backend-ops --test-file. + + Output line layout matches tests/export-graph-ops.cpp: + + ( )* [] + """ + seen: set[tuple] = set() + lines: list[str] = [] + + backend_by_id: dict[int, dict] = {} + for b in self.metadata.get("backends", []): + backend_by_id[b["id"]] = b + + for rec in self.records: + if rec.type != OP_EVENT: + continue + + op_id = GGML_OP_NAMES_TO_ID.get(rec.name, -1) + if op_id < 0: + continue + + if op_id in _EXPORT_SKIP_OPS: + continue + + # --- Build the source list directly from the captured arrays --- + sources: list[tuple[int, list[int], list[int]]] = [] + for i in range(len(rec.ne_src)): + ne_i = rec.ne_src[i] + if not any(v != 0 for v in ne_i): + continue + src_type = rec.type_src[i] if i < len(rec.type_src) and rec.type_src[i] >= 0 else 0 + nb_i = rec.nb_src[i] if i < len(rec.nb_src) else [0, 0, 0, 0] + sources.append((src_type, list(ne_i), list(nb_i))) + + # MUL_MAT_ID needs the ids tensor as src[2]; synthesize one if missing. + if op_id == 30 and len(sources) == 2: + sources.append((24, [sources[1][1][1], 1, 1, 1], [0, 0, 0, 0])) # I32 + + if not sources: + continue + + # --- Output shape --- + if any(v != 0 for v in rec.ne): + ne_out = list(rec.ne) + else: + # Legacy records without captured output ne: fall back to op-specific formula. + src_ne0 = sources[0][1] if len(sources) > 0 else [0, 0, 0, 0] + src_ne1 = sources[1][1] if len(sources) > 1 else [0, 0, 0, 0] + src_ne2 = sources[2][1] if len(sources) > 2 else [0, 0, 0, 0] + computed = _compute_output_ne(op_id, src_ne0, src_ne1, src_ne2) + if computed is None: + continue + ne_out = computed + + # --- Output type --- + out_type = rec.out_type if rec.out_type >= 0 else 0 + + # --- op_params: emit the full 16-int32 block when captured. --- + if any(p != 0 for p in rec.op_params): + op_params = list(rec.op_params) + else: + # Legacy fallback synthesis (best-effort for v2 JSON files). + op_params = [] + if op_id == 30 and len(sources) >= 2: # MUL_MAT_ID + op_params.append(sources[1][1][1]) + elif op_id in (86, 95) and rec.sub_op >= 0: # UNARY, GLU + op_params.append(rec.sub_op) + + bname = "" + if rec.backend_id in backend_by_id: + bname = backend_by_id[rec.backend_id].get("device", "") + if not bname or bname == "unknown": + bname = backend_by_id[rec.backend_id].get("name", "") + + key = (op_id, out_type, tuple(ne_out), tuple(op_params), + tuple((s[0], tuple(s[1]), tuple(s[2])) for s in sources), bname) + if key in seen: + continue + seen.add(key) + + parts: list[str] = [str(op_id), str(out_type), + str(ne_out[0]), str(ne_out[1]), str(ne_out[2]), str(ne_out[3])] + parts.append(str(len(op_params))) + parts.extend(str(p) for p in op_params) + parts.append(str(len(sources))) + for src_type, src_ne, src_nb in sources: + parts.append(str(src_type)) + parts.extend(str(v) for v in src_ne) + parts.extend(str(v) for v in src_nb) + parts.append(rec.name if rec.name else "-") + if bname: + parts.append(bname) + + lines.append(" ".join(parts) + "\n") + + with open(filepath, "w") as f: + f.writelines(lines) + + print(f"Exported {len(lines)} unique ops to: {filepath}") + + def export_html_viewer(self, filepath: str | Path, max_records: int = 0) -> None: + """Export a self-contained interactive HTML timeline viewer using Canvas.""" + import json as json_mod + + metadata_backends = self.metadata.get("backends", []) + backend_by_id: dict[int, dict] = {b["id"]: b for b in metadata_backends} + + backend_names: dict[int, str] = {} + for bid in sorted(set(rec.backend_id for rec in self.records)): + binfo = backend_by_id.get(bid, {}) + name = binfo.get("name", f"Backend {bid}") + device = binfo.get("device", "") + backend_names[bid] = device if device and device != "unknown" else name + + events: list[dict] = [] + cum_us = 0.0 + for rec in self.records: + dur_us = rec.duration_ns / 1000.0 + events.append({ + "n": rec.name, + "d": dur_us, + "s": rec.shape_str, + "b": rec.bytes, + "t": rec.type, + "bid": rec.backend_id, + "start": cum_us, + }) + cum_us += dur_us + total_us = cum_us + + if max_records > 0 and len(events) > max_records: + stride = len(events) // max_records + events = events[::stride][:max_records] + + if total_us == 0: + print("No profiling data to export.") + return + + header_stats = str(len(events)) + ' events | ' + f'{total_us/1000:.1f}' + ' ms' + + # Build backend name map with string keys for JSON + bn_str = {str(k): v for k, v in backend_names.items()} + + # --- HTML --- + html = ( + '\n' + 'ggml Profiler\n\n' + '

ggml Profiler Timeline

' + '' + header_stats + '
\n' + '
' + '' + '' + '' + '' + '' + '' + '
\n' + '
\n' + '
\n' + '
\n' + '
\n' + '
\n' + '
\n' + '' + + with open(filepath, "w") as f: + f.write(html) + + print(f"HTML viewer exported to: {filepath}") + print(f"Open in browser: file://{Path(filepath).resolve()}") + + +def load(filepath: str | Path) -> ProfileData: + """Load a profiler JSON file.""" + return ProfileData.load(filepath) + + +def main() -> None: + import argparse + + parser = argparse.ArgumentParser( + description="llama.cpp profiler analysis tool", + formatter_class=argparse.RawDescriptionHelpFormatter, + epilog=""" +Examples: + python -m tools.profiler.profiler profile.json + python -m tools.profiler.profiler profile.json --chrome-trace trace.json + python -m tools.profiler.profiler profile.json --top-ops 20 + python -m tools.profiler.profiler profile.json --export-ops ops.txt + """, + ) + parser.add_argument("profile", help="Path to profiler JSON file") + parser.add_argument("--chrome-trace", metavar="FILE", + help="Export as Chrome Trace Event format") + parser.add_argument("--html-viewer", metavar="FILE", + help="Export as interactive HTML timeline viewer") + parser.add_argument("--export-ops", metavar="FILE", + help="Export ops in export-graph-ops format (for test-backend-ops --test-file)") + parser.add_argument("--html-max-records", type=int, default=0, + help="Max records in HTML viewer (0=unlimited, set to downsample for huge traces)") + parser.add_argument("--top-ops", type=int, default=0, + help="Show top N operations (0 = show summary)") + parser.add_argument("--top-kernels", type=int, default=0, + help="Show top N longest kernels") + parser.add_argument("--inefficiency", action="store_true", + help="Show inefficiency ranking") + + args = parser.parse_args() + + data = load(args.profile) + + if args.chrome_trace: + data.export_chrome_trace(args.chrome_trace) + + if args.html_viewer: + data.export_html_viewer(args.html_viewer, max_records=args.html_max_records) + + if args.export_ops: + data.export_graph_ops(args.export_ops) + + if args.top_ops > 0: + print(f"\nTop {args.top_ops} operations by total time:\n") + for s in data.top_operations(args.top_ops): + pct = 100.0 * s.total_ns / data.total_ns if data.total_ns > 0 else 0 + print(f" {s.type_name:<5} {s.backend_id:>4} {s.name:<28} {pct:>6.1f}% " + f"{s.count:>6}x {s.total_ms:>10.2f} ms avg={s.avg_us:.2f} us") + print() + + if args.top_kernels > 0: + print(f"\nTop {args.top_kernels} longest kernels:\n") + for rec in data.top_kernels(args.top_kernels): + print(f" {rec.type_name:<5} {rec.backend_id:>4} {rec.name:<28} " + f"{rec.duration_us:>10.2f} us split={rec.split_id}") + print() + + if args.inefficiency: + print("\nInefficiency ranking (time/byte for operations with data):\n") + for s in data.inefficiency_ranking(10): + print(f" {s.name:<28} {s.time_per_byte_ns / 1000:>10.2f} us/byte " + f"{s.count:>6} calls {s.total_bytes / 1e6:.1f} MB") + print() + + if args.top_ops == 0 and args.top_kernels == 0 and not args.inefficiency and not args.chrome_trace and not args.html_viewer and not args.export_ops: + data.summary() + + +if __name__ == "__main__": + main() diff --git a/tools/server/server.cpp b/tools/server/server.cpp index 22378b38c5..b9b0758277 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -554,6 +554,29 @@ int llama_server(common_params & params, int argc, char ** argv) { auto * ll_ctx = ctx_server.get_llama_context(); if (ll_ctx != nullptr) { + // export profiling data if profiling was enabled + if (params.profiling) { + ggml_backend_sched_t sched = llama_context_get_sched(ll_ctx); + if (sched != nullptr) { + if (params.profiling_output.empty()) { + ggml_backend_sched_print_profiling(sched); + } else { + const std::string & path = params.profiling_output; + int ret; + if (path.size() >= 4 && path.compare(path.size() - 4, 4, ".txt") == 0) { + ret = ggml_backend_sched_export_profiling_text(sched, path.c_str()); + } else { + ret = ggml_backend_sched_export_profiling_json(sched, path.c_str()); + } + if (ret == 0) { + SRV_INF("profiling data exported to: %s\n", path.c_str()); + } else { + SRV_ERR("failed to export profiling data to: %s\n", path.c_str()); + } + } + } + } + common_memory_breakdown_print(ll_ctx); } }