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.

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.

On gfx1151 the revert made ROCm0 reject ROCm_Host buffers: token_embd GET_ROWS ran on the CPU and all graph inputs were copied host->device per step (measured, Qwen3.8-27B server load: graph splits 1 -> 2, ROCm0 compute buffer 244 -> 500 MiB target, 672 -> 928 MiB MTP ctx, 991 -> 1079 MiB DFlash drafter). Upstream reverted over ggml-org#27311 (async ubatch input writes racing in-place device reads on UMA APUs). halo 33870a2's scheduler UMA input ring addresses that race and is only detected when the device accepts the host buffer type, so restoring the flag also re-enables the ring (default depth 2). Claude-Session: https://claude.ai/code/session_01MkuKwgPNgpAg7EbxKR2iWC

angt · 2026-09-16T05:33:44Z

I made a special build of this PR (rebased) to make testing easier, you can try with:

curl installama.sh | LLAMA_BUCKET=angt/install.sh LLAMA_VERSION=sched-uma-ring sh

ServeurpersoCom · 2026-09-16T06:58:33Z

I made a special build of this PR (rebased) to make testing easier, you can try with:
curl installama.sh | LLAMA_BUCKET=angt/install.sh LLAMA_VERSION=sched-uma-ring sh

Tested sched-uma-ring build (edf57f7) on 4x Radeon AI PRO R9700 (gfx1201, ROCm 7.13): wikitext PPL identical to master in single GPU and 4 GPU pipeline parallel (Qwen3-0.6B 17.7914, Qwen3.8-27B 6.3354), sanitizer level 2 reports 0 races, pp/tg throughput unchanged.
No repro of the original UMA hazard, as expected on dGPUs where inputs are copied to VRAM rather than computed from host buffers.

angt · 2026-09-16T08:45:38Z

Forgot to revert the revert #28604.. 😞

ServeurpersoCom · 2026-09-16T08:51:16Z

It doesn't matter if we just have to rerun the test; it's scripted:)

JohannesGaessler · 2026-09-16T11:33:19Z

What is the intent for this PR going forward? I am not at all familiar with what is being done here and even just getting up to speed would mean a significant time investment on my end. Will this PR be broken up into smaller chunks that are easier to review?

pwilkin · 2026-09-16T11:51:50Z

@JohannesGaessler this is already split as much as it can be, since it contains #26167 + necessary changes to the sanitizer to cover all spots and the rest is the required logic.

The shortest description of the problem is in #27311 (comment), this is basically trying to fix prop.integrated, which has been broken since forever and therefore has been commented out. Both CUDA and ROCm are affected by it (if we enable prop.integrated), the lazy path I proposed was to add a synchronization, but @ORippler suggested a ring buffer instead, so here we are.

JohannesGaessler · 2026-09-16T12:05:53Z

I previously did some work on the scheduler in #25319 before I went on break. Should that PR in your opinion also be merged into master and if yes in what order relative to this one?

ggerganov · 2026-09-16T12:08:45Z

This PR is not ready for review. Supporting prop.integrated is very low priority and it requires much better CI infrastructure and someone more knowledgable in low-level ggml to design how it should work.

pwilkin · 2026-09-16T12:15:47Z

@JohannesGaessler they are unrelated, in that your work deals with making copies faster, while this PR addresses the case in which there is no copy made so the actual memory is being reused.

@ggerganov but why? I mean, I discussed this with @ORippler pretty extensively (as well as the other guys from NVIDIA), I think I added a pretty detailed explanation of why this is needed and how it works. If @JohannesGaessler can't review this, then who is that "more knowledgeable person"?

ggerganov · 2026-09-16T12:43:28Z

@pwilkin It's low priority because I don't think the impact is significant and warrants such a fundamental change. I didn't say @JohannesGaessler cannot review this. I meant that I would prefer someone who is well-versed in low-level ggml and has demonstrated consistently that they are familiar with the inner workings of the code and how it should be modified/extended. Atm, I don't have this confidence for this PR. If you want to make it work, start with some smaller contributions, improve the existing logic, refactor, improve the CI, etc. Adding a bunch of interfaces to ggml out of nowhere and changing how all user code should work is very difficult to accept.

pwilkin · 2026-09-16T13:21:30Z

All right, I can understand the "trust in me that I know what I'm doing / required changes / performance gains" ratio doesn't quite add up :) I'll park it for now then and try to do some simpler stuff for the time being.

am17an · 2026-09-16T15:19:54Z

can we at least merge the sanitizer PR which catches these issues? It can be run as part of CI in different configurations using small models

pwilkin · 2026-09-16T16:27:09Z

Yeah, I'll extract my sanitizer enhancements from this PR to rebase on your sanitizer PR.

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

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

IMbackK mentioned this pull request Sep 14, 2026

ggml-cuda: avoid direct ROCm_Host compute on HIP integrated GPUs #25863

Closed

pwilkin marked this pull request as draft September 16, 2026 13:21

+                                  }
+                              }
+                              {

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.

IMbackK commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

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

Uh oh!

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

IMbackK commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

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

Uh oh!

Uh oh!

pwilkin commented Sep 8, 2026

Uh oh!

angt commented Sep 10, 2026

Uh oh!

angt commented Sep 16, 2026

Uh oh!

ServeurpersoCom commented Sep 16, 2026

Uh oh!

angt commented Sep 16, 2026

Uh oh!

ServeurpersoCom commented Sep 16, 2026

Uh oh!

JohannesGaessler commented Sep 16, 2026

Uh oh!

pwilkin commented Sep 16, 2026

Uh oh!

JohannesGaessler commented Sep 16, 2026

IMbackK commented Sep 8, 2026 •

edited

Loading

IMbackK commented Sep 8, 2026 •

edited

Loading

am17an commented Sep 16, 2026 •

edited

Loading