Scheduler UMA ring buffer (+ sanitizer and fixes) by pwilkin · Pull Request #27311 · ggml-org/llama.cpp

pwilkin · 2026-08-18T09:24:33Z

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

I have read and agree with the contributing guidelines
AI usage disclosure: Yes

pwilkin · 2026-08-18T09:36:32Z

@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 props.integrated flag can be reenabled on CUDA devices.

pwilkin · 2026-08-19T12:01:34Z

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 · 2026-08-19T14:37:40Z

@aendk Does the behavior here conform to #27258, where we try to add tests to formalize a backend's behavior?

ORippler

Thanks for taking a stab at formalizing this! I'll take a tour on DGX/RTX Spark later on and report back on perf

ORippler · 2026-08-19T14:42:02Z

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


How are these priorities determined?

Augmented the docs to include the entire algorithm.

ORippler · 2026-08-19T14:57:06Z

+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


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

ORippler · 2026-08-19T15:02:24Z

+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`.


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?

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.

ORippler · 2026-08-19T15:17:23Z

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


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

DeMaulwurfn · 2026-08-20T19:51:28Z

I tested #26225 on this box a few days ago and reported the numbers over there, so here is the same again
on its successor. Same machine, same model same method, so everything is directly comparable to what i
posted before. Also nobody seems to be covering the AMD side, the ping in your description only mentions
CUDA integrated boxes. Short version, correct, no speed loss, and the sanitizer stays quiet.

The box is unchanged from last time, a Framework Desktop with Strix Halo, gfx1151, 96 GB unified memory as
a BIOS carve out, Fedora 44 with kernel 7.1.8. Model is still Qwen3.5-122B-A10B as unsloth UD-Q4_K_XL with
-c 262144, flags are -ngl 99 -dio --no-mmap --jinja -fa on --cache-type-k q8_0 --cache-type-v q8_0. Built
from your branchs own .devops/rocm.Dockerfile with one change, AMDGPU_TARGETS=gfx1151 only because i have
exactly one card. Base image rocm/dev-ubuntu-24.04:7.2.1-complete, same as the official recipe. The commit
i tested is b16cc7d3, so after the rebase and after the ring allocation hardening. All throughput numbers
come from the servers own timings object, not from wall clock.

For correctness i use the same needle test as before. A marker KANARIE-<8 hex> goes at a defined position
into a deterministic filler prompt, it has no relation to prompt length or position so it can not be
guessed, and the model has to either echo it or answer NICHTGEFUNDEN, which is german for "not found". Six
prompt lengths times five positions, so 30 cells. Result on your branch is 30 of 30, and it stays 30 of 30
with the sanitizer switched on.

For comparison, the 13 of 30 that i reported in #26225 was measured on build 10454, wich is an ancestor of
your base here. The failing cells were all the ones where the marker sits more then about 1024 tokens away
from the end of the prompt. Nothing, that fixes this, got merged between 10454 and your base, i checked the
log, so the base still has the bug. I did not build the base seperately this time, that is the one gap in
my chain. I did look, at what your base commit actually is though, 60eeeb6 "cuda : skip UMA override for
HIP builds", and it only touches memory reporting, two preprocessor lines, so it cant be what heals this.

The sanitizer is the part i find most interesting, that one is new compared to my last report. First
attempt gave me nothing at all, not even the summary line, and that turned out to be my own fault, the
proxy i normally run in front of llama-server swallows its stderr. So i started llama-server directly
instead and got this:

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.
But zero is zero, the summary line proves the thing was actually running, wich matters because "no races"
and "sanitizer never ran" look exactly the same in a log otherwise, both empty.

Now speed. Prefill in t/s on the same machine and model, 9776 is the last build before #24233 and still
what i run in production because of the corruption.

prompt tokens	9776 (before #24233)	10454 (master, no fix)	your branch b16cc7d3
3373	274,5	380,0	363,1
26357	109,0	318,7	312,4
52592	64,2	265,3	261,6
104777	35,4	196,9	193,1

At 104777 tokens thats 193,1 against 196,9 with no fix at all, so 1,9 percent apart while my measurement
noise on prefill is also 1,9 percent. Your ring buffer keeps the whole scaling. In wall clock, reading a
100k token prompt takes 49 minutes on 9776 and 9 minutes here, so for long context work on an APU this is
the difference between usable and not. Fair warning that 10454 and your base are not the same commit, there
are a couple of weeks of master between them, so please dont read the small gaps at the short end as your
patch.

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.
The build i made for #26225, wich sits on master from end of july, does 11,67 t/s decode at that same
104777 token point. Thats about 21 percent faster then both current master and your branch. Prefill is fine
on all of them, its only decode and only at long context. So something in master between end of july and
now cost that, and your branch just inherits it. Happy to open a seperate issue with the numbers if thats
useful, i did not want to clutter this one.

Anyway from where i sit, this looks good, better then the previous approach because the sanitizer confirms
it structurally and not just by symptoms. If you want anything else run on gfx1151 just say so, the build
pipeline is allready set up so its about an hour turnaround. Test scripts are plain perl and curl if anyone
wants them.

frizikk · 2026-08-24T14:36:41Z

Linux Strix Halo ROCm validation on the exact current candidate: PASS.

parent: bb4caa7540188872173c44d161602d9271386413
head: b3823d816748eeeb6bd9859f8650b3f08ef739c5
AMD Ryzen AI MAX+ 395 / Radeon 8060S, gfx1151, 128 GB UMA
Linux 7.2.0, ROCm 7.2.4, AMD clang 22.0.0
Release HIP build with AMDGPU_TARGETS=gfx1151

Correctness:

head test-backend-sched-ring: PASS
parent test-backend-ops test -b ROCm0: 13,025/13,025 PASS
head test-backend-ops test -b ROCm0: 13,025/13,025 PASS

I also ran a controlled server differential on the exact head with Qwen3.8-27B UD-Q5_K_XL (model SHA-256 176a6a3f034e9cdc447c10cd00329fc9b31002e6589b9295f2ad4f1eefe0f6ab):

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 none

Workload: four distinct prompts submitted concurrently to four slots, three rounds (12 completions), temperature 0, fixed seed, prompt cache disabled.

exact-head mode	stable prompts across 3 rounds	sanitizer races in captured 1-round log
ring enabled (`GGML_SCHED_UMA_RING=2`)	4/4	0
ring disabled (`GGML_SCHED_UMA_RING=1`)	2/4	18

With the ring disabled, the sanitizer reported write-after-read conflicts in ROCm host memory for inp_tokens, K/V indices, the attention mask, sequence-copy input, and positions. With the ring enabled, all 12 completions were stable and no race was reported in the captured control run.

This gives a hardware-backed toggle differential for the UMA input hazard and the proposed ring-buffer fix on Strix Halo.

pwilkin · 2026-08-24T14:46:20Z

@aendk @ORippler WDYT about this one?

aendk · 2026-08-25T12:56:19Z

+                    }
+                }
+
+                {


What is the benefit of this indentation?

aendk · 2026-08-25T15:07:37Z

+                    GGML_ABORT("CUDA graph uid reused after node properties changed - see [TAG_CUDA_GRAPH_UID]");
+                }
+            }
+        }


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.

aendk · 2026-09-07T12:34:08Z

+  skip re-reading the addresses. Rotating moves tensor addresses without re-splitting, so the
+  scheduler re-stamps the split uids whenever it re-points inputs. A backend keeping such a cache
+  must key it on the uid or re-check the addresses itself. Violating this does not degrade
+  gracefully: the cached work replays against stale addresses and the results are quietly wrong.


Am I correct in assuming that rotating input addresses could lead to cache thrashing ?
I believe the cuda graph cache key is the uid.

@aendk nope:

static const void * ggml_cuda_graph_get_key(ggml_cgraph * cgraph) { return cgraph->nodes[0]; }

aendk · 2026-09-08T07:48:30Z

I tested using speedbench with llama-server on a prototype NVIDIA WoA UMA system, and I see no significant performance impact. Needed to test 6x back-to-back so that caching effects do not introduce errors.
Testing suggests that there are spurious slowdowns of 0.3%. That is within noise, I think.

@pwilkin @am17an do you consider the sanitizer merge-ready? Do you think we should merge it as part of this PR?

pwilkin · 2026-09-08T08:40:37Z

@aendk I think so, it's also gated behind a debug flag, so for most users it'll be a no-op anyway, little harm in adding it even if it's imperfect.

Would like to get it merged ASAP because as it stands ROCm on iGPU (Strix Halo, which is a large part of the userbase) is unusable without this.

am17an · 2026-09-08T08:52:44Z

Let's not rush with anything related to the scheduler. I'm not aware of llama.cpp being unusable on the strix halo, I'm running it just fine.

ggerganov · 2026-09-08T09:06:25Z

I am also not yet familiar with the problem that the PR is trying to solve. There is no need to rush - we need to understand very well how the changes work and why they are needed. Atm, I am not confident that this is the case.

am17an · 2026-09-08T09:07:33Z

At least this PR breaks pipeline parallel from what I can tell

am17an · 2026-09-08T09:10:55Z

@@ -0,0 +1,80 @@
+#include "ggml-backend.h"


this test is absolutely useless for this. Doesn't even test the ring.

pwilkin · 2026-09-08T09:14:00Z

Let's not rush with anything related to the scheduler. I'm not aware of llama.cpp being unusable on the strix halo, I'm running it just fine.

Long context on ROCm breaks on Strix Halo without this. Just try running with sanitizer = 2.

0cc4m · 2026-09-08T09:24:36Z

And Vulkan on Strix Halo?

pwilkin · 2026-09-08T09:27:49Z

No, this doesn't touch Vulkan.

0cc4m · 2026-09-08T09:39:18Z

I looked into it, the reason for that is that Vulkan does not try to use host buffers directly, it copies them to a non-host buffer even if it could in theory read it. So IMO the short-term fix is to disable that on ROCm, long term to fix the scheduler.

Assisted-by: Codex

pwilkin · 2026-09-08T09:41:43Z

@ggerganov

I am also not yet familiar with the problem that the PR is trying to solve. There is no need to rush - we need to understand very well how the changes work and why they are needed. Atm, I am not confident that this is the case.

Okay, this is the mechanism:
-> input is assigned to ROCm_Host
-> on iGPU in ROCm, the copy operation to GPU is a no-op - the memory is pinned, so the iGPU reads it directly (Vulkan does a copy, so it's unaffected)
-> the process_ubatch() code calls ggml_backend_sched_graph_compute_async(), so it can start preparing the next ubatch before the previous one is finished
-> the new code allocates buffers that the previous GPU run is still reading

The new test now reproduces this scenario directly. My first idea was just to do a sync, but @ORippler suggested to do a ring buffer instead.

@am17an I wasn't aware this was breaking PP. I'll take a look at this.

aendk · 2026-09-08T09:55:03Z

@pwilkin
I talked to @am17an and I could repro his perplexity findings.

echo tensor parallel
./bin/llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF -f ../scripts/wikitext-2-raw/wiki.test.raw -sm tensor $params
echo "baseline, single GPU"
CUDA_VISIBLE_DEVICES=0 ./bin/llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF -f ../scripts/wikitext-2-raw/wiki.test.raw $params
echo "Pipeline parallel"
./bin/llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF -f ../scripts/wikitext-2-raw/wiki.test.raw $params

yields

  tensor parallel perplexity 2.49.898.752 I Final estimate: PPL = 25.1367 +/- 1.33207
       single GPU perplexity 0.04.420.374 I Final estimate: PPL = 25.1537 +/- 1.33281
pipeline parallel perplexity 0.04.146.302 I Final estimate: PPL = 2954.0289 +/- 186.79159

on a pair of dGPUs. So pipeline parallel is definitely broken in this current PR.

pwilkin · 2026-09-08T10:05:34Z

@aendk @am17an yeah, I'm trying to figure out how to fix pipeline parallel without breaking CUDA graph replay.

am17an · 2026-09-08T10:15:59Z

@pwilkin maybe do what @0cc4m is suggesting rather than making so many changes

IMbackK · 2026-09-08T10:17:25Z

I looked into it, the reason for that is that Vulkan does not try to use host buffers directly, it copies them to a non-host buffer even if it could in theory read it. So IMO the short-term fix is to disable that on ROCm, long term to fix the scheduler.

ie #25863

pwilkin · 2026-09-08T10:19:21Z

@am17an yeah, but that's a performance regression.

IMbackK · 2026-09-08T10:21:06Z

@pwilkin i think the first problem here is that there is no consensus on how to fix the underlying issue, so we just need the stop gap for now

pwilkin · 2026-09-08T10:26:49Z

All right, so we merge the PR you linked?

IMbackK · 2026-09-08T10:28:21Z

might be better to just revert the pr that enabled prop.integrated on hip again.

pwilkin · 2026-09-08T10:29:14Z

Huh, or that :)

pwilkin · 2026-09-08T10:30:33Z

There's a whole discussion in #25863 about this and the only reason I started doing the ring buffer was because @ORippler wanted a proper general solution so we could also enable prop.integrated on CUDA devices that support it. We can revert prop.integrated while we work on this.

IMbackK · 2026-09-08T10:31:34Z

yes lets revert c7d8722

"AMD engineers [...] very likely know better than me." bold assumption

Assisted-by: Codex

pwilkin · 2026-09-08T14:23:01Z

All right, the fix for pipeline parallel is in, I've tested it on both my 5060 and R9700 and there's no regression in graph reuse.

angt · 2026-09-10T15:13:10Z

I've tested this branch with prop.integrated = true re-enabled in HIP on an AMD Ryzen AI system (gfx1151, ROCm 7.2.1), and it works correctly. The output is identical to the output with HIP_LAUNCH_BLOCKING=1.

pwilkin requested review from a team, CISC and ggerganov as code owners August 18, 2026 09:24

github-actions Bot added documentation Improvements or additions to documentation ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Aug 18, 2026

pwilkin requested review from JohannesGaessler, ORippler and am17an August 18, 2026 09:35

ORippler reviewed Aug 19, 2026

View reviewed changes

ORippler mentioned this pull request Aug 19, 2026

ggml: fix backend split scheduler race condition #26040

Merged

aendk reviewed Aug 20, 2026

View reviewed changes

Comment thread ggml/src/ggml-cuda/ggml-cuda.cu

github-actions Bot added the testing Everything test related label Aug 20, 2026

pwilkin force-pushed the sched-uma-ring branch from f046c3e to faa3fdf Compare August 20, 2026 13:04

pwilkin force-pushed the sched-uma-ring branch from b16cc7d to b3823d8 Compare August 21, 2026 13:35

aendk reviewed Aug 25, 2026

View reviewed changes

Comment thread ggml/src/ggml-backend.cpp

}

}

{

aendk Aug 25, 2026

Copy link
Copy Markdown

Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the benefit of this indentation?

aendk reviewed Aug 25, 2026

View reviewed changes

This comment was marked as off-topic.

Sign in to view

DeMaulwurfn mentioned this pull request Aug 25, 2026

HIP backend silently corrupts Qwen3.5-27B (Gated DeltaNet) inference on gfx1151 - oldest context lost; Vulkan correct at identical commit #27556

Open

This comment was marked as duplicate.

Sign in to view

aendk reviewed Sep 7, 2026

View reviewed changes

am17an reviewed Sep 8, 2026

View reviewed changes

ggml: test asynchronous shared-input ring reuse

a51e390

Assisted-by: Codex

IMbackK mentioned this pull request Sep 8, 2026

Revert "ggml-cuda : restore prop.integrated on HIP builds (#24233)" #28604

Merged

ggml: fix pipeline input rotation and preserve graph replay

dced2ff

Assisted-by: Codex

Error00101 mentioned this pull request Sep 8, 2026

Eval bug: HIP/ROCm on gfx1151, wrong logits (not a crash), triggered by prompts longer than n_ubatch. #28211

Open

accemlcc mentioned this pull request Sep 11, 2026

2-node Qwen3.8-Flash-Next (180B) on 2× Strix Halo — recipe kyuz0/amd-strix-halo-vllm-toolboxes#88

Open

+                                  }
+                              }
+                              {

Uh oh!

Conversation

pwilkin commented Aug 18, 2026

Overview

Additional information

Requirements

Uh oh!

pwilkin commented Aug 18, 2026

Uh oh!

pwilkin commented Aug 19, 2026

Uh oh!

ORippler commented Aug 19, 2026

Uh oh!

ORippler left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

DeMaulwurfn commented Aug 20, 2026

Uh oh!

frizikk commented Aug 24, 2026

Uh oh!

pwilkin commented Aug 24, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

This comment was marked as off-topic.

This comment was marked as duplicate.

This comment was marked as duplicate.

aendk Sep 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

aendk commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

am17an commented Sep 8, 2026

Uh oh!

ggerganov commented Sep 8, 2026

Uh oh!

am17an commented Sep 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

0cc4m commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

0cc4m commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

aendk commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

am17an commented Sep 8, 2026

aendk Sep 7, 2026 •

edited

Loading

am17an commented Sep 8, 2026 •

edited

Loading

IMbackK commented Sep 8, 2026 •

edited

Loading

IMbackK commented Sep 8, 2026 •

edited

Loading