Skip to content

sched: pipeline the delivery of a host-resident KV cache - #39

Open
Piggidragon wants to merge 14 commits into
GenerelSchwerz:llama/devfrom
Piggidragon:kv/pipelined-transport
Open

sched: pipeline the delivery of a host-resident KV cache#39
Piggidragon wants to merge 14 commits into
GenerelSchwerz:llama/devfrom
Piggidragon:kv/pipelined-transport

Conversation

@Piggidragon

@Piggidragon Piggidragon commented Aug 26, 2026

Copy link
Copy Markdown

Pipelines the host-to-device delivery of a host-resident KV cache: the transfer is issued one split ahead, on a transfer stream of its own, so a decode token stops paying transfer and attention in series.

Supersedes #38, which targeted beellama/dev; this is the same work rebased onto llama/dev.

What it does

Three pieces, each load-bearing:

  1. A stable prefix. The KV window is not stable for a whole graph - a CPU split writes this ubatch's rows into it between one layer's attention and the next. Everything below the lowest written row is, and at decode depth that is ~99.5% of the bytes. ggml_tensor::stable_prefix records it on the tensor that owns the storage; llama_kv_cache::update_stable_prefixes() sets it from apply_ubatch(), before the graph is built and allocated, so the plan and the deliveries are decided against the same write position even when the graph is reused. build_graph_shift() clears it.
  2. A staging ring outside ggml-alloc's reach. ggml-alloc may recycle a graph-owned input copy after its last graph-level consumer while a look-ahead transfer is still in flight - that is what made an earlier cross-layer prefetch experiment non-exact. The scheduler allocates its own ring and points the staged copies at it before allocation; a ready/release event pair per slot carries the handover each way. Every eligible accelerator gets its own ring, cursor and budget, so a layer-split model pipelines on each device and a device with no room falls back alone.
  3. One range per stream. A window over several streams is one view of a tensor whose streams sit end to end. Treating it as a single flat byte range let the lowest-writing stream cap the stable prefix for every stream above it, and copied the cells between one stream's window and the next that the graph never reads. The prefix is counted within a stream, and such an input is delivered one range per stream. A window over one stream keeps the single flat range it had. See Parallel sequences.
  4. A look-ahead clear of the ring's tail. A delivery L splits ahead recycles the slot of the split L - n_slots back, so n_slots == L + 1 recycles the split just enqueued and still running. The ring keeps 2 slots of margin, deliveries are issued after a split is enqueued, and slot recycling is ordered stream-to-stream, not through the host. Each of those three alone costs the entire gain while still producing correct output - the first working version measured +0.5% and looked like "the copy just doesn't overlap".

--kv-pipeline-depth N (default 1, 0 = ordered path exactly). Only engages where a host-resident cache produces the deliveries; a device-resident run never creates the transport. Deeper look-ahead is worse at every depth measured: at a 19,246-token prompt, 29.83 t/s at N = 1, 28.44 at N = 2, 25.93 at N = 4.

CUDA only. The scheduler enables the ring for a device whose backend registry is named CUDA, that has asynchronous transfers and events, and whose default buffer type the scheduler was configured with. Meta devices are excluded outright, so -sm tensor keeps the ordered path. Everything else - SYCL, WebGPU, Vulkan, the CPU - is untouched and stays ordered.

Measurements

Re-measured on this head, on an RTX 4070 (sm_89, 11,902 MiB usable), driver 610.57.04 / CUDA 13.3, i5-13400F, Qwen3.8-27B-UD-IQ2_M.gguf, -ngl 99 -sm none -mg 0 -t 3 -fa on -ctk q8_0 -ctv q8_0 -b 512 -ub 512, host residency -nkvo --kv-cpu-pinned --recurrent-state-offload, --kv-pipeline-budget 512, everything under taskset -c 0,2,4.

llama-bench, A/B/A/B with reversed arm order (docs/repro/r4-kv-pipeline-ab.sh), both passes shown, plus peak device memory from the context sweep:

context ordered pipelined gain peak device memory ring
4,096 29.980, 29.978 34.802, 34.805 +16.1% +28 MiB 27 MiB
16,384 19.012, 19.009 29.778, 29.788 +56.6% +104 MiB 107 MiB
32,768 12.724, 12.724 15.287, 15.286 +20.1% +206 MiB 213 MiB
65,536 7.635, 7.635 8.659, 8.654 +13.4% +410 MiB 428 MiB

llama-server, greedy, the four 18,432-prefill tasks of the exactness gate at -c 32768:

task prompt ordered pipelined gain
prose 14,821 19.775 30.012 +51.8%
dialogue 15,984 19.155 29.767 +55.4%
records 29,603 13.553 16.396 +21.0%
code 29,670 13.504 16.362 +21.2%

Depth 4 is slower than depth 1 on all four (27.037, 26.584, 15.934, 15.857), which is why N = 1 is the default.

The rows below were taken on earlier heads and are not re-measured here. They are kept because they are the negative and diagnostic results the design rests on, not the headline numbers.

Where the token goes, and why the gain narrows

GGML_SCHED_TRANSPORT_DEBUG=2 is implemented in this PR (the counters existed but nothing accumulated or printed them). Per decode graph behind a 19,246-token prompt:

ordered pipelined
total split-loop time 54.11 ms 31.10 ms
blocked in ordered ggml_backend_tensor_copy 28.30 ms 3.57 ms
blocked waiting for the consumer 25.70 ms 27.31 ms
delivered early / late 0 / 0 MiB 644.0 / 2.3 MiB

644 MiB in the 28.30 ms the ordered arm spends on the same bytes is 22.0 GB/s, and nvidia-smi reports the card at gen4 x16 - about 88% of what the link does in practice. The pipeline came within 5% of the max(copy, compute) ceiling at both 19k and 48k. The narrowing gain is therefore arithmetic: copy grows with context, compute does not, and once copy dominates there is only compute left to hide behind it. =3 names the tensors still on the ordered path.

Two attempts to recover the last 3.57 ms, both recorded in the doc as negative results: issuing the delivery in pieces so a blocking copy can interleave does nothing (the copy engine is FIFO across streams, and small pieces cost throughput), and putting the copy on the consumer's stream moves the time into the consumer wait without changing throughput. That 3.57 ms is almost entirely one 256 KiB graph input queued behind the deliveries - it is bandwidth, not latency.

The lever is bytes, not the link

-ctk q4_0 -ctv q4_0 halves the traffic, and what that is worth depends on which side of the crossover you are:

prompt KV ordered pipelined delivered
19,246 q8_0 17.785 29.833 644.0 MiB
19,246 q4_0 22.904 31.502 343.2 MiB
48,042 q8_0 9.790 11.313 1602.5 MiB
48,042 q4_0 14.146 17.717 850.6 MiB

+5.6% at 19,246 (the consumer wait is 27.31 ms at q8_0 and 27.26 at q4_0 - the pipeline had already reached the compute floor, so the removed bytes were bytes nothing waited for) and +56.6% at 48,042, where copy still dominates.

The ring also wins per MiB against --kv-gpu-layers. At 19,246 with -c 32768, a device-resident layer costs ~68 MiB and the ring ~205 MiB:

no --kv-gpu-layers 4 8
ordered 17.785 20.334
pipelined 29.843 30.263 30.640

Four device-resident layers are worth +14.3% on the ordered path and +1.4% on the pipelined one, for more memory than the ring costs.

Pinning is worth as much as the pipeline and is off by default. Behind a 13,128-token prompt: 21.582 -> 32.252 pinned, 14.945 -> 22.709 unpinned.

Parallel sequences

Measured with llama-batched-bench, 2,048 prompt tokens per sequence, -c 32768 -np 8, generation t/s:

-npl unified d0 unified d1 streams d0 streams d1 before streams d1 after
1 32.69 35.35 32.59 35.32 35.31
2 51.69 58.63 48.68 51.88 61.56
4 72.43 86.52 62.39 65.06 89.59
8 83.90 105.16 70.92 72.18 113.55

The 8-slot row is measured with each arm run on its own. Taken as the last step of a sweep that has already run 1, 2 and 4 slots in the same process, the unified ring is refused for headroom and that arm reads 83.88 instead - the guard reacting to what the process has already allocated rather than to the configuration.

Without the per-stream delivery a non-unified cache pipelines almost nothing: 6.6% of the delivery goes early at 8 slots, because the prefix stops at the lowest stream's head, and the delivery also carries the unread cells between the streams. With it, 8 slots go from 72.18 to 112.70 t/s against 70.61 ordered, and one slot measures 35.31 against 35.32. A non-unified cache is the default, so this is the configuration most -np > 1 runs were in.

Concurrent slots cannot be gated on output the way one sequence can. Their batching varies between runs, so the same build at the same depth produces different greedy output - three runs at N = 0 gave three different hashes. test_transport_multi_stream_ranges stands in for it: every stream's window covered once from its own source offset, the cells between them never moved, and the early and late bytes split as the prefix says.

Device memory

A host-resident cache exists to keep device memory free, so the staging is capped outright by --kv-pipeline-budget (default 128 MiB per device), not by a fraction of what happens to be free. A ring is (N + 2) slots of one attention layer's K+V over the whole context, so it grows linearly with context.

The cap is applied to what the current graph needs, and the warning reports the full-context figure alongside it so the budget can be sized against the number that matters. Enforcing the projection instead was tried and reverted: it refuses the ring for every large -c even when the window never gets near it, which at -c 32768 turned the feature off by default.

A budget or headroom decline is per graph, not latched. The plan is remade for every graph, so a context that outgrows the budget gives the ring back and a later, smaller live window takes it again. The cost is that transient, and it is bounded by the budget the user already authorised. Declining is otherwise free: the ring and the transfer backend's device context are both released, and the transfer backend is created lazily in the first place.

The ring never starves the graph. The rings are laid out and allocated before the graph is, and the configuration is locked by then, so a device that can hold the graph alone but not the graph next to a ring used to fail allocation outright with no way for the caller to retry. If graph reservation fails, the rings are now released and the reservation is retried once on the ordered path, and that scheduler keeps the ordered path from then on. test_transport_releases_ring_for_graph sizes a dummy device to hold the graph or the ring but not both.

The headroom guard was measured and left alone. The ring declines unless it can leave GGML_SCHED_TRANSPORT_HEADROOM (512 MiB) free. The compute buffer on this configuration is 617 MiB and this fork resizes compute buffers at run time, so the guard is about the size of the thing it exists to leave room for. Two candidate changes were checked and neither survives: sizing it against the compute buffer instead of a constant is stricter here (617 against 512), and lowering it is not supported by the cases that looked like evidence for it - run on their own, both the 8-slot unified ring and -d 131072 fit and pipeline at the shipped 512 MiB. The guard refuses only under accumulated pressure inside one process, which is what it is for.

Tried and reverted: planning the ring during reserve so ggml-alloc would not budget blocks for the staged copies. It reclaims nothing, and test_transport_fallback_keeps_allocator_plan pins the invariant: the compute buffer is the same size on both arms.

Validation

  • Byte-identical greedy output at N = 0, N = 1 and N = 4 - all eight tasks (prose, source code, JSON records, dialogue, at 2k and 18k), re-run on this head with a CUDA build: docs/repro/r4-kv-pipeline-exact.sh 0 1 4 exits 0 with every hash matching depth 0.
  • Two independent N = 0 passes agree on all eight, which was not true before this PR. records@18432 used to give different hashes across otherwise identical runs and looked like the pipeline breaking exactness. It is not the task: asked on its own with the prompt cache off it returns the same hash three times running, at -c 32768 and -c 65536. It was the harness - all eight tasks share one server with prompt caching on, and records@18432 is ~29.6k tokens with a task of about the same size ahead of it, so the two do not both fit in a 32,768 cache and placement depended on what was still resident. The harness now sets cache_prompt: false, and gives every task a nonce so no two share a restorable prefix.
  • test-alloc covers entry allocation and bounds against a buffer type whose get_alloc_size exceeds ggml_nbytes, per-byte delivery coverage from the matching source offset, event ordering (nothing waits on an event before it is recorded), the empty graph, the budget decline and recovery, per-backend partial failure, the meta and non-CUDA exclusions, and the ring-starvation fallback. Clean under UBSan.
  • test_transport_multi_stream_ranges covers the per-stream delivery, which the byte-identical gate structurally cannot: that gate is single-sequence, and concurrent slots are not reproducible between runs.
  • Device-resident KV unaffected - the transport is never created, because llama_context passes a depth of 0.

Tensor parallelism

-sm tensor is not pipelined, and nothing in this PR moves it closer. The scheduler excludes meta devices explicitly and requires the CUDA registry name, so a tensor-parallel run keeps the ordered path. Enabling it needs a validated strided head-split write for a host-resident cache, which is not here. An earlier revision of this branch carried three meta-backend prerequisites; they are not in this revision.

Independently of that, -sm tensor together with --no-kv-offload is currently incorrect, which is a pre-existing fork defect and not caused by this PR. Fixed separately in #48; reported upstream as ggml-org#27757.

Same build, same prompt, greedy:

config output hash verdict
-sm layer + -nkvo 3e127464e8b901a7 reference
-sm tensor + device KV 3e127464e8b901a7 correct
-sm tensor + -nkvo 9b82be0158a2fa4d wrong, silently

Cause. TP splits attention by head, but a host-resident cache is one undivided tensor, so the scheduler's copy is classified MIRRORED and the whole window goes to both devices. With 24 query heads split 12/12 and 4 KV heads mirrored, the kernel derives the GQA ratio from the tensors it is handed - 12/4 = 3 instead of 6 - and the second device's queries, renumbered from 0, read the first device's keys. With an uneven split the same fault surfaces as a crash: fattn.cu:371 GGML_ASSERT(Q->ne[2] % K->ne[2] == 0), since 24 split 13/11 is not divisible by 4.

Still to do


Re-run on this head: r4-kv-pipeline-exact.sh (exit 0), r4-kv-pipeline-ab.sh (exit 0), r4-kv-pipeline-context-sweep.sh (exit 0), test-alloc in Release, UBSan and against the CUDA build, git diff --check.

🤖 Generated with Claude Code

https://claude.ai/code/session_01DLsim6XPFPodyRQ9cRVdG1

@GenerelSchwerz

Copy link
Copy Markdown
Owner

Review of head 175236b: changes requested before merge.

The CUDA performance result looks credible, and the declining relative gain is consistent with the overlap ceiling: ordered time is approximately copy + compute, while pipelined time approaches max(copy, compute) plus residual overhead. As context grows, KV traffic grows while the amount of compute available to hide it does not grow at the same rate. NCU is not needed to explain that curve. A short nsys trace is useful after the correctness fixes, but it is not a substitute for them.

Blocking defects

1. The incomplete meta/tensor-parallel path is now reachable

The generic eligibility check selects any non-CPU backend with async tensor set and event interfaces:

for (int i = 0; i < sched->n_backends; i++) {
ggml_backend_t backend = sched->backends[i];
if (backend->iface.set_tensor_async == NULL ||
backend->iface.event_record == NULL ||
backend->iface.event_wait == NULL) {
continue;
}
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
if (dev == NULL || dev->iface.event_new == NULL) {
// Worth naming rather than folding into the generic message below: a backend that
// fans out over several devices -- the meta backend used for tensor parallelism --
// lands here because that layer implements no events, and ordering a transfer stream
// against the consumer is what this is built on. It keeps the ordered path.
if (tr->debug > 0) {
GGML_LOG_INFO("%s: %s cannot order streams with events, staying on the ordered path\n",
__func__, ggml_backend_name(backend));
}
continue;
}
if (ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU) {
continue;
}
// the ring is written through the transfer backend, which only accepts the device's own
// default buffer type; a scheduler configured with anything else keeps the ordered path
if (sched->bufts[i] != ggml_backend_dev_buffer_type(dev)) {
continue;
}
tr->rings[i].eligible = true;
tr->rings[i].alignment = std::max<size_t>(ggml_backend_buft_get_alignment(sched->bufts[i]), 128);

This PR adds those interfaces to the meta backend:

// A meta event is one simple event per device. Recording it records on every simple backend, and
// waiting on it makes every simple backend wait on its own device's event, so a stream is never
// ordered against another device's work by accident.
struct ggml_backend_meta_event_context {
std::vector<ggml_backend_event_t> simple_events;
};
static ggml_backend_event_t ggml_backend_meta_device_event_new(ggml_backend_dev_t dev);
static void ggml_backend_meta_device_event_free(ggml_backend_dev_t dev, ggml_backend_event_t event);
static void ggml_backend_meta_device_event_synchronize(ggml_backend_dev_t dev, ggml_backend_event_t event);
static const ggml_backend_device_i ggml_backend_meta_device_iface = {
/* .get_name = */ ggml_backend_meta_device_get_name,
/* .get_description = */ ggml_backend_meta_device_get_description,
/* .get_memory = */ ggml_backend_meta_device_get_memory,
/* .get_type = */ ggml_backend_meta_device_get_type,
/* .get_props = */ ggml_backend_meta_device_get_props,
/* .init_backend = */ ggml_backend_meta_device_init_backend,
/* .get_buffer_type = */ ggml_backend_meta_device_get_buffer_type,
/* .get_host_buffer_type = */ ggml_backend_meta_device_get_host_buffer_type,
/* .buffer_from_host_ptr = */ nullptr,
/* .supports_op = */ ggml_backend_meta_device_supports_op,
/* .supports_buft = */ ggml_backend_meta_device_supports_buft,
/* .offload_op = */ nullptr,
/* .event_new = */ ggml_backend_meta_device_event_new,
/* .event_free = */ ggml_backend_meta_device_event_free,
/* .event_synchronize = */ ggml_backend_meta_device_event_synchronize,
};
static bool ggml_backend_dev_is_meta(ggml_backend_dev_t dev) {
return dev != nullptr && dev->iface.get_name == ggml_backend_meta_device_iface.get_name;
}
static size_t ggml_backend_meta_dev_n_devs(ggml_backend_dev_t meta_dev) {
GGML_ASSERT(ggml_backend_dev_is_meta(meta_dev));
const ggml_backend_meta_device_context * meta_dev_ctx = (const ggml_backend_meta_device_context *) meta_dev->context;
return meta_dev_ctx->simple_devs.size();
}
static ggml_backend_dev_t ggml_backend_meta_dev_simple_dev(ggml_backend_dev_t meta_dev, size_t index) {
GGML_ASSERT(ggml_backend_dev_is_meta(meta_dev));
const ggml_backend_meta_device_context * meta_dev_ctx = (const ggml_backend_meta_device_context *) meta_dev->context;
GGML_ASSERT(index < meta_dev_ctx->simple_devs.size());
return meta_dev_ctx->simple_devs[index];
}
static ggml_backend_event_t ggml_backend_meta_device_event_new(ggml_backend_dev_t dev) {
const size_t n_devs = ggml_backend_meta_dev_n_devs(dev);
ggml_backend_meta_event_context * event_ctx = new ggml_backend_meta_event_context;
event_ctx->simple_events.reserve(n_devs);
for (size_t i = 0; i < n_devs; i++) {
// null means the simple device has no events, so the meta device has none either
ggml_backend_event_t simple_event = ggml_backend_event_new(ggml_backend_meta_dev_simple_dev(dev, i));
if (simple_event == nullptr) {
for (ggml_backend_event_t e : event_ctx->simple_events) {
ggml_backend_event_free(e);
}
delete event_ctx;
return nullptr;
}
event_ctx->simple_events.push_back(simple_event);
}
ggml_backend_event_t event = new ggml_backend_event;
event->device = dev;
event->context = event_ctx;
return event;
}
static void ggml_backend_meta_device_event_free(ggml_backend_dev_t dev, ggml_backend_event_t event) {
ggml_backend_meta_event_context * event_ctx = (ggml_backend_meta_event_context *) event->context;
for (ggml_backend_event_t simple_event : event_ctx->simple_events) {
ggml_backend_event_free(simple_event);
}
delete event_ctx;
delete event;
GGML_UNUSED(dev);
}
static void ggml_backend_meta_device_event_synchronize(ggml_backend_dev_t dev, ggml_backend_event_t event) {
ggml_backend_meta_event_context * event_ctx = (ggml_backend_meta_event_context *) event->context;
for (ggml_backend_event_t simple_event : event_ctx->simple_events) {
ggml_backend_event_synchronize(simple_event);
}
GGML_UNUSED(dev);

static const ggml_backend_i ggml_backend_meta_i = {
/* .get_name = */ ggml_backend_meta_get_name,
/* .free = */ ggml_backend_meta_free,
/* .set_tensor_async = */ ggml_backend_meta_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_meta_get_tensor_async,
/* .set_tensor_2d_async = */ nullptr,
/* .get_tensor_2d_async = */ nullptr,
/* .cpy_tensor_async = */ nullptr,
/* .synchronize = */ ggml_backend_meta_synchronize,
/* .graph_plan_create = */ nullptr,
/* .graph_plan_free = */ nullptr,
/* .graph_plan_update = */ nullptr,
/* .graph_plan_compute = */ nullptr,
/* .graph_compute = */ ggml_backend_meta_graph_compute,
/* .event_record = */ ggml_backend_meta_event_record,
/* .event_wait = */ ggml_backend_meta_event_wait,
/* .graph_optimize = */ nullptr,

A normal partial-prefix delivery sends the changing tail with a nonzero destination offset:

if (staged && ggml_backend_sched_input_is_staged(sched, split_id, input_id)) {
// whatever prefix was stable went out on the transfer stream earlier; the rest is
// what an earlier split of this graph may still have written, and it is only safe
// to read now that every earlier split has run. It goes on the consumer's own
// stream, where it is already ordered ahead of the kernels and behind the reader
// of whatever occupied this slot before.
const size_t prefix = ggml_backend_sched_input_stable_prefix(input);
const size_t nbytes = ggml_nbytes(input);
if (nbytes > prefix) {
ggml_backend_tensor_set_async(split_backend, input_cpy,
(const char *) input->data + prefix, prefix, nbytes - prefix);
tr->n_bytes_late += nbytes - prefix;
}

Meta hard-asserts that the offset is zero and also requires whole-chunk granularity:

static void ggml_backend_meta_set_tensor_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
const size_t n_backends = ggml_backend_meta_n_backends(backend);
GGML_ASSERT(offset == 0);
GGML_ASSERT(ggml_is_contiguous(tensor));
const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ false);
GGML_ASSERT(split_state.n_segments == 1);
GGML_ASSERT(split_state.nr[0] == 1);
switch (split_state.axis) {
case GGML_BACKEND_SPLIT_AXIS_0:
case GGML_BACKEND_SPLIT_AXIS_1:
case GGML_BACKEND_SPLIT_AXIS_2: {
// Exploit that tensors are contiguous to splice it with simple tensors as "chunks".
const size_t chunk_size_full = tensor->nb[split_state.axis + 1];
GGML_ASSERT(offset % chunk_size_full == 0);
GGML_ASSERT(size % chunk_size_full == 0);
const int64_t i_start = offset /chunk_size_full;
const int64_t i_stop = (offset + size)/chunk_size_full;
size_t offset_j = 0;
for (size_t j = 0; j < n_backends; j++){
ggml_backend_t simple_backend = ggml_backend_meta_simple_backend(backend, j);
ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
const size_t chunk_size_j = simple_tensor->nb[split_state.axis + 1];
if (chunk_size_j == 0) {
continue;
}
ggml_backend_tensor_set_2d_async(simple_backend, simple_tensor, (const char *) data + offset_j, offset, chunk_size_j,
i_stop - i_start, chunk_size_j, chunk_size_full);
offset_j += chunk_size_j;
}
GGML_ASSERT(offset_j == chunk_size_full);
} break;
case GGML_BACKEND_SPLIT_AXIS_MIRRORED: {
for (size_t j = 0; j < n_backends; j++) {
ggml_backend_tensor_set_async(
ggml_backend_meta_simple_backend(backend, j), ggml_backend_meta_buffer_simple_tensor(tensor, j), data, offset, size);
}

The design document says the strided delivery is still missing and the meta work has not been run:

`ggml_backend_meta_buffer_init_tensor` places each per-device tensor at the
same offset from its own buffer's base, taken from the meta tensor's offset
from the placeholder base. What was missing is that `ggml_gallocr_init_tensor`
leaves a tensor that already has `data` alone, so a copy pointed into the ring
never reached `init_tensor` and its per-device tensors were never built. The
plan calls it now. On a plain device buffer this costs nothing: the ring
carries `USAGE_COMPUTE`, which is the case the CUDA `init_tensor` skips.
3. **A transfer-only meta backend.** `ggml_backend_dev_init` on a meta device
runs the whole meta context constructor, which calls `ggml_backend_comm_init`
across every device. The transfer backend only moves tensors and orders
streams, so it asks for `"transfer"` and the constructor skips the collective.
4. **Still missing: the delivery itself.** A host-resident cache reaches attention
permuted, as `[head_dim, n_kv, n_head_kv, 1]` with the heads interleaved inside
each cell, and each device owns one run of heads per cell. That is a strided
write, and `ggml_backend_meta_set_tensor_async` can only splice a contiguous
tensor into whole chunks. It needs the same strided case the synchronous
`set_tensor` needs, which belongs with the fix below.
All of it sits behind a correctness problem that is not this feature's:
**`-sm tensor` together with `--no-kv-offload` currently produces wrong output.**
On one build and one prompt, `-sm layer --no-kv-offload` and `-sm tensor` with a
device-resident cache agree exactly, while `-sm tensor --no-kv-offload` differs.
It does not crash or warn; it generates fluent, different text.
The cause is the GQA head mapping. Tensor parallelism splits attention by head,
but a host-resident cache is one undivided tensor, so the scheduler's copy of it
is classified `MIRRORED` and the whole window goes to every device. With 24 query
heads split 12/12 and 4 KV heads mirrored, the kernel derives the GQA ratio from
the tensors it is handed -- 12/4 = 3 rather than 6 -- and the second device's
queries, renumbered from 0, read the first device's keys. With an uneven split the
same fault surfaces as a crash instead:
`GGML_ASSERT(Q->ne[2] % K->ne[2] == 0)`, because 24 heads split 13/11 is not
divisible by 4.
Head-splitting the copy rather than mirroring it fixes it. That was prototyped
and reproduced the layer-split output byte for byte, and needs four coordinated
changes: classify the scheduler's copy at all (it is a leaf in a compute buffer,
so it never reaches the device's split-state callback), use the head axis for the
permuted `[head_dim, n_kv, n_head_kv, 1]` shape rather than the cache tensor's own
axis, express the granularity in heads aligned to the query split divided by the
GQA ratio, and add a strided write because the heads are interleaved within each
row rather than laid out end to end.
## Future work
- Fix `-sm tensor` with `--no-kv-offload` (above). Until then it should not be
used: it is wrong rather than slow.
- The strided head-split delivery above, and then a measurement. The three meta
backend pieces compile and ship here, and none of them has been run.

This means the documentation claim that tensor parallelism remains ordered does not match the code. A host-KV tensor-parallel run can enter the pipeline and assert.

Fix: explicitly exclude GGML_BACKEND_DEVICE_TYPE_META for now. Prefer moving the untested meta event and transfer groundwork into the later change that implements and validates the strided head-split write. Removing only the offset assertion is not sufficient.

2. Partial per-backend setup failure leaves stale plan membership

Transfer backend or event creation can fail here:

static bool ggml_backend_sched_transport_ensure_backend(ggml_backend_sched_t sched, int backend_id) {
struct ggml_backend_sched_transport * tr = &sched->transport;
struct ggml_backend_sched_transport_ring * r = &tr->rings[backend_id];
if (r->transfer != NULL) {
return true;
}
ggml_backend_dev_t dev = ggml_backend_get_device(sched->backends[backend_id]);
if (dev == NULL) {
return false;
}
// a meta device stands up a collective communicator across every device it wraps when it is
// asked for a backend; this one only transfers and orders streams, so it does not need one
const char * params = ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_META ? "transfer" : NULL;
ggml_backend_t transfer = ggml_backend_dev_init(dev, params);
if (transfer == NULL) {
return false;
}
bool ok = true;
for (int slot = 0; slot < tr->n_slots && ok; slot++) {
r->slots[slot].ready = ggml_backend_event_new(dev);
r->slots[slot].release = ggml_backend_event_new(dev);
ok = r->slots[slot].ready != NULL && r->slots[slot].release != NULL;
}
if (!ok) {
for (int slot = 0; slot < tr->n_slots; slot++) {
ggml_backend_event_free(r->slots[slot].ready);
ggml_backend_event_free(r->slots[slot].release);
r->slots[slot].ready = NULL;
r->slots[slot].release = NULL;
}
ggml_backend_free(transfer);
return false;
}
r->transfer = transfer;
return true;

The caller only sets r->n_staged to zero:

// nothing has been allocated for this ring until here
if (!ggml_backend_sched_transport_ensure_backend(sched, bid)) {
r->n_staged = 0;
continue;
}

It does not clear split_order and input_staged for that backend, unlike the budget, headroom, and allocation failure paths. If another backend succeeds, global staging remains enabled and address assignment reaches the failed backend:

for (int i = 0; i < sched->n_splits; i++) {
if (tr->split_order[i] < 0) {
continue;
}
struct ggml_backend_sched_split * split = &sched->splits[i];
struct ggml_backend_sched_transport_ring * r = &tr->rings[split->backend_id];
char * const ring = (char *) ggml_backend_buffer_get_base(r->buffer);
char * slot = ring + (size_t)(tr->split_order[i] % tr->n_slots) * r->slot_size;
size_t offset = 0;
for (int j = 0; j < split->n_inputs; j++) {
if (!ggml_backend_sched_input_is_staged(sched, i, j)) {
continue;
}
struct ggml_tensor * input_cpy = tensor_copy(split->inputs[j], split->backend_id, sched->cur_copy);
input_cpy->data = slot + offset;
input_cpy->buffer = r->buffer;
// ggml-alloc leaves a tensor that already has data alone, so it never initialises the
// ones pointed into the ring. A buffer that places a tensor as more than an offset --
// a meta buffer builds one sub-tensor per device -- needs this to have happened.
ggml_backend_buffer_init_tensor(r->buffer, input_cpy);
offset += GGML_PAD(ggml_nbytes(split->inputs[j]), r->alignment);
}
GGML_ASSERT(offset <= r->slot_size);

That path calls ggml_backend_buffer_get_base with a null ring buffer.

Fix: factor one backend-decline helper that releases the ring and clears n_staged, every affected split_order entry, and every corresponding input_staged entry. Use it for every failure path. Add failure injection for one backend failing while another succeeds.

3. Crossing the budget permanently disables the feature

Ring eligibility rejects a backend once over_budget is set:

static bool ggml_backend_sched_transport_ring_enabled(ggml_backend_sched_t sched, int backend_id) {
const struct ggml_backend_sched_transport * tr = &sched->transport;
if (tr->depth < 1 || tr->n_slots < 2 || backend_id < 0) {
return false;
}
const struct ggml_backend_sched_transport_ring * r = &tr->rings[backend_id];
return r->eligible && !r->over_budget;
}
static bool ggml_backend_sched_transport_enabled(ggml_backend_sched_t sched) {
for (int i = 0; i < sched->n_backends; i++) {
if (ggml_backend_sched_transport_ring_enabled(sched, i)) {
return true;
}
}
return false;

The budget and headroom paths set that flag permanently:

if (tr->budget > 0 && (ring_size > tr->budget || r->over_budget)) {
r->over_budget = true;
if (!r->reported_no_room) {
GGML_LOG_WARN("%s: transport ring on %s needs %zu MiB now and %zu MiB at the full "
"context, against a %zu MiB budget, staying on the ordered path (raise "
"--kv-pipeline-budget to spend more device memory on it)\n", __func__,
ggml_backend_name(sched->backends[bid]), ring_size >> 20,
ring_size_max >> 20, tr->budget >> 20);
r->reported_no_room = true;
}
ggml_backend_sched_transport_release_ring(sched, bid, true);
// un-stage this ring's inputs: they have no ring to live in
for (int i = 0; i < sched->n_splits; i++) {
if (sched->splits[i].backend_id != bid) {
continue;
}
tr->split_order[i] = -1;
for (int j = 0; j < sched->splits[i].n_inputs; j++) {
tr->input_staged[tr->split_input_ofs[i] + j] = 0;
}
}
continue;

if (dev_free > 0 && ring_size + GGML_SCHED_TRANSPORT_HEADROOM > dev_free) {
if (!r->reported_no_room) {
GGML_LOG_WARN("%s: transport ring on %s would need %zu MiB and leave less than "
"%u MiB of the %zu MiB free, staying on the ordered path\n", __func__,
ggml_backend_name(sched->backends[bid]), ring_size >> 20,
GGML_SCHED_TRANSPORT_HEADROOM >> 20, dev_free >> 20);
r->reported_no_room = true;
}
r->over_budget = true;
ggml_backend_sched_transport_release_ring(sched, bid, true);
for (int i = 0; i < sched->n_splits; i++) {
if (sched->splits[i].backend_id != bid) {
continue;
}
tr->split_order[i] = -1;
for (int j = 0; j < sched->splits[i].n_inputs; j++) {
tr->input_staged[tr->split_input_ofs[i] + j] = 0;
}
}
continue;

The assumption that a context only grows is not valid for a long-running server. The KV view size is derived from the highest currently used cell and can shrink after a sequence is removed:

uint32_t llama_kv_cache::get_n_kv(const slot_info & sinfo) const {
uint32_t result = 0;
// pad the n_kv value so that the graph remains constant across batches and can be reused
// note: this also helps some backends with performance (f.ex https://github.com/ggml-org/llama.cpp/pull/16812#issuecomment-3455112220)
const uint32_t n_pad_cur = std::max(n_pad, 256u);
for (uint32_t s = 0; s < sinfo.n_stream(); ++s) {
const auto & cells = v_cells[sinfo.strm[s]];
result = std::max(std::min(cells.size(), std::max(n_pad_cur, GGML_PAD(cells.used_max_p1(), n_pad_cur))), result);
}
return result;

Once over_budget is set, the initial eligibility check prevents later small graphs from even reaching the calculation that could reconsider the ring. One long request can therefore force all later short requests on the same context onto the ordered path.

Fix: re-evaluate the budget for each new graph. Keep warning suppression separate from eligibility. If allocation churn is a concern, use hysteresis rather than a permanent latch.

4. The public setters can invalidate an allocated graph

Depth is documented as pre-allocation, but the restriction is not enforced. Budget has no timing restriction:

// `depth` is how many splits ahead deliveries run; 0 disables pipelining. The ring holds a
// couple of slots more than that, so that recycling a slot never has to wait for a reader
// that is still running. Requires a destination backend with asynchronous transfers and
// events; where that is missing the setting is ignored. Costs roughly (depth + 2) *
// (largest staged split) of device memory. Must be called before the first graph is
// allocated.
GGML_API void ggml_backend_sched_set_transport_pipeline_depth(ggml_backend_sched_t sched, int depth);
// Hard cap on the staging ring, in bytes. A host-resident cache exists to keep device memory
// free, so the ring is capped outright and not merely against what happens to be free: past
// the cap the scheduler declines and keeps the ordered path. 0 removes the cap. Default 128 MiB.
GGML_API void ggml_backend_sched_set_transport_pipeline_budget(ggml_backend_sched_t sched, size_t bytes);
// Number of staged deliveries and staged bytes issued since the scheduler was created.
GGML_API void ggml_backend_sched_get_transport_pipeline_stats(ggml_backend_sched_t sched, int64_t * n_deliveries, int64_t * n_bytes_early, int64_t * n_bytes_late);

Both setters can free ring storage:

void ggml_backend_sched_set_transport_pipeline_depth(ggml_backend_sched_t sched, int depth) {
GGML_ASSERT(sched);
const char * env = getenv("GGML_KV_PIPELINE_DEPTH");
if (env != NULL) {
depth = atoi(env);
}
depth = std::min(depth, GGML_SCHED_MAX_TRANSPORT_SLOTS - GGML_SCHED_TRANSPORT_MARGIN);
if (depth < 0) {
depth = 0;
}
struct ggml_backend_sched_transport * tr = &sched->transport;
ggml_backend_sched_transport_teardown(sched);
for (int i = 0; i < sched->n_backends; i++) {
tr->rings[i].eligible = false;
tr->rings[i].over_budget = false;
tr->rings[i].reported_no_room = false;
}
tr->depth = depth;
tr->n_slots = depth + GGML_SCHED_TRANSPORT_MARGIN;
if (depth < 1) {
return;

void ggml_backend_sched_set_transport_pipeline_budget(ggml_backend_sched_t sched, size_t bytes) {
GGML_ASSERT(sched);
const char * env = getenv("GGML_KV_PIPELINE_BUDGET_MIB");
if (env != NULL) {
bytes = (size_t) strtoull(env, NULL, 10) * 1024 * 1024;
}
if (sched->transport.budget != bytes) {
sched->transport.budget = bytes;
for (int i = 0; i < sched->n_backends; i++) {
sched->transport.rings[i].reported_no_room = false;
sched->transport.rings[i].over_budget = false;
// a ring in hand may no longer be allowed
ggml_backend_sched_transport_free_ring(sched, i, true);
}
}

Ring freeing does not invalidate the staged input-copy data and buffer fields:

static void ggml_backend_sched_transport_free_ring(ggml_backend_sched_t sched, int backend_id, bool sync_consumers) {
struct ggml_backend_sched_transport_ring * r = &sched->transport.rings[backend_id];
if (r->buffer == NULL) {
return;
}
// nothing may still be reading from or writing into the ring
if (r->transfer) {
ggml_backend_synchronize(r->transfer);
}
if (sync_consumers) {
ggml_backend_synchronize(sched->backends[backend_id]);
}
ggml_backend_buffer_free(r->buffer);
r->buffer = NULL;
r->slot_size = 0;
for (int i = 0; i < GGML_SCHED_MAX_TRANSPORT_SLOTS; i++) {
r->slots[i].release_armed = false;
}
}

Calling either setter after graph allocation can therefore leave the graph pointing into freed ring storage.

Fix: make the configuration immutable once allocation starts and enforce that state. Passing transport configuration at scheduler construction would be safer than mutable exported setters. At minimum, budget needs the same enforced precondition as depth.

5. The claimed ggml_tensor size preservation is false on some 32-bit ABIs

The parent reserves eight trailing bytes:

void * data;
char name[GGML_MAX_NAME];
void * extra; // extra things e.g. for ggml-cuda.cu
char padding[8];
};

The PR replaces that with size_t while claiming sizeof(ggml_tensor) does not change:

void * data;
char name[GGML_MAX_NAME];
void * extra; // extra things e.g. for ggml-cuda.cu
// number of leading bytes of this tensor's storage that are guaranteed not to be
// written during a single graph evaluation. 0 means "not known".
// set on the tensor that owns the storage, by whoever knows what the graph will write;
// a view inherits the part of it that its own byte window covers. read by the backend
// scheduler, which may use it to deliver a host-resident split input to an accelerator
// before the split that reads it runs, see
// ggml_backend_sched_set_transport_pipeline_depth().
// (kept last, in place of the former trailing padding, so that sizeof(struct ggml_tensor)
// does not change)
size_t stable_prefix;
};
static const size_t GGML_TENSOR_SIZE = sizeof(struct ggml_tensor);
// declare that the first nbytes bytes of tensor->data cannot change while a graph that
// reads this tensor is being evaluated. nbytes is clamped to ggml_nbytes(tensor).
// set it on the tensor that owns the storage, not on a view of it, and keep it current:
// it must describe the graph that is about to run, including when that graph is reused.
GGML_API void ggml_set_stable_prefix(struct ggml_tensor * tensor, size_t nbytes);
GGML_API size_t ggml_get_stable_prefix(const struct ggml_tensor * tensor);

That is true on LP64, but on an ILP32 ABI such as i386, size_t is four bytes and the structure can shrink from 256 to 252 bytes.

Fix: preserve an eight-byte storage slot, for example with a fixed-size union or representation, and add layout assertions for representative 32-bit and 64-bit targets.

The memory-cap arithmetic also needs overflow checks. On 32-bit, sufficiently large MiB values can wrap to zero, which means uncapped:

auto create_sched = [&](bool pipeline_parallel) {
sched.reset(ggml_backend_sched_new(
backend_ptrs.data(), backend_buft.data(), backend_ptrs.size(),
max_nodes, pipeline_parallel, cparams.op_offload));
// only a host-resident KV cache produces the deliveries this pipelines
ggml_backend_sched_set_transport_pipeline_budget(sched.get(),
(size_t) cparams.kv_pipeline_budget_mib * 1024 * 1024);
ggml_backend_sched_set_transport_pipeline_depth(sched.get(),
cparams.kv_cpu_pinned || !cparams.offload_kqv ? (int) cparams.kv_pipeline_depth : 0);

void ggml_backend_sched_set_transport_pipeline_budget(ggml_backend_sched_t sched, size_t bytes) {
GGML_ASSERT(sched);
const char * env = getenv("GGML_KV_PIPELINE_BUDGET_MIB");
if (env != NULL) {
bytes = (size_t) strtoull(env, NULL, 10) * 1024 * 1024;
}

Use checked parsing, checked MiB conversion, and checked ring addition/multiplication.

Major merge risks

6. Backend eligibility assumes stronger event semantics than the API provides

Depth 1 is enabled by default for a host-resident cache:

llama.cpp/common/common.h

Lines 581 to 588 in 175236b

bool no_kv_offload = false; // disable KV offloading
bool kv_cpu_pinned = false; // use pinned host buffers for CPU-resident KV cache storage
bool recurrent_state_offload = false; // offload recurrent state independently of attention KV storage
int32_t kv_gpu_layers = 0; // with no_kv_offload, keep this many attention KV layers device-resident
bool phase_aware_workspace = false; // resize compute schedulers between prompt and generation phases
int32_t kv_pipeline_depth = 1; // splits of look-ahead for pipelined delivery of a host-resident KV cache (0 = off)
int32_t kv_pipeline_budget_mib = 128; // hard cap on the device memory that delivery may use (0 = uncapped)
bool warmup = true; // warmup run

The eligibility test uses function-pointer presence as proof of nonblocking stream-to-stream ordering. That is not true for all selected backends. SYCL blocks the host in event->wait():

static void ggml_backend_sycl_event_wait(ggml_backend_t backend, ggml_backend_event_t event) try {
GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
sycl::event* sycl_event = static_cast<sycl::event*>(event->context);
if (ggml_backend_is_sycl(backend)) {
SYCL_CHECK(CHECK_TRY_ERROR(sycl_event->wait()));
} else
GGML_ABORT("fatal error");
} catch (sycl::exception const& exc) {
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
<< ", line:" << __LINE__ << std::endl;
std::exit(1);

WebGPU event wait calls host synchronization:

static void ggml_backend_webgpu_event_record(ggml_backend_t backend, ggml_backend_event_t event) {
ggml_backend_webgpu_context * backend_ctx = (ggml_backend_webgpu_context *) backend->context;
ggml_backend_webgpu_event_context * event_ctx = (ggml_backend_webgpu_event_context *) event->context;
event_ctx->future = backend_ctx->webgpu_ctx->global_ctx->queue.OnSubmittedWorkDone(
wgpu::CallbackMode::AllowSpontaneous, [](wgpu::QueueWorkDoneStatus, wgpu::StringView) {});
event_ctx->recorded = true;
}
static void ggml_backend_webgpu_event_wait(ggml_backend_t backend, ggml_backend_event_t event) {
GGML_UNUSED(backend);
ggml_backend_webgpu_device_event_synchronize(nullptr, event);
}

This contradicts the design requirement that recycling not stop the host and can produce serialization or regressions.

Fix: add a capability that specifically guarantees nonblocking event waits, or keep the feature opt-in until each backend is validated. For the current fork, a temporary CUDA-only gate is safer than generic pointer-based eligibility.

7. The late-tail path does not order against a general async producer

The original path either uses a backend-aware async copy or synchronizes the producer and consumer before a blocking copy:

// 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
if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) {
const int64_t t0 = tr->debug >= 2 ? ggml_time_us() : 0;
ggml_backend_synchronize(input_backend);
if (sched->events[split_backend_id][sched->cur_copy] != NULL) {
ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]);
} else {
ggml_backend_synchronize(split_backend);
}
if (tr->debug >= 2) {
tr->t_sync_us += ggml_time_us() - t0;
}
const int64_t t1 = tr->debug >= 2 ? ggml_time_us() : 0;
ggml_backend_tensor_copy(input, input_cpy);
if (tr->debug >= 2) {
tr->t_copy_us += ggml_time_us() - t1;
tr->n_bytes_ordered += ggml_nbytes(input);
}
if (tr->debug >= 3 && !tr->named_ordered) {
GGML_LOG_INFO("%s: ordered copy %s %zu KiB from %s\n", __func__, input->name,
ggml_nbytes(input) >> 10, ggml_backend_buft_name(input->buffer->buft));
}
}

The staged late-tail path directly calls set_tensor_async from input->data without coordinating with input_backend:

// 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]);
struct ggml_tensor * input = split->inputs[input_id];
struct ggml_tensor * input_cpy = tensor_copy(input, split_backend_id, sched->cur_copy);
if (staged && ggml_backend_sched_input_is_staged(sched, split_id, input_id)) {
// whatever prefix was stable went out on the transfer stream earlier; the rest is
// what an earlier split of this graph may still have written, and it is only safe
// to read now that every earlier split has run. It goes on the consumer's own
// stream, where it is already ordered ahead of the kernels and behind the reader
// of whatever occupied this slot before.
const size_t prefix = ggml_backend_sched_input_stable_prefix(input);
const size_t nbytes = ggml_nbytes(input);
if (nbytes > prefix) {
ggml_backend_tensor_set_async(split_backend, input_cpy,
(const char *) input->data + prefix, prefix, nbytes - prefix);
tr->n_bytes_late += nbytes - prefix;
}
continue;

The current llama KV producer is CPU-oriented, but the new ggml API is generic. An asynchronous producer can still be writing the tail when the destination starts reading it.

Fix: either restrict staging to persistent CPU-produced tensors and document that restriction, or preserve producer ordering for the late region.

8. Inputs with no stable prefix are still redirected into the ring

The public documentation says only inputs carrying a stable prefix are eligible:

// Pipelined delivery of host-resident split inputs.
//
// Without it, a split that reads a host-resident input pays copy + compute in series:
// the transfer is issued on the consumer's own stream immediately before the kernels
// that read it. With it, the scheduler keeps a ring of `depth` staging slots outside
// the graph allocator's reach and issues the stable prefix of a later split's inputs on
// a separate transfer stream while the current split computes, so the transfer retires
// underneath the kernels.
//
// Only inputs that carry a stable prefix (ggml_set_stable_prefix) are eligible: without
// one, the scheduler cannot know that an earlier split of the same graph will not still
// write the bytes it would deliver ahead of time. Everything else keeps the ordered path.
//
// `depth` is how many splits ahead deliveries run; 0 disables pipelining. The ring holds a
// couple of slots more than that, so that recycling a slot never has to wait for a reader
// that is still running. Requires a destination backend with asynchronous transfers and
// events; where that is missing the setting is ignored. Costs roughly (depth + 2) *
// (largest staged split) of device memory. Must be called before the first graph is
// allocated.
GGML_API void ggml_backend_sched_set_transport_pipeline_depth(ggml_backend_sched_t sched, int depth);

Actual ring membership deliberately ignores the prefix:

// Whether a split input belongs in its backend's ring.
//
// Deliberately independent of the stable prefix. Membership decides where an input copy lives,
// which the graph allocator has to know when it reserves -- and at reserve time there is no
// ubatch yet, so no prefix. The prefix decides only how much of a staged input can go early;
// zero means all of it waits for the split, which is the ordered path's timing with the ring's
// storage, and is still correct.
static bool ggml_backend_sched_input_can_stage(
ggml_backend_sched_t sched, struct ggml_backend_sched_split * split, int input_id) {
if (!ggml_backend_sched_transport_ring_enabled(sched, split->backend_id)) {
return false;
}
struct ggml_tensor * input = split->inputs[input_id];
// user inputs must be copied immediately, before the user can overwrite them
if (input->flags & GGML_TENSOR_FLAG_INPUT) {
return false;
}
ggml_backend_buffer_t buf = input->view_src ? input->view_src->buffer : input->buffer;
if (buf == NULL || !ggml_backend_buffer_is_host(buf)) {
return false;
}
// weights take the used-experts path in the split loop, which delivers a subset of the bytes
if (ggml_backend_buffer_get_usage(buf) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) {
return false;
}
return tensor_copy(input, split->backend_id, sched->cur_copy) != NULL;
}

This includes transposed V, which is explicitly assigned a zero prefix:

for (const auto & layer : layers) {
if (layer.k) {
ggml_set_stable_prefix(layer.k, min_row*layer.k->nb[1]);
}
if (layer.v) {
// the transposed V cache scatters each ubatch across the whole tensor, so there is
// no leading region that this ubatch leaves alone
ggml_set_stable_prefix(layer.v, v_trans ? 0 : min_row*layer.v->nb[1]);
}
}

Those inputs consume ring capacity and use the new event path while contributing no overlap. This can cause transposed-V configurations to cross the budget earlier despite the documentation saying they stay untouched.

Fix: use a static transport-candidate annotation for reserve-time membership and a separate per-evaluation stable byte count.

Validation and instrumentation

There are no automated test changes. The exactness harness prints hashes but does not compare them:

DEPTHS=(0 1 4); [ $# -gt 0 ] && DEPTHS=("$@")
rc=0
for D in "${DEPTHS[@]}"; do
echo "== pipeline depth=$D ctx=$CTX prefill lengths=$LENGTHS"
LOG=$(mktemp /tmp/r4-kv-pipeline.XXXX.log)
GGML_KV_PIPELINE_DEPTH=$D taskset -c "$PIN" "$BUILD/bin/llama-server" -m "$MODEL" \
-ngl 99 -sm none -mg 0 -t 3 -nkvo --kv-cpu-pinned --recurrent-state-offload \
-fa on -ctk q8_0 -ctv q8_0 -b 512 -ub 512 -c "$CTX" --parallel 1 \
--host 127.0.0.1 --port "$PORT" --no-warmup > "$LOG" 2>&1 &
SRV=$!
for _ in $(seq 1 600); do
curl -sf "http://127.0.0.1:$PORT/health" >/dev/null 2>&1 && break
sleep 1
done
python3 "$HERE/r4-kv-pipeline-exact.py" "$PORT" "$LENGTHS" || rc=$?
kill $SRV 2>/dev/null; wait $SRV 2>/dev/null
rm -f "$LOG"
done

try:
with urllib.request.urlopen(req, timeout=14400) as r:
d = json.load(r)
except Exception as e:
print(f"{label} REQUEST_FAILED {type(e).__name__}", flush=True)
return False
m = d["choices"][0]["message"]
# reasoning models put most of the generation in reasoning_content; hash both
text = (m.get("reasoning_content") or "") + "\x00" + (m.get("content") or "")
t = d.get("timings", {})
# a reused prefix shows up as a prompt_n far below the prompt actually sent; the hash it
# produces is not comparable to a fresh prefill, so say so rather than reporting it silently
prompt_n = t.get("prompt_n") or 0
reused = prompt_n < want_prefill // 2
print(f"{label:<18} {hashlib.sha256(text.encode()).hexdigest()[:16]} "
f"prompt_n={prompt_n:<7} n={t.get('predicted_n'):<4} "
f"pp={t.get('prompt_per_second'):8.2f} tg={t.get('predicted_per_second'):7.3f}"
f"{' CACHE_REUSE' if reused else ''}", flush=True)
return not reused
ok = True
for length in LENGTHS:
ntok = 256 if length <= 4096 else 128
for name in CORPORA:
prompt = nonce(name, length) + filler(name, length) + "\n\n" + QUESTIONS[name]
ok &= ask(f"{name}@{length}", prompt, ntok, length)
sys.exit(0 if ok else 1)

It only fails on a request error or detected cache reuse. It is a manual measurement harness, not an exactness gate.

Reuse the existing test infrastructure rather than adding a new test file. At minimum cover:

  • partial prefix, zero prefix, and reused-graph prefix changes;
  • one backend failing setup while another succeeds;
  • crossing the budget and then running a smaller graph;
  • setter calls after allocation;
  • explicit meta exclusion;
  • depth 0 preserving the old path;
  • event semantics for every backend that remains enabled.

Some profiling claims are also not reproducible from this head. The document says debug level 3 showed individual copy costs:

The 3.57 ms that remains moves 0.4 MiB, and `GGML_SCHED_TRANSPORT_DEBUG=3` shows
that almost all of it is one copy: `attn_inp_k_rot`, 256 KiB, 18 us on the
ordered path and 3.4 ms behind one split of look-ahead. The 32 KV store copies
cost 353 us between them.
It looks like latency and is not. A blocking copy shares the device's copy engine
with the deliveries and waits for what is already queued there: two staged splits
at 22.0 GB/s is 3.6 ms, which is the number. Two things were tried and neither
helped. Issuing the delivery in pieces so the blocking copy can interleave does
nothing -- the engine is FIFO across streams, `attn_inp_k_rot` stays at 3.4 ms at
every piece size, and small pieces cost throughput (29.80 t/s whole, 28.73 at
4 MiB, 22.01 at 1 MiB). Putting the copy on the consumer's own stream so the host
never blocks moves the time rather than removing it: the ordered copy falls from
3.57 ms to 0.16 ms, the consumer wait rises from 27.31 ms to 31.05 ms, and
throughput does not move (29.808 against 29.834).

The implementation logs only tensor name and size:

const int64_t t1 = tr->debug >= 2 ? ggml_time_us() : 0;
ggml_backend_tensor_copy(input, input_cpy);
if (tr->debug >= 2) {
tr->t_copy_us += ggml_time_us() - t1;
tr->n_bytes_ordered += ggml_nbytes(input);
}
if (tr->debug >= 3 && !tr->named_ordered) {
GGML_LOG_INFO("%s: ordered copy %s %zu KiB from %s\n", __func__, input->name,
ggml_nbytes(input) >> 10, ggml_backend_buft_name(input->buffer->buft));
}

Also, n_stop_recycle increments when a stream wait is enqueued, not when look-ahead actually stops:

// the previous occupant of this slot must be read before the slot is overwritten. This is
// ordered stream to stream rather than through the host: blocking the host here would
// hold back the work it has not enqueued yet, which is what the margin exists to avoid.
if (slot->release_armed) {
ggml_backend_event_wait(r->transfer, slot->release);
slot->release_armed = false;
tr->n_stop_recycle++;
}

The public stats getter exposes only deliveries and early/late bytes, not all the counters described by the document:

void ggml_backend_sched_get_transport_pipeline_stats(
ggml_backend_sched_t sched, int64_t * n_deliveries, int64_t * n_bytes_early, int64_t * n_bytes_late) {
GGML_ASSERT(sched);
if (n_deliveries) { *n_deliveries = sched->transport.n_deliveries; }
if (n_bytes_early) { *n_bytes_early = sched->transport.n_bytes_early; }
if (n_bytes_late) { *n_bytes_late = sched->transport.n_bytes_late; }
}

Restore the instrumentation that produced the documented numbers or narrow the claims.

Commit cleanup

The first two commits use Co-Authored-By for Claude. Repository policy requires Assisted-by for agent assistance:

b6129ae

2c756fe

Later commits use the required trailer. Rewrite the first two before merge.

Recommended development order

  1. Exclude meta and narrow the enabled backend set.
  2. Fix partial-backend cleanup and the permanent budget latch.
  3. Freeze scheduler configuration before allocation.
  4. Preserve 32-bit layout and add checked size arithmetic.
  5. Tighten the candidate/source-ordering contract.
  6. Add automated regression coverage using existing test files.
  7. Make the telemetry match the documentation.
  8. Run the full relevant CI matrix.
  9. Capture a short nsys trace at roughly 4k, 16k, and 32k. Use NCU only if nsys reveals a kernel-side bottleneck.

With --no-kv-offload the attention history lives in host RAM and reaches the
accelerator on every decode token. The scheduler issued that transfer on the
consumer's own stream immediately before the kernels that read it, so a token
cost copy + compute in series.

The bytes and the attention operations are unchanged; only the point at which
the transfer is issued moves. Greedy output is byte-identical to the ordered
path -- verified against a build without these changes, at every look-ahead
tested, single GPU and layer-split across two.

Three pieces, each load-bearing:

- ggml_tensor::stable_prefix records how many leading bytes of a tensor's
  storage the graph about to run will not write. The KV window is not stable for
  a whole graph -- a CPU split writes this ubatch's rows into it between one
  layer's attention and the next -- but everything below the lowest written row
  is, and at decode depth that is essentially all of it. llama_kv_cache sets it
  from apply_ubatch(), before the graph is built and allocated, so the plan and
  the deliveries are decided against the same write position even when the graph
  is reused; build_graph_shift() clears it.

- A staging ring the graph allocator cannot reach. ggml-alloc may recycle a
  graph-owned input copy after its last graph-level consumer while a look-ahead
  transfer is still in flight. The scheduler allocates the ring itself and points
  the staged copies at it before allocation; a ready/release event pair per slot
  carries the handover in each direction. Every eligible accelerator gets its own
  ring, cursor and budget, so a layer-split model pipelines on each device and a
  device with no room falls back alone.

- A look-ahead that stays clear of the ring's tail. A delivery L splits ahead
  recycles the slot of the split L - n_slots back, so n_slots == L + 1 recycles
  the split just enqueued and still running. The ring keeps two slots of margin,
  deliveries are issued after a split is enqueued rather than before, and slot
  recycling is ordered stream to stream rather than through the host. Each of
  those three alone costs the entire gain while still producing correct output.

--kv-pipeline-depth N, default 1, 0 restores the ordered path exactly. It only
engages where a host-resident cache produces the deliveries.

Because a host-resident cache exists to keep device memory free, the staging is
capped outright by --kv-pipeline-budget (default 128 MiB per device) rather than
by a fraction of what happens to be free. A ring is (N + 2) slots of one
attention layer's K and V over the whole context, so it grows with the context:
27 MiB at 4k, 213 MiB at 32k, 1.7 GiB at 256k. Past the cap the scheduler
declines and keeps the ordered path, and declining costs nothing -- the check
runs before anything is allocated, the decision is latched because a context only
grows, and the transfer backend is created lazily and released with the ring.

Single GPU (RTX 4070, Qwen3.8-27B-UD-IQ2_M, -nkvo --kv-cpu-pinned, q8_0 K/V),
A/B/A/B with reversed arm order:

  depth   ordered            pipelined          gain
   4,096  31.7324, 31.7363   37.0889, 37.0741   +16.9%
  16,384  19.6765, 19.6854   31.5352, 31.5807   +60.4%
  32,768  13.0264, 13.0254   15.5325, 15.5329   +19.3%   (needs a raised budget)

Server decode behind an 18,422-token prompt: 18.468 -> 30.685 t/s, +66.2%.

Two GPUs (RTX 4070 + RTX 3060, Qwen3.8-27B-UD-Q5_K_M, -sm layer), both rings
engaged: 13.06 -> 18.05 t/s at 4,096 and 6.86 -> 9.75 t/s at 16,384.

The gain narrows with depth because compute is a shrinking share of the token,
so there is less to hide the copy behind. That is arithmetic, not an
implementation limit, and more look-ahead makes it worse rather than better.

Tensor parallelism keeps the ordered path: the scheduler sees one meta backend
there and the ring is a byte arena, while a meta buffer places tensors as
per-device slices rather than at offsets. It declines rather than staging into
something it cannot address.

docs/kv-transport-pipelining.md carries the design, the numbers and the limits;
docs/repro/ carries the scripts that produced them.

Assisted-by: Claude Opus 5
…cripts

llama-bench does not expose --kv-cpu-pinned or --recurrent-state-offload the way
llama-server does, so two of the reproduction scripts passed flags the binary
rejects. They now probe --help and pass only what it takes.

The feature doc gains what is actually left: why -sm tensor keeps the ordered
path (no events in the meta layer, and a ring that is a byte arena while a meta
buffer places tensors as per-device slices), the correctness problem underneath
it that is not this feature's, and the rest of the open list -- the transient
device-memory peak, the --kv-gpu-layers comparison, and the exactness harness's
dependence on baseline determinism it does not have.

The decline for a backend that cannot record events now says so by name rather
than falling into the generic "no backend supports" line, because the backend it
catches is the meta backend and the next person to look will want to know that.

Assisted-by: Claude Opus 5
The host-time breakdown the feature doc describes had no code behind it: the
counters existed but nothing accumulated or printed them. GGML_SCHED_TRANSPORT_DEBUG=2
now reports the split loop as a mean over each 128 graphs, with the bytes the
ordered path still moves and why the look-ahead stopped; =3 names the tensors
that are still on it. That is what found the rest: 40 blocking copies a token
moving 0.4 MiB, 32 of them the device-to-host KV store.

The budget warning now reports what the ring costs at the full context next to
what it costs now, so --kv-pipeline-budget can be sized against the number that
matters. It is still applied per graph: enforcing the projection would refuse the
ring for every large -c even when the window never gets near it.

llama-bench gains -kvcp and -rso. Without them a host-resident run measures
something else entirely -- 9.02 against 19.43 t/s ordered at 16,384 -- and the
repro scripts had been silently dropping both since llama-bench lost them.

The exactness harness gives every task a nonce derived from its own name and
length, so no two share a prefix the server can restore, and fails a task whose
prompt_n says one was reused anyway.

Assisted-by: Claude Opus 5
…t it does

GGML_SCHED_TRANSPORT_DEBUG=3 now reports what each remaining blocking copy cost,
not just its name. It turns out one of them is almost all of it: attn_inp_k_rot,
256 KiB, 18 us on the ordered path and 3.4 ms behind one split of look-ahead.

That is the copy engine, not latency. A blocking copy waits for the deliveries
already queued on it, and two staged splits at 22.0 GB/s is 3.6 ms. Issuing the
delivery in pieces does not help, the engine is FIFO across streams. Putting the
copy on the consumer's stream so the host never blocks moves the time into the
consumer wait and leaves throughput alone. The doc records both, so the next
person does not spend the afternoon on it again.

Assisted-by: Claude Opus 5
records@18432 was giving different answers across otherwise identical N = 0 runs,
which made it useless as a gate and looked like the pipeline breaking exactness.
It is not the task: asked on its own with the prompt cache off it returns the same
hash three times running, at -c 32768 and at -c 65536.

It is the harness. All eight tasks share one server with prompt caching on, and
records@18432 is about 29.6k tokens with a task of about the same size ahead of
it, so the two do not both fit in a 32,768 cache and placement depended on what
was still resident. The nonce stops a prefix being restored, it does not stop the
pressure. cache_prompt=false does.

Two independent N = 0 passes now agree on all eight tasks, and N = 0, N = 1 and
N = 4 agree on all eight.

Assisted-by: Claude Opus 5
Restrict staging to annotated CUDA inputs, make backend decline complete, re-evaluate budgets, freeze scheduler configuration, and preserve tensor layout. Add regressions for prefix changes and fallback behavior.

Assisted-by: OpenAI Codex
@Piggidragon
Piggidragon force-pushed the kv/pipelined-transport branch from 175236b to 2f70cbe Compare August 28, 2026 17:51
@Piggidragon

Copy link
Copy Markdown
Author

Addressed on head 2f70cbe.

  • Removed the unvalidated meta event and transfer work. Meta and non-CUDA backends are explicitly excluded.
  • Added a static transport annotation. Only persistent host K and non-transposed V tensors are marked; the stable prefix remains per evaluation.
  • Centralized backend-decline cleanup and covered one failed backend alongside one successful backend.
  • Removed the permanent budget latch, so smaller later graphs are reconsidered.
  • Locked depth and budget configuration once graph allocation starts.
  • Preserved the eight-byte ggml_tensor tail slot and added checked MiB conversion and ring-size arithmetic.
  • Restricted eligibility and producer ordering to the validated CUDA host-KV path.
  • Made the exactness harness compare hashes and corrected the telemetry documentation.

The existing test-alloc target now covers partial, zero, and changing prefixes on a reused graph; depth 0; budget recovery; post-allocation setters; partial setup failure; meta exclusion; annotation requirement; and non-CUDA exclusion.

Local validation on the final commit:

  • test-alloc passed
  • llama-cli and llama-server built successfully
  • reproduction shell scripts and Python harness passed syntax checks
  • git diff check passed

The branch is rebased onto llama/dev at 01b141f. The obsolete meta commit was removed and the first two commit trailers now use Assisted-by. CI is queued on the updated head.

I could not capture the requested nsys traces locally because this environment has no working NVIDIA driver. The CUDA CI job is queued, but it does not replace the runtime trace.

@Piggidragon

Copy link
Copy Markdown
Author

Review findings

ggml/src/ggml-backend.cpp:1921 — medium. tr->split_input_ofs[sched->n_splits] = n_inputs_total; runs unconditionally, but the allocation above is guarded by if (tr->plan_capacity < sched->n_splits). On the first planned graph with sched->n_splits == 0 the guard is 0 < 0 → false, so split_input_ofs is still the calloc'd NULL and this is a NULL store. Reachable by any ggml_backend_sched user that opts into the transport and then allocates an empty/node-free graph. Guard the write (or the whole plan) on sched->n_splits > 0.

ggml/src/ggml-backend.cpp:2260 — medium. ggml_backend_sched_transport_plan() is called before the ggml_gallocr_reserve_n fallback at line 2291, so if that fallback fires while a ring is live, galloc re-reserves a graph whose staged input_cpy tensors already carry ring addresses. ggml_gallocr_allocate_node skips anything with data != NULL, so the recorded plan gives them buffer_id = -1 and the compute buffer shrinks by the staged bytes. When the ring is later declined (the context grows past --kv-pipeline-budget, or ggml_backend_buft_alloc_buffer fails), those copies come back with data == NULL, ggml_gallocr_node_needs_realloc returns false for talloc->buffer_id < 0, and the graph is fully re-reserved mid-generation — the "unexpected graph reallocation" this fork instruments with GGML_SCHED_DEBUG_REALLOC, and a compute-buffer size change that the resizable/shared-workspace machinery (sched_buffer_owner, phase_aware_workspace) is not expecting. Clearing the staged copies' data around the reserve and re-pointing them afterwards avoids it.

ggml/src/ggml-backend.cpp:1950 — low. The "reader further down the graph" disqualification scan is O(staged_inputs x total_graph_nodes x GGML_MAX_SRC): for every staged input it walks every node of every split. On a 60-layer model with a host KV cache that is ~100 staged inputs x ~1500 nodes x 10 srcs ≈ 1.5M pointer comparisons per ggml_backend_sched_alloc_graph, i.e. on every prompt-processing ubatch and every decode where the graph is not reused — a couple of milliseconds added to the path the feature exists to shave milliseconds off. The scan result depends only on input_cpy identity, so it can be done in one pass that first collects the staged copies into a set and then sweeps the splits once.

ggml/src/ggml-backend.cpp:2724 — low. GGML_KV_PIPELINE_DEPTH overrides the caller's argument unconditionally, so --kv-pipeline-depth 0 (and the internal 0 that llama_context passes for a device-resident cache) is silently ignored when the variable is set in the environment. Worse, a malformed value makes the setter return false, and llama_context::sched_reserve (src/llama-context.cpp:882) turns that into throw std::invalid_argument("invalid KV transport pipeline configuration") — model load aborts with a message that never mentions the environment variable. Same shape in ggml_backend_sched_set_transport_pipeline_budget for GGML_KV_PIPELINE_BUDGET_MIB. An env var should be a fallback for an unset argument, not an override, and a bad value should warn and be ignored rather than fail context creation.

ggml/src/ggml-backend.cpp:2399 — low. The comment at line 876 says debug >= 3 "names them once", but named_ordered is only set at line 2621, inside the n_graphs % 128 == 0 block. With GGML_SCHED_TRANSPORT_DEBUG=3 every ordered copy is logged on all of the first 128 graphs (hundreds of lines per graph), and a run shorter than 128 graphs never stops. Set named_ordered = true where the naming happens.

docs/repro/r4-kv-pipeline-exact.sh:33 — low. BASE is only assigned inside the if python3 ...; then success branch. If the depth-0 arm fails (server start timeout, CACHE_REUSE, a request error), BASE stays empty and the next depth becomes the baseline — so the gate reports success while only having compared depth 1 against depth 4, never against the ordered path. rc is non-zero in that case, but the diff output claims agreement it never checked. Bind the baseline to the first depth explicitly, or abort when the first arm fails.


🤖 Generated with Claude Code

@GenerelSchwerz

Copy link
Copy Markdown
Owner

AI-assisted review of current head a29866f — I think this needs changes before merge.

  1. The transport ring bypasses the backend allocation contract. Ring entries are sized with ggml_nbytes() and bound by assigning data/buffer directly (1844-1849, 2153-2156). Pristine ggml permits get_alloc_size() to exceed ggml_nbytes(), notably for quantized tensors; CUDA MMQ may then clear that additional padding. Because the ring is marked COMPUTE, a transported quantized input can overwrite the following entry/slot. Please lay entries out with ggml_backend_buft_get_alloc_size() and initialize/bind them through the normal backend tensor allocation path. A dummy backend test where get_alloc_size > nbytes would cover this.

  2. The published 32k result is not reproducible from this head. The documentation says the scripts used --kv-pipeline-budget 512 and notes that 32k requires 204 MiB (docs lines 125-133, 270-272), but neither repro script passes a budget and llama-bench does not parse that option (480-483, 846-879). I confirmed the built binary rejects --kv-pipeline-budget 512. Please expose the budget in llama-bench, include it in output, pass 512 in both scripts, and rerun or revise the table.

  3. The repro scripts fail open. The context sweep converts JSON errors to FAILED and then exits successfully (27-38); with a missing model I got four FAILED arms and exit status 0. The A/B script can likewise hide failures in its first three arms because only the final command determines the inner shell status (29-39). Both also silently omit required kvcp/rso options. Please make option checks and every arm fail closed, with set -euo pipefail and explicit status propagation.

  4. Current tests do not validate transferred data or event ordering. The dummy async copy and event methods only increment counters or do nothing (136-153), while the transport test asserts only statistics (1264-1298). Wrong offsets, overlap, missing initialization, or broken event order can therefore pass. Please add allocation-bound/data-content coverage and a CUDA exactness gate if feasible.

For an eventual pristine llama.cpp port, this also needs to be stacked after the fork's host-KV-storage/accelerator-attention separation. Upstream currently moves the entire attention region to CPU when KV offload is disabled, so this scheduler patch alone has no useful host-to-accelerator path.

Local checks passed: git diff --check, CPU-only release build, test-alloc, Python compilation, bash -n, and shellcheck. CUDA and several portability CI jobs were still pending when reviewed.

The ring laid its entries out with ggml_nbytes() and bound them by writing data
and buffer directly. A buffer type may ask for more than ggml_nbytes() for a
tensor -- CUDA does for a quantized one, and MMQ clears that padding -- so an
entry could reach into the next one. Entries are now sized with
ggml_backend_buft_get_alloc_size() and bound with ggml_backend_tensor_alloc(),
which also gives them the buffer's own initialization and its bounds check.

test-alloc gets a dummy buffer type whose get_alloc_size exceeds ggml_nbytes,
and a two-entry ring test that checks the entries stay inside the ring and out
of each other, that every byte of an entry is delivered once from the matching
source offset, and that nothing waits on an event before it is recorded.

llama-bench takes -kvpb/--kv-pipeline-budget and reports it. The repro scripts
pass 512 and now fail closed: they refuse a build without -kvcp, -rso or -kvpb
instead of dropping the option, and every arm propagates its status. The
llama-bench table in the doc was measured before the budget existed, so it says
so, and the 32,768 row is marked as needing a re-measurement.

Assisted-by: Claude Opus 5
@github-actions github-actions Bot added documentation Improvements or additions to documentation examples testing ggml labels Sep 1, 2026
Piggidragon added a commit to Piggidragon/llama.cpp that referenced this pull request Sep 2, 2026
From the moe-cache-drafting branch: --moe-expert-cache-size,
--moe-expert-cache-l2-pinned-mb, --experimental-logs, the automatic grouped
decode / prefetch / bias residency behaviour, and the layer-split-only and
speculative interactions. Marked as coming from an unmerged branch, like the
PR GenerelSchwerz#57 and GenerelSchwerz#39 material.

Assisted-by: Claude Sonnet 5
Claude-Session: https://claude.ai/code/session_013Bs926e6SdsxnkrqqKq5HP
@GenerelSchwerz

Copy link
Copy Markdown
Owner

Review of head 0933834bb329fd53590f13834889dbffa41e911b: changes requested.

Blocking

1. The optional ring can consume memory required by the graph

The transport planner allocates and retains the ring before graph allocation:

if (r->buffer == NULL || r->slot_size < slot_size[bid]) {
ggml_backend_sched_transport_free_ring(sched, bid, true);
ggml_backend_buffer_type_t buft = sched->bufts[bid];
// The ring is allocated after the graph allocator has reserved its buffers, so it must
// not take the room those buffers may still have to grow into.
ggml_backend_dev_t dev = ggml_backend_get_device(sched->backends[bid]);
size_t dev_free = 0, dev_total = 0;
if (dev != NULL) {
ggml_backend_dev_memory(dev, &dev_free, &dev_total);
}
if (dev_free > 0 && (dev_free <= GGML_SCHED_TRANSPORT_HEADROOM || ring_size > dev_free - GGML_SCHED_TRANSPORT_HEADROOM)) {
if (!r->reported_no_room) {
GGML_LOG_WARN("%s: transport ring on %s would need %zu MiB and leave less than "
"%u MiB of the %zu MiB free, staying on the ordered path\n", __func__,
ggml_backend_name(sched->backends[bid]), ring_size >> 20,
GGML_SCHED_TRANSPORT_HEADROOM >> 20, dev_free >> 20);
r->reported_no_room = true;
}
ggml_backend_sched_transport_decline_backend(sched, bid);
continue;
}
ggml_backend_buffer_t buffer = ggml_backend_buft_alloc_buffer(buft, ring_size);
if (buffer == NULL) {
GGML_LOG_WARN("%s: failed to allocate %zu MiB for the transport ring on %s, "
"pipelining disabled there\n", __func__, ring_size >> 20,
ggml_backend_name(sched->backends[bid]));
ggml_backend_sched_transport_decline_backend(sched, bid);
continue;
}
ggml_backend_buffer_set_usage(buffer, GGML_BACKEND_BUFFER_USAGE_COMPUTE);
r->buffer = buffer;
r->slot_size = slot_size[bid];

Only afterward does graph allocation fall back to ggml_gallocr_reserve_n(). If that allocation fails, the function returns without releasing the optional ring or retrying the ordered path:

// lay out the transport rings and point the staged input copies at them before the graph is
// allocated, so ggml-alloc sees those copies as already allocated and leaves them alone
ggml_backend_sched_transport_plan(sched);
// allocate graph
if (backend_ids_changed || !ggml_gallocr_alloc_graph(sched->galloc, &sched->graph)) {
#ifndef NDEBUG
GGML_LOG_DEBUG("%s: failed to allocate graph, reserving (backend_ids_changed = %d)\n", __func__, backend_ids_changed);
#endif
if (sched->debug_realloc > 0) {
// we are interested only in situations where the graph was reallocated even though its size remained the same [GGML_SCHED_DEBUG_REALLOC]
// example: https://github.com/ggml-org/llama.cpp/pull/17143
const bool unexpected = !backend_ids_changed && sched->debug_prev_graph_size == sched->debug_graph_size;
if (unexpected || sched->debug_realloc > 1) {
GGML_ABORT("%s: unexpected graph reallocation (graph size = %d, nodes = %d, leafs = %d), debug_realloc = %d\n", __func__,
sched->debug_graph_size, sched->graph.n_nodes, sched->graph.n_leafs, sched->debug_realloc);
}
}
// the re-allocation may cause the split inputs to be moved to a different address
// synchronize without ggml_backend_sched_synchronize to avoid changing cur_copy
for (int i = 0; i < sched->n_backends; i++) {
if (sched->transport.rings[i].transfer) {
ggml_backend_synchronize(sched->transport.rings[i].transfer);
}
}
for (int i = 0; i < sched->n_backends; i++) {
ggml_backend_synchronize(sched->backends[i]);
}
ggml_backend_sched_transport_clear_addresses(sched);
if (!ggml_gallocr_reserve_n(sched->galloc, &sched->graph, sched->node_backend_ids, sched->leaf_backend_ids)) {
GGML_LOG_ERROR("%s: failed to allocate graph\n", __func__);
return false;
}
ggml_backend_sched_transport_assign_addresses(sched);
if (!ggml_gallocr_alloc_graph(sched->galloc, &sched->graph)) {
GGML_LOG_ERROR("%s: failed to allocate graph\n", __func__);
return false;

This can turn a graph that fits without pipelining into GGML_STATUS_ALLOC_FAILED. I reproduced it with the existing dummy backend:

  • transported input: 40 MiB
  • default depth 1 ring: 120 MiB, within the default 128 MiB budget
  • graph compute allocation: 640 MiB
  • available device memory: one byte below 760 MiB
  • depth 0: succeeds and uses 640 MiB
  • depth 1: the ring allocation succeeds and passes the 512 MiB headroom check, but graph reservation fails because just under 640 MiB remains

Configuration is then locked here, so the caller cannot disable the ring and retry on the same scheduler:

bool ggml_backend_sched_alloc_graph(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
GGML_ASSERT(sched);
GGML_ASSERT((int)sched->hash_set.size >= graph->n_nodes + graph->n_leafs);
GGML_ASSERT(!sched->is_alloc);
sched->transport.config_locked = true;
sched->cur_copy = sched->next_copy;
sched->next_copy = (sched->next_copy + 1) % sched->n_copies;
ggml_backend_sched_split_graph(sched, graph);
if (!ggml_backend_sched_alloc_splits(sched)) {
return false;
}

The graph-allocation failure path should discard the current graph's optional ring bindings and retry reservation/allocation once on the ordered path.

2. The previously reported zero-input defect remains unresolved

When a split exists but the total number of split inputs is zero, input_staged remains null and is passed to memset with a zero size here:

int n_inputs_total = 0;
for (int i = 0; i < sched->n_splits; i++) {
tr->split_input_ofs[i] = n_inputs_total;
n_inputs_total += sched->splits[i].n_inputs;
}
tr->split_input_ofs[sched->n_splits] = n_inputs_total;
if (tr->input_capacity < n_inputs_total) {
unsigned char * pnew = (unsigned char *) realloc(tr->input_staged, n_inputs_total);
if (pnew == NULL) {
GGML_LOG_WARN("%s: failed to allocate the transport plan, pipelining disabled for this graph\n", __func__);
return;
}
tr->input_staged = pnew;
tr->input_capacity = n_inputs_total;
}
memset(tr->input_staged, 0, n_inputs_total);
tr->plan_n_splits = sched->n_splits;
tr->plan_n_inputs = n_inputs_total;

The new n_splits == 0 guard does not cover this case because the scheduler represents the empty test graph with a split. The PR's own test_transport_empty_graph exercises the path:

llama.cpp/tests/test-alloc.cpp

Lines 1468 to 1478 in 0933834

static void test_transport_empty_graph() {
dummy_backend cuda = dummy_backend_init(SIZE_MAX, 8, true, GGML_BACKEND_DEVICE_TYPE_GPU, "CUDA", false);
dummy_backend cpu = dummy_backend_init(SIZE_MAX, 8, true);
auto graph = make_context();
ggml_backend_t backends[] = { cuda.handle.get(), cpu.handle.get() };
ggml_backend_buffer_type_t bufts[] = { &cuda.buffer_type, &cpu.buffer_type };
ggml_backend_sched_ptr sched(ggml_backend_sched_new(backends, bufts, 2, 128, false, false));
GGML_ASSERT(ggml_backend_sched_set_transport_pipeline_depth(sched.get(), 1));
GGML_ASSERT(ggml_backend_sched_alloc_graph(sched.get(), graph.graph));
GGML_ASSERT(ggml_backend_sched_graph_compute(sched.get(), graph.graph) == GGML_STATUS_SUCCESS);

The test passes in Release, but the exact head fails under UBSan:

ggml/src/ggml-backend.cpp:2052:11: runtime error: null pointer passed as argument 1, which is declared to never be null

This is follow-up to the existing finding, not independent duplicate feedback:

#39 (comment)

Handle n_inputs_total == 0 before the null memset, while initializing every split order to the ordered state.

Will slow the review

1. The PR description materially contradicts the current head

The description still says that three meta-backend prerequisites ship here and that budget decline is latched. The current implementation explicitly excludes meta and requires the exact CUDA registry name:

int n_eligible = 0;
for (int i = 0; i < sched->n_backends; i++) {
ggml_backend_t backend = sched->backends[i];
ggml_backend_dev_t dev = ggml_backend_get_device(backend);
if (dev == NULL || ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_META) {
continue;
}
ggml_backend_reg_t reg = ggml_backend_dev_backend_reg(dev);
if (reg == NULL || strcmp(ggml_backend_reg_name(reg), "CUDA") != 0) {
continue;
}
if (backend->iface.set_tensor_async == NULL ||
backend->iface.event_record == NULL ||
backend->iface.event_wait == NULL) {
continue;
}
if (dev->iface.event_new == NULL) {
continue;
}
if (ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU) {
continue;
}
// the ring is written through the transfer backend, which only accepts the device's own
// default buffer type; a scheduler configured with anything else keeps the ordered path
if (sched->bufts[i] != ggml_backend_dev_buffer_type(dev)) {
continue;
}

The current design document instead says smaller later graphs are reconsidered:

- Under the cap the ring is allocated and the deliveries pipeline.
- Over it the scheduler declines and keeps the ordered path for that graph.
Later graphs are evaluated again, so a smaller live window can use the ring.
- Declining costs nothing in steady state. Both the ring and the transfer
backend's device context are released.
**The cap is applied to what the current graph needs, not to what the full
context would need.** A run whose window stays small keeps the ring whatever
`-n_ctx` says, which is the common case and the reason it is done this way: a
staged input is a view of the cache tensor, so the full-context figure is there
for the asking, but enforcing it would refuse the ring for every large `-c` even
when the window never gets near it. The warning reports both numbers so that
`--kv-pipeline-budget` can be sized against the one that matters.
The cost of deciding per graph is that a context which grows past the budget
allocates a ring for the small early windows and gives it back once it outgrows
them. That transient is bounded by the budget itself, which is the memory the
user already authorised, so it is a property of the cap rather than a defect in

The description's measurement presentation also conflicts with the document, which says the 32k row predates the current budget behavior and needs remeasurement on this head:

> These rows were taken before the budget existed, so they are the uncapped
> numbers. The 32,768 ring is 204 MiB and the default cap is 128 MiB, so that row
> does not reproduce on the current default: it needs `-kvpb 512`, which
> `llama-bench` did not take until now. The scripts pass it, and the row is due a
> re-measurement on the current head.

Update the description so reviewers are evaluating the implementation and validation that actually remain in this revision.

2. Changed prose and public comments extensively violate the no-hard-wrap rule

The repository rule is explicit here:

llama.cpp/AGENTS.md

Lines 79 to 90 in 0933834

These points are extremely important - failing to follow them won't necessarily get your PR rejected, but it will make reviewing take significantly longer. Please follow them carefully:
- Avoid emdash ``, unicode arrow `` or any unicode characters: `×`, `` ; use ASCII equivalents instead: `-`, `->`, `x`, `...`
- Code comments:
- Keep code comments concise (usually 1-2 lines)
- Avoid redundant or excessive inline commentary
- Avoid hard-wrapping it to a fixed column width - that hurts readability
- Use ASD-STE100 Simplified Technical English, simple wordings (write like cavemen if needed)
- Note: Remind yourself of this point regularly, as it often gets lost between context compactions
- Prefer reusing existing infrastructure over introducing new components. Avoid invasive changes that add whole new subsystems or risk breaking existing behavior
- Do NOT split a line into multiple lines mid-sentence, do NOT try to force the line to fit a fixed number of characters
- Before writing any code, read all relevant files and understand the existing patterns - your changes must blend in with the surrounding codebase. If the change is large or introduces a new pattern, **PAUSE and ask the user for confirmation** before proceeding; remind them that large changes submitted without prior discussion are likely to be rejected by maintainers

Examples include:

  • the design-document opening:
    With `--no-kv-offload` (optionally with `--kv-cpu-pinned`), the attention history
    lives in host RAM and has to reach the accelerator on every decode token. The
    backend scheduler used to issue that transfer on the consumer's own stream, right
    before the kernels that read it, so a token cost `copy + compute` in series.
    The transfer and
    the attention arithmetic are the same as before, but the transfer is issued one
    split ahead, on a stream of its own, so the copy engine retires it underneath the
    kernels of the split before it.
  • the public scheduler documentation:
    // Pipelined delivery of host-resident split inputs.
    //
    // Without it, a split that reads a host-resident input pays copy + compute in series:
    // the transfer is issued on the consumer's own stream immediately before the kernels
    // that read it. With it, the scheduler keeps a ring of `depth` staging slots outside
    // the graph allocator's reach and issues the stable prefix of a later split's inputs on
    // a separate transfer stream while the current split computes, so the transfer retires
    // underneath the kernels.
    //
    // Only persistent host inputs marked with GGML_TENSOR_FLAG_TRANSPORT are eligible. Their
    // stable prefix must be current before each evaluation. The producer must be the CPU or the
    // same backend stream that consumes the late region.
    //
    // `depth` is how many splits ahead deliveries run; 0 disables pipelining. The ring holds a
    // couple of slots more than that, so that recycling a slot never has to wait for a reader
    // that is still running. Requires a destination backend with asynchronous transfers and
    // events; where that is missing the setting is ignored. Costs roughly (depth + 2) *
    // (largest staged split) of device memory. Must be called before the first graph is
    // allocated. Returns false after graph allocation starts.
    GGML_API bool ggml_backend_sched_set_transport_pipeline_depth(ggml_backend_sched_t sched, int depth);
    // Hard cap on the staging ring, in bytes. A host-resident cache exists to keep device memory
    // free, so the ring is capped outright and not merely against what happens to be free: past
    // the cap the scheduler declines and keeps the ordered path. 0 removes the cap. Default 128 MiB.
    // Returns false after graph allocation starts. Configuration is immutable then.
    GGML_API bool ggml_backend_sched_set_transport_pipeline_budget(ggml_backend_sched_t sched, size_t bytes);
    // Number of staged deliveries and staged bytes issued since the scheduler was created.
    GGML_API void ggml_backend_sched_get_transport_pipeline_stats(ggml_backend_sched_t sched, int64_t * n_deliveries, int64_t * n_bytes_early, int64_t * n_bytes_late);
  • the public context comments:

    llama.cpp/include/llama.h

    Lines 424 to 431 in 0933834

    uint32_t kv_pipeline_depth; // how many splits ahead the scheduler delivers a host-resident KV cache to the
    // accelerator, so that the transfer runs while the previous split computes.
    // 0 keeps the ordered path, where a decode token pays the transfer and the
    // attention kernels in series. Costs (kv_pipeline_depth + 2) * (largest staged
    // split) of device memory.
    uint32_t kv_pipeline_budget_mib; // hard cap on that device memory, in MiB. Past it the scheduler declines and
    // keeps the ordered path, so a host-resident cache never quietly trades the
    // device memory it exists to save. 0 removes the cap.

This is widespread rather than an isolated nit.

Existing validation concern

I am not duplicating the already-raised request for real CUDA data and ordering validation. It remains unresolved on this exact head: the current CUDA CI job was canceled without running, and commit 0933834b changed ring allocation and binding after the documented runtime validation.

Canceled job:

https://github.com/GenerelSchwerz/llama.cpp/actions/runs/33486643492/job/99788163320

Existing owner feedback:

#39 (comment)

Upstream comparison

No equivalent implementation exists in current pristine upstream. The closest relevant work supports parts of the approach but does not remove the blockers above:

The scope gate otherwise passes. Fork PR #31 documented R4 and its design decision before implementation, and neither upstream ggml-org#21067 nor ggml-org#27311 is a direct duplicate. Upstream issue ggml-org#27757 concerns the separate tensor-parallel host-KV correctness defect.

Verification

  • Release CPU build of test-alloc and llama-bench: passed
  • Release test-alloc: passed
  • UBSan test-alloc: exposed the null memset; no other runtime error appeared when UBSan was allowed to continue
  • Focused default-budget allocation-capacity reproducer: exposed the ring-starvation failure
  • git diff --check: passed
  • Shell syntax, ShellCheck, and Python compilation for the reproduction scripts: passed
  • Benchmark option plumbing for kvcp, kvpd, kvpb, and rso: present
  • CPU, Windows, server, and WebGPU CI jobs passed; CUDA and other self-hosted jobs were canceled while queued
  • No formal review threads exist; earlier feedback is in issue comments

No additional security, lifetime, event-order, public-API, or backend-dispatch finding survived verification.

The rings are laid out and allocated before the graph is, so a device that can
hold the graph alone but not the graph next to a ring turned into
GGML_STATUS_ALLOC_FAILED. The configuration is locked by then, so the caller
could not turn the ring off and retry either. When graph reservation fails the
rings are now released and the reservation is retried once on the ordered path,
and that scheduler keeps the ordered path from then on.

A plan over a split list with no inputs left input_staged unallocated and passed
it to memset, which UBSan reports even at size 0. Such a plan stages nothing, so
it now returns after putting every split back on the ordered path.

test-alloc gets a device capacity on the dummy backend and a test that sizes it
to hold the graph or the ring but not both.

The prose and public comments this branch added were hard-wrapped to a fixed
column, against the repository rule. They are unwrapped, one sentence per line.

Assisted-by: Claude Opus 5
The llama-bench table predated the budget and said so, and the context sweep
and the exactness gate were last run before the ring allocation and binding
changed. All three are re-run on an RTX 4070 with a CUDA build of this head, at
--kv-pipeline-budget 512.

Greedy server output is identical at depth 0, 1 and 4 across all eight tasks.
Throughput is +16.1% at 4,096, +56.6% at 16,384, +20.1% at 32,768 and +13.4% at
65,536, for +28, +104, +206 and +410 MiB of device memory. The 131,072 and
262,144 arms are not re-measured and say so.

The server table is replaced with the four 18,432-prefill tasks of the exactness
gate, which is what this head was actually run on; the copy/compute breakdown
keeps its earlier numbers and says which head they came from.

Assisted-by: Claude Opus 5
… MiB

The context sweep is measured with the budget raised, so its deep rows read
as default behaviour when they are not: past 20,556 rows of window the
default declines and those depths stay ordered. Say so, and add a finer
sweep that puts the peak at 16,384 rows and shows the gain per MiB falling
off as 1/rows^2 above it.

The peak is where copy and compute are equal, and the ring size there works
out to (n_slots / n_attn) * compute * BW - the bytes per row cancel, so the
budget is quant-invariant. That predicts 101 MiB against the 102 MiB
measured, which is what the 128 MiB default is sized against.

Assisted-by: Claude Opus 5
A window over several streams is one view of a tensor whose streams sit
end to end, and both the prefix and the delivery treated it as one flat
byte range. That made the lowest-writing stream cap the stable prefix for
every stream above it, and it copied the cells between one stream's window
and the next, which the graph never reads.

The prefix is now counted within a stream, and a staged input whose last
dimension indexes streams is delivered as one range per stream. A window
over one stream keeps the single flat range it had, so single-sequence
timing and bytes are unchanged.

Behind 8 slots of a non-unified cache this takes decode from 72.18 to
112.70 t/s, against 70.61 ordered. At one slot it measures 35.31 against
35.32 before.

test_transport_multi_stream_ranges pins the delivery: every stream's
window covered once from its own source offset, the unread cells between
them never moved, and the early and late bytes split as the prefix says.
Concurrent slots cannot be gated on output the way one sequence can -
their batching varies between runs, so the same depth gives different
greedy output - which is why this is a unit test.

Assisted-by: Claude Opus 5
The q4_0 arm of the context sweep moves the peak from 16,384 rows to
32,768 and leaves the ring at it at 108 MiB against 102, which is what
"the bytes per row cancel" claims. It was derived before and is measured
now.

Add the parallel numbers and say what the 8-slot row depends on: run on
its own the unified ring fits, and only after a sweep has allocated for
1, 2 and 4 slots in the same process does the headroom guard refuse it.

Assisted-by: Claude Opus 5
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

documentation Improvements or additions to documentation examples ggml testing

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants