Repository navigation
Conversation
Split-K flash decoding for head size 256, GQA 6 and 1-4 query tokens (MTP verify widths) with WMMA and packed in-register dequantization. Used automatically for eligible shapes on RDNA4; GGML_CUDA_DISABLE_PAGED_DECODE falls back to the stock kernels. On an RX 9070 XT it is 2.2-4.6x faster than the stock vec/tile paths at 512-262144 tokens (up to 17x at 2 query tokens), with KL divergence equal to stock within error on Qwen3.8-27B. The kernel addresses KV through an optional page table and supports a streamed ring; this commit uses only the contiguous (identity) case. Adds eval and perf cases for these shapes to test-backend-ops.
ggml_flash_attn_ext_set_pages() lets FLASH_ATTN_EXT read K/V through a page table: k/v are a pool of pages, and each logical page is either a pool page or a page of a separate ring. ggml_flash_attn_ext_set_sync() makes the op wait on a ready flag before reading data another queue writes, and publish a consumed flag when done. The step tag is an input tensor, so one captured graph serves every step. CPU: gathers the logical rows into the work buffer and runs the contiguous path. CUDA (HIP, RDNA4): the decode kernel takes the table (resolved once per 64-token chunk); only blocks that touch ring pages wait. Stock kernels get wait/release kernels around them for synced contiguous K/V. Paged K/V is supported only where the decode kernel applies. test-backend-ops: paged cases with shuffled pool pages and ring pages, against the CPU reference. Paged addressing costs 0-1.9% vs contiguous at 32k-262k.
llama_context_params.kv_resident_size > 0 turns the attention layers' KV cache into a pool of 256-cell pages (llama_kv_stream). Pinned host memory holds pages that leave the pool: - pages are assigned to slots on demand; - a page is written back when evicted and reloaded when needed again; - the moves are async copies on the compute backend's stream, so they are ordered with the graphs. Writes map cells to pool rows through k/v idxs. Attention reads the pool through per-layer page tables (one I32 input, ggml_flash_attn_ext_set_pages) for ubatches the RDNA4 decode kernel takes (up to 4 tokens). Larger ubatches attend a gather of the pages into a scratch slot, with the stock kernels. State save/restore goes through the page map. K-shift is disabled. For now the pool must hold the whole context: streaming the pages that do not fit comes next. Requires -fa on, q8_0 K / q4_0 V, head size 256 and one sequence; hybrid (attention + recurrent) memory only. The MTP draft context does not stream. LLAMA_KV_STREAM_DEBUG=rotate moves every page through host memory into a rotated slot before every ubatch. At 8k with Qwen3.8-27B (greedy, with and without MTP), the generated text is identical to the stock cache: - with the identity layout, at the same speed; - under rotate. The server's RAM prompt cache restores identically through the page map.
Opt-in flags for KV streaming: --kv-stream enables it and requires --kv-resident-mib, the GPU memory for resident KV pages. The speculative draft context never streams.
…tream A small backend extension (ggml-copy-engine.h, get_proc_address "ggml_backend_copy_engine"). It provides: - a queue of host -> device copies; - events on that queue; - coherent mapped host memory that device kernels can poll (the flags of ggml_flash_attn_ext_set_sync). On AMD the copies go to an SDMA engine directly through ROCr, without dependencies. HIP sends the copies of every stream to one engine (0x1 on gfx1201). A copy that waits for its stream's kernels, such as a readback of a graph's output, blocks that engine until they finish. If those kernels wait for this queue's copies, nothing moves. So the queue takes the highest H2D engine (GGML_CUDA_COPY_ENGINE_SDMA overrides). The decode kernel's and the wait kernel's sync waits now trap after 60 s, so a stall fails the process instead of hanging it.
Attended pages that are not resident now stream every ubatch from the host store into two device slots, alternating by layer, through the copy engine: - a copier thread queues a layer's copies once the slot's previous reader has released it; - a publisher thread sets the layer's ready flag when they land; - the attention waits for the flag and releases the slot (ggml_flash_attn_ext_set_sync, with a step tag input). Decode reads streamed pages through the page table, packed at the start of the slot. Larger ubatches get them in place, beside the graph's copies of the resident pages. Placement: - the pages a ubatch writes are resident; - free slots take the newest attended pages; - the oldest page is evicted when a written page needs a slot. A copy job that streams an evicted page waits for its write-back (an event on the compute stream). The page moves are now done when a ubatch is applied for computation, not in apply_ubatch: llama_kv_cache::prepare() applies every ubatch of a batch to find their slots, then undoes them. Stalls: every wait of the copy queue is bounded (LLAMA_KV_STREAM_TIMEOUT, default 30 s) and aborts with the queue state. LLAMA_KV_STREAM_DEBUG=trace logs the queue; LLAMA_KV_STREAM_DEBUG=nocopy runs it (flags, events, waits) without copying anything. Its results are wrong, but a no-MTP step time against a normal run separates the copies' cost from the pipeline's. At 8k with 10 of 40 pages resident (19 of 29 attended pages streamed per decode step), the generated text is identical to the stock cache, with and without MTP: 32.5 vs 33.2 t/s, and 58.2 vs 58.6 t/s with MTP.
When op_params[GGML_FA_PARAM_LSE] holds a device pointer, the MMA kernel writes, for every output row, the log-sum-exp of its attention logits: max + log(rowsum). The scaled and masked logits use the same natural-log units as the normalization. Whole tiles write it in the epilogue; tiles split across stream-k blocks write it in the fixup kernels. This lets a caller attend to parts of K/V separately and merge the normalized results exactly. The vec and tile kernels take the parameter but do not write it. Without the pointer nothing changes.
…-exp ggml_flash_attn_ext_set_page_runs() gives a paged FLASH_ATTN_EXT the page table as runs of consecutive pool or ring pages. The runs live in host memory that the caller keeps current; such a node is never CUDA-graph-captured. On RDNA4 (Q8_0 K / Q4_0 V, head 256), prefill ubatches (more than 4 query tokens) attend each run in chunks of GGML_CUDA_PREFILL_CHUNK tokens (default 6144, where the stock MMA kernel is most efficient). Each chunk is a stock call that writes its log-sum-exp, and the normalized parts are merged in f32. Contiguous K/V longer than a chunk is split the same way. As a result: - the stock kernels' f16 copies of K/V are bounded by a chunk (1 GB of compute buffer at 262k); - so is their f16 P*V accumulation, whose error otherwise grows with the KV length (NMSE vs an f64 reference at 64k: 1.6e-5 chunked vs 1.8e-4 for one stock call). The speed is 0.81-1.05x one stock call at 8k-262k. Ubatches the stock MMA kernel does not take (for example 5-15 query tokens) fall back to the paged decode kernel in 32-row tiles. GGML_CUDA_PAGED_PREFILL forces those tiles for all prefill, to compare. test-backend-ops covers: - prefill shapes (7/33/512 tokens) over shuffled pool pages and ring pages; - a contiguous page layout for perf.
When the backend takes paged K/V for a 512-token ubatch (probed once with ggml_backend_supports_op), prefill uses the same packed ring copies and page table as decode, plus the page runs. The gather into full-layer device slots and the stock kernels' f16 copies are no longer used. Tier 4 (139k tokens, MTP, -c 147456) results: - the largest pool that fits grows from 300 to 1500 MiB; - decode goes from 39.2 to 48.2 t/s; - prefill is 730 t/s vs 736. Tier 3 prefill: 996 t/s vs stock 1004, text identical. Tier 5 now fits at -c 262144.
The chunked prefill (a part's output with its f16 copies of K/V, two log-sum-exp arrays) and the decode kernel's split partials took their scratch from the memory pool at run time. That memory was not part of the reserved compute buffer, so a tight VRAM budget failed in the middle of a prefill (Tier 5, -c 262144, 800 MiB pool) instead of at load. Both now report the scratch through ggml_cuda_flash_attn_ext_get_alloc_size() and use the space after dst, so the graph reserves it.
Two prefill fusions for qwen35-style linear-attention layers, ported from stew675/llama-cpp-rdna-boosts (block 08: norm-gated.cu, gdn-conv.cu; MIT) and made bit-identical to the unfused ops here: - norm-rows: RMS_NORM -> MUL(w) [-> silu-gated MUL] and RMS_NORM -> SCALE on rows of at most 256 columns (per-head norms, the GDN l2 norms) run one wave per row, 8 rows per block, instead of one block per row. The sum repeats rms_norm_f32<256>'s accumulation and reduction order; the square is accumulated as in the stock kernel, because a bare product feeding the tree would be contracted into an fma and round differently. GGML_CUDA_DISABLE_NORM_ROWS=1 turns it off. - gdn-conv: CONCAT(state, x^T) -> SSM_CONV -> SILU reads the state and x directly; the concat only writes the columns the state snapshots copy. The conv runs at the concat node: in the unfused graph x is dead after the concat, and the allocator gives its memory to the conv output (seen in every ubatch) or to tensors created before the conv, so running it at the conv node read clobbered data. In place, each 32-token tile's 3-row halo is saved into unused concat columns first, so the tiles do not race. GGML_CUDA_DISABLE_CONV_FUSION=1 turns it off. Qwen3.8-27B UD-IQ4_XS, RX 9070 XT: perplexity identical (1.7409 +/- 0.04694), same greedy text with MTP; 32k prefill 1162 -> 1188 t/s (+2.2%); GPU time of the 32k prefill graphs 29.26 -> 28.61 s.
The dense FFN down projection in prefill runs as mmq, which first quantizes its activations; those are the SWIGLU output. When that output only feeds the matmul, the mmq quantize now computes silu(gate)*up from the GLU's sources itself, so the GLU kernel and its output round trip go away. The product is rounded on its own (__fmul_rn) as when the GLU kernel stores it, so it is not contracted into the block sums; the q8_1 blocks are bit-identical. Ported from stew675/llama-cpp-rdna-boosts (block 15, FUSE_SWIGLU_MMQ; MIT). ggml_cuda_mul_mat_q takes the GLU as an optional source; the fusion applies only when ggml_cuda_mul_mat would take the mmq path for a non-fp4 weight. GGML_CUDA_DISABLE_SWIGLU_MMQ=1 turns it off. Qwen3.8-27B UD-IQ4_XS, RX 9070 XT: perplexity identical (1.7409 +/- 0.04694), same greedy text with MTP. GPU time of the 32k prefill graphs 28.62 -> 28.21 s; with the norm-rows and conv fusions, 32k prefill 1162 -> 1205 t/s (+3.6%).
When prefill reads the pages in place (paged prefill), a device slot only holds the pages that stream, packed: the attended pages that are not resident. Whenever any page streams the pool is full of attended pages, so at most n_pages - n_slots stream per ubatch. set_backend, which learns whether prefill is paged, now reallocates the two slots to that size (one page when nothing streams); prepare() asserts the bound. Gathers keep full-layer slots. Every page added to the pool now also shrinks both slots by a page. At the Tier 5 shape (-c 262144, MTP, q8_0/q4_0): 832 -> 745 MiB of slots at a 700 MiB pool. That alone does not raise the usable pool: an 800 MiB pool loads but runs out of memory in a long prefill. Same text with full and shrunk slots at one pool.
The compute buffers are sized for a prefill ubatch (398 MiB for the target, 372 MiB for the MTP context at the Tier 5 shape) but decode needs a few MiB. With --kv-resident-mib-decode N, a context that sees a decode batch (at most n_tokens_decode tokens) after a prefill re-reserves its compute buffers at that size and grows the KV stream pool to N MiB; a larger batch shrinks the pool back first and then restores the prefill buffers, so the two never need memory at once. Resizing writes the resident pages back to the host store and reallocates the pool (81 ms at 307 pages). A draft context created against the target (ctx_other) follows it: the target switches its followers to decode before itself, so the MTP buffers shrink before the target pool grows. The first switch waits for a prefill, so a warmup cannot grow the pool before the draft context allocated its buffers. A full state read (a server slot restore) counts as a prefill; partial reads (recurrent checkpoints) do not. common sets n_tokens_decode to the slots times the longest draft plus one, so a verify batch never leaves the decode phase (default 4). Tier 4 shape (-c 147456, MTP, 1400 MiB pool): decode budget up to 2000 MiB (215 -> 307 pages), compute buffers 286 -> 17.4 MiB (target) and 260 -> 5.0 MiB (MTP); decode 48.3 -> 49.4 t/s, which is compute bound there. Same text and draft acceptance with and without the decode budget, also with 5-token drafts (one switch per turn).
A slot file held the target context's state and the context checkpoints, but not the draft context's KV: after a restore, a draft model that attends the whole prompt (MTP) drafted from an empty cache. The checkpoint appendix (version 2) now carries the draft context's full sequence state, and a restore loads it after the checkpoints; version 1 files still load. At 8k with MTP: the restored request gives the same text as a cold one, with 4 prompt tokens reprocessed from the last checkpoint.
Between decode ubatches the host samples and drafts (MTP) while the copy engine idles: 13 ms of a 92 ms cycle at -c 262144 (210k tokens, 607 pages per layer streamed, 1400 MiB decode pool), where the verify is bound by the copies. After a decode ubatch with nothing queued behind it, the copier now copies the first two layers of the next one, with the same pages, into the two slots once their last readers released them. They are published only if the next ubatch streams exactly the same copies; otherwise they are copied again over the prefetched data. A drain (clear, pool resize, state restore) invalidates them. LLAMA_KV_STREAM_DEBUG=noprefetch turns it off. Tier 5 from a restored slot: decode 35.0 -> 37.3 t/s (+6.5%), cycle 92.2 -> 86.1 ms (verify 81.4 -> 74.1 ms; the drafts take 1.1 ms longer beside the copies). 80 of 81 prefetches are used. Same text, reasoning included, with and without, also at 8k where pages move during decode.
The two device slots hold every page that is not resident (809 pages per layer at the Tier 5 shape with a 215-page pool), but decode only streams the attended ones (607 at 210k tokens). --kv-resident-mib-decode now covers the pool and the slots: a decode resize sizes the slots for the pages attended once the next batch is placed plus 16 pages of headroom (LLAMA_KV_STREAM_HEADROOM, for tests), and the pool takes the rest. When the context outgrows the slots, decode() splits the same budget again before the batch and drops the reusable graphs, whose views point into the old buffers. release_pool also frees the slots, so old and new buffers never coexist. Tier 5 from a restored slot (same text, reasoning included, in every run): budget pool pages slot pages decode t/s 1400 + 657 slots 215 809 37.30 (before) 2057 241 597 38.62 2150 258 580 39.59 2250 275 563 40.63 2300 out of memory allocating the slots At 8k with no headroom, the slots grew twice during 420 decode tokens with the same text.
With MTP drafts longer than 3, the verify batches (5+ tokens) left the paged decode kernel: over paged K/V they took the chunked prefill, and over contiguous K/V (the MTP context) the stock kernels, whose f16 copy of the whole K/V needed a 1 GiB compute buffer at -c 262144. Batches of up to 8 query tokens now take the decode kernel: 5 tokens (30 rows) fit one 32-row tile with the 8-token mask staging, 6-8 take two tiles. The KV stream uses the same limit (llama_kv_stream::decode_max_tokens) for its page tables and prefetch. test-backend-ops: 5, 6 and 8 tokens, contiguous and paged. Tier 5 from a restored slot (decode budget reduced by the 150 MiB of recurrent state each extra draft token needs): draft max budget pool pages acceptance decode t/s 3 2200 267 0.729 40.09 4 2050 240 0.650 41.88 5 1900 214 0.594 40.67 6 1750 187 0.494 37.44 Drafts of 4 give the same text as 3; 5 and 6 diverge at one near-tie (coherent both ways), as batch sizes change the rounding.
Add .devops/nix/rocm-rx9070xt.nix with one definition of the qualified gfx1201 configuration: - packages.rocm-rx9070xt: ROCm build for gfx1201 only, with GGML_CUDA_FA_QUANTS including q8_0-q4_0 plus each side's same-type kernel, and the Zen 4 instruction set listed explicitly (SSE4.2, AVX, AVX2, BMI2, FMA, F16C and the AVX-512 extensions). Inside Nix, SOURCE_DATE_EPOCH makes ggml default GGML_NATIVE off and enable no x86 extensions, and the cc-wrapper strips -march=native. - llama-rocm-env: exposes only the gfx1201 device through ROCR_VISIBLE_DEVICES (detected from the KFD topology), so the Ryzen iGPU (gfx1036) is never initialized. - llama-rocm-configure: configures a local build with the same flags. It refuses a build directory whose cached HIP compiler differs, instead of letting CMake reset the cache and drop GGML_HIP. - serve-openai: launches the qualified server. `nix run` uses the Nix build; the dev shell uses ./build-rocm. A model must be given. - devShells.rocm-rx9070xt: ROCM_PATH, HIP_DEVICE_LIB_PATH, rocprofv3, rocminfo, amd-smi, rocm-smi, rocm-bandwidth-test, and python with numpy/matplotlib. package.nix gains fattnQuants and cpuFeatures arguments and uses the current LLAMA_BUILD_UI/LLAMA_USE_PREBUILT_UI options. No ccache: after a change to ggml/src/ggml-cuda/common.cuh, ccache answered 82 of 105 HIP compiles that inline ggml_backend_cuda_context::stream() with objects built against the old header, and the mixed layouts crashed llama-server on startup. llama-rocm-configure passes -DGGML_CCACHE=OFF and the dev shell does not provide ccache. The launcher serves with --kv-stream: a 450 MiB resident pool during prefill and 2100 MiB for the pool and slots between prefills (LLAMA_ARG_KV_RESIDENT_MIB, LLAMA_ARG_KV_RESIDENT_MIB_DECODE), and the draft KV at q8_0/q4_0 like the target's. GPU_MAX_HW_QUEUES defaults to 2: the second HW queue's preemption save area (~140 MiB) is created during the first prompt batch and made a long prefill at full context run out of memory; at Tier 3 and Tier 5 speed is unchanged and the peak is ~85 MiB lower. The launcher passes reasoning_effort=medium as a default template kwarg (LLAMA_REASONING_EFFORT). The model's embedded Qwen template defaults to xhigh, which on a short coding prompt reasoned through the whole 8192-token budget without answering (132 s); medium answered in 840 reasoning tokens (18 s). A request's reasoning_effort still takes precedence. flake.lock is tracked (added with -f; upstream ignores it) to pin nixpkgs 6510408, the ROCm 7.2.3 toolchain, so that a plain `nix develop .#rocm-rx9070xt` builds against it.
A context checkpoint of Qwen3.8-27B with an MTP draft is 208 MiB (149.6 MiB of recurrent state for the target, 57.5 MiB for the draft at 36k tokens), and the server takes three of them for every prompt: at the start of the request, at the last user message, and 4 tokens before the end. Each took about 85 ms, of which only ~10 ms was the copy from the device: resize() zero-filled a fresh allocation at about 3 GB/s, the rate of faulting in 4 KiB pages. Reserve the buffer untouched, madvise(MADV_HUGEPAGE) it, then size it. With THP in madvise mode a checkpoint now takes ~21 ms (resize 8 ms for the target instead of 52 ms), and a short agent turn at 36k tokens (15 new tokens) goes from 606 ms to 415 ms of prompt time. Evicting a checkpoint adds ~5 ms.
LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY saves the parts of a sequence that cannot be rolled back with seq_rm: SWA cells and recurrent states. llama_kv_cache ignored the flag and wrote all of its cells, which is what llama_kv_cache_iswa needs from its SWA cache but not what a plain cache needs: every position of it can be removed. The server's context checkpoints save the draft context with this flag, so with an MTP draft every checkpoint carried the draft's whole KV cache, growing with the context (57.5 MiB at 36k tokens, about 320 MiB at 200k), three checkpoints per prompt and up to 32 kept per slot. A partial save or load of a cache without SWA now writes and reads nothing. After a restore the server removes the draft's positions past the checkpoint, and the ones before it are the same either way. Qwen3.8-27B with MTP at 36k tokens: a checkpoint is 149.6 MiB instead of 208 MiB. A conversation that edits the last user message and regenerates a reply (two checkpoint restores) gives the same text and the same draft acceptance on every turn, and short turns take 15-20 ms less prompt time. A saved slot restores without re-processing. Slot files saved before this change hold a full draft KV per checkpoint; the server's test-load drops that draft data and keeps the checkpoints.
RDNA4 runs mul_mat_vec_q with one wave per block for ncols_dst > 1, and calc_rows_per_block gave it one row per block, so every wave reloaded the q8_1 activations of all columns for a single row. Four rows per block share them. Each row's sum keeps the same per-lane order and warp reduction, so results are bit-identical. With more than one row per block, the last block of a row count that is not a multiple of the block's rows read the weights of the rows past the end (only the stores were guarded). The kernel now gets nrows_x and those rows read the last row again; their results are still not stored. RX 9070 XT, 17408 x 5120, 4 distinct weight matrices per op (streamed from VRAM), change in time with 4 rows per block, ncols_dst = 2 / 4 / 8: Q8_0 -2 / -2 / -28% Q4_K -5 / -27 / -29% Q4_0 -2 / -4 / -50% Q5_K -1 / -24 / -28% IQ4_XS -1 / -7 / -42% Q6_K 0 / -23 / -33% IQ3_S -16 / -32 / -40% Q3_K -15 / -20 / -27% IQ3_XXS -14 / -33 / -34% Q2_K -19 / -19 / -23% IQ2_S -13 / -32 / -34% IQ2_XS -14 / -30 / -32% No measured type or column count (2-8) is slower. 2 rows gains less; 8 rows gains a little more for Q4_K/Q5_K and loses for the IQ2/IQ3 types and Q3_K. Other types (Q4_1, Q5_0, Q5_1, IQ1_*, IQ2_XXS, IQ4_NL, MXFP4) get the same setting unmeasured. Qwen3.8-27B UD-IQ4_XS with MTP (4-token verify): verify 40.0 -> 36.8 ms at 8k, decode 61.7 -> 67.0 t/s at 8k and 63.9 -> 69.5 t/s at 36k, same greedy text. test-backend-ops: MUL_MAT 2446/2446 and MUL_MAT_ID 933/933 pass, as before the change; excluded are the two huge-N f16 MUL_MAT cases and the MUL_MAT_ID dup cases, which fail on ROCm without it.
pd_combine divided the weighted sum of the splits by the sum of their weights. When every key of a row is masked, all splits have a maximum of -inf, every weight is 0, and the row came out as 0/0 = NaN. The CPU reference and pd_merge give 0 for such a row; pd_combine now does too.
GGML_CUDA_PAGED_PREFILL sent prefill ubatches to the paged decode kernel's row tiles, as a prototype to measure against the stock kernels; the chunked paged prefill replaced it. The ablate argument of pd_partial skipped the K loads, the V loads or the next chunk's fetch to measure their cost, and every launch passed 0.
…ebug LLAMA_KV_STREAM_DEBUG=trace, nocopy and noprefetch logged the copy queue, skipped the copies and disabled the prefetch, for measurements; rotate stays, as it tests the streaming paths with a pool that holds the whole context. LLAMA_KV_STREAM_HEADROOM set the decode slots' headroom for tests and nothing sets it; the headroom is now a constant of 16 pages. With separate prefill and decode budgets, every prompt switches phase twice, and each switch logged the reserve, the compute buffer sizes and the pool and slot sizes at info level. A phase switch now reserves quietly (sched_reserve(true)) and the pool resize logs at debug level; the reserve at startup still logs at info.
After a decode ubatch, the copier copies the first layers of the next one ahead, and the next job uses them if it streams the same copies. The prefetch was matched by tag, but the tag only advances for ubatches that stream pages. A ubatch in between that streamed nothing left the tag unchanged, so the job after it matched the prefetch, although that ubatch may have evicted pages and written them back after the prefetch had copied them: the attention could read stale K/V. Jobs now carry a count of every prepared ubatch, and a prefetch is used only by the job of the ubatch right after the one it was copied for.
A KV streaming ubatch of more than 8 tokens got the page runs, so its attention took the chunked paged prefill: for every 6,144-token chunk of the context the stock kernel converts K/V to f16, which a few query tokens do not amortize. n-gram drafts (ngram-mod) verify up to 65 tokens and paid it on every step: 1.79 ms per attention layer at 9 tokens and 64k context, against 0.56 ms at 8. Without page runs, paged K/V of any size takes the decode kernel's row tiles, the faster up to 96 query tokens. Per attention layer, chunked prefill / row tiles: query tokens 4k 16k 64k 17 155 / 79 us 580 / 273 us 2463 / 1086 us 65 362 / 310 1425 / 1202 5906 / 4921 96 368 / 355 1448 / 1375 5977 / 5436 128 332 / 414 1328 / 1595 5400 / 6299 Decode t/s with --spec-type draft-mtp,ngram-mod (n_max 16, n_min 12) at the launcher's config, before -> after (MTP alone: 91 at 8k, 81 at 71k): copy 80 lines, 8k 186 -> 236 copy with 3 renames, 8k 152 -> 181 copy 80 lines, 71k 58 -> 146 The copy and edit outputs are byte-identical before and after. On reasoning ngram-mod rarely drafts, about -2% at 71k. Batches of 9-96 tokens now differ numerically from before (a request's first batch, the template prefix, is one), as with any kernel change: KLD against ubatch-512 logits at ubatch 32/64/96 is 0.0038/0.0035/0.0035, against 0.0036/0.0000/0.0039 before, with perplexity unchanged. The paged test attaches page runs as llama does; 145 Q8_0/Q4_0 FLASH_ATTN_EXT cases pass, paged 9-96 tokens on row tiles and 97 and 512 chunked.
--spec-type draft-mtp,ngram-mod with --spec-ngram-mod-n-max 16 and --spec-ngram-mod-n-min 12 (n_match stays 24). ngram-mod drafts instead of MTP when the output repeats the context, as when code is copied or edited; on reasoning it rarely drafts. Decode t/s at the launcher's config, without a desktop session, MTP alone / with ngram-mod: copy 80 lines, 8k 91 / 236 copy with 3 renames, 8k 91 / 181 copy 80 lines, 71k 81 / 146 explain (thinking), 8k 63 / 62 explain (thinking), 71k 56 / 55 n_max 64 decoded 364 t/s on the 8k copy but 118 at 71k, where a 65-token verify costs 4.9 ms per attention layer. n_match 16 was 6% faster than 24 on edits and 1% slower on reasoning. --spec-draft-p-min 0.5-0.9 cost MTP 12-30% on reasoning, so it stays 0. The decode buffers grow by about 40 MiB; the peak at 71k was 16,140 of 16,304 MiB. LLAMA_SPEC_TYPE=draft-mtp turns the n-gram drafts off.
The stream created the copy engine only when the initial (prefill) pool could not hold the context. A decode budget smaller than the context streams too, so a prefill budget that held the context with a smaller decode budget asserted at the first decode ubatch that streamed: -c 16384 --kv-resident-mib 1000 --kv-resident-mib-decode 120 llama-kv-stream.cpp:366: GGML_ASSERT(ce && "llama_kv_stream: no copy engine") failed The cache now passes the pages the decode budget holds to the stream, which creates the engine when either phase cannot hold every page (the decode budget also holds at least one ring page), and fails at load if the device has none. The load log states what each phase keeps on the device: llama_kv_stream: 69 of 1024 pages resident (17664 tokens), ... llama_kv_stream: between prefills, the decode budget holds up to 323 of 1024 pages (82688 tokens) llama_kv_stream: the other pages stream from the host store into the ubatches that attend them A decode budget smaller than the prefill budget is accepted as given, with a warning: decode has the smaller compute buffers, so such a budget only keeps fewer pages resident between prefills. The repro above now runs and writes the same text as with a decode budget that holds the context (89.6 vs 90.2 t/s). The launcher's config writes the same text at the same speed as before.
…ocess A phase switch writes the resident pages back to the host store and frees the pool before it allocates the next phase's, so that the two never coexist. When the next pool did not fit (another process took the VRAM after the server started), the allocation threw out of llama_decode with the cache left without a pool, and the server later aborted at exit: E srv update_slots: decode() failed: KV streaming: failed to allocate a pool of 322 pages ggml-backend.cpp:345: GGML_ASSERT(buffer) failed (llama_kv_cache::memory_breakdown) The budget is the user's, so the pool is never shrunk to fit. Instead resize_kv_stream logs what does not fit and why, and leaves the cache without a pool, with every page in the host store: resize_kv_stream: KV streaming: the decode pool of 322 pages (2093 MiB of the 2100 MiB decode KV budget) does not fit: ROCm0 has 1928 MiB free of 16304 MiB. Lower the decode KV budget, or free VRAM that other processes hold. The KV cache stays in host memory, and the next batch tries again. decode returns -2 before anything of the batch runs: the server reports the request as a compute error, clears the slot as for any fatal decode error, and keeps running. The next batch allocates its phase's pool again. No graph is reserved while the pool tensors are detached (the allocator would place them in the compute buffer). The cache's buffer walks skip the missing pool. Tested with a process holding all but 60 MiB of the VRAM free while the server idles in the prefill phase: requests fail with the message above while it holds the memory, the server exits cleanly without a pool, and once the memory is released the next request writes the same text as without the interruption. The launcher's config decodes at the same speed (explain 68.72/68.65 vs 68.68/68.60 t/s, copy with ngram-mod 185.8/187.1 vs 186.9/187.0) with the same text.
The default, made explicit so that the banner and the comment state it. In four agent sessions on the same task (24k-82k decoded tokens each), the false n-gram drafts cost 1.0-2.6 s per session (under 0.5% of decode) with n_match 24 and 32 alike, and 32 did not draft less often per token (4.7 vs 4.5 drafts per 1k tokens). The longest coding session, 17k-127k context, gained an estimated 4.4% of decode time with 24.
…tmuls mmvq quantizes its activations to q8_1 whatever the weight type, so the matrix-vector products that follow on the same activations (q/k/v, the GDN projections, gate/up) quantized them again: 497 quantize launches in a 4-token verify of Qwen3.8-27B, one per matmul, for 257 distinct activation tensors. The backend context now keeps the last copy in a persistent buffer, and an mmvq on the same activations (same data, shape and strides) reads it instead. The graph evaluation drops the copy at the start of every graph, while another stream runs, and after any node whose output overlaps the activations (an in-place op, or memory the allocator gave to a later tensor). The buffer is allocated once, outside a CUDA graph capture, so captured graphs keep a fixed address. MUL_MAT_ID keeps the pool path. HIP gets the mappings for cudaStreamIsCapturing. Launcher config, RX 9070 XT, alternated with the previous build, same text in every run: explain (MTP) 68.72/68.44 -> 69.86/69.73 t/s (+1.7%) copy (MTP + ngram-mod) 186.8/186.9 -> 188.2/188.5 t/s (+0.8%) test_mul_mat_siblings runs three matmuls on the same in-graph activations in one evaluation, with a node between them, and in its inplace variant an in-place op rewriting them before the third. With the invalidation disabled the inplace cases fail (ERR 0.59-0.75); with it all 12 pass. MUL_MAT (2445 cases without the huge-N ones) and MUL_MAT_ID (932) pass as before.
In a 4-token speculative verify, each linear-attention layer of qwen35 concatenates its conv state [3, C] with the new columns x^T into [T + 3, C], copies four 3-column views of the concat into the recurrent cache (the conv state as of each of the last n_rs_seq + 1 tokens, for rollback), and runs the conv on the concat. The concat takes the generic non-contiguous kernel, one 256-thread block per channel for 7 values (13 us for 287 KB at C = 10240), and the copies are four more launches. One kernel now writes the snapshots from the state and x, and skips the copy nodes. In batches of fewer than 256 tokens it also writes the concat, and the conv runs at its own node: 5 launches per layer become 1, 240 -> 48 per verify of Qwen3.8-27B. In prefill ubatches the direct conv of 7866103 reads the state and x at the concat node, so nothing reads the concat and it is no longer written: the kernel that wrote the columns the copies read (gdn_concat_tail) goes, and with it the check that kept the in-place halos below them. The direct conv was also meant for the verify batches on RDNA4 but never matched them: the allocator gives the conv output the memory of the state (T = 3) or of the state and x (T = 4), and at T = 2 the in-place check rejected a halo that a single tile does not use. It now takes prefill ubatches only. The conv's node asks the same question as the concat's (gdn_conv_plan), so it is skipped only when it ran there. Launcher config, RX 9070 XT, alternated with the previous build, same text in every run: explain (MTP) 69.73/69.75 -> 71.73/71.67 t/s (+2.8%) copy (MTP + ngram-mod) 188.3/188.9 -> 191.6/191.7 t/s (+1.6%) sqlite_64k, -c 262144 68.4-68.8 -> 70.1-70.4 t/s decode (+2.4%), prefill, peak VRAM unchanged test_gdn_conv_state builds the graph of build_conv_state (the concat, K snapshot copies into a cache, conv + silu) and checks the snapshots and the conv output: T = 1, 2, 3, 4, 16 take the snapshot kernel with the concat and 256 the snapshot kernel with the direct conv. A wrong snapshot fails every case with K > 1, a wrong direct conv output the T = 256 case. SSM_CONV, SSM_CONV_BIAS_SILU, CONCAT and CPY pass as before.
Every layer of qwen35 adds its residual and then normalizes the sum
(attn_residual -> attn_post_norm, and l_out -> the next layer's
attn_norm): ADD -> RMS_NORM -> MUL(w), with the sum still read by the
next residual add. That is 128 add kernels per 4-token verify of
Qwen3.8-27B, and as many per prefill ubatch.
One kernel now forms the sum, writes it, and computes the norm and the
product from it: rms_norm_f32<block_size, true> with the add folded in.
The sum is formed as the add kernel forms it and the square accumulated
as rms_norm_f32 accumulates it, so the sum and the product are
bit-identical to the unfused ops. The allocator puts the sum in place
over one input and the product over the other (both dead after the add);
that aliasing is allowed, since a block reads its row of the inputs
before the reduction and writes the product after it, and any partial
overlap is not.
Launcher config, RX 9070 XT; this build and the previous one each ran
alternated with the step-1 build (explain 69.73-69.77 t/s, copy
187.2-188.9 t/s in both runs); same text in every run:
explain (MTP) 71.73/71.67 -> 72.44/72.40 t/s (+1.0%)
copy (MTP + ngram-mod) 191.6/191.7 -> 193.4/193.7 t/s (+1.0%)
8k prefill ~1273 -> ~1283 t/s
sqlite_64k, -c 262144 70.1-70.4 -> 70.7-71.0 t/s decode, 960 -> 965
t/s prefill, peak VRAM unchanged
test_add_rms_norm_mul checks the sum and the product for rows of 1000
and 5120 columns (both block sizes of the norm), 1, 4 and 512 rows, and
an in-place add; a wrong product fails it. ADD_RMS_NORM, RMS_NORM,
RMS_NORM_MUL_ADD and ADD pass as before.
The linear-attention layers of qwen35 form their gate as
softplus(alpha + dt) * a, with dt and a one row each: three element-wise
kernels on a [48, n_tokens] tensor, 144 per 4-token verify of
Qwen3.8-27B. The existing UNARY + MUL fusion skips it, as it needs both
MUL operands to have the same shape.
One kernel now computes the three ops for the pattern that
test_add_unary_mul describes (bias and scale one row each, F32,
softplus), with the same op_softplus, so the gate is bit-identical to
the unfused ops. As in the model graph the allocator may give the gate
the memory of alpha, the output may be alpha itself, but share no other
memory with the inputs.
Launcher config, RX 9070 XT; this build and the previous one each ran
alternated with the step-1 build (explain 69.73-69.77 t/s, copy
187.2-188.9 t/s in both runs); same text in every run:
explain (MTP) 72.44/72.40 -> 72.96/72.97 t/s (+0.7%)
copy (MTP + ngram-mod) 193.4/193.7 -> 194.6/194.8 t/s (+0.6%)
sqlite_64k, -c 262144 70.7-71.0 -> 71.1-71.4 t/s decode, prefill
and peak VRAM unchanged
test_add_unary_mul (ADD_SOFTPLUS_MUL) now runs the fused kernel on CUDA:
its six fusing cases fail with a wrong gate and the must-not-fuse cases
pass unchanged; ADD_SILU_MUL, ADD_SIGMOID_MUL, ADD and MUL pass as
before.
… mmvq
After the residual add and norm, a layer's projections (q/k/v or the
linear-attention qkv, z, alpha and beta; the FFN gate and up) multiply
the normalized activations with mmvq, which first quantizes them to
q8_1: 128 of the 257 quantize launches of a 4-token verify of
Qwen3.8-27B. For batches that mmvq multiplies (at most
MMVQ_MAX_BATCH_SIZE rows), the fused add + norm now also writes that
copy, as quantize_q8_1 makes it: in its output loop the lanes of a warp
hold the 32 values of one q8_1 block, so the block's max and sum come
from the same warp reductions, and the copy is bit-identical.
The copy goes into the context's mmvq buffer (ggml_cuda_mmvq_q8_1_cache,
now reserved through reserve() by both), and becomes valid after the
evaluator's written() pass over the fused nodes: the product is often
where the elided RMS_NORM node's result would have been, so marking it
earlier would drop it at once.
Launcher config, RX 9070 XT; this build and the previous one each ran
alternated with the step-1 build (explain 69.73-69.78 t/s, copy
187.2-188.9 t/s in both runs); same text in every run:
explain (MTP) 72.96/72.97 -> 73.26/73.17 t/s (+0.3%)
copy (MTP + ngram-mod) 194.6/194.8 -> 195.4/195.0 t/s (+0.3%)
sqlite_64k, -c 262144 71.1-71.4 -> 71.3-71.7 t/s decode, prefill
and peak VRAM unchanged
test_add_rms_norm_mul_mat multiplies the fused product by two Q4_K,
IQ4_XS or Q8_0 matrices for 1, 4, 8 (the fused copy) and 9 rows (mmvq's
own quantize); with a wrong scale in the copy the 1-8 row cases fail and
the 9-row case passes. MUL_MAT_ID passes as before, MUL_MAT (without the
huge-N cases) but for a q5_0 k=32 case that fails at random on the
previous build too (2 of 150 runs there, 0 of 150 here).
… between its passes
A decode batch gives add_rms_norm_mul_f32 a few rows, one 1024-thread
block each, so the kernel waits on load latency: the second pass read
back the sums it had just written and loaded the weight. Rows of up to
8 * block_size columns now keep each thread's sums and weights in
registers from the first pass (the same values, so the outputs stay
bit-identical); longer rows take the two-pass loop as before.
Per 4-token verify of Qwen3.8-27B (127 kernels): 716 -> 585 us of
kernel time. Launcher config, RX 9070 XT; this build and the previous
one each ran alternated with the step-1 build (explain 69.72-69.78 t/s,
copy 187.9-188.6 t/s in both runs); same text in every run:
explain (MTP) 73.26/73.17 -> 73.45/73.34 t/s (+0.25%)
copy (MTP + ngram-mod) 195.4/195.0 -> 195.6/195.2 t/s (+0.1%)
sqlite_64k, -c 262144 71.3-71.7 -> 71.5-71.8 t/s decode, prefill
and peak VRAM unchanged
test_add_rms_norm_mul gets a 10000-column case (the two-pass loop).
build_rs gathers a linear-attention layer's recurrent state from the
cache into the graph (GET_ROWS of the sequence's row) before the gated
delta net reads it: 3 MB per layer of Qwen3.8-27B (128 x 128 x 48
floats), 48 gathers per verify, ~7 us each. The snapshots already go
the other way without a copy (ggml_cuda_try_gdn_cache_fusion).
With one sequence the kernel now reads the state from row ids[0] of the
cache itself, and the gather is skipped with the no-op view after it
(ggml_cuda_gdn_state_rows_plan; the GDN and the gather node ask the
same plan). Each thread reads the state elements it later writes as
snapshots, so the row may be one the snapshots overwrite, which is the
usual case (the newest snapshot replaces the state it started from).
The fusion requires the snapshot fusion (the GDN node runs in the fusion
path), no other reader of the gather but empty views (build_rs's extra
states) and no write to the cache between the gather and the GDN, and
skips batches that the MMA path may take (more tokens than snapshot
slots: n-gram verifies and prefill keep the gather).
Launcher config, RX 9070 XT, alternated with the previous build, same
text in every run:
explain (MTP) 73.48/73.40 -> 74.52/74.37 t/s (+1.3%)
copy (MTP + ngram-mod) 195.5/195.3 -> 196.3/195.8 t/s (+0.3%)
sqlite_64k, -c 262144 71.5-71.8 -> 72.4-72.7 t/s decode (+1.2%),
prefill and peak VRAM unchanged
test_gated_delta_net_state_rows builds the graph of build_rs and
build_recurrent_attn for one sequence (the state gathered from a row of
the cache the snapshots are copied into) and checks the attention output
and the snapshots, with the state in row 0 (overwritten by the newest
snapshot) or another row; reading a wrong row fails every case of 1, 2
and 4 tokens, and the 5-token cases (the gather kept) pass.
GATED_DELTA_NET, GATED_DELTA_NET_CACHE_FUSION and GET_ROWS pass as
before.
The narrow-row norms of 7866103 (one wave per row for rows of at most 256 columns, from 4096 rows) had no test: no test case reached 4096 rows. test_rms_norm_rows builds qwen35's per-head norms, RMS_NORM -> MUL(w), the gated delta net's l2 norm (RMS_NORM -> SCALE, also on strided views as of the conv output) and the gated form (-> silu-gated MUL), at decode and verify sizes (the stock kernels) and prefill sizes. With a wrong output in rms_rows_f32 the four prefill cases fail, one per mode and the strided one; the others pass. Two findings while writing it: - In qwen35's graph the gated form never matches: the z projection's matmul runs between the norm and the silu gate, so the norm cannot take the gate. The decode trace runs the narrow-row kernel only as MUL and SCALE. - Below 4096 rows the narrow-row kernel is slower than the stock norm: in a 4-token verify of Qwen3.8-27B (176 per-head norms) 617 us against 339 us, so the threshold stays; a comment at the check now says so.
…ir own on HIP too The fused residual add + RMS norm writes the q8_1 copy of its product for mmvq. On HIP the compiler contracted the product into the first add of each q8_1 block sum (v_fmac_f32 in the gfx1201 ISA), so the sum took the unrounded product, while quantize_q8_1 sums the stored value. The quants and the scale were exact; the half-precision block sum could differ in its last bit, which mmvq reads only for q4_0, q4_1, q5_0 and q5_1 weights. __fmul_rn does not prevent this on HIP: without OCML_BASIC_ROUNDED_OPERATIONS, clang's __clang_hip_math.h defines it as a plain product. ggml_cuda_fmul_rn rounds a product on its own (on HIP under #pragma clang fp contract(off), else __fmul_rn). The fused norm and the prefill SWIGLU -> mmq quantize use it; the latter was not contracted, and its ISA is unchanged. The fused norm's lane code moves to quantize_q8_1_lane (common.cuh), for the next q8_1 writer. gfx1201 ISA: the q8_1 variants of add_rms_norm_mul_f32 have 9 fewer FMAs; the variants without the copy are unchanged. ADD_RMS_NORM_MUL 26/26 and SWIGLU 36/36 pass.
In a 4-token verify of Qwen3.8-27B the FFN down projection (mmvq) first
quantizes the SWIGLU output to q8_1: 64 of the 130 quantize launches
left per verify. For batches that mmvq multiplies (at most
MMVQ_MAX_BATCH_SIZE rows, rows a multiple of MATRIX_ROW_PADDING), the
gated unary kernel now also writes that copy into the context's mmvq
buffer, as quantize_q8_1 makes it (quantize_q8_1_lane on the product
rounded on its own), so the copy is bit-identical. This applies to
every F32 GLU op of that kernel (SWIGLU, GEGLU, REGLU, ...). The
evaluator now applies a pending copy after unfused nodes too.
Kernel trace at 8k: 1,526 -> 1,462 kernels per verify, quantize_q8_1
130 -> 66 (256 -> 113 us), the SWIGLU kernel ~130 -> 179 us.
Launcher config, RX 9070 XT, alternated with the previous build; same
text and acceptance in every run:
explain (MTP) 74.46/74.35 -> 74.87/74.64 t/s (+0.5%)
copy (MTP + ngram-mod) 195.9/195.8 -> 194.6/195.7 t/s (noise: its
n-gram verifies are 17 rows)
sqlite_64k, -c 262144 72.42-72.78 -> 72.65-72.94 t/s decode,
prefill and peak VRAM unchanged
test_glu_mul_mat multiplies a GLU output by a Q4_K, IQ4_XS or Q8_0
matrix. With a 10% scale error in the copy the 1-8 row cases fail and
the 9 and 64 row cases pass: those take the prefill SWIGLU -> mmq
quantize, which had no test before (a scale error there fails them).
The GLU ops, MUL_MAT 2444/2444 and MUL_MAT_ID 931/931 (without the
huge-N cases) pass.
…8 tokens)
In a verify batch of Qwen3.8-27B each linear-attention layer computes
MUL_MAT(ssm_alpha) -> RESHAPE -> ADD(dt) -> SOFTPLUS -> MUL(a) -> RESHAPE
and MUL_MAT(ssm_beta) -> RESHAPE -> SIGMOID: two Q8_0 [5120, 48] mmvq
launches of ~10 us each (12 blocks), the fused gate and the sigmoid, 4
launches per layer. For Q8_0 weights and batches of 2-8 tokens that both
matmuls would run as mmvq, gdn_gate_beta_q8_0 computes both matrices,
one row per block (96 blocks), and writes the gate and beta.
Each dot product is summed as mul_mat_vec_q sums it for that column
count (the same lane mapping, vec_dot_q8_0_q8_1, cross-warp order and
warp reduction; rows do not depend on each other), and the element-wise
ops are those of the unfused kernels (op_sigmoid and op_softplus now
use ggml_cuda_op_{sigmoid,softplus}_single), so both outputs are
bit-identical. Decode (1 token) is unchanged. ggml_cuda_mul_mat_takes
reports the kernel ggml_cuda_mul_mat would choose (it replaces
ggml_cuda_mul_mat_takes_mmq), and mmvq_src1_q8_1 gets the activations'
q8_1 copy for mmvq and the new kernel.
llama-perplexity logits at -ub 2, 4 and 8 (SQLite, 2048 tokens) are
byte-identical to the previous build; with a 10% error in beta the -ub 4
logits differ, so the kernel runs there. Kernel trace at 8k: 1,462 ->
1,318 kernels per verify, the replaced kernels 1,107 -> 398 us.
Launcher config, RX 9070 XT, alternated with the previous build; same
text and acceptance in every run:
explain (MTP) 74.73/74.67 -> 76.32/76.18 t/s (+2.1%)
copy (MTP + ngram-mod) 195.8/195.9 -> 197.3/197.4 t/s (+0.7%)
sqlite_64k, -c 262144 72.63-73.03 -> 73.90-74.19 t/s decode (+1.6%),
prefill and peak VRAM unchanged
test_gdn_gate_beta builds the pattern in the model's node order: with a
10% error in the kernel's beta or gate the Q8_0 cases of 2-8 tokens fail
(on SIGMOID or MUL); the 1 and 9 token and Q4_0 cases pass. MUL_MAT
2444/2444, MUL_MAT_ID 931/931, GATED_DELTA_NET, SIGMOID, SOFTPLUS and
ADD_SOFTPLUS_MUL pass.
…DS read and the V scales from a compact copy With 3-4 query tokens (two 16-row tiles) the RDNA4 paged decode kernel ran at ~380-405 GB/s against ~590 GB/s for 1-2 tokens. Timing ablations at 131k context put most of the cost in the LDS reads of the K and V dequantization, not in the WMMAs, the conversion math or KV streaming. - K: a fragment's 8 Q8_0 quants start at an even byte of the row whose offset mod 8 depends on the block alone (34-byte blocks), so they come from the two aligned 8-byte words around them with a constant shift: one read instead of four 2-byte reads. - V: storing the raw rows to LDS also copies each block scale (byte 2 * v of 16-byte piece v) into a [block][token] array, so the 8 token scales of a fragment are one 16-byte read instead of eight. LDS grows by 1 KB to 40.6 KB, still three blocks per WGP; no variant spills (the 8-row one spilled 3 VGPRs before). The same bytes give the same values in the same WMMA order: llama-perplexity logits at -ub 2, 4 and 8 are byte-identical to the previous build. With a broken K shift they differ (PPL 2.2261 -> 4.4713); a wrong V scale fails 96 of the 100 paged FLASH_ATTN_EXT cases. test-backend-ops perf, paged Q8_0/Q4_0, us per call for 1 / 3 / 4 query tokens: 32k 93 / 143 / 145 -> 84 / 120 / 123 131k 372 / 548 / 559 -> 370 / 462 / 473 262k 724 / 1084 / 1097 -> 722 / 905 / 930 sqlite_64k, -c 262144, launcher config: decode 73.91-74.21 -> 75.33-75.65 t/s (+1.9%), prefill and peak VRAM unchanged, same text and acceptance in all 12 replies. Tried and dropped: a compact copy of the K scales (5 spilled VGPRs, slower) and pairing the V byte reads across lanes with DPP (the swaps cost more than the reads they saved).
…e it starts without waiting for them On the RX 9070 XT a 4-token mmvq loses ~9-10 us per launch to its ramp-up and tail. Once running, IQ4_XS streams at ~618 GB/s (time against matrix size, 8k to 70k rows), but every launch adds the fixed cost, and a verify step launches ~400 of them. Matmuls of the same activations that run back to back (gate and up; q, k and v) do not depend on each other, so each may start while the ones before it finish. - Direct launches: a sibling whose call launches only its matmul (q8_1 copy from the cache) is launched with hipExtAnyOrderLaunch, i.e. without the AQL barrier bit. In a harness of back-to-back IQ4_XS matmuls this removes the fixed cost: 17408x5120 85.8 -> 78.2 us, 6144x5120 37.0 -> 27.8 us, 1024x5120 12.8 -> 4.9 us. - Graph capture drops that flag (ROCm 7.2.3: capturehipExtLaunchKernel does not store it). In a capture a sibling runs on stream 1 with exactly the capture dependencies of its group's first matmul kernel (hipStreamGetCaptureInfo_v2, hipStreamUpdateCaptureDependencies; a plain event wait would keep the stream's node from an earlier fork as a second dependency) and joins back. CLR replays such a branch without a barrier on the same queue only with DEBUG_HIP_FORCE_ASYNC_QUEUE=1 and DEBUG_HIP_FORCE_GRAPH_QUEUES=1; otherwise on a second HW queue, which gained nothing. So captures fork only when both are set. - A sibling's src0 must be a graph leaf (weights) and its output must not overlap the group's: the group then only reads memory that stays live while the sibling writes. GGML_CUDA_DISABLE_ANY_ORDER turns it off. CUDA builds are unchanged. Same math in the same order: llama-perplexity logits at -ub 2, 4 and 8 are byte-identical to the previous build, with and without the CLR settings. test-backend-ops MUL_MAT 2451/2451, without the two huge-N f16 cases that fail on ROCm before this change too. The new adjacent-sibling cases overlap in a trace, and do not with GGML_CUDA_DISABLE_ANY_ORDER. 8k, launcher config, decode t/s explain / copy, two rounds each, with the next commit's node order: any-order off 76.18-76.35 / 198.37; on with the CLR settings 78.59-78.60 / 199.70-199.86 (+3.0% / +0.7%); on without them 76.37-76.47 / 198.28-198.88. Same text and drafts throughout.
The gated delta net layers computed z right before the gated norm that uses it, after the delta net, and the attention layers ran Q's norm and RoPE between the Q and K projections. Expanding the matmuls where they are built puts qkv and z, and Q, K and V, back to back, so ggml-cuda on HIP lets them overlap (previous commit). Graph order only: same math, logits byte-identical. With HIP graphs off and direct any-order launches, explain decode at 8k went 76.49 -> 77.73 t/s with this order (only gate/up and K/V were adjacent before).
ggml-cuda forks matmuls of the same activations in captured HIP graphs only when DEBUG_HIP_FORCE_ASYNC_QUEUE=1 and DEBUG_HIP_FORCE_GRAPH_QUEUES=1, the mode in which CLR replays the fork on one queue without a barrier: +3.0% decode on reasoning at 8k (76.3 -> 78.6 t/s), +0.7% on copies. The banner prints both and GGML_CUDA_DISABLE_ANY_ORDER.
…ves so an iteration's loads go out together On the RX 9070 XT a 4-token IQ4_XS mmvq launch costs ~9 us beyond its streaming time. A per-block timeline (start and end of every block) splits it into ~2.1 us between the previous launch's last block and this launch's first, ~1.5 us to start 1280 waves, and the rest inside the kernel: once the GPU drains at the end of a launch each wave is slow, because each loop iteration waits on memory 3-4 times in a row. To fit 12 waves per SIMD (120 VGPRs; gfx1201 allocates in 24s) the compiler issued the q8_1 loads, waited, then the weights, then their scales. amdgpu_waves_per_eu(10, 10) on the multi-column kernel lets it issue all of an iteration's loads before the first wait (IQ4_XS 4 columns: 124 -> 138 VGPRs, still 10 waves per SIMD). Per type and column count, from a sweep of 2-8 columns: - IQ4_XS and Q5_K, 2-5 columns: 10. Past 5 columns the cap makes the compiler serialize the loads to save registers (92-96 VGPRs) and the kernel gets up to 2x slower. - IQ3_S, 2-7 columns: 8; it spills at 10, and at 8 columns at 8. - IQ3_XXS does not gain. 8 for IQ4_XS loses overall: a 5120-row launch (1280 blocks) no longer fits on the GPU at once. The same schedule from -mllvm --amdgpu-sched-strategy=max-ilp or relaxed occupancy for the whole file runs at 8 waves and made decode slower. Isolated, 4 columns, us per matmul: IQ4_XS 17408x5120 86.4 -> 84.9, 5120x17408 84.2 -> 82.6; Q5_K 17408x5120 115.1 -> 113.0, 1024x5120 16.4 -> 14.7; IQ3_S 17408x5120 86.2 -> 83.2. IQ4_XS fixed cost per launch (time against size, 8k-70k rows) 8.8 -> 7.3 us at the same ~614 GB/s. Only the 14 targeted instantiations change in the ISA; other architectures and CUDA keep __launch_bounds__. Same math in the same order: llama-perplexity logits at -ub 2, 4 and 8 are byte-identical to the previous build. test-backend-ops MUL_MAT 2451/2451 (without the two huge-N f16 cases that fail on ROCm before this too), GLU_MUL_MAT 19/19. 8k, launcher config, alternated with the previous build, decode t/s: explain 78.55/78.54 -> 79.44/79.48 (+1.2%), copy 199.69/199.58 -> 199.71/199.92. Same text and drafts throughout. This rests on the compiler's scheduling (ROCm 7.2.3): re-measure after a compiler update or a change to these loops.
…byte-wise SIMD ops HIP has no byte-wise SIMD compare or subtract; __vcmpne4 and __vsub4 are emulated byte by byte, and every 4 grid bytes of IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS and IQ3_S take one of each. ggml_cuda_apply_signs does the same in about half the instructions: m = (((s & sel) + 0x7F7F7F7F) >> 7) & 0x01010101 is 1 in the bytes to negate, (m << 8) - m is 0xFF there, and (g ^ 0xFF) + 1 = -g does not carry into the next byte because the grid bytes of these types are all nonzero. IQ2_S and IQ3_S now pick their sign bits with the same 0x08040201 / 0x80402010 selectors as the other types, and on HIP the sign byte is broadcast with a v_perm instead of a multiply (quarter rate on RDNA). CUDA keeps __vcmpne4 and __vsub4. The 1-column IQ3_S loop goes from 191 to 119 VALU per iteration. Isolated on the RX 9070 XT, us per matmul, 4 columns: IQ3_S 17408x5120 83.9 -> 74.8, 5120x17408 81.8 -> 77.5; IQ3_XXS 84.4 -> 76.4, 82.4 -> 74.4; IQ2_S 17408x5120 87.3 -> 76.1; IQ2_XS 83.4 -> 76.1; IQ2_XXS 74.3 -> 66.5. 1 column: IQ3_S 17408x5120 77.1 -> 73.8, 5120x17408 78.7 -> 77.7, IQ2_XS 62.3 -> 59.4, IQ2_XXS 62.2 -> 57.1, but IQ3_S at 8192-12288 rows is 5-15% slower (10240x5120 48.2 -> 53.0). Its loop now waits for the signs between the grid loads; reordering the source or an occupancy target does not change that. Plain decode (every matmul 1 column, tg128) still gains: 36.47 -> 36.62 t/s. The same idea for Q3_K (vil - vih) was 1-3% slower and is not used. Same results: llama-perplexity logits at -ub 2, 4 and 8 are byte-identical to the previous build. test-backend-ops MUL_MAT 2451/2451 (without the f16 cases with n=2105344 and 8400000 that fail on ROCm before this too), MUL_MAT_ID 931/931. 8k, launcher config, alternated with the previous build, decode t/s: explain 79.45/79.51 -> 80.40/80.39 (+1.2%), copy 199.82/200.22 -> 200.77/200.82. Same text and drafts throughout.
…nd a ring page keeps the ring The decode resize gave the pool everything and the ring one page only when the budget held n_need pool pages and a ring page, but n_need - 1 pool pages and the ring page already hold the n_need pages. A budget in between took the split branch, which then gave the pool all n_need pages and the ring 0, and set_pool reads a ring of 0 as one for every page that is not resident: at 2100 MiB on Qwen3.8-27B, a switch to decode at 78337 to 78592 cells asked for 569 MiB of slots on top of the pool and failed with an OOM. Each budget whose remainder over the pool page size is below a ring page has one such context window.
…t step The catch-up pass ran the MTP layer over every token of the verify batch right after the target's verify decode, before acceptance was known, and the first draft step then ran it again for the new token. process() now keeps the verify rows (token, position and paired hidden row) instead; accept() says how many to keep (the sampled token and the accepted drafts), and draft() decodes them in its first step ahead of the new token. That is one MTP pass instead of two per cycle, and rejected drafts never enter the draft KV. A process() call that comes first decodes the kept rows below its own first position, in a decode of their own, and drops the rest. That covers: - the replay after an n-gram checkpoint restore, which decodes the rows again; - a turn that ends right after a verify; - a prompt that diverges; - the verify of an n-gram draft, after which the MTP draft() did not run. The rows only go onto the draft state they were deferred from, whose last position must be right before them. Shared-memory and chained-head MTP keep the old path. Two things the merged step needs: - The MTP draft context's decode batch is one token per slot longer than a verify. A merged step of a full verify plus the new token otherwise switches the draft context to its 372 MiB prefill buffers mid-decode, which runs out of memory at the 262k launcher config. - With masked nextn embeddings the context copied the first n_outputs rows of t_h_nextn, but most MTP graphs (qwen35, qwen35moe, deepseek32, glm-dsa, ...) set it before narrowing it to the output rows. With rows ahead of the output in the batch, the next draft step read the wrong hidden state (explain: 254/170 -> 315/149 drafted/accepted). The copy now takes the output rows when the tensor still holds every row of the ubatch. 8k, launcher config, 3 rounds alternated with the previous build, time per verify cycle: explain 37.55 -> 37.34 ms (73.50 -> 74.72 t/s; drafted/accepted 565/325 -> 559/327, the MTP rows run in a different batch size); copy unchanged (n-gram drafts dominate). Same text throughout, same peak VRAM. A five-turn conversation with MTP and ngram-mod gives the same text and draft counts as before: - a turn that ends at max_tokens; - follow-ups that reuse it; - a diverging one; - a cached repeat.
…eir own A single-head MTP draft step that also decodes the rows kept from the last verify has several tokens with outputs, and it shared the graph arena of the single-token draft steps after it. Every cycle the next step then rebuilt its graph, ggml-cuda found the graph's properties changed and ran it eagerly, and the step after that captured the CUDA graph again. With an arena of its own, the single-token steps keep theirs: the first one rebuilds at the same addresses and replays the captured graph, the second reuses it. Host time to launch the draft steps of a cycle (8k explain): steps 1 and 2 319 + 349 -> 265 + 216 us. 8k explain, launcher config, 3 rounds, time per verify cycle: 37.34 -> 37.23 ms (74.72 -> 74.95 t/s), same text; copy unchanged. Together with the MTP catch-up merge, at the agent sessions' sampling (temperature 1.0, top-k 20, top-p 0.95, min-p 0.05): explain 37.84 -> 37.47 ms per cycle.
sched_reserve dropped every graph arena and the reserve graph, and a new llm_graph_result zero-fills its graph metadata: ~11 MB for Qwen3.8-27B (27712 nodes), ~3 ms. Each switch between the prefill and decode phases reserves, so every turn recreated them, in the MTP draft context as well. They are now reset when the node budget is the same. At the launcher config: 3 fewer arena creations (~9 ms) per phase switch per context. The 8k benchmarks switch once per run and show no change (explain 37.23 -> 37.24 ms per cycle).
Agent requests sample at temperature 1.0 (top-k 20, top-p 0.95, min-p 0.05 in every request of the recent sessions), where greedy MTP drafts are accepted less often than in greedy benchmarks. --spec-draft-sampling probabilistic samples the drafts and verifies them by rejection sampling, which keeps the target's output distribution. 8k, launcher config, that sampling, seeds 1-3, decode t/s: explain 63.7/71.0/68.5 -> 74.0/76.9/79.8 (+13.6% on the mean), acceptance ~48% -> ~60%, the same time per cycle; copy 231.3/231.1 -> 231.4/231.7 on seeds 2-3 (n-gram drafts). Greedy requests draft greedily either way.
Each call of the Q4_K and Q5_K vec dots took 2 ints of quants, 16 bytes apart, and 4 single ints of q8_1 per column, and unpacked the 6-bit scales for them: twice the work per weight byte of IQ4_XS, and only narrow loads. A call now takes 16 contiguous bytes of a 32-byte group (VDR 4), whose low and high nibbles pair with the same 16 values of two q8_1 blocks: one 16-byte weight load, 16-byte q8_1 loads, half the calls. The integer dot products are unchanged; the float sums group differently, so the results are not bit-identical. Q5_K no longer gets the RDNA4 cap of 10 waves: its loop now needs 142 VGPRs and spilled under it, and gains as much without it. Isolated, 4 columns, us per matmul (rows x K): - Q4_K: 6144x5120 41.6 -> 39.6, 5120x6144 42.3 -> 37.0, 10240x5120 63.4 -> 57.6, 1024x5120 14.3 -> 11.2, 5120x17408 101.4 -> 90.4, 17408x5120 100.5 -> 91.1; also faster at 1, 2 and 5 columns (up to -27%). - Q5_K: 1024x5120 16.2 -> 14.2, 5120x6144 46.5 -> 45.1, 5120x17408 116.5 -> 110.2, 17408x5120 115.8 -> 111.5, the output head 129006x5120 758.9 -> 749.6. Weighted by the matrices of one verify: about -0.27 ms (Q4_K) and -0.14 ms (Q5_K). In situ, on fixed input (llama-perplexity at ubatch 4: every step a 4-token verify, first 2048-token pass): 15.62/ 15.62/15.67 s at HEAD -> 15.47 s, -0.35 ms per verify (-1.1%). At the launcher config (8k, 3 rounds) the text changes, and with it the drafts: explain 37.24 -> 36.31 ms per cycle, 74.9 -> 80.4 t/s, but accepted/drafted 327/559 -> 336/522, so most of that is the new text; at the agent sessions' sampling 63.5 -> 66.9 t/s; copy 59.46 -> 59.23 ms per cycle with the same drafts. KLD at ubatch 4 (every matmul a 4-column mmvq), SQLite source, 8 x 2048 tokens, against HEAD: mean 0.0031, 99.9% 0.20, top-1 same 99.0%, PPL ratio 0.9989 +/- 0.0013 (ubatch 1: 0.0031, 98.97%, 0.9990 +/- 0.0019). HEAD's own ubatch 512 against its ubatch 4 differs more: mean 0.0048, 99.9% 0.31, top-1 same 98.7%. Both at ubatch 4 against HEAD's ubatch-512 logits as a third reference, HEAD -> this change: mean KLD 0.00478 -> 0.00463, top-1 same 98.69% -> 98.89%, PPL ratio 1.0015 -> 0.9997. MUL_MAT 2451/2451, MUL_MAT_ID 931/931.
Support gate/up + GLU fusion in mmvq on RDNA4 when the column count fits in registers (up to 4 columns for IQ4_XS and Q4_K).
Document adaptive KV streaming, RDNA 4 optimizations, setup instructions, full-context benchmark matrix, and real-world multi-turn agent performance.
Includes a deterministic 1024-case LSE page-run merge oracle.
The RDNA2 TILE fallback requires page-run metadata even when a decode batch is <= 8 tokens. Without it, the paged prefill/decode selection fails for these batches. Tested on RX 6950 XT ROCm 10.2 Preview: 4K decode 15.28 -> 49.64 tok/s; 40K decode 1.80 -> 27.11 tok/s, both with external MTP2 and Q8_0/Q4_0 KV.
RX 6950 XT ROCm 10.2 Preview + Swift MTP2, corrected page-runs. Report 40K, 112K and 127.9K prompt outcomes, with known performance limits and confounders.
achyudh
force-pushed
the
feature/adaptive-kv-stream-rocm
branch
from
October 22, 2026 02:54
7048fd6 to
3ef1275
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Add a correctness-first RDNA2 (gfx1030) fallback for the adaptive KV streamer's paged FlashAttention path.
The specialized paged decoder requires RDNA4 WMMA. For RDNA2, this branch:
No KVMem integration or KV ownership changes.
AMD RX 6950 XT measurements
Debian 13, ROCm 10.2 Preview, Swift Qwen3.8-27B IQ3_S with external MTP2, 131072 configured context, K/V q8_0/q4_0, 256 MiB resident KV, 96 fixed generated tokens:
Later requests reused a previous prompt prefix; prefill speeds apply only to newly evaluated suffixes. After the missing short-token page-run propagation was fixed, 112K decoding improved from 0.75 to 13.74 tok/s. A 1,024-case deterministic host numerical merge oracle and targeted HIP compilation also passed.
Full benchmark methods: https://github.com/PourrezJ/llama.cpp/blob/rdna2/paged-attention/docs/rdna2-gfx1030-benchmarks.md
Limitations and review requests
Draft — not ready to merge.
Please review RDNA2 dispatch selection, page-run metadata semantics and LSE merge correctness.