[briansp2020]

Thanks for r19, and for the TODO #42 write-ups - the §14 analysis of the compute-buffer margin was a good read.

This is a small, separate wip/ item in the same form as #104 and #106, so you can review, re-validate and fold it in however suits you: one patch to common.cuh that makes the wave32 warp reductions use DPP lane exchanges on RDNA3/RDNA4 instead of __shfl_xor (which is an LDS ds_bpermute per step on gfx1201).

What it changes: warp_reduce_sum (float, float2, int) and warp_reduce_max do the same XOR butterfly - same partner at the same offset, same order 16, 8, 4, 2, 1 - via permlanex16 and update_dpp row_xmask. The results are bit-identical. Other targets keep __shfl_xor, and -DGGML_HIP_NO_DPP_XOR restores it on RDNA too.

Results on r19 (2 x R9700, ROCm 10.0, same session, details in the README):

• HC_MIX BF16 (Flash-Next, n_embd 2560): 45.0 -> 33.0 us at 4 tokens, 88.6 -> 53.7 us at 8; bandwidth-bound matmuls unchanged.
• Server, Flash-Next GSQ-RCO IQ3_XXS + MTP, -sm layer, 256K, no other local changes: greedy +4 %, sequences +3-4.5 %, 2 concurrent +11 %. Greedy output unchanged in every run, 0 faults.
• test-backend-ops: ~9,700 cases over 16 ops that use the reductions, all OK.

A data point for your §14, in case it's useful: we hit the same grow-in-place reallocation with all experts in VRAM under -sm tensor (with the #106 patches), at a 259.6k-token prompt. The reserve (n_kv 262144) packed into 1572.9 MiB per device; the real end-of-prefill graph (n_kv ~260k, so the F16 K/V copies are 254/508 MiB instead of 256/512) packed into 1579.3 and then 1592.5 MiB, though its live set was no larger. Each growth freed and reallocated the buffer, and a 2 ms sysfs sampler showed card 0 at 32.45-32.47 of 32.62 GB for the ~20-30 ms of each swap. We patched it the same way you did, padding the allocation in ggml_vbuffer_alloc; on this config 3 % was enough (the overrun was ~1.25 %), and the 259.6k peak went from 32,466 to 31,273 MiB with no reallocation. Since your margin covers it, I've left our version out of this PR.

Also, in case it helps when you get to #106: its nine patches rebase onto r19 without conflicts, and we're running them on r19 in production now.

Happy to re-run anything or rework this if you'd rather it land differently. No rush on our side.

(Claude-assisted, as before.)

[stew675]

Thank you @briansp2020

I'll get to all your PRs soon. I just got a real breakthrough in my fight with the whole MoE cache system for both tensor and layer split.

It's finally able to handle multiple concurrent sessions doing both prefill and decode with MTP, and not tie itself into knots, AND be able to dynamically scale the GPU cache arena as required.

This design I'm about to push also fixes a whole class of potential GPU allocation issues.

I've had a lot of big wins today and it's time to sleep. I'll put the polishing touches on the code tomorrow and then I'll check out your PRs

[briansp2020]

Looking forward to it. With your caching and my updates, Flash Next will be run very fast and useful local model!

[stew675]

Looking forward to it. With your caching and my updates, Flash Next will be run very fast and useful local model!

r22 about to ship out. I'm not calling it done, done. More like a beta, but it's the first time that the new dynamic slab allocator makes an appearance, and this is the correct path moving forwards. It allows the prefill and decode paths to co-exist and dynamically shift memory back and forth between them. Without the slab allocator everything relies on the HIP allocator, and let me tell you, that thing makes a mess!

Basically what happens when -ncmoe is > 0, we now grab as much VRAM as possible and manage it ourselves. The prefill work-pools and the MoE caching arena share this slab with a floating pointer. When the prefill engine needs more space to work it evicts enough entries from the MoE cache to satisfy its needs and moves the boundary pointer up. This means that some MoE cache still remains, but decode slows down of course (less experts able to be cached).

When the prefill is done, it relinquishes its memory back to the MoE cache pool, and then decode can then ramp up to full speed again.

It's not perfect, and it's a real juggling act, and almost certainly there's still some scenarios where the balls get dropped (ie. server crash). We'll work through those as we go.

[stew675]

#108 (comment)

The new slab allocator is a massive win!

[stew675]

Thanks, and thanks for the #106 follow-up detail as well. I reviewed this against the current tip (v16-a55e952b8-r25). The inner patch applies cleanly to ggml/src/ggml-cuda/common.cuh, and the approach checks out:

• row_xmask for offsets 1 to 8 is DPP control 0x160 | offset, and offset 16 is permlanex16. Both move the same partner value at the same offset, and the order 16, 8, 4, 2, 1 matches the __shfl_xor loop, so the sums and maxes are bit-identical.
• WARP_SIZE is hard-coded to 32 in common.cuh, and the only explicit reduction widths in the tree are 32, so the static_assert(width <= 32) cannot fire on the RDNA3/RDNA4 path.
• The GGML_USE_HIP and RDNA3/RDNA4 guards match the vendor macros in ggml/src/ggml-cuda/vendors/hip.h.

One point for the fold: the switch is a build-time define (-DGGML_HIP_NO_DPP_XOR), not a runtime env var. The delivery default-on policy uses runtime env vars as kill switches, so if I fold this I will likely convert it to a runtime gate (or document it alongside the existing build-time gates).

The HC_MIX numbers (45.0 to 33.0 us at 4 tokens, 88.6 to 53.7 us at 8) and the server gains are a solid case. Per the promotion rules, individual validations do not carry over, so I will re-run the combination gates on the folded tip before release.

Plan: fold into block 15, re-validate, and confirm the exact text once that is done.

[stew675]

Folded into delivery block 15 and shipped in v16-a55e952b8-r26, with the build-time gate
-DGGML_HIP_NO_DPP_XOR. I measured the DPP benefit directly here: BF16 HC_MIX on gfx1201 / ROCm
7.14 went 58.84 -> 45.26 us at 4 tokens and 100.00 -> 62.55 us at 8 (+23 % / +37 %), with
test-backend-ops 961/961 on MUL_MAT_ID/HC_MIX/HC_COMBINE and the same-seed coherence hash
unchanged.

One deviation from the PR as submitted: the runtime kill switch is not included. I implemented
a device-side runtime gate (per-TU static __device__ flag + setter kernel + host registrar, since
HIP without RDC gives each TU its own device symbols) and it reproducibly wedged the FA prefill at
~4096 tokens with MTP, with the env var set to either 1 or 0; the plain patch and
-DGGML_HIP_NO_DPP_XOR both complete the gate. The butterfly code itself is fine, so the
build-time gate is what ships. The campaign is archived under
archive/work/r26-rdna4-dpp-butterflies/. Thank you.

[briansp2020]

Thanks for folding this in and for the direct measurement, which came out stronger than ours. Interesting find on the runtime gate wedging FA prefill; the build-time gate is great for us. We'll move to r26 and drop our local copy.