ggml : --prefetch-experts-slots — lookahead H2D prefetch of host-resident MoE experts

[leshchukandrej]

--prefetch-experts-slots: lookahead H2D prefetch of host-resident MoE expert weights

Branch: port/prefetch-experts-slots (head: leshchukandrej/beellama.cpp fork) — commit on top of current master
Scope: 8 files, +327 / −0, single feature
Flags: --prefetch-experts-slots N (CLI) / llama_context_params.prefetch_experts_slots / ggml_backend_sched_set_prefetch_experts_slots()

Problem

When MoE expert weights are not GPU-resident (--n-cpu-moe N / --cpu-moe, or
host pages when the model does not fully fit VRAM), the scheduler uploads each
used expert tensor host→device right before the MUL_MAT_ID split that consumes
it. In the common single-context case (n_copies == 1) that H2D transfer is not
overlapped: each split boundary synchronizes, so large prefill batches wait on
PCIe/NVLink instead of computing. Long-prompt prefill TTFT is dominated by this
serial upload cost.

What this does

During prefill, expert weight uploads are issued one split ahead of need
(1-deep lookahead) through a second backend instance on the same device
(a separate CUDA stream), into rotating staging buffers. The consuming split
performs a per-split cross-stream event wait that is already satisfied by launch
time — so the H2D bytes arrive during compute, not after it.

split i        :  [ MUL_MAT_ID compute ]        uploads for split i+2 fired here
prefetch stream:     |====== H2D expert(i+2) → staging slot =====|
split i+2      :  [ wait ready[i+2] (no-op) ] [ MUL_MAT_ID compute ]

Design

• slots == 0 (default): feature fully off — no state, no allocations, the
  scheduler behaves exactly as before.
• slots >= 2: pipeline on. Staging memory = slots * max_expert_tensor,
  lazily allocated on first use. 3 is recommended; capped at 4.
• Fires only for splits whose first node is GGML_OP_MUL_MAT_ID with
  host-resident weight inputs (GGML_BACKEND_BUFFER_USAGE_WEIGHTS, is_host),
  and only at prefill scale (ids->ne[0]*ids->ne[1] >= 2*n_expert). At
  decode-scale batches the stock per-split copy of used experts is kept —
  routing ids carry information there, and decode is unaffected by construction
  (callback_eval mode never fires).
• Requires device async + events caps (CUDA); on any allocation/cap failure
  or with an eval callback installed the feature disables itself and the regular
  copy path is used. Correctness never depends on prefetch.
• Safety: the input-copy tensor is only re-pointed at the staging slot for the
  duration of its split and restored right after graph launch, and the
  ready-event wait orders the copy before the kernels. Gated per-split (not
  per-graph) so tool_choice semantics are preserved.

Lossless

Prefetch changes when bytes land on device, never what is computed: the same
host weights are copied, the consuming graph/kernels are unchanged, and the
event wait guarantees completion before launch. slots = 0 leaves every code
path untouched.

Empirically (temp 0, same prompt/seed, --n-cpu-moe 20 on a 24B A3B, q4_0 KV):

prefetch	prompt	TTFT	output
OFF	2762 tok	3501 ms	682 chars
ON (slots 3)	2762 tok	3786 ms	682 chars
byte-identical:			True

Measured effect

Large-prefill TTFT is the target; decode is flat (untouched by design).
Benchmarks below measured on a downstream (beellama v0.4.5) build carrying this
byte-identical code, on an RTX 5070 Ti (12 GB) with host-expert configs:

config	prompt	TTFT off	TTFT on	recall/output
24B A3B, -ncmoe 20, q4_0 KV	~42k tok	15.71 s	13.97 s (−11%)	identical both
24B A3B, -ncmoe 20, q4_0 KV	200 tok	0.51 s	0.41 s	identical both
21.8 GB 35B A3B Q4_K_M auto-offload	~42k tok	21.19 s	16.46 s (−22%)	identical both

Notes: with -ncmoe 20 ~20 layers' experts are host-resident; the 35B row used
no -ncmoe — auto-offload alone leaves half the experts host-resident, the
realistic deployment target. The feature is off by default, so no-one pays the
(fixed, ~prefill-only) staging cost unless they opt in.

Files

file	change
common/arg.cpp	flag
common/common.{h,cpp}	common_params plumbing
include/llama.h	llama_context_params field
src/llama-cparams.h	cparams field (default 0)
src/llama-context.cpp	wire into sched_reserve()
ggml/include/ggml-backend.h, ggml/src/ggml-backend.cpp	scheduler prefetch state machine

[ggml-gh-bot]

Hi @leshchukandrej, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[pwilkin]

/bot review

[ggml-gh-bot]

Automated code review

Review of PR #28414 (--prefetch-experts-slots). Checklists run: scope/quick-reject gate, security, approach/design (new scheduler subsystem in ggml/), ggml/backend, public API, general conventions.

Blocking

(point 1) Staging state is per-scheduler but devices are per-split: multi-GPU layer-split is silently broken. prefetch_backend, the slot buffers, and the events are created once, from the first split backend that fires (ggml_backend_sched_prefetch_init, ggml/src/ggml-backend.cpp:1696-1728). With --split-mode layer + 2 GPUs and host-resident experts, MoE splits on GPU1 will fire through the GPU0 prefetch backend, with slot buffers on GPU0 (ggml_backend_get_default_buffer_type(split_backend) is only used at first init). The consuming split on GPU1 then launches compute reading a GPU0 pointer with no peer access enabled - silent wrong results, not an error. The feature needs a device guard: only fire when the split backend's device matches the prefetch device (or keep per-device prefetch state), and otherwise fall back to the regular copy path.

(point 2) Disable/realloc while a lookahead fire is outstanding leaves dangling slot pointers. Once try_fire_prefetch repoints input_cpy->buffer/data at a slot for a future split, two things can free that slot before the future split consumes it: (a) ggml_backend_sched_prefetch_disable called from a failed inline prefetch_init (OOM or cap failure on a later split), and (b) the slot-realloc branch in prefetch_init that frees an old, too-small buffer. The consuming split's block (ggml/src/ggml-backend.cpp:1846-1860) only checks lookahead[split_id].slot != -1 - it never re-checks sched->prefetch_experts or that the slot buffer is alive - so it skips the regular input copy, waits an event recorded before the free, and launches compute reading freed device memory. Also, prefetch_disable synchronizes only the one split_backend passed in, not other backends whose kernels may still be reading slots. At minimum, consumption must verify the prefetch is still armed, restore the saved buffer/data, and fall back to the normal copy; ideally disable should not be reachable while fires are outstanding.

(point 3) Missing ggml_backend_buffer_is_host() check in both fire paths. The PR description says the feature engages only for "host-resident" weights, and the existing used-experts copy path it sits next to requires ggml_backend_buffer_is_host(input->buffer) (ggml/src/ggml-backend.cpp:1953), but try_fire_prefetch and the inline block check only GGML_BACKEND_BUFFER_USAGE_WEIGHTS. A WEIGHTS-usage, non-host input would be handed to ggml_backend_cuda_set_tensor_async as an cudaMemcpyHostToDevice source - invalid for a device pointer. Add the is_host check to match the description and the sibling code.

(point 4) Comment content and non-ASCII characters. The repeated "mindcontrol-port" / "mindcontrol prefetch-wait A/B verdict" / "the only mode that preserves tool_choice semantics" / "tool_calls" references (also in the public headers ggml/include/ggml-backend.h:357 and include/llama.h:410) are meaningless outside the source project and read like leftover port notes; per AGENTS.md, comments must be concise, self-contained, and ASCII-only. Concretely: ggml/src/ggml-backend.cpp:1927 and :2057 contain em-dashes (—), which are explicitly prohibited. All "mindcontrol"/"tool_choice" references must be rewritten as plain descriptions of what the code does and why.

(point 5) Scope gate: this is a new scheduler subsystem with no prior discussion. This adds ~300 lines of cross-stream pipelining machinery to the core scheduler, plus new public surface (ggml_backend_sched_set_prefetch_experts_slots, a llama_context_params field, a CLI flag). Per CONTRIBUTING.md, a change of this size should start with an issue/discussion before the PR; the PR description references none. Please link or open one.

Will slow the review

(point 6) Dead debug machinery should be removed. prefetch_wait_mode >= 2 is unreachable (the setter only ever sets 0 or 1), so pending_prefetch_slots, last_prefetch_split_backend, and the end-of-graph sync block (ggml/src/ggml-backend.cpp:2107-2113) are dead code. Same for the inline prefetch block (:1897-1934): with LOOKAHEAD=1, priming + per-split refire covers every eligible split, and the inline block's conditions are identical to try_fire_prefetch's, so it cannot fire in practice. Drop both (or unify with try_fire_prefetch), and hardcode the lookahead/mode values instead of storing them as struct fields. This alone removes ~80 lines and several struct members.

(point 7) Unconditional per-step heap allocation. std::vector<prefetch_pending> lookahead(sched->n_splits) (ggml/src/ggml-backend.cpp:1799) is constructed on every compute_splits call - including plain decode with the feature off (the default), adding a malloc/free per token for all users. Gate all of this scratch state behind sched->prefetch_experts.

(point 8) Leak when the slot count is reduced. In the i >= 2 fallback in prefetch_init (ggml/src/ggml-backend.cpp:1739-1744), prefetch_n_slots is lowered to i, but sched->prefetch_slots[i] (the old, too-small buffer), plus the events and buffers for indices >= the new count, are then skipped by the ggml_backend_sched_free loop (:2221-2227, iterates i < prefetch_n_slots). Free the leftovers when shrinking.

(point 9) Per-fire rescan. ggml_backend_sched_prefetch_init calls ggml_backend_sched_sched_prefetch_max_size (O(splits x inputs)) on every fire. Compute the max once per compute_splits call (or cache it) instead.

(point 10) Benchmark data is confusing. The "Lossless" table shows ON being slower (3501 ms off vs 3786 ms on) while the "Measured effect" table shows improvements, and all numbers come from a downstream fork build ("beellama v0.4.5"), not this tree. Maintainers will ask for vanilla-build numbers, so having them ready (and clarifying the first table) will save a round trip.

Nits

(point 11) std::stoi in the new arg handler is fine (the parser catches std::invalid_argument), but the help string is very long; consider trimming it and putting detail in the PR/docs.

(point 12) LOOKAHEAD as an uppercase local is unusual for the surrounding style; the rest of the function uses lowercase locals.

The overall design (second backend instance + rotating staging slots + ready/free event pairs) is a reasonable shape for the problem, and the off-by-default gating plus the per-split event wait are the right instincts. But the multi-GPU staging-device issue (point 1) and the disable-while-outstanding UAF (point 2) are correctness bugs that must be fixed, and points 4/6 alone will roughly halve the diff.

This review was generated automatically by pi coding agent using zai-org/GLM-5.3. It may contain mistakes. Maintainers make the final call.

[XBold]

Cool approach. The 1-deep lookahead with a second CUDA stream is a clean way to hide PCIe latency without changing the graph topology. Keeping it prefill-only means decode performance is guaranteed flat.

A couple of things I'm curious about:

1. Why cap at 4 slots? For very long prefill batches (your ~42k tok test showed 22% improvement), the expert routing set is broader and the gap between compute and transfer is largest. Did you test with 8+ slots and find a point of diminishing returns, or is it a conservative cap to keep staging memory bounded?

2. Self-disable on allocation failure is smart for correctness, but in practice on a card like the 5070 Ti (12 GB), what's the failure mode? Is it not enough VRAM for staging, or CUDA stream and event creation failing? Worth documenting the typical failure conditions so users know when it silently falls back.

3. The ids->ne[0]*ids->ne[1] >= 2*n_expert prefill-scale gate is a heuristic to avoid firing on short prompts where the prefetch overhead outweighs the benefit, or is there a deeper correctness concern?

Byte-identical output with slots=0 is the right design. Curious if you've seen this help with models beyond qwen4exp. Mixtral-style routing where the expert set is smaller but more frequently repeated is one I'm wondering about.

[leshchukandrej]

I have tested it on 35b moe only. 4 as a cap is found on different tests to be optimal as after 3, rge gain is much smaller but the vram usage bigger. 30% is a real gain, after that it bis almost non visible. Probably limitation of the card as it is almost 100% saturated on 3 comparing to 50-60% when feature off

…

On Mon, Sep 7, 2026, 13:35 XBold ***@***.***> wrote: *XBold* left a comment (ggml-org/llama.cpp#28414) <#28414 (comment)> Cool approach. The 1-deep lookahead with a second CUDA stream is a clean way to hide PCIe latency without changing the graph topology. Keeping it prefill-only means decode performance is guaranteed flat. A couple of things I'm curious about: 1. *Why cap at 4 slots?* For very long prefill batches (your ~42k tok test showed 22% improvement), the expert routing set is broader and the gap between compute and transfer is largest. Did you test with 8+ slots and find a point of diminishing returns, or is it a conservative cap to keep staging memory bounded? 2. *Self-disable on allocation failure* is smart for correctness, but in practice on a card like the 5070 Ti (12 GB), what's the failure mode? Is it not enough VRAM for staging, or CUDA stream and event creation failing? Worth documenting the typical failure conditions so users know when it silently falls back. 3. *The ids->ne[0]*ids->ne[1] >= 2*n_expert prefill-scale gate* is a heuristic to avoid firing on short prompts where the prefetch overhead outweighs the benefit, or is there a deeper correctness concern? Byte-identical output with slots=0 is the right design. Curious if you've seen this help with models beyond qwen4exp. Mixtral-style routing where the expert set is smaller but more frequently repeated is one I'm wondering about. — Reply to this email directly, view it on GitHub <#28414?email_source=notifications&email_token=AIUMH6YXMBY2AP3UNN6THUT5N4EXNA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKNJXGQ4TINBRGAYKM4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-5574944100>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AIUMH62VFBKJVIYEVGNYVDD5N4EXNAVCNFSNUABFKJSXA33TNF2G64TZHM3DCMRTGU2DOOBUHNEXG43VMU5TKMZVGI4DONZVGYZ2C5QC> . You are receiving this because you were mentioned.Message ID: ***@***.***>

[1jeffchristensen]

Confirming blocking point 1 from the automated review with an end-to-end repro on CUDA. The failure is silent wrong output, not a crash, and it is easy to mistake for a model or quant problem.

Setup: Windows, llama.cpp master 465e49b9c plus this PR, CUDA 13.3, 2x RTX 5060 Ti (sm_120) and 1x RTX 3060 (sm_86). Model sh0wie/Qwen3.8-Flash-Next-REAP-288-GGUF Q4_K_M (qwen4exp, 288 experts). Served with -sm layer and -ncmoe, so the experts are host resident.

Symptom: short answers are correct, then any generation past a few hundred tokens becomes a solid run of / for the whole token budget. With --prefetch-experts-slots 0 the identical command is correct every time.

llama-server -m <model> -ngl 99 -c 8192 -sm layer -ts 3,1 -mg 0 -ncmoe 26 \
    --cache-type-k q8_0 --cache-type-v q8_0 --prefetch-experts-slots 3
# "What is 17*23? Answer with just the number." -> "391", correct
# "Write at least 900 words about ..."          -> 1400 tokens, every one of them "/"
# same command with --prefetch-experts-slots 0  -> 1363 tokens of correct prose

Two things made this hard to attribute:

1. It is masked by allocation failure. When the staging buffer cannot be allocated the feature disables itself and output is correct, so the configurations with the least free VRAM were the ones that looked healthy, and freeing VRAM broke them. At ctx 163840 with -ncmoe 26 the 478 MiB staging allocation failed and every task passed. The same two GPUs at the same context with -ncmoe 30, which leaves about 3.5 GiB more free on device 0, allocated successfully and produced only /. That is the reverse of how an out of memory line in the log usually reads.
2. A short smoke test does not see it. The arithmetic answer is correct, and even a 180 token paragraph is correct. Detection needs a few hundred generated tokens.

It reproduces with two devices, so it is not specific to three, nor to the mixed sm_86 card. Configurations tested, all -sm layer with host resident experts: CUDA0+CUDA1, CUDA0+CUDA2, and all three together, at ctx 8192, 32768, 65536 and 163840. Every one that did not hit the staging allocation failure produced /.

On this model the feature was also not a win. A 96k token prefill measured 135.0 tok/s with --prefetch-experts-slots 3 against 140.4 tok/s with it off, same placement, same build.

[leshchukandrej]

It suppose to give the gain only for the moe and only if the gpu is not oversaturelsted already. Also, the values to use are 2, 3, 4. 1 suppose to be treated as off or it may corrupt the output if not

…

On Mon, Sep 7, 2026, 14:03 1jeff ***@***.***> wrote: *1jeffchristensen* left a comment (ggml-org/llama.cpp#28414) <#28414 (comment)> Confirming blocking point 1 from the automated review with an end-to-end repro on CUDA. The failure is silent wrong output, not a crash, and it is easy to mistake for a model or quant problem. Setup: Windows, llama.cpp master 465e49b plus this PR, CUDA 13.3, 2x RTX 5060 Ti (sm_120) and 1x RTX 3060 (sm_86). Model sh0wie/Qwen3.8-Flash-Next-REAP-288-GGUF Q4_K_M (qwen4exp, 288 experts). Served with -sm layer and -ncmoe, so the experts are host resident. Symptom: short answers are correct, then any generation past a few hundred tokens becomes a solid run of / for the whole token budget. With --prefetch-experts-slots 0 the identical command is correct every time. llama-server -m <model> -ngl 99 -c 8192 -sm layer -ts 3,1 -mg 0 -ncmoe 26 \ --cache-type-k q8_0 --cache-type-v q8_0 --prefetch-experts-slots 3 # "What is 17*23? Answer with just the number." -> "391", correct # "Write at least 900 words about ..." -> 1400 tokens, every one of them "/" # same command with --prefetch-experts-slots 0 -> 1363 tokens of correct prose Two things made this hard to attribute: 1. It is masked by allocation failure. When the staging buffer cannot be allocated the feature disables itself and output is correct, so the configurations with the least free VRAM were the ones that looked healthy, and freeing VRAM broke them. At ctx 163840 with -ncmoe 26 the 478 MiB staging allocation failed and every task passed. The same two GPUs at the same context with -ncmoe 30, which leaves about 3.5 GiB more free on device 0, allocated successfully and produced only /. That is the reverse of how an out of memory line in the log usually reads. 2. A short smoke test does not see it. The arithmetic answer is correct, and even a 180 token paragraph is correct. Detection needs a few hundred generated tokens. It reproduces with two devices, so it is not specific to three, nor to the mixed sm_86 card. Configurations tested, all -sm layer with host resident experts: CUDA0+CUDA1, CUDA0+CUDA2, and all three together, at ctx 8192, 32768, 65536 and 163840. Every one that did not hit the staging allocation failure produced /. On this model the feature was also not a win. A 96k token prefill measured 135.0 tok/s with --prefetch-experts-slots 3 against 140.4 tok/s with it off, same placement, same build. — Reply to this email directly, view it on GitHub <#28414?email_source=notifications&email_token=AIUMH6ZELSDHUJYBCS2KPZD5N4IAFA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKNJXGUYTMMJUGI22M4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-5575161425>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AIUMH653GODVQVWPAIVK3DT5N4IAFAVCNFSNUABFKJSXA33TNF2G64TZHM3DCMRTGU2DOOBUHNEXG43VMU5TKMZVGI4DONZVGYZ2C5QC> . You are receiving this because you were mentioned.Message ID: ***@***.***>

[leshchukandrej]

What was the b/ub param set? On Mon, Sep 7, 2026, 15:14 Andrii Leshchuk ***@***.***> wrote:

…

It suppose to give the gain only for the moe and only if the gpu is not oversaturelsted already. Also, the values to use are 2, 3, 4. 1 suppose to be treated as off or it may corrupt the output if not On Mon, Sep 7, 2026, 14:03 1jeff ***@***.***> wrote: > *1jeffchristensen* left a comment (ggml-org/llama.cpp#28414) > <#28414 (comment)> > > Confirming blocking point 1 from the automated review with an end-to-end > repro on CUDA. The failure is silent wrong output, not a crash, and it is > easy to mistake for a model or quant problem. > > Setup: Windows, llama.cpp master 465e49b plus this PR, CUDA 13.3, 2x > RTX 5060 Ti (sm_120) and 1x RTX 3060 (sm_86). Model > sh0wie/Qwen3.8-Flash-Next-REAP-288-GGUF Q4_K_M (qwen4exp, 288 experts). > Served with -sm layer and -ncmoe, so the experts are host resident. > > Symptom: short answers are correct, then any generation past a few > hundred tokens becomes a solid run of / for the whole token budget. With --prefetch-experts-slots > 0 the identical command is correct every time. > > llama-server -m <model> -ngl 99 -c 8192 -sm layer -ts 3,1 -mg 0 -ncmoe 26 \ > --cache-type-k q8_0 --cache-type-v q8_0 --prefetch-experts-slots 3 > # "What is 17*23? Answer with just the number." -> "391", correct > # "Write at least 900 words about ..." -> 1400 tokens, every one of them "/" > # same command with --prefetch-experts-slots 0 -> 1363 tokens of correct prose > > Two things made this hard to attribute: > > 1. It is masked by allocation failure. When the staging buffer cannot > be allocated the feature disables itself and output is correct, so the > configurations with the least free VRAM were the ones that looked healthy, > and freeing VRAM broke them. At ctx 163840 with -ncmoe 26 the 478 MiB > staging allocation failed and every task passed. The same two GPUs at the > same context with -ncmoe 30, which leaves about 3.5 GiB more free on > device 0, allocated successfully and produced only /. That is the > reverse of how an out of memory line in the log usually reads. > 2. A short smoke test does not see it. The arithmetic answer is > correct, and even a 180 token paragraph is correct. Detection needs a few > hundred generated tokens. > > It reproduces with two devices, so it is not specific to three, nor to > the mixed sm_86 card. Configurations tested, all -sm layer with host > resident experts: CUDA0+CUDA1, CUDA0+CUDA2, and all three together, at ctx > 8192, 32768, 65536 and 163840. Every one that did not hit the staging > allocation failure produced /. > > On this model the feature was also not a win. A 96k token prefill > measured 135.0 tok/s with --prefetch-experts-slots 3 against 140.4 tok/s > with it off, same placement, same build. > > — > Reply to this email directly, view it on GitHub > <#28414?email_source=notifications&email_token=AIUMH6ZELSDHUJYBCS2KPZD5N4IAFA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKNJXGUYTMMJUGI22M4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-5575161425>, > or unsubscribe > <https://github.com/notifications/unsubscribe-auth/AIUMH653GODVQVWPAIVK3DT5N4IAFAVCNFSNUABFKJSXA33TNF2G64TZHM3DCMRTGU2DOOBUHNEXG43VMU5TKMZVGI4DONZVGYZ2C5QC> > . > You are receiving this because you were mentioned.Message ID: > ***@***.***> >

[1jeffchristensen]

Defaults on every run in that report: -b 2048 -ub 512. Neither flag was passed. Slots was 3, never 1, so the "1 may corrupt" caveat is not what I hit.

Agreed on both of your points about where the feature applies, and I think they actually sharpen the bug rather than explain it away. It is a MoE, and it is served with -ncmoe so the experts are host resident, which is the case the feature targets. And "only if the GPU is not oversaturated" is exactly the failing regime here. Same two GPUs, same model, ctx 163840, -sm layer -ts 3,1 -mg 0, q8_0 KV, default -b 2048 -ub 512, --prefetch-experts-slots 3, the only variable being -ncmoe:

-ncmoe 26   device 0 at 14672 MiB   478.12 MiB staging alloc FAILS   output correct
-ncmoe 30   device 0 at 11070 MiB   staging alloc succeeds           output is solid '/'

So when the card is full enough that the staging allocation fails, the feature disables itself and everything is fine. Giving it about 3.5 GiB more headroom on device 0 is what breaks it. That is consistent with the static finding: the slots are allocated once from the first split's device, and a split on the other device then reads a foreign pointer. Single device never hits it, which fits your 35B results.

One correction to my earlier comment, which I stated more flatly than my data supports. I said the throughput comparison was on the same build. It was not. The 135.0 tok/s with the flag was one binary and the 140.4 without it was another, the second adding sm_86 to the arch list and PR #27044. Both at default -b 2048 -ub 512, same placement, same model, same prompt. Neither difference should touch sm_120 prefill and the 4% gap is outside my 1.4% cold/warm spread, but a cross-build pair cannot settle a 4% question. Please read that line as "no gain observed on this model", not as "slower". I did not run a clean same-build A/B because I had already turned the flag off for correctness. The correctness finding does not depend on it either way.

For what it is worth, on this box the prefill lever that did pay was ubatch: -ub 1024 -b 4096 took the same 96k prefill from 142.3 to 269.8 tok/s at no decode cost. If a same-build perf A/B on a single device with slots 2/3/4 at -ub 1024 would be useful to you I am happy to run it, since at -ub 512 there may simply not be enough in flight for the prefetch to hide anything.

[pwilkin]

Given my prior tests with prefetching experts, I'm not considering any PRs of this sort unless someone can clearly show me that this beats purely using --mmap with the page cache for performance.

[leshchukandrej]

@pwilkin again, this pr boosts the partial resident in RAM, not SSD. It allows the CPU to compute n chunks without sequential chunk by chunk allowing to saturate the VRAM more and works only if VRAM is undersaturated. Im playing with the similar approach but for the ssd async reading right now.

@1jeffchristensen regarding the multigpu setup, I dont have it so I cannot test the behavior unfortunately. Most likely there is an issue with that

[leshchukandrej]

@1jeffchristensen yes, reread the implementation and it will not work in the current implementation for the multigpu units, unfortunately. You might need to do some work on top of it to make ot work

[briansp2020]

Thanks for putting this together — a one-split lookahead on a second stream with rotating staging slots is exactly the shape we arrived at independently last week on a single AMD card, so flagging our numbers and one thing I noticed in the diff, in case they're useful. Nothing is urgent on our side.

Setup. Radeon AI PRO R9700 32 GB (gfx1201), ROCm 10.0, PCIe 5.0 x16, Ryzen 9 9900X, 128 GB DDR5. Qwen3.8-Flash-Next UD-Q4_K_XL (512 experts, top-10, 48 MoE layers, ~72 GB of expert weights), -ncmoe 48 -b 4096 -ub 4096 -c 196608, q8_0 KV. Our tree is not master — rdna-boosts fork point 790cf51aa plus #27861 (96 cache slots) with a staging ring added in the CUDA backend (ab0b1e69c in our tree, env-gated, default path byte-identical). The mechanism is independent of both, so I'd expect it to carry.

Prefill t/s (warm, single request):

	27k	64k	150k	190k
stock copy path	1,014	1,079	—	917
3 staging slots	1,518	1,473	1,329	1,232
4 or 6 slots	1,555–1,563	1,597–1,600	—	—

Decode unchanged (24.9–25.6 t/s); greedy output identical to the stock path over 5 prompts × 3 repetitions. 3 → 4 slots is +8% and 4 → 6 is nothing, which matches your "diminishing after 3". Two ROCm-specific data points: a hipMemcpy microbench here gives 53 GB/s H2D from pageable memory as well as pinned (the runtime pins on the fly), so the overlap works with plain mmap on ROCm — the CUDA "pageable = host-synchronous" caveat may deserve a line in the description. And your prefetch_max_size sizing is the right call because Unsloth UD quants mix types per layer (ffn_down_exps is Q5_1 on 43 layers and Q8_0 on 5 here, 629 vs 891 MB); it means the VRAM cost is N × the largest expert tensor — 3 slots was ~2.0 GB on this model.

One thing in the diff to double-check. The slot is allocated with ggml_backend_buft_alloc_buffer(buft, ggml_nbytes(input)) and gets the default USAGE_ANY. The CUDA buffer type pads quantized tensors whose ne[0] % MATRIX_ROW_PADDING != 0 (ggml_backend_cuda_buffer_type_get_alloc_size), and mmq.cu clears that padding before the kernel only when src0->buffer usage is COMPUTE ("If src0 is a temporary compute buffer, clear any potential padding"). While a tensor is pointed at a slot its buffer is the slot buffer, so neither happens, and the kernel reads ggml_row_size(type, 512 − ne0 % 512) bytes past the allocation — init_tensor's comment ("initialize padding to 0 to avoid possible NaN values") is the reason that region is normally zeroed. It's live for ffn_down_exps on this model (ne0 = 640) and on the 35B-A3B you tested (n_ff_exp = 768); whether it shows up depends on what's in the slack past the allocation. Sizing the slot with ggml_backend_buft_get_alloc_size(buft, input) and zeroing the tail once at allocation (or giving the slot buffer COMPUTE usage so MMQ clears it) should close it.

Smaller data point: at prefill scale the ids read-back that decides which experts to copy is itself a device sync per MoE layer; copying the whole tensor asynchronously without it was worth ~+11% on its own before any overlap, consistent with your ≥ 2·n_expert gate skipping it.

Happy to re-run anything on this ROCm single-GPU setup if you revise the branch — we can't help with the multi-GPU case, unfortunately.

(Disclosure: the measurements and our implementation were done with Claude Code assisting, and this comment was drafted with it and reviewed by me.)

[briansp2020]

Follow-up with data from running your branch itself on the ROCm box (ported onto our tree at abb8b628e, same model and flags as above, 96 cache slots from #27861), in case it helps — and one correction to my earlier padding suggestion.

Correctness: greedy outputs identical to the stock copy path on every run (5 prompts × repetitions), at 2, 3 and 4 slots. No issues.

Speed on ROCm, mmap'd weights (prefill t/s, 27k / 64k): stock 1,014 / 1,079 → 2 slots 1,059 / 1,051 → 3 slots 1,319 / 1,292 → 4 slots 1,318 / 1,291. Decode unchanged. With --load-mode none (pinned host memory), 3 slots: 1,459–1,532 / 1,484, i.e. parity with our staging-ring numbers above. So on ROCm the mmap case leaves ~13% on the table, and I think I can say why.

The timeline (GPU timestamps around every large H2D copy and every split's kernels; the same per-layer pattern repeats in every layer, times in ms):

big split (down_{L-1} + attention) ends   3289.00
copy down_L 629 MB                          3289.01 – 3300.82
gate_L                                      3289.00 – 3291.49
up_L                                        3300.82 – 3303.95   <- starts exactly when the down copy ends;
                                                                   its own copy finished at 3241

up_L's data has been resident for ~50 ms, but its kernels wait for copy(down_L), which the lookahead enqueued on the copy stream before the host issued up_L's event_wait. A small HIP microbench reproduces the mechanism: stream A does copy1, hipEventRecord(e1), copy2; stream B does hipStreamWaitEvent(e1). If B's wait is issued before copy2 is enqueued, B is released after copy1 (1.3 ms); if it is issued after copy2 is enqueued and the source is pageable, B is released after copy2 (22 ms). With pinned host memory both orders release after copy1. So on ROCm an event recorded after a pageable async copy is resolved at wait time against the copy stream's tail rather than the record point — which is exactly the order your lookahead produces. CUDA should be unaffected; it may explain why you couldn't see it.

Two cheap ways around it, if you want ROCm+mmap to get the full effect: issue the consumer's ggml_backend_event_wait(split_backend, prefetch_ready[...]) for split i+1 right after graph_compute_async(i) and before try_fire_prefetch(i+2) (the wait is then always ahead of the next enqueue), or note in the description that ROCm wants --no-mmap. Happy to test either here.

Correction on the padding point: zeroing the slot's tail once at allocation is not enough, because slots rotate between tensors of different sizes (a 629 MB down_exps landing in a slot that held an 891 MB one has the old bytes in its padding tail). What worked for us is ggml_backend_buffer_set_usage(slot, GGML_BACKEND_BUFFER_USAGE_COMPUTE) after allocation — then mmq.cu clears the padding before each use, exactly as it does for the scheduler's own input copies — plus sizing with ggml_backend_buft_get_alloc_size.

One more VRAM note: since every slot is sized to the largest tensor (891 MB here), 3 slots cost ~2.7 GB and 4 didn't fit next to the 96-slot cache; it doesn't affect the design, just the sizing advice.

(As before: measurements and the port were done with Claude Code assisting; drafted with it and reviewed by me.)