Add tensor name to JSON output

* opencl: allow loading binary kernel

* opencl: add libdl.h

* ggml-backend-dl is in ggml, which depends backend libs, thus
  ggml-opencl cannot depend on ggml-backend-dl
* add libdl.h to break cyclic dep

* opencl: allow loading bin kernel lib

* opencl: load `gemm_moe_mxfp4_f32_ns` from kernel lib if available

* opencl: load q8_0 gemm from kernel lib

* opencl: load q4_0 moe gemm from kernel lib

* opencl: load q4_1 moe gemm from kernel lib

* opencl: load q4_k moe gemm from kernel lib

* opencl: always declare `get_adreno_bin_kernel_func_t`

* opencl: rephrase message

* opencl: fix for rebase

* opencl: update doc

common/arg.cpp

@@ -496,13 +496,15 @@ void common_models_handler_apply(common_models_handler & handler, common_params
     }
 
     // handle hf_plan tasks
-    auto add_tasks = [&opts, &tasks](const hf_cache::hf_files & model_files, common_params_model & model) {
+    auto add_tasks = [&opts, &tasks](const hf_cache::hf_files  & model_files,
+                                    const hf_cache::hf_file    & primary,
+                                    common_params_model        & model) {
         for (size_t i = 0; i < model_files.size(); ++i) {
             auto & model_file = model_files[i];
-            bool is_first = (i == 0);
-            tasks.emplace_back(model_file, opts, [&, is_first]() {
-                if (is_first) {
-                    // only use first part as model path
+            bool is_primary = (model_file.path == primary.path);
+            tasks.emplace_back(model_file, opts, [&, is_primary]() {
+                if (is_primary) {
+                    // the primary file is the first split (00001-of), use it as model path
                     model.path = hf_cache::finalize_file(model_file);
                 } else {
                     hf_cache::finalize_file(model_file);
@@ -511,7 +513,7 @@ void common_models_handler_apply(common_models_handler & handler, common_params
         }
     };
     if (!plan.model_files.empty()) {
-        add_tasks(plan.model_files, params.model);
+        add_tasks(plan.model_files, plan.primary, params.model);
     }
     if (!plan.mmproj.local_path.empty()) {
         tasks.emplace_back(plan.mmproj, opts, [&]() {
@@ -539,12 +541,12 @@ void common_models_handler_apply(common_models_handler & handler, common_params
 
     // handle plan_spec (e.g. --spec-draft-hf)
     if (!plan_spec.model_files.empty()) {
-        add_tasks(plan_spec.model_files, params.speculative.draft.mparams);
+        add_tasks(plan_spec.model_files, plan_spec.primary, params.speculative.draft.mparams);
     }
 
     // handle vocoder plan (e.g. --hf-repo-v)
     if (!plan_voc.model_files.empty()) {
-        add_tasks(plan_voc.model_files, params.vocoder.model);
+        add_tasks(plan_voc.model_files, plan_voc.primary, params.vocoder.model);
     }
 
     // run all tasks in parallel
@@ -1238,6 +1240,21 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.completion = true;
         }
     ));
+    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"),
@@ -2841,6 +2858,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}));
+    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),

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"
@@ -1295,6 +1296,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);
 }
 

common/common.h

@@ -5,6 +5,7 @@
 #include "llama-cpp.h"
 
 #include "ggml-opt.h"
+#include "ggml-profiler.h"
 #include "ggml.h"
 
 #include <set>
@@ -475,6 +476,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<size_t> fit_params_target = std::vector<size_t>(llama_max_devices(), 1024 * 1024*1024);
@@ -709,6 +711,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;
 

docs/backend/OPENCL.md

@@ -1,16 +1,26 @@
 # llama.cpp for OpenCL
 
-- [Background](#background)
-- [OS](#os)
-- [Hardware](#hardware)
-- [DataType Supports](#datatype-supports)
-- [Model Preparation](#model-preparation)
-- [CMake Options](#cmake-options)
-- [Android](#android)
-- [Windows 11 Arm64](#windows-11-arm64)
-- [Linux](#Linux)
-- [Known Issue](#known-issues)
-- [TODO](#todo)
+- [llama.cpp for OpenCL](#llamacpp-for-opencl)
+  - [Background](#background)
+    - [Llama.cpp + OpenCL](#llamacpp--opencl)
+  - [OS](#os)
+  - [Hardware](#hardware)
+    - [Adreno GPU](#adreno-gpu)
+  - [DataType Supports](#datatype-supports)
+  - [Model Preparation](#model-preparation)
+  - [Binary Kernel Library](#binary-kernel-library)
+  - [CMake Options](#cmake-options)
+  - [Android](#android)
+    - [I. Setup Environment](#i-setup-environment)
+    - [II. Build llama.cpp](#ii-build-llamacpp)
+  - [Windows 11 Arm64](#windows-11-arm64)
+    - [I. Setup Environment](#i-setup-environment-1)
+    - [II. Build llama.cpp](#ii-build-llamacpp-1)
+  - [Linux](#linux)
+    - [I. Setup Environment](#i-setup-environment-2)
+    - [II. Build llama.cpp](#ii-build-llamacpp-2)
+  - [Known Issues](#known-issues)
+  - [TODO](#todo)
 
 ## Background
 
@@ -34,11 +44,13 @@ The llama.cpp OpenCL backend is designed to enable llama.cpp on **Qualcomm Adren
 
 **Verified devices**
 
-| Adreno GPU                           | Status  |
-|:------------------------------------:|:-------:|
-| Adreno 750 (Snapdragon 8 Gen 3)      | Support |
-| Adreno 830 (Snapdragon 8 Elite)      | Support |
-| Adreno X85 (Snapdragon X Elite)      | Support |
+| Adreno GPU                            | Status  |
+|:-------------------------------------:|:-------:|
+| Adreno 750 (Snapdragon 8 Gen 3)       | Support |
+| Adreno 830 (Snapdragon 8 Elite)       | Support |
+| Adreno 840 (Snapdragon 8 Elite Gen 5) | Support |
+| Adreno X1-85 (Snapdragon X Elite)     | Support |
+| Adreno X2-90 (Snapdragon X2 Elite)    | Support |
 
 > A6x GPUs with a recent driver and compiler are supported; they are usually found in IoT platforms.
 However, A6x GPUs in phones are likely not supported due to the outdated driver and compiler.
@@ -47,42 +59,43 @@ However, A6x GPUs in phones are likely not supported due to the outdated driver
 
 | DataType               | Status                     |
 |:----------------------:|:--------------------------:|
+| Q1_0                   | Support                    |
 | Q4_0                   | Support                    |
-| Q6_K                   | Support, but not optimized |
+| Q4_1                   | Support                    |
+| Q5_0                   | Support                    |
+| Q5_1                   | Support                    |
 | Q8_0                   | Support                    |
+| Q4_K                   | Support                    |
+| Q5_K                   | Support                    |
+| Q6_K                   | Support                    |
 | MXFP4                  | Support                    |
+| IQ4_NL                 | Support                    |
 
 ## Model Preparation
 
-You can refer to the general [llama-quantize tool](/tools/quantize/README.md) for steps to convert a model in Hugging Face safetensor format to GGUF with quantization.
+Since common quantizations are supported now, it is recommanded to download GGUF models directly from Huggingface.
 
-Currently we support `Q4_0` quantization and have optimized for it. To achieve best performance on Adreno GPU, add `--pure` to `llama-quantize` (i.e., make all weights in `Q4_0`). For example,
+## Binary Kernel Library
 
-```sh
-./llama-quantize --pure ggml-model-qwen2.5-3b-f16.gguf ggml-model-qwen-3b-Q4_0.gguf Q4_0
-```
+A prebuilt binary kernel library has been introduced for Adreno GPUs.
+It currently targets X2 GPUs (X2-90, X2-85 and X2-45) found in Snapdragon X2 SoC.
+The library currently contains kernels for MUL_MAT_ID with Q4_0, Q4_1, Q4_K, MXFP4.
+The library must be manually downloaded from https://softwarecenter.qualcomm.com/catalog/item/Adreno_Kernel_Library_GGML.
 
-Since `Q6_K` is also supported, `Q4_0` quantization without `--pure` will also work. However, the performance will be worse compared to pure `Q4_0` quantization.
+To allow using the kernel library, add `-DGGML_OPENCL_USE_ADRENO_BIN_KERNELS=ON` when configuring with CMake.
+Then, extract `adreno-opencl-kernels.dll` from the zip file downloaded from the above URL and put it alongside the executables.
+If kernels compatible with the current GPU are found in the library, they will be loaded and used.
 
-### `MXFP4` MoE Models
-
-OpenAI gpt-oss models are MoE models in `MXFP4`. The quantized model will be in `MXFP4_MOE`, a mixture of `MXFP4` and `Q8_0`.
-For this quantization, there is no need to specify `--pure`.
-For gpt-oss-20b model, you can directly [download](https://huggingface.co/ggml-org/gpt-oss-20b-GGUF) the quantized GGUF file in `MXFP4_MOE` from Hugging Face.
-
-Although it is possible to quantize gpt-oss-20b model in pure `Q4_0` (all weights in `Q4_0`), it is not recommended since `MXFP4` has been optimized for MoE while `Q4_0` is not. In addition, accuracy should degrade with such pure `Q4_0` quantization.
-Hence, using the default `MXFP4_MOE` quantization (see the link above) is recommended for this model.
-
-> Note that the `Q4_0` model found [here](https://huggingface.co/unsloth/gpt-oss-20b-GGUF/blob/main/gpt-oss-20b-Q4_0.gguf) is a mixture of `Q4_0`, `Q8_0` and `MXFP4` and gives better performance than `MXFP4_MOE` quantization.
 
 ## CMake Options
 
 The OpenCL backend has the following CMake options that control the behavior of the backend.
 
-| CMake options                     | Default value  | Description                               |
-|:---------------------------------:|:--------------:|:------------------------------------------|
-| `GGML_OPENCL_EMBED_KERNELS`       | `ON`           | Embed OpenCL kernels into the executable. |
-| `GGML_OPENCL_USE_ADRENO_KERNELS`  | `ON`           | Use kernels optimized for Adreno.         |
+| CMake options                        | Default value  | Description                               |
+|:------------------------------------:|:--------------:|:------------------------------------------|
+| `GGML_OPENCL_EMBED_KERNELS`          | `ON`           | Embed OpenCL kernels into the executable. |
+| `GGML_OPENCL_USE_ADRENO_KERNELS`     | `ON`           | Use kernels optimized for Adreno.         |
+| `GGML_OPENCL_USE_ADRENO_BIN_KERNELS` | `OFF`          | Allow using binary kernel lib for Adreno. |
 
 ## Android
 
@@ -277,6 +290,5 @@ ninja
 
 ## TODO
 
-- Optimization for Q6_K
-- Support and optimization for Q4_K
 - Improve flash attention
+- Improve OpenCL C kernels performance

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

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

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

ggml/include/ggml-profiler.h

@@ -0,0 +1,134 @@
+#pragma once
+
+#include "ggml-backend.h"
+#include "ggml.h"
+
+#include <stdint.h>
+
+#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

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

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" {
@@ -144,6 +145,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 {

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 <assert.h>
 #include <limits.h>
@@ -20,6 +21,7 @@
 #include <stdlib.h>
 #include <string.h>
 #include <algorithm>
+#include <string>
 #include <vector>
 
 #ifdef __APPLE__
@@ -231,6 +233,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);
 }
 
@@ -825,6 +836,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<ggml_profile_record> copy_records;        // copy events recorded by the scheduler
+    std::vector<ggml_profile_record> 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<backend_meta>        profiling_backend_meta;
 };
 
 #define hash_id(tensor) ggml_hash_find_or_insert(&sched->hash_set, tensor)
@@ -1538,6 +1563,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;
@@ -1546,11 +1604,28 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s
     std::vector<int32_t> ids;
     std::vector<ggml_bitset_t> used_ids;
 
+    // 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;
         ggml_backend_t split_backend = sched->backends[split_backend_id];
 
+        // 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]);
@@ -1564,7 +1639,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) {
@@ -1621,12 +1714,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<size_t>(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,
@@ -1635,6 +1730,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++;
@@ -1658,9 +1758,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) {
@@ -1668,7 +1793,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));
+                        }
                     }
                 }
             }
@@ -1721,6 +1884,32 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s
         }
     }
 
+    // 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;
 }
 
@@ -1787,6 +1976,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);
 
@@ -1797,6 +2004,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]);
@@ -2369,3 +2603,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<op_stats> 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<op_stats> 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<backend_stats> 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;
+}

ggml/src/ggml-blas/ggml-blas.cpp

@@ -1,6 +1,7 @@
 #include "ggml-impl.h"
 #include "ggml-blas.h"
 #include "ggml-backend-impl.h"
+#include "ggml-profiler.h"
 
 #include <future>
 #include <vector>
@@ -25,6 +26,11 @@ struct ggml_backend_blas_context {
 #ifndef GGML_USE_OPENMP
     std::vector<std::future<void>> tasks;
 #endif
+
+    // Profiling state
+    bool                             profiling_enabled  = false;
+    int                              profiling_split_id = -1;
+    std::vector<ggml_profile_record> profiling_records;
 };
 
 static void ggml_backend_blas_mul_mat(ggml_backend_blas_context * ctx, struct ggml_tensor * dst) {
@@ -232,6 +238,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);
@@ -241,16 +259,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;
@@ -286,10 +312,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)
@@ -302,6 +329,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;
 }
 

ggml/src/ggml-cann/ggml-cann.cpp

@@ -3035,7 +3035,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;
 }

ggml/src/ggml-cpu/ggml-cpu.c

@@ -6,6 +6,7 @@
 #include "traits.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"
@@ -1172,8 +1173,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;
@@ -1263,9 +1264,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);
@@ -1474,8 +1475,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;
@@ -1542,8 +1543,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));
@@ -3046,17 +3047,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(&params, 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
@@ -3067,14 +3106,15 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
             ggml_compute_forward(&params, 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);
+            }
         }
     }
 

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<ggml_profile_record> 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;
 }
 

ggml/src/ggml-cuda/common.cuh

@@ -1388,6 +1388,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;
@@ -1499,6 +1502,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 {

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"
@@ -89,6 +90,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<cudaEvent_t> start_events;
+    std::vector<cudaEvent_t> end_events;
+    std::vector<uint64_t>    cpu_timestamps;  // CPU-side timestamps for global ordering
+    int event_count = 0;
+
+    std::vector<ggml_profile_record> records;
+    std::vector<int> 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); }); }
 
@@ -4398,8 +4485,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));
                 }
@@ -4470,6 +4572,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;
@@ -4521,6 +4636,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;
 }
 
@@ -5709,12 +5833,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;
 }
 

ggml/src/ggml-cuda/vendors/hip.h

@@ -64,8 +64,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

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

ggml/src/ggml-hexagon/ggml-hexagon.cpp

@@ -3783,6 +3783,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);

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

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 <Foundation/Foundation.h>
 
@@ -79,6 +80,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) {
@@ -450,6 +464,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
@@ -609,6 +651,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<MTLCommandBuffer> cmd_buf = ctx->cmd_bufs[n_cb].obj;
+                if (cmd_buf) {
+                    [cmd_buf waitUntilCompleted];
+                }
+            }
+            for (int i = 0; i < n_cb; ++i) {
+                id<MTLCommandBuffer> 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;
@@ -660,6 +762,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);
@@ -704,6 +810,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) {

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
 //
@@ -293,6 +304,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
 //

ggml/src/ggml-metal/ggml-metal-device.m

@@ -8,6 +8,7 @@
 #include <Metal/Metal.h>
 
 #include <stdatomic.h>
+#include <mach/mach_time.h>
 
 #ifndef TARGET_OS_VISION
 #define TARGET_OS_VISION 0
@@ -517,6 +518,21 @@ void ggml_metal_encoder_end_encoding(ggml_metal_encoder_t encoder) {
     [encoder->obj endEncoding];
 }
 
+//
+// MTLCounterSampleBuffer wrapper
+//
+
+struct ggml_metal_sample_buf {
+    id<MTLCounterSampleBuffer> 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<MTLDevice> mtl_device;
 
@@ -1362,6 +1378,139 @@ const struct ggml_metal_device_props * ggml_metal_device_get_props(ggml_metal_de
     return &dev->props;
 }
 
+//
+// Profiling helpers
+//
+
+// Look up the MTLCommonCounterSetTimestamp counter set, or nil if the device doesn't expose it.
+static id<MTLCounterSet> ggml_metal_device_get_timestamp_counter_set(ggml_metal_device_t dev) {
+    if (@available(macOS 10.15, iOS 14.0, *)) {
+        NSArray<id<MTLCounterSet>> * sets = [dev->mtl_device counterSets];
+        for (id<MTLCounterSet> 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<MTLCounterSet> 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<MTLCounterSampleBuffer> 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
 //

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -77,6 +77,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());
@@ -100,6 +105,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;
 
@@ -144,6 +155,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;
@@ -501,12 +522,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);
     }

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:
 //

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"
@@ -9,6 +10,7 @@
 
 #include <mutex>
 #include <string>
+#include <vector>
 
 #define GGML_METAL_NAME "MTL"
 #define GGML_METAL_MAX_DEVICES 16
@@ -20,6 +22,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<ggml_profile_record> 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
 ////////////////////////////////////////////////////////////////////////////////
@@ -589,6 +644,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;
@@ -606,10 +703,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;
 }
 
@@ -700,10 +800,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);

ggml/src/ggml-opencl/CMakeLists.txt

@@ -31,6 +31,11 @@ if (GGML_OPENCL_EMBED_KERNELS)
     target_include_directories(${TARGET_NAME} PRIVATE "${CMAKE_CURRENT_BINARY_DIR}/autogenerated")
 endif ()
 
+if (GGML_OPENCL_USE_ADRENO_BIN_KERNELS)
+    message(STATUS "OpenCL will use precompiled binary kernels for Adreno (improved performance on some platforms)")
+    add_compile_definitions(GGML_OPENCL_USE_ADRENO_BIN_KERNELS)
+endif ()
+
 function(ggml_opencl_add_kernel KNAME)
     set(KERN_HDR ${CMAKE_CURRENT_BINARY_DIR}/autogenerated/${KNAME}.cl.h)
     set(KERN_SRC ${CMAKE_CURRENT_SOURCE_DIR}/kernels/${KNAME}.cl)

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -13,6 +13,22 @@
 #include "ggml-backend-impl.h"
 #include "ggml.h"
 
+#ifdef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+#include "libdl.h"
+#ifdef _WIN32
+#define KERNEL_LIB_NAME "adreno-opencl-kernels.dll"
+#else
+#define KERNEL_LIB_NAME "libadreno-opencl-kernels.so"
+#endif // _WIN32
+#endif // GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+
+typedef const void * (*get_adreno_bin_kernel_func_t)(
+    const char * name,
+    const char * gpu_name,
+    const char * compiler_ver,
+    size_t     * out_size
+);
+
 #include <CL/cl.h>
 
 #include <inttypes.h>
@@ -476,6 +492,8 @@ struct ggml_backend_opencl_context {
 
     bool adreno_has_large_buffer;
     bool adreno_use_large_buffer;
+    bool adreno_use_bin_kernels;
+    get_adreno_bin_kernel_func_t get_adreno_bin_kernel_func = nullptr;
     ggml_cl_compiler_version adreno_cl_compiler_version;
 
     std::string kernel_compile_opts;  // cached for lazy-compiled kernels.
@@ -718,15 +736,15 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_gated_delta_net_f32[4][2][2] = {};
 
     cl_kernel kernel_timestep_embedding;
-    cl_kernel kernel_gemv_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns;
-    cl_kernel kernel_gemv_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns;
+    cl_kernel kernel_gemv_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns_bin;
+    cl_kernel kernel_gemv_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns_bin;
     cl_kernel kernel_gemv_moe_q5_0_f32_ns, kernel_gemm_moe_q5_0_f32_ns;
     cl_kernel kernel_gemv_moe_q5_1_f32_ns, kernel_gemm_moe_q5_1_f32_ns;
-    cl_kernel kernel_gemv_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns;
+    cl_kernel kernel_gemv_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns_bin;
     cl_kernel kernel_gemv_moe_q5_k_f32_ns, kernel_gemm_moe_q5_k_f32_ns;
     cl_kernel kernel_gemv_moe_q6_k_f32_ns, kernel_gemm_moe_q6_k_f32_ns;
     cl_kernel kernel_gemv_moe_mxfp4_f32, kernel_gemm_moe_mxfp4_f32;
-    cl_kernel kernel_gemv_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns;
+    cl_kernel kernel_gemv_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns_bin;
     cl_kernel kernel_moe_reorder_b;
     cl_kernel kernel_moe_histogram, kernel_moe_scan, kernel_moe_fill, kernel_moe_scatter;
     cl_kernel kernel_mul_mv_id_q4_0_f32_8x_flat;
@@ -870,6 +888,20 @@ struct ggml_backend_opencl_context {
 #endif
     }
 
+    const void * get_adreno_bin_kernel(const std::string &kernel_name, size_t *bin_size) const {
+        if (!get_adreno_bin_kernel_func) {
+            return nullptr;
+        }
+
+        size_t sz;
+        const void * kernel_bin = get_adreno_bin_kernel_func(
+            kernel_name.c_str(), device_name.c_str(), driver_version.c_str(), &sz);
+        if (bin_size) {
+            *bin_size = sz;
+        }
+        return kernel_bin;
+    }
+
 #ifdef GGML_OPENCL_USE_ADRENO_KERNELS
     // Transpose kernels
     cl_program program_transpose;
@@ -891,7 +923,7 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_gemv_noshuffle_q4_0_f32_32000_1_4096;
     cl_kernel kernel_gemv_noshuffle_q4_1_f32;
     cl_kernel kernel_gemm_noshuffle_q4_1_f32;
-    cl_kernel kernel_gemm_noshuffle_q8_0_f32;
+    cl_kernel kernel_gemm_noshuffle_q8_0_f32, kernel_gemm_noshuffle_q8_0_f32_bin;
     cl_kernel kernel_gemv_noshuffle_q8_0_f32;
     cl_kernel kernel_gemm_noshuffle_q1_0_f32;
     cl_kernel kernel_gemv_noshuffle_q1_0_f32;
@@ -988,6 +1020,32 @@ static cl_program build_program_from_source(cl_context ctx, cl_device_id dev, co
     return build_program_from_source_ex(ctx, dev, program_buffer, compile_opts, /*fatal=*/true);
 }
 
+static cl_program build_program_from_binary(cl_context ctx, cl_device_id dev, const char* program_buffer, const std::string &compile_opts, size_t bin_size = 0) {
+    cl_program p;
+    char *program_log;
+    size_t log_size;
+    int err;
+
+    p = clCreateProgramWithBinary(ctx, 1, &dev, &bin_size, (const unsigned char**)&program_buffer, NULL, &err);
+    if(err < 0) {
+        GGML_LOG_ERROR("OpenCL error creating program from binary");
+        exit(1);
+    }
+
+    err = clBuildProgram(p, 0, NULL, compile_opts.c_str(), NULL, NULL);
+    if(err < 0) {
+        clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
+        program_log = (char*) malloc(log_size + 1);
+        program_log[log_size] = '\0';
+        clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, log_size + 1, program_log, NULL);
+        GGML_LOG_ERROR("ggml_opencl: kernel compile error:\n\n%s\n", program_log);
+        free(program_log);
+        exit(1);
+    }
+
+    return p;
+}
+
 static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) {
     // compiler options for general kernels
     auto opencl_c_std =
@@ -1014,6 +1072,17 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) {
     }
 }
 
+static bool use_adreno_bin_kernels(ggml_backend_opencl_context * backend_ctx) {
+#ifndef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+    return false;
+#else
+    if (backend_ctx->gpu_family != GPU_FAMILY::ADRENO) {
+        return false;
+    }
+    return backend_ctx->adreno_use_bin_kernels;
+#endif // GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+}
+
 static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
     if (backend_ctx->kernels_loaded) {
         return;
@@ -3323,6 +3392,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         GGML_LOG_CONT(".");
     }
 
+    // gemm_noshuffle_q8_0_f32_bin
+    {
+        size_t bin_size = 0;
+        backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin = nullptr;
+
+        if (use_adreno_bin_kernels(backend_ctx)) {
+            const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q8_0_f32_ila", &bin_size);
+            if (kernel_bin && bin_size > 0) {
+                cl_program prog =
+                    build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, compile_opts, bin_size);
+
+                CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin = clCreateKernel(prog, "kernel_gemm_noshuffle_q8_0_f32_ila", &err), err));
+                CL_CHECK(clReleaseProgram(prog));
+                GGML_LOG_CONT(".");
+            }
+        }
+    }
+
     // gemv_noshuffle_general_q8_0_f32
     {
         std::string CL_gemv_compile_opts = std::string("-cl-std=") + opencl_c_std +
@@ -3424,6 +3511,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         GGML_LOG_CONT(".");
     }
 
+    // gemm_moe_q4_1_f32_ns_bin
+    {
+        size_t bin_size = 0;
+        backend_ctx->kernel_gemm_moe_q4_1_f32_ns_bin = nullptr;
+
+        if (use_adreno_bin_kernels(backend_ctx)) {
+            const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_q4_1_f32_ns_ila", &bin_size);
+            if (kernel_bin && bin_size > 0) {
+                cl_program prog =
+                    build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size);
+
+                CL_CHECK((backend_ctx->kernel_gemm_moe_q4_1_f32_ns_bin = clCreateKernel(prog, "kernel_gemm_moe_q4_1_f32_ns_ila", &err), err));
+                CL_CHECK(clReleaseProgram(prog));
+                GGML_LOG_CONT(".");
+            }
+        }
+    }
+
     // gemv_moe_mxfp4_f32
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -3490,6 +3595,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         GGML_LOG_CONT(".");
     }
 
+    // gemm_moe_q4_0_f32_ns_bin
+    {
+        size_t bin_size = 0;
+        backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin = nullptr;
+
+        if (use_adreno_bin_kernels(backend_ctx)) {
+            const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_q4_0_f32_ns_ila", &bin_size);
+            if (kernel_bin && bin_size > 0) {
+                cl_program prog =
+                    build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size);
+
+                CL_CHECK((backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin = clCreateKernel(prog, "kernel_gemm_moe_q4_0_f32_ns_ila", &err), err));
+                CL_CHECK(clReleaseProgram(prog));
+                GGML_LOG_CONT(".");
+            }
+        }
+    }
+
     // gemv_moe_q5_0_f32_ns
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -3592,6 +3715,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         GGML_LOG_CONT(".");
     }
 
+    // gemm_moe_q4_k_f32_ns_bin
+    {
+        size_t bin_size = 0;
+        backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin = nullptr;
+
+        if (use_adreno_bin_kernels(backend_ctx)) {
+            const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_q4_k_f32_ns_ila", &bin_size);
+            if (kernel_bin && bin_size > 0) {
+                cl_program prog =
+                    build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size);
+
+                CL_CHECK((backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin = clCreateKernel(prog, "kernel_gemm_moe_q4_k_f32_ns_ila", &err), err));
+                CL_CHECK(clReleaseProgram(prog));
+                GGML_LOG_CONT(".");
+            }
+        }
+    }
+
     // gemv_moe_q5_k_f32_ns
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -3689,9 +3830,27 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         cl_program prog =
             build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), CL_moe_compile_opts);
 
-        CL_CHECK((backend_ctx->kernel_gemm_moe_mxfp4_f32_ns = clCreateKernel(prog, "kernel_gemm_moe_mxfp4_f32_ns", &err), err));
-        CL_CHECK(clReleaseProgram(prog));
-        GGML_LOG_CONT(".");
+            CL_CHECK((backend_ctx->kernel_gemm_moe_mxfp4_f32_ns = clCreateKernel(prog, "kernel_gemm_moe_mxfp4_f32_ns", &err), err));
+            CL_CHECK(clReleaseProgram(prog));
+            GGML_LOG_CONT(".");
+    }
+
+    // gemm_moe_mxfp4_f32_ns_bin
+    {
+        size_t bin_size = 0;
+        backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin = nullptr;
+
+        if (use_adreno_bin_kernels(backend_ctx)) {
+            const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_moe_mxfp4_f32_ns_ila", &bin_size);
+            if (kernel_bin && bin_size > 0) {
+                cl_program prog =
+                    build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, CL_moe_compile_opts, bin_size);
+
+                CL_CHECK((backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin = clCreateKernel(prog, "kernel_gemm_moe_mxfp4_f32_ns_ila", &err), err));
+                CL_CHECK(clReleaseProgram(prog));
+                GGML_LOG_CONT(".");
+            }
+        }
     }
 
     // moe_reorder_b
@@ -4770,6 +4929,27 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
     backend_ctx->adreno_use_large_buffer = getenv("GGML_OPENCL_ADRENO_USE_LARGE_BUFFER") != nullptr &&
                                            backend_ctx->gpu_family == GPU_FAMILY::ADRENO;
 
+#ifdef GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+    // try loading adreno binary kernels if enabled
+    // if fails to load, builtin kernels will be used
+    {
+        dl_handle * kernel_lib_handle = dl_load_library(KERNEL_LIB_NAME);
+        backend_ctx->adreno_use_bin_kernels = false;
+
+        if (kernel_lib_handle) {
+            backend_ctx->get_adreno_bin_kernel_func = (get_adreno_bin_kernel_func_t)dl_get_sym(kernel_lib_handle, "get_adreno_kernels");
+            if (backend_ctx->get_adreno_bin_kernel_func) {
+                GGML_LOG_INFO("ggml_opencl: loaded bin kernel library %s\n", KERNEL_LIB_NAME);
+                backend_ctx->adreno_use_bin_kernels = true;
+            } else {
+                GGML_LOG_INFO("ggml_opencl: bin kernel library %s is invalid, will use builtin kernels\n", KERNEL_LIB_NAME);
+            }
+        } else {
+            GGML_LOG_INFO("ggml_opencl: failed to load %s, will use builtin kernels\n", KERNEL_LIB_NAME);
+        }
+    }
+#endif // GGML_OPENCL_USE_ADRENO_BIN_KERNELS
+
     cl_int err;
 
     // A local ref of cl_context for convenience
@@ -6044,7 +6224,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;
@@ -9209,6 +9390,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;
@@ -14972,6 +15154,99 @@ static void ggml_cl_mul_mat_q8_0_f32_adreno(ggml_backend_t backend, const ggml_t
         CL_CHECK(clReleaseMemObject(b_img));
         CL_CHECK(clReleaseMemObject(b_sub_buf));
     } else {
+        // use bin kernel if available
+        if (backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin) {
+            int K_pad = K;
+
+            cl_mem b_sub_buf = nullptr;
+            cl_mem d_sub_buf = nullptr;
+
+            cl_mem a_img = nullptr;
+            cl_mem s_img = nullptr;
+            cl_mem b_img = nullptr;
+            cl_mem d_img = nullptr;
+
+            // subbuffer for activations
+            region.origin = offset1;
+            region.size = K_pad * N * sizeof(float);
+            CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
+
+            // Create subbuffer and image1d_buffer for dst
+            region.origin = (extrad->offset); // + dst->view_offs;
+            region.size = M * N * sizeof(float);
+            CL_CHECK((d_sub_buf = clCreateSubBuffer((extrad->data_device), 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
+
+            // create an image for A
+            img_fmt = { CL_R, CL_FLOAT};
+            memset(&img_desc, 0, sizeof(img_desc));
+            img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            img_desc.image_width = M * K / 4;    // Divide by 4 for char -> float
+            img_desc.buffer = extra0_q8_0->q;
+            CL_CHECK((a_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+            // create an image for Scale
+            img_fmt = { CL_R, CL_HALF_FLOAT};
+            memset(&img_desc, 0, sizeof(img_desc));
+            img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            img_desc.image_width = M * K / 32;    // Block size is 32
+            img_desc.buffer = extra0_q8_0->d;
+            CL_CHECK((s_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+            // create an image for B from sub_buffer
+            img_fmt = {CL_R, CL_FLOAT};
+            memset(&img_desc, 0, sizeof(img_desc));
+            img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            img_desc.image_width = K_pad * N;
+            img_desc.buffer = b_sub_buf;
+            CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+            // img for d
+            img_fmt = {CL_R, CL_FLOAT};
+            memset(&img_desc, 0, sizeof(img_desc));
+            img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            img_desc.image_width = M * N;
+            img_desc.buffer = d_sub_buf;
+            CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+            // gemm
+            kernel = backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin;
+
+            bool layoutA_Mfirst = true;
+            bool layoutS_Mfirst = true;
+            bool layoutB_Nfirst = false;
+            bool layoutC_Mfirst = true;
+
+            cl_uint lineStrideMatrixAinBytes = layoutA_Mfirst ? M * 4 : K;                // int8
+            cl_uint lineStrideMatrixSinBytes = layoutS_Mfirst ? M * 2 : (K / 32) * 2;     // fp16
+            cl_uint lineStrideMatrixBinBytes = layoutB_Nfirst ? N * 4 : K_pad * 4;        // fp32
+            cl_uint lineStrideMatrixCinBytes = layoutC_Mfirst ? M * 4 : N * 4;            // fp32
+
+            CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem), &a_img));
+            CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_mem), &s_img));
+            CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem), &b_img));
+            CL_CHECK(clSetKernelArg(kernel,  3, sizeof(int),    &extra1->offset));
+            CL_CHECK(clSetKernelArg(kernel,  4, sizeof(cl_mem), &d_img));
+            CL_CHECK(clSetKernelArg(kernel,  5, sizeof(int),    &extrad->offset));
+            CL_CHECK(clSetKernelArg(kernel,  6, sizeof(int),    &K));
+            CL_CHECK(clSetKernelArg(kernel,  7, sizeof(int),    &lineStrideMatrixAinBytes));
+            CL_CHECK(clSetKernelArg(kernel,  8, sizeof(int),    &lineStrideMatrixSinBytes));
+            CL_CHECK(clSetKernelArg(kernel,  9, sizeof(int),    &lineStrideMatrixBinBytes));
+            CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int),    &lineStrideMatrixCinBytes));
+
+            size_t global_work_size[] = { 64, (size_t)CEIL_DIV(M, 64), (size_t)CEIL_DIV(N, 64)};
+            size_t local_work_size[]  = { 64, 2, 2 };
+
+            backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+
+            CL_CHECK(clReleaseMemObject(b_sub_buf));
+            CL_CHECK(clReleaseMemObject(d_sub_buf));
+            CL_CHECK(clReleaseMemObject(a_img));
+            CL_CHECK(clReleaseMemObject(s_img));
+            CL_CHECK(clReleaseMemObject(b_img));
+            CL_CHECK(clReleaseMemObject(d_img));
+            return;
+        }
+
         cl_mem b_sub_buf = nullptr;
         cl_mem b_sub_buf_trans = nullptr;
         cl_mem b_img = nullptr;
@@ -17825,6 +18100,9 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
 
                 } else { // for gemm
                     kernel = backend_ctx->kernel_gemm_moe_q4_0_f32_ns;
+                    if (backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin) {
+                        kernel = backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin;
+                    }
 
                     // Reorder router if called from test-backend-ops or when new router is generated.
                     // Otherwise reuse the reordered result from previous mul_mat_id call.
@@ -17870,6 +18148,11 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     cl_image_desc image_desc_buf_src1;
                     image_format_buf_src1 = {CL_RGBA, CL_FLOAT};
                     image_desc_buf_src1 = {CL_MEM_OBJECT_IMAGE1D_BUFFER, static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size / 4), 0,0,0,0,0,0,0, {buf_src1_reordered}};
+                    if (backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin) {
+                        // bin kernel uses slightly different image format
+                        image_format_buf_src1 = {CL_R, CL_FLOAT};
+                        image_desc_buf_src1.image_width = static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size);
+                    }
                     image_src1_reordered = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &image_format_buf_src1, &image_desc_buf_src1, NULL, &status);
                     CL_CHECK(status);
 
@@ -18042,6 +18325,9 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
 
                 } else { // for gemm
                     kernel = backend_ctx->kernel_gemm_moe_q4_1_f32_ns;
+                    if (backend_ctx->kernel_gemm_moe_q4_1_f32_ns_bin) {
+                        kernel = backend_ctx->kernel_gemm_moe_q4_1_f32_ns_bin;
+                    }
 
                     // Reorder router if called from test-backend-ops or when new router is generated.
                     // Otherwise reuse the reordered result from previous mul_mat_id call.
@@ -18087,6 +18373,11 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     cl_image_desc image_desc_buf_src1;
                     image_format_buf_src1 = {CL_RGBA, CL_FLOAT};
                     image_desc_buf_src1 = {CL_MEM_OBJECT_IMAGE1D_BUFFER, static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size / 4), 0,0,0,0,0,0,0, {buf_src1_reordered}};
+                    if (backend_ctx->kernel_gemm_moe_q4_1_f32_ns_bin) {
+                        // bin kernel uses slightly different image format
+                        image_format_buf_src1 = {CL_R, CL_FLOAT};
+                        image_desc_buf_src1.image_width = static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size);
+                    }
                     image_src1_reordered = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &image_format_buf_src1, &image_desc_buf_src1, NULL, &status);
                     CL_CHECK(status);
 
@@ -18648,6 +18939,9 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
 
                 } else { // for gemm
                     kernel = backend_ctx->kernel_gemm_moe_q4_k_f32_ns;
+                    if (backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin) {
+                        kernel = backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin;
+                    }
 
                     // Reorder router if called from test-backend-ops or when new router is generated.
                     // Otherwise reuse the reordered result from previous mul_mat_id call.
@@ -18689,6 +18983,11 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     CL_CHECK(status);
                     cl_image_format image_format_buf_src1 = {CL_RGBA, CL_FLOAT};
                     cl_image_desc image_desc_buf_src1 = {CL_MEM_OBJECT_IMAGE1D_BUFFER, static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size / 4), 0,0,0,0,0,0,0, {buf_src1_reordered}};
+                    if (backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin) {
+                        // bin kernel uses slightly different image format
+                        image_format_buf_src1 = {CL_R, CL_FLOAT};
+                        image_desc_buf_src1.image_width = static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size);
+                    }
                     image_src1_reordered = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &image_format_buf_src1, &image_desc_buf_src1, NULL, &status);
                     CL_CHECK(status);
 
@@ -19172,6 +19471,9 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
 
                 } else { // for gemm
                     kernel = backend_ctx->kernel_gemm_moe_mxfp4_f32_ns;
+                    if (backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin) {
+                        kernel = backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin;
+                    }
 
                     // Reorder router if called from test-backend-ops or when new router is generated.
                     // Otherwise reuse the reordered result from previous mul_mat_id call.
@@ -19218,6 +19520,11 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
                     cl_image_desc image_desc_buf_src1;
                     image_format_buf_src1 = {CL_RGBA, CL_FLOAT};
                     image_desc_buf_src1 = {CL_MEM_OBJECT_IMAGE1D_BUFFER, static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size / 4), 0,0,0,0,0,0,0, {buf_src1_reordered}};
+                    if (backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin) {
+                        // bin kernel uses slightly different image format
+                        image_format_buf_src1 = {CL_R, CL_FLOAT};
+                        image_desc_buf_src1.image_width = static_cast<size_t>(ne00 * max_post_router_tile * n_tile_size);
+                    }
                     image_src1_reordered = clCreateImage(backend_ctx->context, CL_MEM_READ_ONLY, &image_format_buf_src1, &image_desc_buf_src1, NULL, &status);
                     CL_CHECK(status);
 

ggml/src/ggml-opencl/libdl.h

@@ -0,0 +1,79 @@
+#pragma once
+
+#ifdef _WIN32
+#   define WIN32_LEAN_AND_MEAN
+#   ifndef NOMINMAX
+#       define NOMINMAX
+#   endif
+#   include <windows.h>
+#   include <winevt.h>
+#else
+#    include <dlfcn.h>
+#    include <unistd.h>
+#endif
+#include <filesystem>
+
+namespace fs = std::filesystem;
+
+#ifdef _WIN32
+
+using dl_handle = std::remove_pointer_t<HMODULE>;
+
+struct dl_handle_deleter {
+    void operator()(HMODULE handle) {
+        FreeLibrary(handle);
+    }
+};
+
+static inline dl_handle * dl_load_library(const fs::path & path) {
+    // suppress error dialogs for missing DLLs
+    DWORD old_mode = SetErrorMode(SEM_FAILCRITICALERRORS);
+    SetErrorMode(old_mode | SEM_FAILCRITICALERRORS);
+
+    HMODULE handle = LoadLibraryW(path.wstring().c_str());
+
+    SetErrorMode(old_mode);
+
+    return handle;
+}
+
+static inline void * dl_get_sym(dl_handle * handle, const char * name) {
+    DWORD old_mode = SetErrorMode(SEM_FAILCRITICALERRORS);
+    SetErrorMode(old_mode | SEM_FAILCRITICALERRORS);
+
+    void * p = (void *) GetProcAddress(handle, name);
+
+    SetErrorMode(old_mode);
+
+    return p;
+}
+
+static inline const char * dl_error() {
+    return "";
+}
+
+#else
+
+using dl_handle = void;
+
+struct dl_handle_deleter {
+    void operator()(void * handle) {
+        dlclose(handle);
+    }
+};
+
+static inline dl_handle * dl_load_library(const fs::path & path) {
+    dl_handle * handle = dlopen(path.string().c_str(), RTLD_NOW | RTLD_LOCAL);
+    return handle;
+}
+
+static inline void * dl_get_sym(dl_handle * handle, const char * name) {
+    return dlsym(handle, name);
+}
+
+static inline const char * dl_error() {
+    const char *rslt = dlerror();
+    return rslt != nullptr ? rslt : "";
+}
+
+#endif

ggml/src/ggml-openvino/ggml-openvino.cpp

@@ -703,6 +703,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;

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 <stdio.h>
+#include <string.h>
+
+#ifdef _WIN32
+#    define WIN32_LEAN_AND_MEAN
+#    ifndef NOMINMAX
+#        define NOMINMAX
+#    endif
+#    include <windows.h>
+#else
+#    include <time.h>
+#    include <unistd.h>
+#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;
+}

ggml/src/ggml-rpc/ggml-rpc.cpp

@@ -785,7 +785,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;
 }

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -6200,7 +6200,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;

ggml/src/ggml-virtgpu/ggml-backend.cpp

@@ -65,6 +65,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;

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -1,4 +1,5 @@
 #include "ggml-vulkan.h"
+#include "ggml-profiler.h"
 #include <vulkan/vulkan_core.h>
 #if defined(GGML_VULKAN_RUN_TESTS) || defined(GGML_VULKAN_CHECK_RESULTS)
 #include <chrono>
@@ -1900,8 +1901,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;
@@ -2091,6 +2092,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<ggml_profile_record> records;
+    std::vector<uint64_t>            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;
 
@@ -2151,8 +2167,9 @@ struct ggml_backend_vk_context {
     topk_moe_mode fused_topk_moe_mode {};
     bool fused_topk_moe_scale {};
 
-    // for GGML_VK_PERF_LOGGER
-    std::unique_ptr<vk_perf_logger> perf_logger;
+    // Profiling
+    std::unique_ptr<vk_perf_logger> perf_logger;   // legacy env-var profiler
+    ggml_vk_profiler_state * profiler_state = nullptr;
     vk::QueryPool query_pool;
     std::vector<const char *> query_fusion_names;
     std::vector<int> query_fusion_node_count;
@@ -14589,10 +14606,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.
@@ -15136,7 +15157,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);
     }
 }
@@ -16258,7 +16279,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) {
@@ -16286,6 +16309,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;
@@ -16536,7 +16563,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
@@ -16544,6 +16571,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;
@@ -16577,7 +16607,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();
@@ -16591,15 +16621,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<uint64_t> 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");
+
+        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];
@@ -16614,10 +16669,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) {
@@ -16973,13 +17058,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;
 }
 

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
 
 #if defined(DATA_A_F32)
 #define QUANT_K 1
@@ -1756,7 +1757,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();
@@ -1793,11 +1796,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

ggml/src/ggml-webgpu/ggml-webgpu.cpp

@@ -4004,6 +4004,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;
 }

ggml/src/ggml-zdnn/ggml-zdnn.cpp

@@ -503,7 +503,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;

ggml/src/ggml-zendnn/ggml-zendnn.cpp

@@ -467,7 +467,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;

include/llama.h

@@ -555,6 +555,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);
 

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"
@@ -2446,6 +2447,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;
@@ -3618,6 +3624,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,

tests/test-backend-ops.cpp

@@ -9605,7 +9605,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
     return test_cases;
 }
 
-static std::vector<std::unique_ptr<test_case>> make_test_cases_from_file(const char * path) {
+static std::vector<std::unique_ptr<test_case>> make_test_cases_from_file(const char * path, const char * backend_name = nullptr) {
     std::ifstream f(path);
 
     if (!f.is_open()) {
@@ -9663,6 +9663,17 @@ static std::vector<std::unique_ptr<test_case>> 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)));
     }
 
@@ -9702,7 +9713,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);
@@ -9965,12 +9976,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 <n> runs tests using <n> 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;
@@ -9981,12 +9993,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];
@@ -10048,6 +10064,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();
 
@@ -10071,7 +10091,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++;
@@ -10084,11 +10105,10 @@ int main(int argc, char ** argv) {
         ggml_backend_reg_t reg = ggml_backend_dev_backend_reg(dev);
         auto ggml_backend_set_n_threads_fn = (ggml_backend_set_n_threads_t) ggml_backend_reg_get_proc_address(reg, "ggml_backend_set_n_threads");
         if (ggml_backend_set_n_threads_fn) {
-            // TODO: better value for n_threads
             ggml_backend_set_n_threads_fn(backend, 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),

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 <set>
 #include <fstream>
 #include <iostream>
-#include <random>
 
 // 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<int32_t> op_params;
     std::vector<input_tensor> 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<test_object> & tests) {
+static void extract_graph_ops(ggml_cgraph * cgraph, const char * label, std::set<test_object> & 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;
 
@@ -195,19 +208,21 @@ int main(int argc, char ** argv) {
 
     std::set<test_object> 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());
 

tools/cli/cli.cpp

@@ -672,6 +672,29 @@ int llama_cli(int argc, char ** argv) {
     ctx_cli.ctx_server.terminate();
     inference_thread.join();
 
+    // Export profiling data if profiling was enabled
+    if (params.profiling) {
+        ggml_backend_sched_t sched = llama_context_get_sched(ctx_cli.ctx_server.get_llama_context());
+        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) {
+                    console::log("\nProfiling data exported to: %s\n", path.c_str());
+                } else {
+                    console::error("\nFailed to export profiling data to: %s\n", path.c_str());
+                }
+            }
+        }
+    }
+
     // bump the log level to display timings
     common_log_set_verbosity_thold(LOG_LEVEL_INFO);
     common_memory_breakdown_print(ctx_cli.ctx_server.get_llama_context());

tools/completion/completion.cpp

@@ -988,6 +988,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);
 

tools/profiler/__init__.py

@@ -0,0 +1 @@
+# llama.cpp profiler analysis tools