public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-19 07:08:57 +02:00
Code Issues Packages Projects Releases Wiki Activity

metal : single-source fusion table + fusion debug rework (llama/28164)

Browse Source 
* metal : rework fusion patterns into a single table

All fusable op patterns for the Metal backend are now declared once in a
fusion table (ggml-metal-fuse.cpp) and consumed by both the graph optimizer
(ggml_metal_fuse_max, packing) and the op encoders (ggml_metal_fuse_next,
compute). The two phases share the same pattern table plus ggml_can_fuse_subgraph_ext
for the structural checks, and differ only in the mode used for the pattern
check (STRUCTURAL at optimize time, since tensors are not allocated yet, and
FULL at compute time, including Metal buffer placement). This also protects the
snake activation (MUL + SIN + SQR + MUL + ADD) from being reordered during graph
optimization, which was previously unprotected.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : fix absolute output indices in fusion patterns

ggml_can_fuse_subgraph_ext expects the outputs array to contain absolute graph
node indices (it indexes cgraph->nodes[outputs[i]]), but the fusion table query
was passing a relative index (n_ops - 1). As a result the last node of every
pattern was not recognized as an output and was subjected to the elidable
use-count check, which failed for essentially all fusions. This silently
disabled the norm/MUL fusion and caused a ~5% token-generation regression.

Pass the absolute graph index of the last node instead.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : fuse gated_delta_net with cache cpy

Add GGML_METAL_FUSE_GDN_CACHE to the fusion table: when the gated_delta_net
kernel is followed by a cpy that scatters its recurrent state snapshots into
the KV cache, the kernel writes the snapshots straight into the cache buffer
and the trailing cpy is elided.

The gdn output has other consumers (the attn scores view), so unlike the
elision-chain patterns this is not a simple chain: a 'raw' flag on the fusion
pattern skips the generic chain/shape and ggml_can_fuse_subgraph_ext checks,
making the pattern-specific check callback the sole validator. Packing
(ggml_metal_fuse_max) now matches on the same view-transparent node sequence
that the compute phase uses, so the gdn + cache cpy group is packed along with
any intermediate views and stays adjacent through the reorder.

The fused cpy is a view consumer of the gdn (it writes the cache directly),
so its mem-range is skipped in the encoder; the skip is restricted to CPY
nodes consuming the previous fused node through a view so other fusions are
unaffected.

Add test_gated_delta_net_cache_fusion and register 5 cases.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : drop is_view_consumer mem-range skip

The is_view_consumer skip was carried over from the upstream gated_delta_net
cache-fusion draft, but it is not needed: keeping the elided cpy's mem-range in
the concurrency tracker only ever adds a (conservative) memory barrier at the
fusion point. It can never remove a barrier, so it cannot introduce a race. The
worst case is one spurious barrier per gdn+cache-cpy fusion, which is within
run-to-run noise on Qwen3.5-0.8B Q8_0.

Dropping the check keeps the mem-range loop uniform for all fused groups.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : rename gated_delta_net fused state output args

Rename the fused cache-write kernel argument to match the rest of the kargs:
state_out_stride -> nb_out (and widen it to uint64_t), and the local buffer id
bid_state_out -> bid_out.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : rename raw fusion flag to unsafe

raw did not convey that the flag opts a fusion pattern out of the generic
elision-chain safety net (ggml_can_fuse_subgraph_ext + chain/shape checks).
rename it to 'unsafe' to make explicit that the pattern's check callback is the
sole validator and must re-establish the safety guarantees itself.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : tidy fusion pattern checks and table

- const-correct ggml_metal_fuse_outputs buffer
- annotate unused check-callback parameters
- drop a redundant size_t cast
- align the ops/table initializers and add blank-line separation

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : add generic fusion stats via ad-hoc proc-address API

Add a device-owned fusion context that lets a test tool count how many
times each fusion pattern fires and toggle fusion. It is exposed through
the ad-hoc ggml_backend_reg_get_proc_address mechanism with generic names
so the testing tool is backend-agnostic:

- ggml_backend_fusion_stats_init: start collecting fusion stats; when a
  context is created afterwards it registers the labels/counters and
  encodes single-threaded (n_cb == 0) so the counters are race-free
- ggml_backend_fusion_stats_reset / _get_stats / _set_enabled

The context lives on the metal device (not on the last backend context),
so counters accumulate across contexts and reads are always consistent.
The enable/disable toggle is initialized from GGML_METAL_FUSION_DISABLE
and can be overridden by the test through set_enabled. Labels are
synthesized from the fuse table via ggml_metal_fuse_label (e.g.
"GATED_DELTA_NET+CPY").

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : add fusion count regression test with per-backend baseline

test-fusion runs every dummy model generated by test-llama-archs on a
single backend (single-threaded encoding, n_cb == 0) with fusion enabled
and disabled, and for each mode (prefill / decode) reports the per-fusion
counters and the NMSE between the fused and unfused logits, plus the NMSE
against a CPU reference.

A fusion pattern that silently stops matching (or fires when it should
not) is caught as a regression by comparing the counters against a
committed per-backend TSV baseline:

- --record writes the golden baseline, --check (default) validates it
- the unfused run doubles as a control: its counters must be all-zero
- NMSE is skipped when it is NaN or the arch is already broken on the
  device (e.g. plamo2 on Metal), so the count check is the hard gate
- baseline counts depend only on graph structure, not weights (verified
  stable across weight seeds)
- the fusion stats API is resolved through the ad-hoc get_proc_address
  mechanism with generic names; a backend that does not export it makes
  the test fail with an error

The committed MTL0.tsv baseline covers 110 dummy archs (298 rows).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : rename fusion api helpers to match stats_init signature

Align the test with the ad-hoc fusion stats API: fusion_stats_init no
longer takes an enable bool (stats are turned on by calling it), so the
proc-address wrappers and typedefs are renamed to the api_* convention.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : rename backend to device in fusion test CLI

The fusion test operates on a compute device (e.g. MTL0), not a backend,
so rename the --backend argument to --device and the backend_name
variable to device_name. Keep "backend" where it refers to the ggml
backend interface (the ad-hoc proc-address mechanism).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : add --model and --help to fusion test

--model FILE runs the fusion regression test over a single model file
instead of enumerating a --models DIR. --models and --model are mutually
exclusive. Also add a --help/-h option that prints the usage.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : use backend base name for fusion baseline output

The fusion test is invoked with a specific device name (e.g. MTL0), but
its output - the recorded baseline and the header it writes - should be
named after the backend base name (e.g. MTL, via ggml_backend_reg_name),
since the counters depend on the backend, not on the specific device
index. Rename the committed baseline to MTL.tsv.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : run fusion test from ci instead of ctest

The fusion test needs Metal and generates a lot of dummy models, so it
does not belong in the generic ctest suite. Move it to ci/run.sh as
gg_run_test_fusion, gated on GG_BUILD_METAL like
gg_run_test_llama_archs_tensor_split: it generates the dummy models with
test-llama-archs -o and then validates the fusion counts against the
committed baseline. test-fusion.cpp is still built (llama_build) but no
longer registered as a ctest.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : align fusion baseline TSV columns

Pad the TSV fields to fixed widths so the columns line up regardless of
the variable arch and fusion-label lengths, and trim each field on parse
so the padded file is still accepted. Regenerate the committed MTL.tsv
baseline in the padded format (data unchanged, verified identical modulo
padding).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : widen label column and align fusion TSV header

Give the label column more room (28 chars) and fix the column header
widths so they match the data rows (moe/mode/label), keeping the header
aligned with the values. Regenerate the MTL.tsv baseline in the new
format (data unchanged).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : switch fusion baseline from TSV to CSV

Use comma-separated values like the rest of the project, keeping the
padded, aligned columns. Split on ',' and trim on parse. Rename the
committed baseline to MTL.csv (data unchanged, verified identical modulo
padding/separator). Update the ci/run.sh check path accordingly.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* cont : rebase + update MTL stats

* tests : avoid graph reallocations for some archs

* metal : tidy fusion debugging context and op init

- simplify the shared fusion debugging context comments
- shorten the ggml_metal_fusion struct comment
- align the ggml_metal_fuse struct fields and comments
- move the fusion parameter of ggml_metal_op_init right after dev

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : dedup fusion baseline into any mode

prefill and decode always produce the same per-graph fusion count, so
store a single row per label with mode = "any" and the per-graph count
instead of two rows. this halves the baseline size and keeps the check
stable.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* ci : move fusion model generation to a separate step

the dummy models generated by test-llama-archs are reused by other tests,
so generate them once in their own step instead of inside test_fusion.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : bump nmse thold

* models : fix plamo2 graph

* tests : remove "skip" logic from test-fusion

* tests : set qwen3tts dummy vocab to codec head size

the dummy qwen3tts model used a vocab of 4096 while the codec head is
3072, so the graph padded the output with -inf which made the NMSE in
test-fusion produce NaN. use the exact codec head size instead so the
padding is not generated at all.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : regen fusion baseline

reflect the plamo2 graph fix, which changed its fusion pattern split
(RMS_NORM+MUL 11->10, RMS_NORM+MUL+ADD 3->4; same total).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* ci : skip dummy model generation on OpenVINO

test-llama-archs does not build on the OpenVINO platform, so do not try
to generate the dummy models there.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* cont : minor

* tests : enable test-llama-archs on windows

* cont : disable on windows + workaround

* metal : naming nits

* test-fusion : add instructions to update baseline

* context : fix Kimi-K3 graph reserve

* fusion : update MTL

* cont : fix naming

* metal : rework fusion info storage

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* metal : align fusion info API

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* metal : use opaque fusion handle in ad-hoc API

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* ci : move fusion test to dedicated workflow

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* cont : run only on ggml changes

* cont : simplify

* fusion : remove multi-output stuff for now

* ci : fix typo
This commit is contained in:
Georgi Gerganov
2026-09-14 20:45:06 +03:00
parent 59cca2c760
commit 9d70ac2f1f
13 changed files with 844 additions and 232 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-metal/CMakeLists.txt
+1
View File
@@ -10,6 +10,7 @@ ggml_add_backend_library(ggml-metal
ggml-metal-device.cpp
ggml-metal-common.cpp
ggml-metal-context.m
ggml-metal-fusion.cpp
ggml-metal-ops.cpp
ggml-metal-tuning.cpp
)
ggml/src/ggml-metal/ggml-metal-common.cpp
+10 -37
View File
@@ -1,4 +1,5 @@
#include "ggml-metal-common.h"
#include "ggml-metal-fusion.h"
#include "ggml.h"
#include "ggml-impl.h"
@@ -390,59 +391,31 @@ static std::vector<int> ggml_metal_graph_optimize_reorder(const std::vector<node
}
void ggml_graph_optimize(ggml_cgraph * gf) {
constexpr int MAX_FUSE = 16;
const int n = gf->n_nodes;
enum ggml_op ops[MAX_FUSE];
std::vector<node_info> nodes;
nodes.reserve(gf->n_nodes);
// fuse nodes:
// we don't want to make reorders that break fusing, so we first pack all fusable tensors
// and perform the reorder over the fused nodes. after the reorder is done, we unfuse
//
// the fusable sequences are declared in the fusion table (ggml-metal-fuse.cpp), so the
// packing here is driven by the same patterns that the op encoders will later use
for (int i = 0; i < n; i++) {
node_info node = {
/*.node =*/ gf->nodes[i],
/*.fused =*/ {},
};
// fuse only ops that start with these operations
// can be expanded when needed
if (node.op() == GGML_OP_ADD ||
node.op() == GGML_OP_NORM ||
node.op() == GGML_OP_RMS_NORM) {
ops[0] = node.op();
const int f = ggml_metal_fusion_max(gf, i);
int f = i + 1;
while (f < n && f < i + MAX_FUSE) {
// conservatively allow fusing only these ops
// can be expanded when needed
if (gf->nodes[f]->op != GGML_OP_ADD &&
gf->nodes[f]->op != GGML_OP_MUL &&
gf->nodes[f]->op != GGML_OP_NORM &&
gf->nodes[f]->op != GGML_OP_RMS_NORM) {
break;
}
ops[f - i] = gf->nodes[f]->op;
f++;
}
// add the fused tensors into the node info so we can unfuse them later
for (int k = 1; k < f; k++) {
++i;
f -= i;
for (; f > 1; f--) {
if (ggml_can_fuse(gf, i, ops, f)) {
break;
}
}
// add the fused tensors into the node info so we can unfuse them later
for (int k = 1; k < f; k++) {
++i;
// the .dst() becomes the last fused tensor
node.add_fused(gf->nodes[i]);
}
// the .dst() becomes the last fused tensor
node.add_fused(gf->nodes[i]);
}
nodes.push_back(std::move(node));
ggml/src/ggml-metal/ggml-metal-context.h
+1
View File
@@ -33,6 +33,7 @@ ggml_metal_event_t ggml_metal_get_ev_cpy(ggml_metal_t ctx);
void ggml_metal_set_n_cb (ggml_metal_t ctx, int n_cb);
void ggml_metal_set_abort_callback (ggml_metal_t ctx, ggml_abort_callback abort_callback, void * user_data);
bool ggml_metal_supports_family (ggml_metal_t ctx, int family);
void ggml_metal_capture_next_compute(ggml_metal_t ctx);
ggml/src/ggml-metal/ggml-metal-context.m
+33 -22
View File
@@ -6,6 +6,7 @@
#import "ggml-metal-impl.h"
#import "ggml-metal-common.h"
#import "ggml-metal-ops.h"
#import "ggml-metal-fusion.h"
#import <Foundation/Foundation.h>
@@ -36,15 +37,12 @@ struct ggml_metal {
// additional, inference-time compiled pipelines
ggml_metal_pipelines_t pipelines_ext;
bool use_fusion;
bool use_concurrency;
bool use_graph_optimize;
int debug_graph;
int debug_fusion;
// how many times a given op was fused
uint64_t fuse_cnt[GGML_OP_COUNT];
struct ggml_metal_fusion_info * finfo;
// capture state
int capture_compute;
@@ -139,7 +137,6 @@ ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) {
res->d_queue = dispatch_queue_create("ggml-metal", DISPATCH_QUEUE_CONCURRENT);
res->use_fusion = getenv("GGML_METAL_FUSION_DISABLE") == nil;
res->use_concurrency = getenv("GGML_METAL_CONCURRENCY_DISABLE") == nil;
{
@@ -147,20 +144,19 @@ ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) {
res->debug_graph = val ? atoi(val) : 0;
}
{
const char * val = getenv("GGML_METAL_FUSION_DEBUG");
res->debug_fusion = val ? atoi(val) : 0;
}
res->use_graph_optimize = true;
if (getenv("GGML_METAL_GRAPH_OPTIMIZE_DISABLE") != NULL) {
res->use_graph_optimize = false;
}
memset(res->fuse_cnt, 0, sizeof(res->fuse_cnt));
res->finfo = ggml_metal_device_get_fusion_info(dev);
if (ggml_metal_fusion_info_stats(res->finfo)) {
ggml_metal_fusion_info_labels_init(res->finfo);
res->n_cb = 0;
}
GGML_LOG_INFO("%s: use fusion = %s\n", __func__, res->use_fusion ? "true" : "false");
GGML_LOG_INFO("%s: use fusion = %s\n", __func__, ggml_metal_fusion_info_enabled(res->finfo) ? "true" : "false");
GGML_LOG_INFO("%s: use concurrency = %s\n", __func__, res->use_concurrency ? "true" : "false");
GGML_LOG_INFO("%s: use graph optimize = %s\n", __func__, res->use_graph_optimize ? "true" : "false");
@@ -222,15 +218,18 @@ void ggml_metal_free(ggml_metal_t ctx) {
ctx->pipelines_ext = nil;
}
if (ctx->debug_fusion > 0) {
if (ggml_metal_fusion_info_debug(ctx->finfo) > 0) {
GGML_LOG_DEBUG("%s: fusion stats:\n", __func__);
for (int i = 0; i < GGML_OP_COUNT; i++) {
if (ctx->fuse_cnt[i] == 0) {
const int n_fusions = ggml_metal_fusion_info_n_fusions(ctx->finfo);
for (int i = 0; i < n_fusions; i++) {
const uint64_t count = ggml_metal_fusion_info_count(ctx->finfo, i);
if (count == 0) {
continue;
}
// note: cannot use ggml_log here
GGML_LOG_DEBUG("%s: - %s: %" PRIu64 "\n", __func__, ggml_op_name((enum ggml_op) i), ctx->fuse_cnt[i]);
GGML_LOG_DEBUG("%s: - %s: %" PRIu64 "\n", __func__, ggml_metal_fusion_info_label(ctx->finfo, i), count);
}
}
@@ -481,10 +480,17 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph *
@autoreleasepool {
ctx->gf = gf;
ctx->n_nodes_0 = MIN(n_main, gf->n_nodes);
ctx->n_nodes_1 = gf->n_nodes - ctx->n_nodes_0;
if (ctx->n_cb == 0) {
// single-threaded encoding: the whole graph is encoded by one command buffer
ctx->n_nodes_0 = gf->n_nodes;
ctx->n_nodes_1 = 0;
ctx->n_nodes_per_cb = 0;
} else {
ctx->n_nodes_0 = MIN(n_main, gf->n_nodes);
ctx->n_nodes_1 = gf->n_nodes - ctx->n_nodes_0;
ctx->n_nodes_per_cb = (ctx->n_nodes_1 + ctx->n_cb - 1) / ctx->n_cb;
ctx->n_nodes_per_cb = (ctx->n_nodes_1 + ctx->n_cb - 1) / ctx->n_cb;
}
if (ctx->capture_compute >= 0) {
ctx->capture_compute--;
@@ -682,6 +688,12 @@ ggml_metal_event_t ggml_metal_get_ev_cpy(ggml_metal_t ctx) {
}
void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) {
// when fusion stats are collected the graph must be encoded by a single thread so the
// counters are race-free; override whatever the caller requested
if (ggml_metal_fusion_info_stats(ctx->finfo)) {
n_cb = 0;
}
if (ctx->n_cb != n_cb) {
ctx->n_cb = MIN(n_cb, GGML_METAL_MAX_COMMAND_BUFFERS);
@@ -717,13 +729,12 @@ void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) {
ctx->dev,
cmd_buf,
ctx->gf,
ctx->finfo,
idx_start,
idx_end,
ctx->use_fusion,
ctx->use_concurrency,
ctx->capture_compute,
ctx->debug_graph,
ctx->debug_fusion);
ctx->debug_graph);
for (int idx = 0; idx < ggml_metal_op_n_nodes(ctx_op); ++idx) {
const int res = ggml_metal_op_encode(ctx_op, idx);
ggml/src/ggml-metal/ggml-metal-device.h
+5
View File
@@ -325,6 +325,11 @@ 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);
struct ggml_metal_fusion_info;
// the device-owned fusion debugging context (NULL unless fusion debugging is enabled)
struct ggml_metal_fusion_info * ggml_metal_device_get_fusion_info(ggml_metal_device_t dev);
//
// device buffers
//
ggml/src/ggml-metal/ggml-metal-device.m
+17
View File
@@ -1,4 +1,5 @@
#import "ggml-metal-device.h"
#import "ggml-metal-fusion.h"
#import "ggml-impl.h"
#import "ggml-backend-impl.h"
@@ -896,6 +897,9 @@ struct ggml_metal_device {
struct ggml_metal_device_props props;
// shared fusion debugging context
struct ggml_metal_fusion_info * finfo;
// virtual address for GPU memory allocations
atomic_uintptr_t addr_virt;
};
@@ -1274,6 +1278,13 @@ ggml_metal_device_t ggml_metal_device_init(int device, int n_devices) {
dev->props.max_working_set_size = dev->mtl_device.maxBufferLength;
}
{
const char * val = getenv("GGML_METAL_FUSION_DEBUG");
dev->finfo = ggml_metal_fusion_info_init(
getenv("GGML_METAL_FUSION_DISABLE") == nil,
val ? atoi(val) : 0);
}
snprintf(dev->props.name, sizeof(dev->props.name), "%s%d", "MTL", device);
const char * gpu_name = [[dev->mtl_device name] UTF8String];
if (n_devices > 1) {
@@ -1348,6 +1359,8 @@ void ggml_metal_device_free(ggml_metal_device_t dev) {
assert(dev != NULL);
@autoreleasepool {
ggml_metal_fusion_info_free(dev->finfo);
ggml_metal_rsets_free(dev->rsets);
ggml_metal_library_free(dev->library);
@@ -1935,6 +1948,10 @@ static void ggml_metal_device_disable_tensor(ggml_metal_device_t dev) {
dev->props.has_tensor = false;
}
struct ggml_metal_fusion_info * ggml_metal_device_get_fusion_info(ggml_metal_device_t dev) {
return dev->finfo;
}
//
// device buffers
//
ggml/src/ggml-metal/ggml-metal-fusion.cpp
+502
View File
@@ -0,0 +1,502 @@
#include "ggml-metal-fusion.h"
#include "ggml-backend-impl.h"
#include "ggml-metal-device.h"
#include <algorithm>
#include <string>
#include <vector>
// ---- helpers -------------------------------------------------------------
// true if two tensors live in the same Metal buffer
static bool ggml_metal_fusion_same_buffer(const ggml_tensor * a, const ggml_tensor * b) {
if (!a || !b) {
return false;
}
ggml_backend_buffer_t ba = a->view_src ? a->view_src->buffer : a->buffer;
ggml_backend_buffer_t bb = b->view_src ? b->view_src->buffer : b->buffer;
ggml_metal_buffer_t ca = (ggml_metal_buffer_t) ba->context;
ggml_metal_buffer_t cb = (ggml_metal_buffer_t) bb->context;
return ggml_metal_buffer_get_id(ca, a).metal == ggml_metal_buffer_get_id(cb, b).metal;
}
// ---- pattern checks ------------------------------------------------------
// NORM/RMS_NORM + MUL + ADD: the weight/bias of each fused step must match the norm input
// width, be contiguous rows, and the fused outputs must stay F32
static bool ggml_metal_fusion_check_norm(
const ggml_metal_fusion * fusion,
const ggml_tensor * const * nodes,
ggml_metal_fusion_mode mode) {
GGML_UNUSED(mode);
GGML_ASSERT(fusion->n_ops >= 2);
for (int j = 1; j < fusion->n_ops; j++) {
// the fused MUL/ADD must read the previous node as src0
if (nodes[j]->src[0] != nodes[j - 1]) {
return false;
}
// the weight/bias must have the same row width as the norm input
if (nodes[j]->src[1]->ne[0] != nodes[0]->ne[0]) {
return false;
}
if (!ggml_is_contiguous_rows(nodes[j]->src[1])) {
return false;
}
if (nodes[j]->type != GGML_TYPE_F32) {
return false;
}
}
return true;
}
// ADD x N: each ADD reads the previous ADD as src0, and all addends must share layout
// (and, in FULL mode, live in the same Metal buffer)
static bool ggml_metal_fusion_check_add_chain(
const ggml_metal_fusion * fusion,
const ggml_tensor * const * nodes,
ggml_metal_fusion_mode mode) {
GGML_ASSERT(fusion->n_ops >= 2);
for (int j = 1; j < fusion->n_ops; j++) {
if (nodes[j]->src[0] != nodes[j - 1]) {
return false;
}
if (!ggml_are_same_layout(nodes[j]->src[1], nodes[j - 1]->src[1])) {
return false;
}
if (mode == GGML_METAL_FUSION_FULL) {
if (!ggml_metal_fusion_same_buffer(nodes[j]->src[1], nodes[0]->src[1])) {
return false;
}
}
}
return true;
}
// GATED_DELTA_NET + CPY: the trailing cpy scatters the gdn state snapshots into the recurrent
// cache, so the gdn kernel writes them straight to the cache and the cpy is elided.
// mirrors ggml_metal_op_can_fuse_gdn_cache (PR #25788). the gdn output has other consumers (the
// attn scores view), so unlike the other patterns this is not an elision chain: the structural
// checks live entirely in this callback (unsafe = true).
static bool ggml_metal_fusion_check_gdn_cache(
const ggml_metal_fusion * fusion,
const ggml_tensor * const * nodes,
ggml_metal_fusion_mode mode) {
GGML_UNUSED(fusion);
const ggml_tensor * gdn = nodes[0];
const ggml_tensor * cpy = nodes[1];
// the kernel skips the snapshot tail, so the gdn output must not be a graph output
if (gdn->type != GGML_TYPE_F32 || (gdn->flags & GGML_TENSOR_FLAG_OUTPUT)) {
return false;
}
if (cpy->op != GGML_OP_CPY || (cpy->flags & GGML_TENSOR_FLAG_OUTPUT)) {
return false;
}
const int64_t S_v = gdn->src[2]->ne[0];
const int64_t H = gdn->src[2]->ne[1];
const int64_t n_tokens = gdn->src[2]->ne[2];
const int64_t n_seqs = gdn->src[2]->ne[3];
const int64_t K = ggml_get_op_params_i32(gdn, 0);
const size_t tail_off = ggml_row_size(GGML_TYPE_F32, S_v * H * n_tokens * n_seqs);
const int64_t D = S_v * S_v * H;
const int64_t n_written = std::min<int64_t>(n_tokens, K);
const ggml_tensor * src = cpy->src[0]; // gdn snapshot tail view
const ggml_tensor * dst = cpy->src[1]; // cache view
// src must be this gdn's snapshot tail (contiguous, at the tail offset)
if (src->op != GGML_OP_VIEW || src->view_src != gdn ||
src->view_offs != tail_off || !ggml_is_contiguous(src)) {
return false;
}
const int64_t expected_ne[GGML_MAX_DIMS] = { D, n_seqs, n_written, 1 };
if (dst->type != GGML_TYPE_F32 ||
!std::equal(expected_ne, expected_ne + GGML_MAX_DIMS, dst->ne) ||
dst->nb[0] != ggml_type_size(GGML_TYPE_F32) ||
dst->nb[1] != ggml_row_size(GGML_TYPE_F32, D)) {
return false;
}
if (mode == GGML_METAL_FUSION_FULL) {
// the cache must be allocated so the kernel can write straight to its buffer
if (dst->data == nullptr) {
return false;
}
}
return true;
}
// MUL + SIN + SQR + MUL + ADD (snake activation)
static bool ggml_metal_fusion_check_snake(
const ggml_metal_fusion * fusion,
const ggml_tensor * const * nodes,
ggml_metal_fusion_mode mode) {
GGML_UNUSED(fusion);
GGML_UNUSED(mode);
const ggml_tensor * mul0 = nodes[0];
const ggml_tensor * sin_node = nodes[1];
const ggml_tensor * sqr = nodes[2];
const ggml_tensor * mul1 = nodes[3];
const ggml_tensor * add = nodes[4];
// x carries the full activation shape, a is the broadcast operand
const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1];
const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0];
// mul1 reads sqr and inv_b in either operand order
const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0];
// closure check: the trailing add reads the same x as the leading mul
const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0];
// x is in the supported whitelist and every chain intermediate shares x's type.
// a and inv_b bind as device const float * in the kernel, so they stay F32.
const bool types_ok =
(x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) &&
(a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) &&
(mul0->type == x->type) && (sin_node->type == x->type) &&
(sqr->type == x->type) && (mul1->type == x->type) &&
(add->type == x->type);
// a / inv_b collapse to [1, C, 1, 1], x and add stay 2D
const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];
const bool dim_ok =
(x->ne[2] == 1) && (x->ne[3] == 1) &&
(add->ne[2] == 1) && (add->ne[3] == 1) &&
(a->ne[2] == 1) && (a->ne[3] == 1) &&
(inv_b->ne[2] == 1) && (inv_b->ne[3] == 1);
// kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous
const bool contig_ok =
ggml_is_contiguous(x) && ggml_is_contiguous(add) &&
ggml_is_contiguous(a) && ggml_is_contiguous(inv_b);
return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x;
}
// ---- patterns ------------------------------------------------------------
static const ggml_op ops_norm_mul[] = { GGML_OP_NORM, GGML_OP_MUL };
static const ggml_op ops_norm_mul_add[] = { GGML_OP_NORM, GGML_OP_MUL, GGML_OP_ADD };
static const ggml_op ops_rms_norm_mul[] = { GGML_OP_RMS_NORM, GGML_OP_MUL };
static const ggml_op ops_rms_norm_mul_add[] = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD };
static const ggml_op ops_add_2[] = { GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_add_3[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_add_4[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_add_5[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_add_6[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_add_7[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
static const ggml_op ops_snake[] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD };
static const ggml_op ops_gdn_cache[] = { GGML_OP_GATED_DELTA_NET, GGML_OP_CPY };
static const ggml_metal_fusion ggml_metal_fusions[] = {
{ GGML_METAL_FUSION_NORM_MUL, ops_norm_mul, 2, false, ggml_metal_fusion_check_norm },
{ GGML_METAL_FUSION_NORM_MUL_ADD, ops_norm_mul_add, 3, false, ggml_metal_fusion_check_norm },
{ GGML_METAL_FUSION_NORM_MUL, ops_rms_norm_mul, 2, false, ggml_metal_fusion_check_norm },
{ GGML_METAL_FUSION_NORM_MUL_ADD, ops_rms_norm_mul_add, 3, false, ggml_metal_fusion_check_norm },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_2, 2, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_3, 3, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_4, 4, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_5, 5, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_6, 6, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_7, 7, false, ggml_metal_fusion_check_add_chain },
{ GGML_METAL_FUSION_SNAKE, ops_snake, 5, false, ggml_metal_fusion_check_snake },
{ GGML_METAL_FUSION_GDN_CACHE, ops_gdn_cache, 2, true, ggml_metal_fusion_check_gdn_cache },
};
const ggml_metal_fusion * ggml_metal_fusion_all(int * n) {
*n = (int) sizeof(ggml_metal_fusions) / sizeof(ggml_metal_fusions[0]);
return ggml_metal_fusions;
}
// ---- shared fusion info ---------------------------------------------------
static std::string ggml_metal_fusion_label(const ggml_metal_fusion * fusion) {
GGML_ASSERT(fusion != nullptr);
std::string label;
for (int j = 0; j < fusion->n_ops; j++) {
if (j > 0) {
label += '+';
}
label += ggml_op_name(fusion->ops[j]);
}
return label;
}
struct ggml_metal_fusion_info {
std::vector<std::string> labels;
std::vector<uint64_t> counts;
bool enabled;
bool stats;
bool labels_set;
int debug;
};
struct ggml_metal_fusion_info * ggml_metal_fusion_info_init(bool enabled, int debug) {
ggml_metal_fusion_info * finfo = new ggml_metal_fusion_info;
finfo->enabled = enabled;
finfo->stats = debug > 0;
finfo->labels_set = false;
finfo->debug = debug;
if (finfo->stats) {
ggml_metal_fusion_info_labels_init(finfo);
}
return finfo;
}
void ggml_metal_fusion_info_free(struct ggml_metal_fusion_info * finfo) {
delete finfo;
}
bool ggml_metal_fusion_info_enabled(const struct ggml_metal_fusion_info * finfo) {
return finfo->enabled;
}
bool ggml_metal_fusion_info_stats(const struct ggml_metal_fusion_info * finfo) {
return finfo->stats;
}
int ggml_metal_fusion_info_debug(const struct ggml_metal_fusion_info * finfo) {
return finfo->debug;
}
int ggml_metal_fusion_info_n_fusions(const struct ggml_metal_fusion_info * finfo) {
return (int) finfo->labels.size();
}
const char * ggml_metal_fusion_info_label(const struct ggml_metal_fusion_info * finfo, int idx) {
GGML_ASSERT(idx >= 0 && idx < (int) finfo->labels.size());
return finfo->labels[idx].c_str();
}
uint64_t ggml_metal_fusion_info_count(const struct ggml_metal_fusion_info * finfo, int idx) {
GGML_ASSERT(idx >= 0 && idx < (int) finfo->counts.size());
return finfo->counts[idx];
}
void ggml_metal_fusion_info_count_fusion(struct ggml_metal_fusion_info * finfo, const struct ggml_metal_fusion * fusion) {
if (!finfo->stats || fusion == nullptr) {
return;
}
int n = 0;
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
int idx = -1;
for (int i = 0; i < n; i++) {
if (&all[i] == fusion) {
idx = i;
break;
}
}
if (idx >= 0 && idx < (int) finfo->counts.size()) {
finfo->counts[idx]++;
}
}
void ggml_metal_fusion_info_set_enabled(struct ggml_metal_fusion_info * finfo, bool enabled) {
finfo->enabled = enabled;
}
void ggml_metal_fusion_info_labels_init(struct ggml_metal_fusion_info * finfo) {
if (finfo->labels_set) {
return;
}
int n = 0;
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
finfo->labels.clear();
finfo->counts.assign(n, 0);
finfo->labels.reserve(n);
for (int i = 0; i < n; i++) {
finfo->labels.emplace_back(ggml_metal_fusion_label(&all[i]));
}
finfo->labels_set = true;
}
void ggml_metal_fusion_info_stats_init(struct ggml_metal_fusion_info * finfo) {
finfo->stats = true;
ggml_metal_fusion_info_labels_init(finfo);
}
void ggml_metal_fusion_info_stats_reset(struct ggml_metal_fusion_info * finfo) {
std::fill(finfo->counts.begin(), finfo->counts.end(), 0);
}
int ggml_metal_fusion_info_stats_get(const struct ggml_metal_fusion_info * finfo, const char ** labels, uint64_t * counts, int n) {
const int n_fusions = (int) finfo->labels.size();
if (labels == nullptr) {
return n_fusions;
}
const int n_fill = std::min(n, n_fusions);
for (int i = 0; i < n_fill; i++) {
labels[i] = finfo->labels[i].c_str();
if (counts != nullptr) {
counts[i] = finfo->counts[i];
}
}
return n_fill;
}
// ---- queries -------------------------------------------------------------
// find the longest pattern matching the node sequence starting at idx
// (idx is a position in node_idxs, which maps to graph node indices)
const ggml_metal_fusion * ggml_metal_fusion_next(
const ggml_cgraph * gf,
const int * node_idxs,
int n_idxs,
int idx,
ggml_metal_fusion_mode mode,
int * n_out) {
int n = 0;
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
const ggml_metal_fusion * res = nullptr;
int best = 1;
for (int i = 0; i < n; i++) {
const ggml_metal_fusion * fusion = &all[i];
// only look for a longer match than the current best
if (fusion->n_ops <= best) {
continue;
}
if (idx + fusion->n_ops > n_idxs) {
continue;
}
const ggml_tensor * nodes[GGML_METAL_FUSION_MAX];
// the op sequence must match exactly
bool ok = true;
for (int j = 0; j < fusion->n_ops; j++) {
nodes[j] = gf->nodes[node_idxs[idx + j]];
if (nodes[j]->op != fusion->ops[j]) {
ok = false;
break;
}
}
if (!ok) {
continue;
}
if (!fusion->unsafe) {
// common element-wise chain constraints: each node reads the previous one,
// and all nodes have the same shape
for (int j = 1; j < fusion->n_ops && ok; j++) {
if (nodes[j]->src[0] != nodes[j - 1] && nodes[j]->src[1] != nodes[j - 1]) {
ok = false;
break;
}
if (!ggml_are_same_shape(nodes[j], nodes[j - 1])) {
ok = false;
break;
}
}
if (!ok) {
continue;
}
// all current fusions are single-output elision chains, so the last node is the only output
// TODO: multi-output fusions: store pattern-relative offsets in the table and translate them here
int outputs_buf[1];
outputs_buf[0] = node_idxs[idx + fusion->n_ops - 1];
// structural subgraph checks (op sequence, elidable uses, view containment)
if (!ggml_can_fuse_subgraph_ext(gf, node_idxs + idx, fusion->n_ops, fusion->ops, outputs_buf, 1)) {
continue;
}
}
// pattern-specific checks (the sole validator for unsafe patterns)
if (fusion->check && !fusion->check(fusion, nodes, mode)) {
continue;
}
best = fusion->n_ops;
res = fusion;
}
*n_out = best;
return res;
}
// optimize phase: maximum number of nodes starting at idx (a raw sequential graph index) that
// could be fused, chaining patterns back-to-back. matching runs on the same filtered (view
// transparent) node sequence that the compute phase uses, so the returned count is the raw index
// span from idx to the last matched node (intermediate views are packed along).
int ggml_metal_fusion_max(const ggml_cgraph * gf, int idx) {
// an empty/view node cannot start a pattern - pack it alone
if (ggml_op_is_empty(gf->nodes[idx]->op) || ggml_is_empty(gf->nodes[idx])) {
return 1;
}
// collect the non-empty node indices starting at idx
int idxs[GGML_METAL_FUSION_MAX];
int n_idxs = 0;
for (int i = idx; i < gf->n_nodes && n_idxs < GGML_METAL_FUSION_MAX; i++) {
if (!ggml_op_is_empty(gf->nodes[i]->op) && !ggml_is_empty(gf->nodes[i])) {
idxs[n_idxs++] = i;
}
}
if (n_idxs == 0) {
return 1;
}
int total = 0;
int i_f = 0;
while (i_f < n_idxs && total < GGML_METAL_FUSION_MAX) {
int len = 1;
const ggml_metal_fusion * fusion = ggml_metal_fusion_next(gf, idxs, n_idxs, i_f, GGML_METAL_FUSION_STRUCTURAL, &len);
if (!fusion || total + len > GGML_METAL_FUSION_MAX) {
break;
}
total += len;
i_f += len;
}
if (i_f == 0) {
return 1;
}
// map the matched non-empty nodes back to the raw index span (views are included)
return std::min(GGML_METAL_FUSION_MAX, idxs[i_f - 1] - idx + 1);
}
ggml/src/ggml-metal/ggml-metal-fusion.h
+104
View File
@@ -0,0 +1,104 @@
// single source of truth for the fusions supported by the Metal backend
//
// every fusable subgraph is declared exactly once as a ggml_metal_fusion entry in
// the table in ggml-metal-fusion.cpp. both the graph optimizer (ggml_metal_fusion_max)
// and the op encoders (ggml_metal_fusion_next) consult this same table, so the two
// phases can never disagree about what can be fused.
#pragma once
#include "ggml-impl.h"
#include <stdbool.h>
#ifdef __cplusplus
extern "C" {
#endif
// the maximum number of nodes that can be fused in a single kernel
// (also the maximum length of a packed fusion group during graph optimization)
#define GGML_METAL_FUSION_MAX 16
typedef enum ggml_metal_fusion_mode {
// structural checks only; used by the graph optimizer, at which point the graph
// tensors are not allocated yet, so buffer placement cannot be verified
GGML_METAL_FUSION_STRUCTURAL = 0,
// full checks, including buffer placement; used by the op encoders
GGML_METAL_FUSION_FULL,
} ggml_metal_fusion_mode;
// identifier of each fusion pattern so the op encoders know which kernel to use
typedef enum ggml_metal_fusion_id {
GGML_METAL_FUSION_NONE = 0,
GGML_METAL_FUSION_NORM_MUL, // NORM/RMS_NORM + MUL
GGML_METAL_FUSION_NORM_MUL_ADD, // NORM/RMS_NORM + MUL + ADD
GGML_METAL_FUSION_ADD_CHAIN, // ADD x N (N in [2, 7])
GGML_METAL_FUSION_SNAKE, // MUL + SIN + SQR + MUL + ADD
GGML_METAL_FUSION_GDN_CACHE, // GATED_DELTA_NET + CPY (write snapshots into the recurrent cache)
} ggml_metal_fusion_id;
struct ggml_metal_fusion {
ggml_metal_fusion_id id;
const enum ggml_op * ops; // op sequence (fixed length)
int n_ops; // number of ops
// if unsafe: the generic chain/shape + ggml_can_fuse_subgraph checks are skipped and the
// check callback below is the sole validator (used for patterns that are not elision chains,
// e.g. the gdn + cache-cpy write-through fusion)
bool unsafe;
// extra backend constraints on top of ggml_can_fuse_subgraph
// nodes[j] is the j-th node of the pattern
bool (*check)(const struct ggml_metal_fusion * fusion,
const struct ggml_tensor * const * nodes,
ggml_metal_fusion_mode mode);
};
typedef struct ggml_metal_fusion ggml_metal_fusion;
// the single table of all fusions supported by the Metal backend
const ggml_metal_fusion * ggml_metal_fusion_all(int * n);
// ---- shared fusion info ---------------------------------------------------
// shared fusion debugging context, owned by the device; newly created backend contexts for that
// device register with it so the fusion counters are race-free and accumulate across contexts.
struct ggml_metal_fusion_info; // defined in ggml-metal-fusion.cpp
struct ggml_metal_fusion_info * ggml_metal_fusion_info_init(bool enabled, int debug);
void ggml_metal_fusion_info_free(struct ggml_metal_fusion_info * finfo);
bool ggml_metal_fusion_info_enabled(const struct ggml_metal_fusion_info * finfo);
bool ggml_metal_fusion_info_stats (const struct ggml_metal_fusion_info * finfo);
int ggml_metal_fusion_info_debug (const struct ggml_metal_fusion_info * finfo);
int ggml_metal_fusion_info_n_fusions(const struct ggml_metal_fusion_info * finfo);
const char * ggml_metal_fusion_info_label (const struct ggml_metal_fusion_info * finfo, int idx);
uint64_t ggml_metal_fusion_info_count (const struct ggml_metal_fusion_info * finfo, int idx);
void ggml_metal_fusion_info_count_fusion(struct ggml_metal_fusion_info * finfo, const struct ggml_metal_fusion * fusion);
void ggml_metal_fusion_info_set_enabled (struct ggml_metal_fusion_info * finfo, bool enabled);
void ggml_metal_fusion_info_stats_init ( struct ggml_metal_fusion_info * finfo);
void ggml_metal_fusion_info_stats_reset( struct ggml_metal_fusion_info * finfo);
int ggml_metal_fusion_info_stats_get (const struct ggml_metal_fusion_info * finfo, const char ** labels, uint64_t * counts, int n);
void ggml_metal_fusion_info_labels_init( struct ggml_metal_fusion_info * finfo);
// compute phase: longest fusion starting at idx (a position in node_idxs) that matches in `mode`.
// returns the matching pattern (nullptr if no fusion) and sets *n_out to the number of nodes consumed.
const ggml_metal_fusion * ggml_metal_fusion_next(
const struct ggml_cgraph * gf,
const int * node_idxs,
int n_idxs,
int idx,
ggml_metal_fusion_mode mode,
int * n_out);
// optimize phase: maximum number of nodes starting at idx (a raw sequential graph index) that
// could be fused, chaining patterns back-to-back. returns at least 1.
int ggml_metal_fusion_max(const struct ggml_cgraph * gf, int idx);
#ifdef __cplusplus
}
#endif
ggml/src/ggml-metal/ggml-metal-impl.h
+1
View File
@@ -985,6 +985,7 @@ typedef struct {
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
uint64_t nb_out; // 0 => snapshots are appended after the attn scores (unfused)
} ggml_metal_kargs_gated_delta_net;
typedef struct {
ggml/src/ggml-metal/ggml-metal-ops.cpp
+107 -167
View File
@@ -7,6 +7,7 @@
#include "ggml-metal-impl.h"
#include "ggml-metal-common.h"
#include "ggml-metal-device.h"
#include "ggml-metal-fusion.h"
#include "ggml-metal-tuning.h"
#include <cassert>
@@ -31,24 +32,22 @@ struct ggml_metal_op {
ggml_metal_device_t dev,
ggml_metal_cmd_buf_t cmd_buf,
ggml_cgraph * gf,
ggml_metal_fusion_info * finfo,
int idx_start,
int idx_end,
bool use_fusion,
bool use_concurrency,
bool use_capture,
int debug_graph,
int debug_fusion) {
int debug_graph) {
this->dev = dev;
this->lib = ggml_metal_device_get_library(dev);
this->enc = ggml_metal_encoder_init(cmd_buf, use_concurrency);
this->mem_ranges = ggml_mem_ranges_init(debug_graph);
this->finfo = finfo;
this->idx_start = idx_start;
this->idx_end = idx_end;
this->use_fusion = use_fusion;
this->use_concurrency = use_concurrency;
this->use_capture = use_capture;
this->debug_graph = debug_graph;
this->debug_fusion = debug_fusion;
this->gf = gf;
idxs.reserve(gf->n_nodes);
@@ -78,15 +77,24 @@ struct ggml_metal_op {
return ggml_graph_node(gf, idxs[i]);
}
bool can_fuse(int i0, const ggml_op * ops, int n_ops) const {
assert(use_fusion);
// consult the fusion table for the longest pattern starting at i0
// returns the matching pattern (nullptr if no fusion) and sets *n_out to the number of nodes
const ggml_metal_fusion * can_fuse(int i0, enum ggml_metal_fusion_mode mode, int * n_out) const {
assert(use_fusion());
assert(i0 >= 0 && i0 < n_nodes());
if (i0 + n_ops > n_nodes()) {
return false;
}
return ggml_metal_fusion_next(gf, idxs.data(), (int) idxs.size(), i0, mode, n_out);
}
return ggml_can_fuse_ext(gf, idxs.data() + i0, ops, n_ops);
// whether to attempt fusion; the toggle lives in the shared fusion debugging context owned
// by the device (initialized from GGML_METAL_FUSION_DISABLE, overridable by the test)
bool use_fusion() const {
return ggml_metal_fusion_info_enabled(finfo);
}
// record that a fusion fired, indexed by the matching table entry
void count_fusions(const ggml_metal_fusion * fusion) const {
ggml_metal_fusion_info_count_fusion(finfo, fusion);
}
ggml_metal_device_t dev;
@@ -94,12 +102,13 @@ struct ggml_metal_op {
ggml_metal_encoder_t enc;
ggml_mem_ranges_t mem_ranges;
bool use_fusion;
// shared fusion debugging context
ggml_metal_fusion_info * finfo;
bool use_concurrency;
bool use_capture;
int debug_graph;
int debug_fusion;
private:
ggml_cgraph * gf;
@@ -115,24 +124,22 @@ ggml_metal_op_t ggml_metal_op_init(
ggml_metal_device_t dev,
ggml_metal_cmd_buf_t cmd_buf,
ggml_cgraph * gf,
ggml_metal_fusion_info * finfo,
int idx_start,
int idx_end,
bool use_fusion,
bool use_concurrency,
bool use_capture,
int debug_graph,
int debug_fusion) {
int debug_graph) {
ggml_metal_op_t res = new ggml_metal_op(
dev,
cmd_buf,
gf,
finfo,
idx_start,
idx_end,
use_fusion,
use_concurrency,
use_capture,
debug_graph,
debug_fusion);
debug_graph);
return res;
}
@@ -1868,6 +1875,8 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
ggml_metal_library_t lib = ctx->lib;
ggml_metal_encoder_t enc = ctx->enc;
const bool use_fusion = ctx->use_fusion();
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
@@ -1880,6 +1889,31 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
auto pipeline = ggml_metal_library_get_pipeline_gated_delta_net(lib, op);
// when fused with the trailing cache cpy, the snapshots are written straight into the
// recurrent cache and the cpy is skipped (see GGML_METAL_FUSION_GDN_CACHE)
ggml_metal_buffer_id bid_out = ggml_metal_get_buffer_id(op);
uint64_t nb_out = 0;
int n_fuse = 1;
if (use_fusion) {
int n = 1;
const ggml_metal_fusion * fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
if (fusion && fusion->id == GGML_METAL_FUSION_GDN_CACHE) {
const ggml_tensor * dst_cache = ctx->node(idx + 1)->src[1]; // cache view
bid_out = ggml_metal_get_buffer_id(dst_cache);
nb_out = dst_cache->nb[2]/sizeof(float);
n_fuse = 2;
ctx->count_fusions(fusion);
if (debug_fusion > 1) {
GGML_LOG_DEBUG("%s: fuse: GATED_DELTA_NET + CPY\n", __func__);
}
}
}
int ida = 0;
ggml_metal_kargs_gated_delta_net args = {
@@ -1918,23 +1952,25 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
/*.nb1 =*/ nb1,
/*.nb2 =*/ nb2,
/*.nb3 =*/ nb3,
/*.nb_out =*/ nb_out,
};
ggml_metal_encoder_set_pipeline(enc, pipeline);
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), ida++);
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), ida++); // args
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), ida++); // q
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), ida++); // k
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[2]), ida++); // v
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[3]), ida++); // gate
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[4]), ida++); // beta
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[5]), ida++); // state
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), ida++); // dst
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), ida++); // dst (attn)
ggml_metal_encoder_set_buffer (enc, bid_out, ida++); // state_out
const int nsg = pipeline.nsg;
ggml_metal_encoder_dispatch_threadgroups(enc, op->src[2]->ne[0]/nsg, op->src[2]->ne[1], op->src[2]->ne[3], 32, nsg, 1);
return 1;
return n_fuse;
}
int ggml_metal_op_solve_tri(ggml_metal_op_t ctx, int idx) {
@@ -3718,56 +3754,20 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
return 1;
}
// Snake activation autofuse: mul -> sin -> sqr -> mul -> add
static bool ggml_metal_op_can_fuse_snake(ggml_metal_op_t ctx, int idx) {
static constexpr ggml_op snake_ops[5] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD };
if (ctx->node(idx)->op != GGML_OP_MUL || !ctx->can_fuse(idx, snake_ops, 5)) {
return false;
}
const ggml_tensor * mul0 = ctx->node(idx + 0);
const ggml_tensor * sin_node = ctx->node(idx + 1);
const ggml_tensor * sqr = ctx->node(idx + 2);
const ggml_tensor * mul1 = ctx->node(idx + 3);
const ggml_tensor * add = ctx->node(idx + 4);
// x carries the full activation shape, a is the broadcast operand
const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1];
const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0];
// mul1 reads sqr and inv_b in either operand order
const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0];
// closure check: the trailing add reads the same x as the leading mul
const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0];
// x is in the supported whitelist and every chain intermediate shares x's type.
// a and inv_b bind as device const float * in the kernel, so they stay F32.
const bool types_ok =
(x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) &&
(a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) &&
(mul0->type == x->type) && (sin_node->type == x->type) &&
(sqr->type == x->type) && (mul1->type == x->type) &&
(add->type == x->type);
// a / inv_b collapse to [1, C, 1, 1], x and add stay 2D
const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];
const bool dim_ok =
(x->ne[2] == 1) && (x->ne[3] == 1) &&
(add->ne[2] == 1) && (add->ne[3] == 1) &&
(a->ne[2] == 1) && (a->ne[3] == 1) &&
(inv_b->ne[2] == 1) && (inv_b->ne[3] == 1);
// kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous
const bool contig_ok =
ggml_is_contiguous(x) && ggml_is_contiguous(add) &&
ggml_is_contiguous(a) && ggml_is_contiguous(inv_b);
return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x;
}
int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
if (ctx->use_fusion && ggml_metal_op_can_fuse_snake(ctx, idx)) {
return ggml_metal_op_snake_fused(ctx, idx);
int n_fuse = 1;
const ggml_metal_fusion * fusion = nullptr;
if (ctx->use_fusion()) {
int n = 1;
fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
n_fuse = n;
// snake activation autofuse: mul -> sin -> sqr -> mul -> add
if (fusion && fusion->id == GGML_METAL_FUSION_SNAKE) {
ctx->count_fusions(fusion);
return ggml_metal_op_snake_fused(ctx, idx);
}
}
ggml_tensor * op = ctx->node(idx);
@@ -3775,9 +3775,9 @@ int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
ggml_metal_library_t lib = ctx->lib;
ggml_metal_encoder_t enc = ctx->enc;
const bool use_fusion = ctx->use_fusion;
const bool use_fusion = ctx->use_fusion();
const int debug_fusion = ctx->debug_fusion;
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
@@ -3822,57 +3822,19 @@ int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
/*.o1 =*/ { bid_src1.offs },
};
ggml_op fops[8];
int n_fuse = 1;
// c[0] = add(a, b[0])
// c[1] = add(c[0], b[1])
// c[2] = add(c[1], b[2])
// ...
if (use_fusion) {
fops[0] = GGML_OP_ADD;
fops[1] = GGML_OP_ADD;
fops[2] = GGML_OP_ADD;
fops[3] = GGML_OP_ADD;
fops[4] = GGML_OP_ADD;
fops[5] = GGML_OP_ADD;
fops[6] = GGML_OP_ADD;
fops[7] = GGML_OP_ADD;
// note: in metal, we sometimes encode the graph in parallel so we have to avoid fusing ops
// across splits. idx_end indicates the last node in the current split
for (n_fuse = 0; n_fuse <= 6; ++n_fuse) {
if (!ctx->can_fuse(idx + n_fuse, fops + n_fuse, 2)) {
break;
}
ggml_tensor * f0 = ctx->node(idx + n_fuse);
ggml_tensor * f1 = ctx->node(idx + n_fuse + 1);
if (f0 != f1->src[0]) {
break;
}
// b[0] === b[1] === ...
if (!ggml_are_same_layout(f0->src[1], f1->src[1])) {
break;
}
// only fuse ops if src1 is in the same Metal buffer
ggml_metal_buffer_id bid_fuse = ggml_metal_get_buffer_id(f1->src[1]);
if (bid_fuse.metal != bid_src1.metal) {
break;
}
//ctx->fuse_cnt[ops[n_fuse + 1]->op]++;
args.o1[n_fuse + 1] = bid_fuse.offs;
if (use_fusion && fusion && fusion->id == GGML_METAL_FUSION_ADD_CHAIN) {
// the offsets of the fused addends are relative to the start of the src1 buffer
for (int i = 1; i < n_fuse; i++) {
args.o1[i] = ggml_metal_get_buffer_id(ctx->node(idx + i)->src[1]).offs;
}
++n_fuse;
ctx->count_fusions(fusion);
if (debug_fusion > 1 && n_fuse > 1) {
if (debug_fusion > 1) {
GGML_LOG_DEBUG("%s: fuse: ADD x %d\n", __func__, n_fuse);
}
}
@@ -4080,9 +4042,9 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
ggml_metal_library_t lib = ctx->lib;
ggml_metal_encoder_t enc = ctx->enc;
const bool use_fusion = ctx->use_fusion;
const bool use_fusion = ctx->use_fusion();
const int debug_fusion = ctx->debug_fusion;
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
@@ -4110,8 +4072,6 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
/*.nbf3 =*/ { nb03 },
};
ggml_op fops[8];
int n_fuse = 1;
ggml_metal_buffer_id bid_fuse[2] = { bid_src0, bid_src0 };
@@ -4120,55 +4080,35 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
// d[1] = mul(d[0], b)
// d[2] = add(d[1], c)
if (use_fusion) {
fops[0] = op->op;
fops[1] = GGML_OP_MUL;
fops[2] = GGML_OP_ADD;
int n = 1;
const ggml_metal_fusion * fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
for (n_fuse = 0; n_fuse <= 1; ++n_fuse) {
if (!ctx->can_fuse(idx + n_fuse, fops + n_fuse, 2)) {
break;
if (fusion && (fusion->id == GGML_METAL_FUSION_NORM_MUL || fusion->id == GGML_METAL_FUSION_NORM_MUL_ADD)) {
n_fuse = n;
ctx->count_fusions(fusion);
for (int i = 1; i < n_fuse; i++) {
const ggml_tensor * fn = ctx->node(idx + i);
bid_fuse[i - 1] = ggml_metal_get_buffer_id(fn->src[1]);
args.nef1[i] = fn->src[1]->ne[1];
args.nef2[i] = fn->src[1]->ne[2];
args.nef3[i] = fn->src[1]->ne[3];
args.nbf1[i] = fn->src[1]->nb[1];
args.nbf2[i] = fn->src[1]->nb[2];
args.nbf3[i] = fn->src[1]->nb[3];
}
ggml_tensor * f0 = ctx->node(idx + n_fuse);
ggml_tensor * f1 = ctx->node(idx + n_fuse + 1);
if (f0 != f1->src[0]) {
break;
}
if (f1->src[1]->ne[0] != op->ne[0]) {
break;
}
if (!ggml_is_contiguous_rows(f1->src[1])) {
break;
}
if (f1->type != GGML_TYPE_F32) {
break;
}
//ctx->fuse_cnt[f1->op]++;
bid_fuse[n_fuse] = ggml_metal_get_buffer_id(f1->src[1]);
args.nef1[n_fuse + 1] = f1->src[1]->ne[1];
args.nef2[n_fuse + 1] = f1->src[1]->ne[2];
args.nef3[n_fuse + 1] = f1->src[1]->ne[3];
args.nbf1[n_fuse + 1] = f1->src[1]->nb[1];
args.nbf2[n_fuse + 1] = f1->src[1]->nb[2];
args.nbf3[n_fuse + 1] = f1->src[1]->nb[3];
}
++n_fuse;
if (debug_fusion > 1 && n_fuse > 1) {
if (n_fuse == 2) {
GGML_LOG_DEBUG("%s: fuse: %s + MUL\n", __func__, ggml_op_name(op->op));
}
if (n_fuse == 3) {
GGML_LOG_DEBUG("%s: fuse: %s + MUL + ADD\n", __func__, ggml_op_name(op->op));
if (debug_fusion > 1) {
if (n_fuse == 2) {
GGML_LOG_DEBUG("%s: fuse: %s + MUL\n", __func__, ggml_op_name(op->op));
}
if (n_fuse == 3) {
GGML_LOG_DEBUG("%s: fuse: %s + MUL + ADD\n", __func__, ggml_op_name(op->op));
}
}
}
}
ggml/src/ggml-metal/ggml-metal-ops.h
+4 -3
View File
@@ -8,17 +8,18 @@ extern "C" {
typedef struct ggml_metal_op * ggml_metal_op_t;
struct ggml_metal_fusion; // forward decl (ggml-metal-device.h)
ggml_metal_op_t ggml_metal_op_init(
ggml_metal_device_t dev,
ggml_metal_cmd_buf_t cmd_buf,
struct ggml_cgraph * gf,
struct ggml_metal_fusion_info * finfo,
int idx_start,
int idx_end,
bool use_fusion,
bool use_concurrency,
bool use_capture,
int debug_graph,
int debug_fusion);
int debug_graph);
void ggml_metal_op_free(ggml_metal_op_t ctx);
ggml/src/ggml-metal/ggml-metal.cpp
+42
View File
@@ -4,6 +4,7 @@
#include "ggml-backend-impl.h"
#include "ggml-metal-device.h"
#include "ggml-metal-fusion.h"
#include "ggml-metal-context.h"
#include "ggml-metal-ops.h"
#include "ggml-metal-tuning.h"
@@ -906,6 +907,30 @@ static const char * ggml_backend_metal_tuning_device_token(ggml_backend_dev_t de
return ggml_metal_device_id_token(ggml_metal_device_get_props(ctx_dev)->device_id);
}
// generic fusion debugging API (ad-hoc proc-address mechanism): the test resolves the device
// fusion context once and passes that opaque handle to the rest of the functions
typedef void * ggml_backend_fusion_t;
static ggml_backend_fusion_t ggml_backend_metal_fusion_get(ggml_backend_dev_t dev) {
return ggml_metal_device_get_fusion_info((ggml_metal_device_t)dev->context);
}
static void ggml_backend_metal_fusion_stats_init(ggml_backend_fusion_t finfo) {
ggml_metal_fusion_info_stats_init((struct ggml_metal_fusion_info *) finfo);
}
static void ggml_backend_metal_fusion_stats_reset(ggml_backend_fusion_t finfo) {
ggml_metal_fusion_info_stats_reset((struct ggml_metal_fusion_info *) finfo);
}
static int ggml_backend_metal_fusion_stats_get(ggml_backend_fusion_t finfo, const char ** labels, uint64_t * counts, int n) {
return ggml_metal_fusion_info_stats_get((struct ggml_metal_fusion_info *) finfo, labels, counts, n);
}
static void ggml_backend_metal_fusion_set_enabled(ggml_backend_fusion_t finfo, bool enabled) {
ggml_metal_fusion_info_set_enabled((struct ggml_metal_fusion_info *) finfo, enabled);
}
static void * ggml_backend_metal_get_proc_address(ggml_backend_reg_t reg, const char * name) {
if (strcmp(name, "ggml_backend_get_features") == 0) {
return (void *)ggml_backend_metal_get_features;
@@ -928,6 +953,23 @@ static void * ggml_backend_metal_get_proc_address(ggml_backend_reg_t reg, const
if (strcmp(name, "ggml_backend_metal_tuning_device_token") == 0) {
return (void *)ggml_backend_metal_tuning_device_token;
}
// generic fusion debugging API (ad-hoc proc-address mechanism, not part of the official
// ggml backend interface yet; a backend that adopts it exports these exact names)
if (strcmp(name, "ggml_backend_fusion_get") == 0) {
return (void *)ggml_backend_metal_fusion_get;
}
if (strcmp(name, "ggml_backend_fusion_stats_init") == 0) {
return (void *)ggml_backend_metal_fusion_stats_init;
}
if (strcmp(name, "ggml_backend_fusion_stats_reset") == 0) {
return (void *)ggml_backend_metal_fusion_stats_reset;
}
if (strcmp(name, "ggml_backend_fusion_stats_get") == 0) {
return (void *)ggml_backend_metal_fusion_stats_get;
}
if (strcmp(name, "ggml_backend_fusion_set_enabled") == 0) {
return (void *)ggml_backend_metal_fusion_set_enabled;
}
return NULL;
ggml/src/ggml-metal/kernels/gated_delta_net.metal
+17 -3
View File
@@ -15,6 +15,7 @@ kernel void kernel_gated_delta_net_impl(
device const char * b,
device const char * s,
device char * dst,
device char * dst_fuse,
uint3 tgpig[[threadgroup_position_in_grid]],
uint3 tpitg[[thread_position_in_threadgroup]],
uint3 ntg[[threads_per_threadgroup]]) {
@@ -65,6 +66,12 @@ kernel void kernel_gated_delta_net_impl(
// per-(seq,head) offset within a slot
const uint state_out_base = (i23*args.ne21 + i21)*S_v*S_v + i20*S_v;
// when fused with the cache cpy, write the snapshots straight into the cache buffer using
// the slot stride; otherwise append them after the attn scores (nb_out == 0)
const bool fused = args.nb_out > 0;
const device float * state_out = fused ? (device float *)dst_fuse : (device float *)dst + attn_size;
const uint slot_stride = fused ? (uint)args.nb_out : state_size_per_snap;
for (short t = 0; t < args.ne22; t++) {
float s_k = 0.0f;
@@ -116,7 +123,7 @@ kernel void kernel_gated_delta_net_impl(
if (K > 1) {
const int target_slot = (int)args.ne22 - 1 - (int)t;
if (target_slot >= 0 && target_slot < (int)K) {
device float * dst_state = (device float *) (dst) + attn_size + (uint)target_slot * state_size_per_snap + state_out_base;
device float * dst_state = (device float *)state_out + (uint)target_slot * slot_stride + state_out_base;
FOR_UNROLL (short j = 0; j < NSG; j++) {
const short is = tx*NSG + j;
dst_state[is] = ls[j];
@@ -126,7 +133,7 @@ kernel void kernel_gated_delta_net_impl(
}
if (K == 1) {
device float * dst_state = (device float *) (dst) + attn_size + state_out_base;
device float * dst_state = (device float *)state_out + state_out_base;
FOR_UNROLL (short j = 0; j < NSG; j++) {
const short is = tx*NSG + j;
dst_state[is] = ls[j];
@@ -158,6 +165,7 @@ kernel void kernel_gated_delta_net_impl(
device const char * b,
device const char * s,
device char * dst,
device char * dst_fuse,
uint3 tgpig[[threadgroup_position_in_grid]],
uint3 tpitg[[thread_position_in_threadgroup]],
uint3 ntg[[threads_per_threadgroup]]) {
@@ -230,7 +238,13 @@ kernel void kernel_gated_delta_net_impl(
dst_attn += args.ne21*S_v;
}
device float * dst_state = (device float *) (dst) + args.ne23*args.ne22*args.ne21*S_v + (i23*args.ne21 + i21)*S_v*S_v + i20;
// when fused with the cache cpy, write the snapshots straight into the cache buffer using
// the slot stride; otherwise append them after the attn scores (nb_out == 0)
const bool fused = args.nb_out > 0;
const device float * state_out = fused ? (device float *)dst_fuse : (device float *)dst + args.ne23*args.ne22*args.ne21*S_v;
const uint slot_stride = fused ? (uint)args.nb_out : S_v*S_v;
device float * dst_state = (device float *)state_out + (i23*args.ne21 + i21)*slot_stride + i20;
device T * dstt_state = (device T *) (dst_state);
FOR_UNROLL (short j = 0; j < NSG; j++) {