Repository navigation
Conversation
* ui : move models manager to its own route, keep Discover as a dialog Models now lives at #/models instead of a dialog. Discover opens as a standalone dialog over it, and leaving the page returns to the chat it was opened from. Assisted by: Claude (Opus 5.5 Medium) * ui : polish the models page layout and filters Give the models page a page header with the Discover action, calmer column labels, tinted section and family headers, and an empty state. Load the model list on a direct visit, slide the detail pane over the table below xl, and draw the filter toggles as one box matching the context select, with room for their shadow. Assisted by: Claude (Opus 5.5 Medium) * ui : round the models table family headers Assisted by: Claude (Opus 5.5 Medium) * ui : fix models page navigation and Discover gating Record the models page's way back after every navigation, so Back, a typed URL or a reload no longer return to a stale chat. Escape closes the detail pane, then leaves the page. Gate every Discover entry point on one shared check, recheck the API key on /models, skip the duplicate router fetch on a direct load, and share the route load between pages. Assisted by: Claude (Opus 5.5 Medium) * ui : show model details in a drawer Open the models manager's detail pane in a drawer, from the right on desktop and from the bottom on a phone, instead of an inline side panel and a full-screen cover. Pad the pane on both sides, drop its close button on a phone, and assert that filtered rows are gone from the page rather than merely hidden. Assisted by: Claude (Opus 5.5 Medium) * ui : make whole models table rows clickable A click anywhere on a model or quant row opens it, unless it lands on one of the row's own controls (load toggle, actions menu). The name button stays the keyboard path to the same action. Assisted by: Claude (Opus 5.5 Medium)
The src1 power-of-two rescale in kernel_mul_mm_id is superseded by the per expert-token amax scaling done in the compute graph (ref: ggml-org#30861), which keeps the MoE down-projection activations within f16 range before the op runs. Drop the two-stage reduction kernels, their scratch buffer and the GGML_PREC_F32 dispatch added in ggml-org#26223, as well as the test cases that drive MUL_MAT_ID src1 out of f16 range. Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* server_role * cont
…org#31055) * hexagon: KV-head team split for HMX flash attention in multicore * hexagon: merge FA_HEAD_SPLIT and FA_SPLIT into single GGML_HEXAGON_FA_SPLIT flag 0 = off (token-split only); 1 = head-split, plus team-split when n_kv_heads < N_cores so stranded cores still participate (e.g. gemma-4 with 2 KV heads on 4 cores). Where heads divide evenly across cores, plain head-split is used.
* jinja : fix select and selectattr argument handling * jinja : address review feedback * jinja : Simplified test attr
* model : add DeepSeek-V4.1-Flash (deepseek41) * deepseek41 : remove MXFP8, convert the Engram table to BF16 * deepseek41 : compute the compressed-KV read layers once in the cache * deepseek41 : use get_cells() for the shared-history check, drop llama_kv_cache::seq_pos_tok_le * deepseek41 : gather the Engram rows in-graph like qwen4exp, drop the host gather * deepseek41 : keep the Engram hash constants in hparams.ple_*, one module per Engram layer * deepseek41 : drop the LLAMA_DSV41_KV_MIRROR toggle * deepseek41 : keep il last in build_hc_mixes * deepseek41 : sort the top-k rows before the sparse gather * tests : add deepseek41 to the Metal fusion baseline --------- Co-authored-by: Daniel Han <danielhanchen@gmail.com>
…l-org#30242) * chat : reject JSON nested more than 128 levels before walking it The schema builder, the jinja engine and the decision code walk request JSON with recursion. A request nested deep enough overflows the stack of the HTTP thread and llama-server dies (SIGBUS on macOS, SIGSEGV on Linux): for a tool schema from a few hundred levels on macOS and about 10000 on Linux. Check the depth first and answer 400 past 128 levels: the chat inputs in common_chat_templates_apply (tools, tool-call arguments, chat_template_kwargs, json_schema), the json_schema of /completion, and the whole /v1/systemone body. A gemma4 tool result that is too deep stays text. * fix the root, not the call side --------- Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* opencl: add `mul_mv_q2_k_f32` * opencl: add q2_k SOA mv and tiled mm * opencl: add q2_k noshuffle gemm/gemv for adreno * opencl: add q2_k dp4a gemm for adreno Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com> * opencl: tune q2_k kernels for adreno --------- Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: add topk * opencl: optimize topk * opencl: add fp lightning indexer * opencl: optimize lightning indexer * opencl: refactor * opencl: refactor tiled topk kernel * opencl: adjust format
…ps (ggml-org#31071) * hex-tools: support for --compare option in the profile.py * hex-unary: improve collapse logic to also do splits if nrows is small * hex-colapse: handle all colapsing at the host and make it generic * hex-colapse: remove leftover unary helper
…rg#30106) WM of the small tiles was 32, or 8 for 8-wide subgroups, whatever the tile's BLOCK_SIZE and WARP. With 16-wide subgroups (Mali) the second subgroup of a workgroup wrote into the next tile's columns. With 8-wide subgroups some _id tiles left rows uncomputed. Derive WM so there is one warp tile per subgroup. Tested on my Pixel 9 Android phone. Before the change there were 20 MUL_MAT and 6 MUL_MAT_ID failures. After the change they all pass. Fixes ggml-org#28637 Assisted-by: Claude Opus 5.5
…9920) * vulkan: Add a larger tile size for conv3d (coopmat2 only) This adds a larger tile size, in practice it is only for coopmat2 where we don't need shared memory to store the output (would use more than 64KB). Add some additional representative backend tests. At least, this is what codex came up with from looking at sd.cpp Wan VAE. * only apply large tile size for newer driver versions
…p rows (ggml-org#29948) * cuda: fuse FFN gate/up and GLU in MMQ with interleaved gate/up rows For prompt processing, run the gate and up matmuls of a dense FFN as one mul_mat_q launch and apply the GLU in the write-back. - The loaders take the gate weights as a second pointer and fill the weight tile with up and gate rows interleaved in chunks of 8, so each thread holds an up value and the matching gate value (rows r and r + 8 of an MMA tile). - The MMA loop is unchanged: same tile, warp layout, vec_dot and one sum[] per thread. The fused write-back computes the GLU from registers and writes half the rows. - The fused path uses the regular tiling (no stream-k) with the column tiles innermost so that neighbouring blocks share weight tiles. - Enabled for Q4_K, Q5_K, Q6_K, Q8_0 and IQ4_XS on NVIDIA Turing and newer. * cuda: assert the fused MMQ gate argument and trim redundant comments * cuda: move the fused MMQ hint check to should_fuse and add TODOs * cuda: pass prec_src1 to the fused MMQ write-back * cuda: fuse only with efficient tiling and no fallback, merge the GLU write-back * cuda: require has_fusion explicitly, drop the fusion-only src1 precision check * cuda: launch the fused MMQ kernel through the stream-k kernel with one block per tile * cuda: inline the MMQ tiling-efficiency check * tests: one fused gate/up MMQ case per type and GLU op * cuda: always launch the fused MMQ kernel with one block per tile * tests: shorten the fused gate/up MMQ test comment * cuda: select the fused MMQ loader and write-back in ggml_cuda_mmq_get_util_funcs * cuda: run the fused MMQ kernel on the plain tiling path * cuda: put the column tiles innermost in the MMQ plain tiling path
Co-authored-by: Shreya Agrawal <shreyaagrawal@Shreyas-MacBook-Air.local>
* llama : respect no_alloc in memory caches Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB) * llama : TODO: move no_alloc out of llama_hparams Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB) * llama : TODO: move no_alloc out of llama_hparams Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
…based expert detection (ggml-org#30820) * llama-quant : fix quantizing 4D tensors nrows_total was ne[1]*ne[2], so for a 4D tensor only the first ne[2] slices were quantized and the rest of the rows were dropped. Use ggml_nrows(tensor) to count all rows, including ne[3]. Fixes ggml-org#29916 * llama-quant : detect expert tensors by name and fix imatrix size for 4D tensors The MXFP4_MOE path used ne[2] > 1 to detect MoE expert tensors. This is wrong for non-expert tensors with 3 or more dims, e.g. MLA projections (BailingMoE3) and 4D conv kernels. Use the tensor name instead: all tensors used by GGML_OP_MUL_MAT_ID are named *_exps or *_chexps. This also removes the BailingMoE3 special case. The imatrix holds one vector of ne[0] values per 2D slice, so the expected size is ne[0]*ne[2]*ne[3], not ne[0]*ne[2]. This matches how the imatrix is indexed per row for 4D tensors in llama_tensor_quantize_impl. Assisted-by: Claude Opus 5.5
* qwen: merge tensors on convert * convert: use generator functions * qwen4exp: remove conts and pre-scale
…ml-org#31153) * common, grammar: add recursive depth limit, make it a requirement * catch more recursive functions * common : tidy up peg parser changes * grammar : reject candidates in a loop, without a depth limit The depth limit in llama_grammar_reject_candidates_for_stack counted the code points of a vocab token, so any token of 256 chars or more was always rejected under a grammar, such as the 512-space token of GLM-4.7-Flash or the long space tokens of Laguna-S-2.1. The rejection now walks the code points with its own stack of levels instead of recursing: there is no limit and no stack risk, whatever the length of a token from the GGUF, and the masks match master. * grammar : track rejected candidates as indices into the input At a given code point a candidate only differs from its input by the position of its code point pointer, so the levels now carry indices into the input candidates instead of copies. The first level needs no special case and the rejects map back to the input directly. Grammar sampling gets faster than master on GLM-4.7, Laguna, Qwen3.6 and gemma-4 with the json, list and spaces grammars, same masks. * nit : shorten the comment of llama_grammar_reject_step * common : add depth_guard for schema walks * grammar : move GRAMMAR_RECURSIVE_LIMIT into the .cpp * common : reduce recursion limit for peg parser * improve --------- Co-authored-by: Alde Rojas <hello@alde.dev> Co-authored-by: Pascal <admin@serveurperso.com>
* cuda : use std::string instead of VLA for archName
Modernize the code (old C style) and fix the following warnings:
```
ggml/src/ggml-cuda/ggml-cuda.cu:180:19: warning: variable length arrays in C++ are a Clang extension [-Wvla-cxx-extension]
180 | char archName[archLen + 1];
| ^~~~~~~~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:180:19: note: read of non-const variable 'archLen' is not allowed in a constant expression
ggml/src/ggml-cuda/ggml-cuda.cu:179:9: note: declared here
179 | int archLen = strlen(devName);
| ^
```
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
* Add a comment about archName
Signed-off-by: Adrien Gallouët <adrien@gallouet.fr>
---------
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
Signed-off-by: Adrien Gallouët <adrien@gallouet.fr>
* kv-cache : use ggml_backend_tensor_copy to refill the DSV4 K mirrors * llama : move the Engram pad id to llama_hparams and the token map to llama_model The saver no longer needs the deepseek41 model class. * deepseek41 : drop the F32 src1 request on the HC mixes matmul * ggml : add softplus TODOs in the backends that still use log(1 + exp(x)) * convert : make the DSV4/DSV4.1 modify_tensors generators, drop the redundant vision filter * deepseek41 : clean up comments * deepseek41 : drop the F32 src1 requests on the router and indexer matmuls * convert : use format_tensor_name directly in DeepseekV4Model * s/comp_kv_src/comp_il_src/g * cont : fix TODOs for log1p --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* ui : let the expanded sidebar be resized by dragging its edge Bring back the desktop sidebar resize handle that was dropped with the shadcn sidebar primitive. The width is clamped between 18rem and 32rem, persists across reloads, and resets on a double-click of the handle. The tab bar follows the same width so a wider sidebar never covers it. Assisted by: Claude (Opus 5.5 Medium) * ui : rubber-band the sidebar edge past its bounds and drag it shut Past the max width the edge resists harder the further it is pulled, and the same happens below the min width until the pull passes 9rem, where it snaps to the pointer and letting go hides the sidebar. A release outside the bounds springs the width back to the nearest one. Assisted by: Claude (Opus 5.5 Medium) * ui : show the sidebar edge strip faintly on hover and solid while dragging The strip is clipped to the sidebar's rounded corners through its own overlay, so the sidebar itself is never clipped. Assisted by: Claude (Opus 5.5 Medium) * ui : remember whether the desktop sidebar is open across reloads A phone opens the sidebar as an overlay, so it always starts closed there and its toggling never overwrites the desktop preference. Assisted by: Claude (Opus 5.5 Medium) * ui : fade the sidebar labels and conversation list with its width The fade follows the live width, so it tracks both the toggle's width transition and a drag below the minimum width. Assisted by: Claude (Opus 5.5 Medium) * ui : persist the sidebar width on release and clamp a restored one A drag now updates only the live width and writes it to localStorage once when released. A width restored from localStorage is clamped to the current bounds, so a value saved under older bounds does not stick. Assisted-by: Claude (Opus 5.5 Medium)
Assisted-by: Claude
* speculative : save and restore MTP hidden state Assisted-by: Codex * tests : remove MTP state test Assisted-by: Codex * speculative : simplify MTP checkpoint state Save only the boundary hidden row; verification rows are rebuilt on the supported restore paths. Assisted-by: GPT-6 --------- Co-authored-by: zouyongxing <zouyongxing@bytedance.com>
* hexagon: add support for DSV4_HC_PRE * hexagon: hsv4-hc-pre plain/gated split kernel and l2 prefetch optimization * hex-dsv4-hc-pre: double-buffered DMA for prefill * hex-dsv4-hc-pre: refactor DMA block setup * hex-dsv4-hc-pre: precompute kernel params * hex-dsv4-hc-pre: support for n_hc != 4 * hexagon: add support for DSV4_HC_POST * hex-dsv4-hc-pre: refactor gated PRE prefill DMA to dst writeback after each block * hex-dsv4-hc-pre: remove in-loop vector spills and redundant device-side checks * hex-run: add support for test-llama-archs tool * hex-dsv4: refactoring and optimizations Remove non-dma paths. Clean up DMA pipelines and dma_addr_t casts. Add dims collapsing support. Tested with test-llama-archs -a qwen4ext and test-backend-ops --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
…gml-org#31596) Playwright's bundled Chromium lags behind Chrome security releases (1.56.1 ships 141, and even 1.64.1 ships 156.0.8078.4, before the 156.0.8078.73 fixes), and Playwright launches it with --no-sandbox by default. Outside CI, use the installed, auto-updated Chrome with the sandbox enabled. CI keeps the bundled Chromium, so the self-hosted runner's preinstalled 1.56.1 browsers still work. Assisted-by: Claude (Opus 5.5 Medium)
* hexagon: faster TOP_K for the backend sampler (skip tiles, then block scan) Skip the bitonic sort for tiles whose max cannot reach the running k-th best value, then replace the tile sort with top_k_scan(): reject 256-float blocks with one HVX max-reduce and insert only surviving elements into a sorted k-entry list. Used for k in 2..64 and rows wider than 1024, in both the single-row and multi-row paths. Ties keep the first-seen index. NaN is treated as -INF. * hex-top-k: eliminate scalar reads from vtcm --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
…ml-org#31356) * Optimizing Gemma4 graph for prompt processing TODO: - Currently complete microbatch needs to contain tokens which does not require the logits; need to optimize this further - There are a few more operations that are performed after the kv cache is calculated, need to safely optimize them. * Not filtering the rows when they are unmasked * Cropping out the attention mask too. * gemma-4: improved pruning, mask cropping and PLE scaling Prune before the first KV-shared layer with the single-stream, causal, non-embedding safety gates. Use cropped masks only during shared-layer attention. Preserve requested layer-input embeddings in the KV prefix. Remove the PLE contiguous materialization and fold equivalent PLE scaling into RMSNorm parameters/weights. * spec-mtp: Enable masked h_nextn extraction only when the draft context shares the target KV cache This lets Gemma 4 prune unused prompt rows before its shared layers so that we keep the full perf P.P gains with MTP. Avoid extracting target hidden states for output-free prompt chunks in that shared-MTP path. When a batch has a selected final token, extract only that row unless the whole verification batch is selected. Keep ordinary MTP and classic speculative decoding behavior unchanged: they still use unmasked h_nextn because the draft context must catch up on all target rows. speculative-simple needs an update to select the final prompt row only for shared-KV MTP, so the optimization works there without changing other speculative modes. And it looks like the server already selects its final prompt row for first-token sampling. * spec-mtp: make sure to not drop first draft round * gemma4: enable pruning based on static cparams.kv_unified/n_seq_max=1 * gemma4-mtp: simplify prunning logic and avoid graph branches/re-allocs * gemma4: update pruning gates to use only n_seq_max==1, kv_unified can now be used for multi-seq * spec-simple: revert spec-simple changes for now * gemma4: remove causal_attn from pruning gate (last source of reallocs) * cont : adapt to llm_attn_mask changes * cont : add TODO --------- Co-authored-by: Meet Amin <meetamin378@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* hexagon: add TOP_K_MASK op * Simplify Hexagon TOP_K_MASK kernel and do not duplicate host configuration
* add chatterbox TTS support Co-authored-by: Pascal <admin@serveurperso.com> * add built-in voice support * rm redundant code * clean up * clean up 2 * Apply batched suggestions from code review Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * improve --------- Co-authored-by: Pascal <admin@serveurperso.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* hexagon: initial support for LIGHTNING_INDEXER op * hex-lght-idxr: hmx for prefill * hex-lght-idxr: q8_0 decode gather transpose, two-block decode compute * hex-lght-idxr: pipelined hmx path for decode and prefill * hex-lid: add support for K (f32) * hex-lid: add support for K (f32) * hex-lid: add support for q4_0 * hex-lid: remove scalar VTCM reads in the HMX q4_0 K prep and remove vector spills * hex-lid: host/kernel kparams alighnment and remove sw devices in kernel * hex-lit: update dma pipeline to use canonical dbl buffering in all paths * hex-lit: add GGML_HEXAGON_LI_SELECT to force hmx/hvx paths * hex-lit: add support for dims collapse * hex-lit: make kparams constistent with other ops --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* common : fix -hf tag matching inside longer quant names ":F16" picked a BF16 file when the repo had both, since "BF16" contains "F16". The same applied to ":Q4_XS" vs "IQ4_XS". The tag now has to start at a token boundary in the path. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * common : accept underscore as -hf tag boundary Some repos separate the quant with an underscore, e.g. stepfun-ai/Step-3.5-Flash-GGUF-Q4_K_S ships step3p5_flash_Q4_K_S-00001-of-00012.gguf, which :Q4_K_S no longer matched. Add _ to the token boundary and pin the tag and default lookup cases. --------- Co-authored-by: Pascal <admin@serveurperso.com>
…g#31144) The gguf writer staged each tensor into a temporary vector before writing it. For host-buffer tensors the data pointer is directly writable, so skip the staging copy and write it straight to the file. This avoids the per-tensor allocation, zero-init and memcpy. Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml: bound views by the bytes their strides reach A view was checked against its source with the size of a contiguous tensor of the same shape, computed before its strides were set. Views whose rows overlap failed that check although they stay inside their source, and views with wider strides passed it while reaching past the end. The view functions now pass their strides to the tensor constructor, which checks the bytes they actually reach, like ggml_nbytes. The conv windows of delta_net are created with all their rows at once instead of being widened after the check. * ggml: assert that a strided view starts from contiguous elements The view functions set nb[0] to the type size, so the source must have contiguous elements along dim 0 for the view to read them.
* opencl: improve q2_0 decoding and small f32 matmuls Assisted-by: Codex * opencl: optimize small f32 matmuls on Adreno 8xx Assisted-by: Codex
Co-authored-by: Your Name <you@example.com>
The RISC-V vector intrinsic specification requires __riscv_v_intrinsic to be defined when the compiler supports the API, even when the vector extension is not selected. Clang therefore defines it for rv64gc, causing ggml to enter RVV branches where vector types are rejected because the required ISA extensions are absent. Use __riscv_v for optional RVV implementations; this macro represents the selected target ISA. Retain __riscv_v_intrinsic checks in SpaceMIT sources, where they validate mandatory compiler API support rather than select optional vector code. Closes: ggml-org/ggml#1535
…ant (ggml-org#31635) * llama : follow ctx_other devices and flash attention in Gemma4 assistant * llama : include <algorithm> for std::any_of
* ggml : bump version to 0.29.0 * authors : update
* llama.cpp : bump version to 0.9.0 * authors : update
Brings the fork from upstream b11553 to the v0.9.0 release (557 upstream commits, through v0.7.0 and v0.8.0). It replaces the v0.7.0 merge offered in mxxm-t#22 and keeps that merge's resolutions. Every fork feature is kept and combined with upstream's change, with one exception: DeepSeek-V4.1 now uses upstream's implementation. DeepSeek-V4.1: upstream added its own V4.1 (ggml-org#31086, v0.8.0) under the fork's architecture name deepseek41 but with a different GGUF layout, so one tree cannot load both. Upstream's implementation and format are taken; its converter (conversion/deepseek41.py) comes with this merge. The fork's tensor-parallel routing for V4.1 (design-B head split, shared-expert split) is kept. Upstream's version already has the DSpark query norm, the SwiGLU clamp order, the dead layer-0 mix and the verify-batch block padding that the fork fixed. The host-side Engram gather and the LLAMA_DSV41_* diagnostic gates are dropped. GGUFs converted for the fork's V4.1 (deepseek41.engram.layer_ids, blk.N.engram_embed) stop at load with a missing-key error and must be reconverted. LLAMA_PLE_SHARD is limited to qwen4exp, because V4.1 keeps its Engram tables in the same tensor. FEATURES.md and README.md are updated. Textual conflicts (27 files): - common/speculative.cpp: the fork's deferred DFlash prompt encode and upstream's device-buffer hidden features (ggml-org#31506) exclude each other, and the deferred encode is decided first. The device path reads the fork's per-sequence rows, so combined server batches work. Upstream's LiliCorr setup (ggml-org#30506) comes before the fork's mask-token fallback, and a LiliCorr drafter stays off the deferred path. Upstream's gemma4 MTP row selection is combined with the fork's staged deferred-prefill input. - ggml/src/ggml-backend-meta.cpp: upstream's TOP_K_MASK joins the replicated DeepSeek-V4 ops. - ggml/src/ggml-cuda/common.cuh: the fork's context fields are kept with upstream's initializers and graph-eviction timers (ggml-org#30662). The fork's HIP kernel-set record is kept per graph chunk. - ggml/src/ggml-cuda/gated_delta_net.cu: both the fork's chunk kernel header and upstream's MMA header. - ggml/src/ggml-cuda/ggml-cuda.cu: upstream captures a graph in chunks (ggml-org#30804). The fork's HIP kernel-set check (hipGraphExecUpdate swaps kernel functions without validating launch geometry) runs per chunk. Its out-of-memory fallback now resumes direct evaluation at the first node of the chunk that did not fit, because the earlier chunks already ran. Upstream's graph invalidation on pool frees (ggml-org#26771) is added to the fork's fallible pool allocation. Upstream's multi-node NCCL comm_init_rank (ggml-org#28967) sits beside the fork's custom all-reduce. Both gate fusions are kept: the fork's DeepSeek-V4 MUL/ADD/SIGMOID/SCALE and upstream's qwen4exp hc_post (ggml-org#30602). - ggml/src/ggml-cuda/mmq.cu: the fork's gfx906 tile-width rule moves into upstream's ggml_cuda_mmq_get_J_best. The fork's column and token chunking under workspace pressure is kept and passes upstream's fused gate/up/GLU arguments (ggml-org#29948, ggml-org#31188) on every chunk. - ggml/src/ggml-cuda/norm.cu: upstream's float4 rms_norm variants and dispatcher (ggml-org#29720) take the fork's clamped-grid parameter. - ggml/src/ggml-cuda/topk-moe.cu: both sides added the 384-expert case; one is kept. - src/llama-arch.cpp: the fork's sm-tensor replication helper is kept with upstream's llama_arch_* wrappers. The fork's V4.1 keys and tensors are removed. - src/llama-context.cpp: upstream's early return for device-buffer layer inputs comes before the fork's accumulation. - src/llama-context.h, src/llama-ext.h: the fork's MTP pre-norm accumulation and upstream's layer-input device buffer, events and embedding-row API are both kept. - src/llama-graph.cpp, src/llama-graph.h: the fork's snapshot-ring inputs are kept with upstream's s_slot_rev (ggml-org#30878), now written through ggml_backend_tensor_set. Upstream's llm_attn_mask (ggml-org#29510) is used, with the fork's rule that a KV-only graph leaves the mask unallocated. The rotation constants stay on the device (ggml-org#31073). The fork's replica-aware build_get_k is kept with upstream's MLA V view. - src/llama-hparams.h, src/models/models.h, src/models/deepseek41.cpp, src/models/dflash.cpp: upstream's DeepSeek-V4.1, as above. dflash keeps the fork's optional sliding-window pattern for sidecars. - src/llama-kv-cache-dsv4.cpp, src/llama-kv-cache-dsv4.h: upstream's file (V4.1 tiers, compressed-KV mirrors, rs_max ggml-org#31368, no_alloc) with the fork's V4 replica sync, per-sequence replica clearing and phase-aware prefill snapshot gate. - src/llama-kv-cache.cpp: the fork's KV replicas and upstream's kv-rows functions are both kept. - src/llama-memory-recurrent.cpp: the fork's ring-aware seq_rm is kept. Outside ring mode the rollback bound is upstream's cell.rs_max (ggml-org#31368). - src/llama-model.cpp: the fork's design-B head split and general shared-expert split are kept, and upstream's GLM-5.3-next head split is added. memory_breakdown keeps the fork's skip with upstream's no_alloc(). The fork's hybrid memory filters still leave MiniMax-01 out, and upstream's CLOVER_MOE is added. mxxm-t#22's resolution (no per-file buffer map; load_tensors_hadamard, then the PLE prefault) is kept. - src/models/deepseek4.cpp: upstream's file plus the fork's two V4 changes: the layer_inp taps are expanded at the end of the graph, and the HCA compressed cache is accessed through build_get_k/build_cpy_k. - src/models/qwen4exp.cpp: the fork's ring conv-state writes, with upstream's ggml-org#31443 (no ggml_cont before the copy) applied to them. - tools/gguf-split/gguf-split.cpp: the fork's chunked copy with upstream's truncation check on each chunk. - tools/server/server-context.cpp: release() drops the prompt in the fork's cases and for upstream's realtime, TTS and audio-stream tasks. Speculative verify uses upstream's local spec_i_batch (ggml-org#27196) with the fork's per-row token probabilities. Merged without a textual conflict but changed here: - ggml/src/ggml-cuda/gated_delta_net.cu (from mxxm-t#22): upstream's MMA chunked GDN kernel (ggml-org#29353) writes only the final state, so it is skipped when the snapshot ring keeps more than one state. - ggml/src/ggml-cuda/mmq.cu: mmq_args gained a trailing fusion member. The fork's chunked launches built mmq_args without it, which compiles and silently drops a requested fusion. Both now pass it. - src/llama-context.*, src/llama-graph.*: the fork's per-layer input taps and upstream's device buffer were both named layer_inp_dev. The fork's are now layer_inp_tap and are not allocated in device-buffer mode. - common/speculative.cpp: in upstream's device-feature mode the host extraction that fills the fork's deferred accumulation is skipped, so the deferred replay would have injected zeros. This is why the two modes exclude each other. - tools/server/server-context.cpp: upstream moves slot.spec_i_batch into a local before verification, so the fork's later copy was always empty and every speculative token lost its probabilities. - src/llama-graph.cpp: the fork's MTP KV-only store and its KV shift with replicas used the old single rotation. They now use the per-layer rotation (ggml-org#31073), and a replica uses its owner's. - src/llama-kv-cache-dsv4.cpp: after an ubatch whose snapshots the prefill gate skipped, rs_max is 0, so a rollback into it is refused instead of restoring stale state. - src/llama-arch.h, src/models/models.h, src/llama-model.cpp: the duplicate V4.1 enumerators, struct and case label are removed, as is a fork engram_token_map member that shadowed upstream's. - src/llama-model.cpp: LLAMA_PLE_SHARD is limited to qwen4exp. - tests/llama-context/cases-rs.cpp: upstream's rollback cases assume only the last ubatch's snapshots survive. For the four architectures with the fork's snapshot ring (one sequence), rollback_bound expects the ring's bound (n_rs_seq within the decoded history), and rollback expects the ring to accept a rollback across the replayed single-token ubatches. - tests/test-llama-archs.cpp: the fork's deepseek41 skip is removed; upstream's fixture builds the Engram tables. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…onts The grouped top-k MoE fusion (ggml-org#29181) runs one 32-lane row per WARP_SIZE, so on GCN and CDNA two rows share a 64-lane wavefront. group_mask_warp_inplace() took the kept groups from __ballot_sync(0xFFFFFFFF, ...), and HIP's __ballot_sync ANDs the wavefront's ballot with its mask. The second row of each wavefront therefore read the first row's bits, and selected its experts from the first row's groups. On HIP the ballot is now taken over the whole wavefront and shifted down to the row's first lane. CUDA is unchanged. test-backend-ops -o TOPK_MOE on gfx906 (Radeon Pro Vega II, ROCm 10.0): upstream v0.9.0 fails 75 of 672 cases, and the merge without this change 74 to 78 per die, all with n_expert_groups > 1 and more than one row. With this change: 672/672 on each of four dies, and 672/672 on ten further runs with fresh random inputs. Single-row calls were not affected. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Brings in the column-thread gated delta net kernel (cdc81da), which landed on master while this branch was open. The only conflict is the include block of ggml/src/ggml-cuda/gated_delta_net.cu: the column-thread header is added next to upstream's MMA header. The kernel is combined unchanged. It is tried first in launch_gated_delta_net, and upstream's MMA kernel is still skipped whenever the snapshot ring keeps more than one state. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
nbritton
pushed a commit
to exabit-io/llama.cpp-gfx906-tuning
that referenced
this pull request
Nov 6, 2026
reports/2026-11-06-v090-rebase.md: merge of llama.cpp v0.9.0 into mxxm-t's fork (PR mxxm-t/mx-llama.cpp#23), the rebased Exabit master and candidates, the 0.9.0 package, DeepSeek-V4.1 and Nemotron 3.5 Bolt runs. RCCL topology is now system configuration: tools/rccl-topology/ (generator + systemd unit) writes /var/cache/rccl/topology.xml at every boot and /etc/rccl.conf points RCCL at it, so the scripts no longer set NCCL_TOPO_FILE (the old /root file named PCI addresses that moved with the 2026-10-23 card install, and RCCL had been failing over to the internal all-reduce). RELEASE-RUNBOOK.md: state after v0.9.0 and the cycle's silent-break classes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Author
|
Results of the model tests mentioned in the description, on the same box (4x gfx906, ROCm 10.0, 125 W caps). The branch includes your DeepSeek-V4.1-Flash-1017, UD-IQ1_S (unsloth, upstream's GGUF format:
The automatic fit cannot measure a DFlash/DSpark drafter before it has a target context, so it fills the GPUs with the target and the 4.9 GiB drafter then fails to allocate. NVIDIA Nemotron 3.5 Bolt 74B-A9B 24GB (
The two builds' texts diverge within the first 300 characters (batched verification arithmetic), so the acceptance numbers are not directly comparable. Both builds feed the drafter through the same eager host-copy path: upstream keeps LiLiCorr off the device-buffer path too. On gfx906 the drafter costs more than it saves in both builds (about 18-19 tok/s with it, against 39-41 without). One more note for your own testing: |
nbritton
pushed a commit
to exabit-io/mx-llama.cpp
that referenced
this pull request
Nov 6, 2026
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
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.
Merges the llama.cpp v0.9.0 release (
cd959ba, b12110) intomaster, plus one HIP fix for an upstream bug that v0.9.0 brings in. The branch is the merge commit onc4d12e2c0, then the fix, then a merge of your currentmaster(cdc81dae4, the column-thread GDN kernel, which landed while the PR was being prepared). It fast-forwards frommasterand contains no other changes.This replaces #22, which merged v0.7.0 and was not merged. v0.9.0 contains v0.7.0 and v0.8.0, and this merge keeps #22's resolutions.
GGML_HIP_RCCLstaysOFFas in your tree.Review first: DeepSeek-V4.1 now uses upstream's implementation
Upstream merged its own DeepSeek-V4.1 in v0.8.0 (ggml-org#31086, then ggml-org#31137, ggml-org#31292, ggml-org#31300). It uses your architecture name,
deepseek41, but a different GGUF layout, so one tree cannot load both without a format switch. This merge takes upstream's implementation and format. That is what upstream's converter (conversion/deepseek41.py, part of this merge) writes, and what the published GGUFs use:unsloth/DeepSeek-V4.1-Flash-1017-GGUF, for example, carriesdeepseek41.attention.kv_source_layersand theple.*Engram keys.attn_q_b,attn_output_a/attn_output_band the sinks, and the shared-expert split.hc_prenode (5c8b240), and stable pooled block counts for verify batches (eefc4e7; upstream llama : remove simple batch splitting for DSV4 ggml-org/llama.cpp#31300 fixes the same thing and adds a regression test).LLAMA_DSV41_NO_ENGRAM,NO_COMPRESS,QNORMandHOST_INPUT_THREADSgates. Upstream gathers the Engram rows in the graph from the lazily readper_layer_token_embd. Your 16K prefill gain from the host gather (162 -> 588 t/s at-sm tensor -tps 2) has not been measured on upstream's path.deepseek41.engram.layer_ids,blk.N.engram_embed) stop at load with a missing-key error. They need to be reconverted withconvert_hf_to_gguf.py.LLAMA_PLE_SHARDnow applies to qwen4exp only. V4.1 keeps all its Engram tables in the same tensor, and the shard's head map would split them wrongly.FEATURES.mdandREADME.mdare updated to match.What v0.8.0 and v0.9.0 bring
There are 557 upstream commits since b11553 (441 since v0.7.0). The ones that touch code this fork changes:
rms_normvariants. cuda: potential fix for Q8_1 correction sum ggml-org/llama.cpp#31247: the Q8_1 correction sum is computed from the quantized values.flash_attn_ext_rows. kv-cache : keep rotation constants on the GPU ggml-org/llama.cpp#31073: KV rotation constants stay on the device.How the merge was resolved
Every fork feature is kept and combined with upstream's change, apart from the DeepSeek-V4.1 choice above. There are 27 conflicted files; the merge commit message lists each one. The main combinations:
hipGraphExecUpdateswaps kernel functions without validating launch geometry) now runs per chunk. Your out-of-memory fallback now resumes direct evaluation at the first node of the chunk that did not fit, because the earlier chunks have already run.ggml_cuda_mmq_get_J_best. Your column and token chunking under workspace pressure is kept and passes upstream's fusion arguments on every chunk.rms_norm: upstream's float4 variants and dispatcher take your clamped-grid parameter.seq_rmis kept. Outside ring mode, the rollback bound is upstream'scell.rs_max(memory: refuse rollbacks past the snapshots of the last ubatch ggml-org/llama.cpp#31368).spec_i_batch(server: Support logprobs with speculative decoding ggml-org/llama.cpp#27196) is used, with your per-row token probabilities.Your newest commit, the column-thread GDN kernel (cdc81da), is combined unchanged; the only conflict was the include block of
gated_delta_net.cu. The kernel is tried first inlaunch_gated_delta_net, and upstream's MMA kernel is still skipped whenever the snapshot ring keeps more than one state.These merged without a textual conflict but would have been wrong, so the merge commit changes them:
mmq_argsgained a trailing fusion member. Your chunked launches builtmmq_argswithout it, which compiles and silently drops a requested fusion.layer_inp_dev. Your taps are nowlayer_inp_tap, and they are not allocated in device-buffer mode.slot.spec_i_batchinto a local before verification, so your later copy was always empty and every speculative token lost its probabilities.LLAMA_DSV4_NO_PREFILL_SNAPSHOT),rs_maxis 0. A rollback into that ubatch is refused instead of restoring stale state.engram_token_mapmember that shadowed upstream's.test-llama-context'srollback_boundcase expects your snapshot ring's bound for the four ring architectures.The second commit: HIP fix for the grouped top-k fusion
The grouped-experts top-k fusion (ggml-org#29181) runs one 32-lane row per
WARP_SIZE. On GCN and CDNA, two rows therefore share a 64-lane wavefront, and HIP's__ballot_syncANDs the wavefront's ballot with its mask. The second row of each wavefront read the first row's group bits and picked its experts from the wrong groups. The fix takes the whole wavefront's ballot on HIP and shifts it to the row; CUDA is unchanged.On gfx906, stock v0.9.0 fails 75 of 672
TOPK_MOEcases, all withn_expert_groups > 1and more than one row. With the fix, 672/672 pass on each of four dies, and again on ten further runs with fresh random inputs. It is a separate commit so you can take or drop it on its own.Points to check
flash_attn_ext_rows(llama + CUDA: flash_attn_ext_rows (fix unified kv for multi-seq) ggml-org/llama.cpp#29510) has only an NVIDIA MMA kernel. gfx906 reports it as unsupported, so the auto probe leaves it off.LLAMA_DRAFT_PIPELINEgate skips the pipeline ring for contexts withctx_other. Upstream now setsctx_otheron every DFlash context, so self-contained DFlash drafters skip it too.-sm tensorsplit rules don't cover the qwen4exp tensors that the new converter merges (qwen: merge tensors on convert ggml-org/llama.cpp#31028). Upstream has the same gap; GGUFs converted before qwen: merge tensors on convert ggml-org/llama.cpp#31028 are not affected.Testing
All tests ran on a Mac Pro 7,1 with 2x Radeon Pro Vega II Duo (4x gfx906, 32 GB each), ROCm 10.0.0-gfx906+20261018172427 (HIP 7.15.26333) and kernel 7.0, GPUs capped at 125 W. The build used the tree's own defaults (
-DGGML_HIP=ON -DAMDGPU_TARGETS=gfx906, RCCL off).delta-net-base.cpp)test-backend-ops, every op-sm tensor-sm layer -lm dio, greedy 64 tokensllama-server --spec-type draft-mtp --spec-draft-n-max 2, 256 tokensn_probs: 5top_logprobs-np 4, four concurrent requests-sm tensor --n-cpu-moe 41 -lm mlock,LLAMA_PLE_SHARD=1test-llama-context, dummy models of every architecture, one workerllama-gguf-split: split into 5 parts, then mergednemotron_h_moe), one die--spec-type draft-dflash)The build,
test-backend-ops, perplexity,-sm layer, MTP (256 tokens) and Flash-Next rows were measured on the final tree, includingcdc81dae4; perplexity is 5.6126 with and without it. The other rows were measured on the same branch beforecdc81dae4was merged in. The MTP rows are single requests (n=1).test-llama-contextas ctest runs it aborts on this branch, and in the same two ways on #22's tree, so this predates the merge. First, withGGML_SCHED_DEBUG_REALLOC=1it stops on graph reallocations that come from the fork's scheduler. Second, with four contexts in parallel on one device it hits a ROCm error inggml_backend_cuda_cpy_tensor_async. Stock v0.9.0 passes both. This PR does not address them.Not tested (no hardware or models here): CUDA and NVIDIA builds (including the fused MMQ paths and the multi-node NCCL path), AMD architectures other than gfx906, DeepSeek-V4, DFlash, gemma, MiniMax, and
-tpspipeline stages. DeepSeek-V4.1 with DSpark, and Nemotron 3.5 Bolt with its LiLiCorr drafter, are being tested on this box now; I will post the results here.🤖 Generated with Claude Code