Scheduler UMA ring buffer (+ sanitizer and fixes) - #27311
Conversation
github-actions
Bot
added
documentation
|
@ggerganov this should enable proper host memory usage on UMA devices other than Metal (i.e. ones that actually pin host memory), once this is verified the |
|
As suggested by @am17an I've asked Claude to add a post-mortem doc about my chases of the various bugs during this PR here: https://pwilkin.github.io/llama-scheduler/ring.html |
ORippler
left a comment
ORippler
left a comment
There was a problem hiding this comment.
Thanks for taking a stab at formalizing this! I'll take a tour on DGX/RTX Spark later on and report back on perf
| 2. **Op support and operand location.** Otherwise the highest priority backend that supports the | ||
| op is used, preferring the backend holding the operands. Ops reading tensors in a buffer marked | ||
| `GGML_BACKEND_BUFFER_USAGE_WEIGHTS` prefer that buffer's backend, so that weights are not | ||
| copied. |
There was a problem hiding this comment.
How are these priorities determined?
There was a problem hiding this comment.
Augmented the docs to include the entire algorithm.
| says the memory is dead - so a later split on the owning backend can be given the same memory and | ||
| overwrite it. | ||
|
|
||
| Such tensors are pinned for the lifetime of the graph with `ggml_gallocr_pin_tensor()`, keeping |
There was a problem hiding this comment.
For my understanding. We are talking about pinning the memory on backend A for a node later on consumed by backend B. This will hold also for the path A -> C -> A -> B (which is why ggml_set_output is insufficient, as it garuantuees only within single cgraph execution).
| That move is only made when the target backend supports the op. Support can be conditional on | ||
| the tensor types - CUDA runs `GGML_OP_SET` only for F32 and I32 - so the backend owning the | ||
| aliased memory is not guaranteed to be able to run the op writing into it. There is no correct | ||
| placement in that case: the scheduler copies operands into a split, never results out of one, so | ||
| whichever backend runs the op, the write cannot reach the aliased memory. The node is left where | ||
| the earlier passes put it, which is what the scheduler did before this rule existed, and the | ||
| reason is logged under `GGML_SCHED_DEBUG`. |
There was a problem hiding this comment.
so the backend owning the
aliased memory is not guaranteed to be able to run the op writing into it. There is no correct
placement in that case: the scheduler copies operands into a split, never results out of one, so
whichever backend runs the op, the write cannot reach the aliased memory.
Wouldn't the correct behavior be to check for this during node-placement + expansion time?
There was a problem hiding this comment.
It might be, but I didn't want to do a full scheduler refactor for this. One is probably due anyway since the work in this PR outlined quite a few issues with the scheduler and quite a few of the current solutions (esp. regarding the "special" input / output tensors) seem really hacky.
| It maintains a vector clock per actor - the host thread and each backend - and a shadow map of | ||
| which actor last read or wrote every byte of every buffer. Synchronization points (backend | ||
| synchronize, event record, event wait, event synchronize, async copies) advance those clocks. When | ||
| an access conflicts with a previous one and no happens-before edge connects them, it reports: |
There was a problem hiding this comment.
Can we make this respect inter-backend events? AFAIK it's not current POR, but CUDA offers cross-stream events like such
CUDA_CHECK(cudaEventRecord(cuda_ctx_src->copy_event, cuda_ctx_src->stream()));
// wait on dst stream for the copy to complete
CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx_dst->stream(), cuda_ctx_src->copy_event, 0));
f046c3e to
faa3fdf
Compare
|
I tested #26225 on this box a few days ago and reported the numbers over there, so here is the same again The box is unchanged from last time, a Framework Desktop with Strix Halo, gfx1151, 96 GB unified memory as For correctness i use the same needle test as before. A marker KANARIE-<8 hex> goes at a defined position For comparison, the 13 of 30 that i reported in #26225 was measured on build 10454, wich is an ancestor of The sanitizer is the part i find most interesting, that one is new compared to my last report. First ggml-sched-sanitize: 0 race(s) reported That was over four cells at 2000 and 4000 tokens, not the full grid, so its four requests and not thirty. Now speed. Prefill in t/s on the same machine and model, 9776 is the last build before #24233 and still
At 104777 tokens thats 193,1 against 196,9 with no fix at all, so 1,9 percent apart while my measurement Decode is unchanged as well, 9,62 t/s at 104777 tokens against 9,61 on unpatched master One thing that is not about your PR at all, but you might want to know since it showed up in the same runs. Anyway from where i sit, this looks good, better then the previous approach because the sanitizer confirms |
b16cc7d to
b3823d8
Compare
|
Linux Strix Halo ROCm validation on the exact current candidate: PASS.
Correctness:
I also ran a controlled server differential on the exact head with Qwen3.8-27B UD-Q5_K_XL (model SHA-256 GGML_SCHED_SANITIZE=1 GGML_SCHED_SANITIZE_NONFATAL=1 \
GGML_SCHED_UMA_RING=<2 or 1> \
llama-server -m model.gguf -ngl 999 -fa on -c 16384 \
-np 4 --kv-unified -b 2048 -ub 2048 -lm noneWorkload: four distinct prompts submitted concurrently to four slots, three rounds (12 completions), temperature 0, fixed seed, prompt cache disabled.
With the ring disabled, the sanitizer reported write-after-read conflicts in ROCm host memory for This gives a hardware-backed toggle differential for the UMA input hazard and the proposed ring-buffer fix on Strix Halo. |
| } | ||
| } | ||
|
|
||
| { |
There was a problem hiding this comment.
What is the benefit of this indentation?
| GGML_ABORT("CUDA graph uid reused after node properties changed - see [TAG_CUDA_GRAPH_UID]"); | ||
| } | ||
| } | ||
| } |
There was a problem hiding this comment.
Is this addition debug code, or should similar things be rolled out to all other UMA backends?
Also, my suggestion would be to move this specific code to ggml_cuda_graph_verify_uid() itself, for better structure / smaller general functions.
Independent validation on gfx1151 (Strix Halo) — fixes silent output corruption, no measurable throughput costI tested this PR on a Ryzen AI MAX+ 395 (gfx1151 / Radeon 8060S, 128 GB unified memory). It fixes silent numerical corruption in the HIP backend that affects both perplexity and ordinary generation. Throughput is unchanged. All comparisons are single-variable: Environment
Symptom 1 — perplexity is garbage whenever
|
| build | -c |
-b |
-ub |
chunks | PPL |
|---|---|---|---|---|---|
| bb4caa7 (base) | 512 | 2048 | 512 | 100 | 872,006.5 |
| bb4caa7 (base) | 4096 | 4096 | 512 | 20 | 1.336e10 |
| pr27311 | 512 | 2048 | 512 | 100 | 63.59 |
| pr27311 | 4096 | 4096 | 512 | 20 | 50.28 |
| Vulkan (reference) | 512 | 2048 | 512 | 100 | 64.10 |
| Vulkan (reference) | 4096 | 4096 | 512 | 20 | 50.66 |
The PR's values land on the Vulkan reference. The base's do not.
The trigger is the batch split, not small ubatches. On the base build, holding -b at 4096 and varying -ub (chunks=8):
-c |
-b |
-ub |
PPL | |
|---|---|---|---|---|
| 4096 | 4096 | 4096 | 56.55 | ok |
| 4096 | 4096 | 2048 | 757.85 | corrupt |
| 4096 | 4096 | 1024 | NaN | corrupt |
| 4096 | 4096 | 512 | 6.97e9 | corrupt |
and with ub == b, varying the size:
-c |
-b |
-ub |
PPL | |
|---|---|---|---|---|
| 512 | 512 | 512 | 73.30 | ok |
| 1024 | 1024 | 1024 | 84.52 | ok |
| 2048 | 2048 | 2048 | 43.75 | ok |
| 512 | 2048 | 512 | 816,259 | corrupt |
-ub 512 is fine when -b 512. 8/8 on the rule ub < b ⇒ corrupt.
Not model-specific. -c 2048 -b 2048, chunks=8, base build:
| model | Vulkan ub512 |
HIP ub512 |
HIP ub2048 |
|---|---|---|---|
| Llama-3.1-8B-Instruct Q4_K_M | 5.89 | 3,703.86 | 5.87 |
| Nemotron-3-Nano-4B Q4_K_M | 9.81 | 947.09 | 9.77 |
| phi-4 (14B) Q4_K_M | 4.88 | 5,315.35 | 4.86 |
| gpt-oss-20b MXFP4 | 643.56 | 1,881.46 | 641.75 |
| Qwen3-Coder-30B-A3B Q4_K_M | 6.85 | 68,548.58 | 6.85 |
| gemma-4-26B-A4B UD-Q4_K_M | 5,994.41 | 135,285.65 | 5,792.47 |
(The last two rows have high baselines on both backends — an unrelated issue I'm looking into separately. The HIP/Vulkan divergence is what matters here.)
Symptom 2 — ordinary generation is degenerate once the prompt exceeds n_ubatch
This is the same root cause but is worth calling out because llama-perplexity does not catch it, and it is what users actually hit: any RAG or long-context request silently returns nonsense.
llama-cli, Llama-3.1-8B-Instruct-Q4_K_M, --temp 0 --seed 1, ~3,500-token prompt (a wikitext passage) ending in QUESTION: In one sentence, who is Robert Boulter?\nANSWER:
| build | -b |
-ub |
output |
|---|---|---|---|
| bb4caa7 | 4096 | 512 | The of the best the of the of the of the of the of the… |
| bb4caa7 | 2048 | 2048 | The , the , the , and also , the , and also , the… |
| pr27311 | 4096 | 512 | Robert Boulter is an English film, television, and theatre actor. |
| pr27311 | 2048 | 2048 | Robert Boulter is an English film, television, and theatre actor. |
| Vulkan (reference) | 2048 | 512 | Robert Boulter is an English film, television, and theatre actor. |
Here the trigger is prompt_tokens > n_ubatch. On the base build:
| prompt | -b |
-ub |
result |
|---|---|---|---|
| ~1.5k tok | 2048 | 512 | corrupt |
| ~1.5k tok | 1024 | 1024 | corrupt |
| ~1.5k tok | 1024 | 512 | corrupt |
| ~3.5k tok | 4096 | 512 | corrupt |
| ~3.5k tok | 4096 | 4096 | ok |
| ~3.5k tok | 2048 | 2048 | corrupt |
| ~3.5k tok | 8192 | 512 | corrupt |
7/7 — the only passing configuration is the one where the whole prompt fits in a single ubatch. Flash attention on or off makes no difference.
Confirming the root cause directly
To check this is the integrated path and not something else, I made the flag runtime-switchable in ggml_cuda_init() on an otherwise unmodified tree and flipped it on one binary (patch verified present in the built libggml-hip.so):
| test | integrated ON |
integrated OFF |
|---|---|---|
PPL, -c 512 -b 2048 -ub 512 |
850,121 | 72.80 |
gen, ~3.5k prompt, b4096 ub512 |
corrupt | correct |
gen, ~3.5k prompt, b2048 ub2048 |
corrupt | correct |
gen, ~1.5k prompt, b2048 ub512 |
corrupt | correct |
Both symptoms have the single root cause this PR addresses.
Throughput cost: none
llama-bench, gemma-4-E4B-it-Q4_K_M, -ngl 999 -fa 1 -r 3, three independent runs per config; mean and (max−min)/mean:
| pp512 | pp4096 | tg128 | |
|---|---|---|---|
| bb4caa7 (base) | 1980.25 ±0.68% | 1742.21 ±0.38% | 55.17 ±0.06% |
| pr27311 | 1976.79 ±0.94% | 1728.65 ±2.10% | 55.13 ±0.01% |
| delta | −0.17% | −0.78% | −0.07% |
All deltas are inside the run-to-run spread; pp4096 varies 2–3% between identical runs on this machine, so single measurements are not informative at this granularity. Consistent with the 97.7% figure reported earlier in this thread.
For reference I also applied #25863's ggml-cuda.cu diff onto this same base and measured it identically: pp512 1957.75 ±0.89%, pp4096 1722.25 ±3.00%, tg128 55.05 ±0.18%. The two PRs are indistinguishable on throughput here.
Rebased onto current master
The branch is 63 commits behind master, so I rebased it to check it still applies and still works. All 17 commits replay onto eab8ee41f (master as of 2026-08-25) with no conflicts, and the diffstat is unchanged (1348 insertions / 18 files).
On that rebase:
- Builds clean for HIP/gfx1151.
- The PR's own new tests pass:
test-backend-sched-ring(rc=0) andtest-alloc(test_pinned_view_root_no_inplace PASSED). - Master itself is still affected: PPL 986,205 at
-c 512 -b 2048 -ub 512, and the long-prompt generation case still returnsThe of the best the of the of the of.... - Rebased PR: PPL 79.64 on the same config, generation correct.
- Throughput vs master, 3 runs each:
| pp512 | pp4096 | tg128 | |
|---|---|---|---|
master eab8ee41f |
1972.67 ±0.58% | 1733.65 ±2.20% | 55.19 ±0.15% |
| this PR, rebased | 1971.46 ±0.64% | 1745.74 ±1.23% | 55.16 ±0.23% |
No measurable cost. Happy to share the rebased branch if that saves you the merge.
Comparison with #25863
I also built and tested #25863 (ce82541ac) on the same hardware. It fixes both symptoms as well:
| test | #25863 | #27311 |
|---|---|---|
PPL c512 b2048 ub512 |
63.67 | 63.59 |
PPL c4096 b4096 ub512 |
50.17 | 50.28 |
gen, long prompt b4096 ub512 |
correct | correct |
gen, long prompt b2048 ub2048 |
correct | correct |
I have no correctness evidence favouring one over the other on this hardware. Either resolves the corruption; the choice is whether you want the targeted mitigation or the general scheduler fix.
Reproducing
cmake -S llama.cpp -B build-rocm \
-DCMAKE_BUILD_TYPE=Release -DGGML_HIP=ON -DGGML_NATIVE=ON -DGGML_CUDA_FA=ON \
-DAMDGPU_TARGETS=gfx1151 -DCMAKE_HIP_ARCHITECTURES=gfx1151
cmake --build build-rocm -j"$(nproc)" --target llama-perplexity llama-cli
# symptom 1 - corrupt on base, sane on this PR
./build-rocm/bin/llama-perplexity -m gemma-4-E4B-it-Q4_K_M.gguf \
-f wiki.test.raw -ngl 999 -fa 1 -c 512 -b 2048 -ub 512 --chunks 100
# symptom 2 - needs a prompt longer than n_ubatch, and needs generation to show up
./build-rocm/bin/llama-cli -m Meta-Llama-3.1-8B-Instruct-Q4_K_M.gguf \
-ngl 999 -fa 1 -c 8192 -b 4096 -ub 512 --temp 0 --seed 1 -n 30 -st -f long_prompt.txtHappy to run additional configurations on this hardware if it would help move the PR along.
|
@nabe2030 asked whether the ring buffer fixes the MTP collapse from #27572, so I ran the comparison directly on the affected hardware before anything else got built on top of it. Three builds were tested under identical conditions on AMD Strix Halo (gfx1151, 128GB unified, ROCm 7.2.2, Qwen3.8-27B-UD-Q4_K_XL + mmproj-F16): variant A is current master Concurrent reproducer — the result that matters:
All three variants pass the sequential test, confirming the race requires concurrency as described in #27572. The stock-master round-1 failures produced completions of 7, 33, and 5 tokens with empty content after short reasoning — the exact symptom from the original report. Throughput parity (temp 0, 3-run median):
The stock-master benchmark is omitted since the server is racing and the numbers are meaningless. A and B are within 1–4% on prefill and 0–1% on generation, which is within run-to-run noise on this rig — the ring buffer costs nothing measurable on a single-GPU spec-decode workload. Conclusion: this PR resolves #27572 in addition to the multi-slot corruption and the #27579 tool-array corruption that @nabe2030 already verified against it, making it the fourth issue closed by the input ring buffer. The sched-synchronize variant remains a valid minimal backport for release branches, but the ring buffer is the better fix architecturally since it removes the host stall entirely rather than extending it to single-GPU setups, and it enables host/GPU overlap that the synchronize approach structurally cannot provide. The b10581-era sequential regression (acceptance 0.0 even at |
|
Results from gfx1151 on this PR, focused on the cells that were not yet covered here. Environment: AMD Ryzen AI Max+ 395 / Radeon 8060S (gfx1151), 128 GB unified, Two caveats up front. "ring off" below means 1. Sanitizer race counts under realistic load
@frizikk reported 18 races over 4 prompts × 3 rounds earlier in this thread. Different 2. Multi-slot output corruption, concurrency sweep (MTP off)Generated text captured and checked, n=20 per cell:
Ring off reproduces the known signature (repeated 3. Tool-array corruption (#27579), single slot29,355-token prompt, 17 tools,
The stock-master row is the control for the caveat above: ring off and stock master sit at This failure mode is worth separating from the degenerate-text one. The output stays 4. MTP acceptance (#27572) — independent confirmation@ByungHyun21 already settled this above with a cleaner three-way comparison. Ours agrees,
5. vLLM control on the same siliconWorth stating since "gfx1151 is just broken" comes up whenever these reports appear: on One environment note that may be useful to others reporting here: our "ROCm 7.14" builds @ByungHyun21 referred to the multi-slot and tool-array results as already verified; the |
Two problems, both reached by test-opt, which builds a scheduler with the backend under test in front of the full backend list - so for the CPU device that is a CPU backend in front of the CPU backend. The detection asked whether any backend other than the last accepts the last backend's buffer type. A CPU backend trivially accepts CPU memory, so a CPU-only scheduler was treated as a device computing on host memory and given a ring it has no use for. Skip CPU devices: the ring exists for a non-CPU device reading memory the host writes. That alone was a segfault rather than wasted memory. The tensor aliased by the ring sat at slot cur_copy, so which tensor occupied a given leaf index changed from one allocation to the next. ggml_gallocr_needs_realloc() compares node and leaf counts and sizes but never identity, so it saw no reason to reallocate, and a leaf_alloc recorded for a pre-allocated tensor (buffer_id -1) was applied to a placeholder that still needed allocating, indexing galloc->buffers out of bounds. Keep the aliased tensor at slot 0 so leaf identity is stable, and rotate purely by re-pointing the inputs, which both paths now do explicitly. This is what the graph reuse path already relied on; the allocate path was getting rotation as a side effect of where split_graph placed the alias. Also stop GGML_SCHED_UMA_RING from switching the ring on where it was not detected. It is an escape hatch for a scheduler that has one, not a way to impose one on a caller that does not meet its contract - ggml_opt keeps its inputs across separate allocations and does not. Sanitizer race reports are unchanged: 94/0 with the ring off and on, 0 at partial offload, reuse path still rotating, output correct at full and partial offload. Assisted-By: Claude Opus 5
The detection excluded CPU devices, but BLAS is an accelerator device that computes on host memory, so a BLAS plus CPU scheduler was still given a ring. ggml_opt keeps its inputs across separate allocations and does not meet the ring's contract, so test-opt went to 73/118 with it enabled. Excluding device types one at a time is the wrong shape. The hazard is a reader that is still reading after graph_compute returns, so require the backend to be asynchronous. That covers CPU and BLAS together, and anything else synchronous: their work is finished before control comes back, and nothing of theirs can be reading the inputs the host is about to overwrite. Verified against a BLAS build of test-opt, which is the configuration that failed: 118/118, ring never enabled. The ring is still enabled for ROCm and its race reports are unchanged at 94/0 with it off and on, 0 at partial offload. Assisted-By: Claude Opus 5
Pass 4 assigns a node that aliases its source to the backend owning that source, so the write reaches the aliased memory instead of a discarded copy. It did so unconditionally. Support for these ops can be conditional on the tensor types - CUDA runs GGML_OP_SET only for F32 and I32, and OpenCL has no GGML_OP_ACC at all - so the backend that owns the aliased memory is not guaranteed to be able to run the op writing into it. Forcing the node there hands the backend an op it rejects, which aborts in ggml_backend_graph_compute. There is no placement that computes such a graph correctly: operands are copied into a split, results are never copied out, so the write cannot reach the aliased memory from any other backend. Check support before moving, and leave the node where the earlier passes put it otherwise - the behaviour before this rule existed. The declined move is logged under GGML_SCHED_DEBUG. Measured on Qwen3.5-4B at -ngl 16, the case the rule was added for: 13 moves before and after, 0 declined, output unchanged. Races still 0 with the ring on and 1 with it off, and test-opt passes 46/46. Assisted-By: Claude Opus 5
Assisted-by: Codex
Assisted-by: Codex
Assisted-by: Codex
2cdd638 to
47744e5
Compare
pwilkin
removed
Vulkan
|
Sheesh... sorry, bad rebase. |
11 participants
Overview
As per discussion in #25863 , implement the ring buffer mechanism for input tensors, on top of @am17an 's sanitizer (#26167), plus additional hardening for the sanitizer and extra scheduler fixes (there was an error with duplicating pinned memory that was a view's base).
Additional information
Makes host buffers viable again, ping @ORippler for feedback / tests on CUDA integrated boxes. No measurable efficiency losses.
Supersedes #25863 , #26167 , #26225
Moved all scheduler documentation to a dedicated doc.
Requirements