vulkan: support sparse Flash Attention by 0cc4m · Pull Request #28105 · ggml-org/llama.cpp

0cc4m · 2026-08-31T17:38:40Z

Overview

Vulkan support for #27970

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES, Claude was used to assist.

On gfx1151 RADV the gather loses at every depth (11k -4.4, 115k -2.7 t/s vs masked, draft-mtp n-max 2): masked FA already skips fully-masked tiles, the per-block bias avoids the mask upload the gather exists to dodge, and gather mode forces the per-cell bias whose upload costs more than the gather saves. QWEN4EXP_QSA_GATHER=1 re-enables it for A/B. Revisit when Vulkan sparse FA (upstream ggml-org#28105) lands.

LynxPDA · 2026-09-11T10:42:49Z

@0cc4m Hi, I tried this PR on a Strix Halo machine with the Vulkan backend. Short contexts seem to work fine, but with long contexts global attention appears to break. The generation stays locally coherent, but the model intermittently loses parts of the previous history/messages.
A simple repro is to ask it to analyze a known-working code file of about 20-30k tokens. Without this PR, it gives a detailed analysis and confirms the code works. With this PR, it claims the code is broken, stitched together from drafts/hallucinations, and contains many truncated functions and unrelated fragments. It looks like some parts of the long context are not being attended to properly. Happy to provide more details if needed.

0cc4m · 2026-09-11T11:28:01Z

Please check with latest version. I'll try to reproduce it.

0cc4m · 2026-09-11T12:36:43Z

I tested it and didn't see an issue, it gives correct responses with long context input.

LynxPDA · 2026-09-11T13:49:39Z

@0cc4m it seems that draft had prefill-sparse. There was an issue with qwen4exp that I described earlier.
I can confirm that the issue is not present in the current PR.

0cc4m · 2026-09-11T16:19:04Z

qwen4exp is not supported by sparse FA yet, as far as I know.

jeffbolznv

I haven't gone through all the code yet, but wanted to get some early feedback out..

Port PR ggml-org#28105's sparse flash-attention compaction and wire it to the qwen4exp QSA prefill mask. The mask of a QSA layer is exactly the top-k selection intersected with causality, so only n_kv_max (= top-k width) cells per row are finite; the backend now compacts those positions per mask row and flash attention reads K/V/mask through the per-row index list instead of scanning the whole cache. - ggml_flash_attn_ext_set_sparse stores the per-row finite bound in op_params[5] (op_params[4] stays the fork's n_kv_raw); CUDA fattn passes the hint through for reference. - flash_attn_sparse_compact.comp builds the per-row index list with a deterministic subgroup-ballot scan (ascending position order, -1 padded). Upstream's atomic slot assignment is a race: the list order is the softmax accumulation order, so the run-to-run bits differ; the scan makes the sparse path bit-stable and identical between the cache on/off arms, which the A/B harness requires. - vulkan FA pipelines gain USE_SPARSE (bit 16) and the fork's DYNAMIC_KV moves to bit 32; cm2's sparse-only tensor-layout updates and gather offsets stay behind USE_SPARSE so the dense specialization keeps its codegen (an unguarded runtime KV select halved dense throughput on gfx1151). - The sparse gate follows the tiling contract of the shader: one index list and one mask row are resolved per DISPATCH TILE, so every row of a tile has to be the same query. That holds when the rows are the gqa heads of one token (gqa_ratio > 1, i.e. decode); large-N shapes run with gqa_ratio == 1 and correctly decline to dense. It also declines when the cache is under max(4096, min_ratio * n_kv_max) cells. FA_SPARSE_DISABLE reverts to dense for A/B. - Extend flash_attn_union/gather_union with a KV-head dimension and a batch offset, plus a grouped prefill driver (64-row groups, opt-in via GGML_VK_FA_TOPK_UNION_GQA): one compact set per group with the scratch reused per group. Inert by default; the per-row sparse path measured ahead of any shared-set compaction. That pp512 measurement was the broken configuration: it took the shared-tile path, which is fast and wrong. With the tiling fixed (see the following commit), the sparse path serves decode (gqa_ratio > 1) and prefill declines to dense. Micro model pp512 and A/B figures above therefore do not describe this commit as merged; re-measure before quoting them.

jeffbolznv · 2026-09-14T13:14:19Z

-                    const uint v_row = j * Bc + row;
+                    uint32_t vcol;
+                    bool kv_active = fa_kv_index(j * Bc + row, vcol);
+                    const uint v_row = USE_SPARSE ? vcol : (j * Bc + row);


Can this just use vcol unconditionally? It doesn't follow the same pattern as other replacements.

Yes, that wasn't necessary. Fixed.

jeffbolznv · 2026-09-14T13:15:35Z

+    const int r = data_sparse[sparse_base + blockCoords[0]];
+    if (r < 0) { return f16vec4(0); }
+    const uint32_t o = g_k_off_elem + uint(r) * k_stride + blockCoords[1] * FA_GATHER_BS + coordInBlock[1];
+    return f16vec4(data_kf16[o], data_kf16[o + 1], data_kf16[o + 2], data_kf16[o + 3]);


I think it should be possible to declare an f16vec4 binding and just do one load.

jeffbolznv · 2026-09-14T13:18:50Z

        const bool k_use_decode = (bs_k > 1u);
-        if (k_use_decode) {
+        if (USE_SPARSE) {
+            coopMatLoadTensorNV(K_T, data_k, k_offset, sliceTensorLayoutNV(tensorLayoutK, j * Bc, Bc, 0, HSK_pad), tensorViewTranspose FAGATHERK);


Should FAGATHERK just be called FADECODEK now and remove the branch?

They use different addressing, so I don't think that's possible here.

I don't understand. I'm just talking about renaming the #define, and then combining line 392 with line 394. I think this should work?

They are distinguished by a spec constant flag in the same shader compile, I can only separate them if I compile them separately.

* upstream/master: (72 commits) HIP: Enable AllReduce for ROCm (ggml-org#27825) opencl: choose the MoE expert matmul by batch size for speculative decoding/MTP (ggml-org#27637) ci: build MUSA for only 1 arch (ggml-org#28944) docs: Rule of thumb for AI review time [no ci] (ggml-org#28945) rpc : hash-cache only weights (ggml-org#28789) cuda: support row-contiguous SUM_ROWS (ggml-org#26308) models : move build_arch_graph() after graph() template specialization (ggml-org#28934) vulkan: support sparse Flash Attention (ggml-org#28105) OpenVINO: optimize stateful decode and GPU MoE inference (ggml-org#28638) opencl: add generic ssm_scan (ggml-org#28881) ci: bump kleidiai runners from 22.04 to 24.04 (ggml-org#28885) metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3) (ggml-org#28599) ci: Bump CUDA Windows x64 builds to 13.4.1 (ggml-org#28930) ci : fix android release (ggml-org#28936) cuda : enable i16 and i32 for DUP (ggml-org#28897) cmake : use PROJECT_SOURCE_DIR instead of CMAKE_SOURCE_DIR (ggml-org#28771) webui: stop re-probing disabled /tools endpoint on every message (ggml-org#28646) ci : reuse build tag name when used instead of safe one (ggml-org#28911) CI: hip-quality-check: ignore spill added in bfdc321 (ggml-org#28909) HIP: fattn-mma: use fp32 accumulation on MFMA devices (ggml-org#28576) ...

* vulkan: add sparse Flash Attention support for DSV4/GLM * tune implementation * add tests * avoid nondeterministic atomicAdd * add cm2 decode vector support * simplify logic and make variable names more consistent * add cm2 f16vec4 binding for decode vector

* server : add missing headers (#28795) Signed-off-by: Adrien Gallouët <angt@huggingface.co> * ggml-cuda: hip add specific config table for AMD GCN (#27841) * server : allow model downloads at model limit fix issue #26809 (#28530) * ui : add cache (#28802) Signed-off-by: Adrien Gallouët <angt@huggingface.co> * cmake: leave the timestamp out of precompiled headers on clang (#28816) Clang stores the modification time of the precompiled header sources inside the header and refuses the header when they differ. A cached header restored from another checkout carries the timestamps of that checkout, so the build fails. The option covers the compilers ccache treats as MSVC while they are clang underneath, clang-cl and the Intel LLVM drivers. * jinja : support dot property integer literals (#28817) * common : implement common_schema internal representation for JSON schemas (#28736) * common : implement common_schema types * common : implement a json schema optimizer * common : reduce optimizations * common : refactor json-schema-to-grammar to use common_schema * common : use common_trie * common/schema : implement type/kind resolution * cont : cleanup * cont : remove common_chat_tool_parameters * cont : simplify schema resolution * cont : pass common_schema through the json-schema-to-grammar builder * cont : cleanup * cont : move enums under common_schema and add type enum * cont : reduce test cases * cont : clean up * cont : clean up * refactor : rename common_schema_parse to common_schema_from_json * tests : fix gcc dangling-reference warning in test-json-schema * tests : take the schema label as const char * to satisfy gcc dangling-reference * refactor : rename common_schema_builder parse_* methods to build_* * cont : fix may_be_string * cont : properly handle empty tool parameters * cont : add tests for empty $ref * cont : remove dead code * cont : update docs * cont : make "{}" mean any object for json_object as well * cont : restore (min|max)Length to imply string type * cont : rename common_schema to common_chat_schema * common: add LOG_JSON macro to log structured data (#28586) * add LOG_JSON macro * fit: add demo LOG_JSON * chat : improve parsing of complex types in qwen3-coder (#28742) * chat : improve schema support in qwen3 parser * cont : clean up grammar a bit * opencl: apply the noshuffle row-alignment rule to q4_K, q5_K and q8_0, not just q6_K (#28575) * vulkan: workaround NV queuesubmit driver bug (#28830) There is a driver bug where two queues on the same VkDevice simultaneously submitting can break some internal synchronization. Until it's fixed, add a mutex around queuesubmit. * ci : cap test-backend-ops parallel jobs at 2 and add a 3600s timeout (#28833) - Clamp the -j parallelism to min(nproc, 2) so a single-core runner uses -j 1 and multi-core runners use at most -j 2, instead of unconditionally using $(nproc). - Add a 3600s timeout to both test-backend-ops runs (the high-perf CPU path and the default path) so a hung test cannot stall CI indefinitely. - Note a TODO to reduce the timeout to 1800s in the future. Assisted-by: pi:llama.cpp/Qwen3.8-27B * ci : remove leftover command (#28839) * tests : reduce FA test sizes (#28842) * pi : prefer PI_MODEL_NAME env var for model disclosure (#28853) Assisted-by: pi:llama.cpp/Qwen3.8-27B * ci : run editorconfig and code-style checks on ubuntu-slim (#28854) Move the EditorConfig Checker and Code Style Checker workflows from the `[self-hosted, fast]` runners to `ubuntu-slim`, which is an established runner label in the repo. Assisted-by: pi:llama.cpp/Qwen3.8-27B * sycl : Fix get mem error (#28227) * fix for unsupport zes API * optimize the code * adjust the log level * rm unused head files * Update docs/backend/SYCL.md Co-authored-by: Titaniumtown <titaniumtown@proton.me> * fix the error to detect level zero SDK/dev package, stop build after detect the error * update the message * fix the build error when missed to install level zero dev package * rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions --------- Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com> Co-authored-by: Titaniumtown <titaniumtown@proton.me> Co-authored-by: Neo Zhang <NA> * tests : fix typo in test-quant-type-selection for nemotron 3 nano (#28835) Corrects a typo in `tests/test-quant-type-selection` for the Nvidia Nemotron 3 Nano 30B A3B model, which was referred to as *nvidia-nemotron-nano-3-30b-a3b*. The error made the test skip that test case, rather than failing the test. [no release] * ggml-cpu(s390x): guard VXE-only repack helpers (#28775) * models : guard the expert FFN size fallback in nemotron-h against a zero divisor (#28779) The NextN/MTP tail loop derives the expert FFN size as n_ff/n_expert_used when expert_feed_forward_length gives nothing for the layer. Both values come from per-layer arrays that legitimately hold 0 on layers that are not MoE, so a checkpoint whose predict layers hold 0 in both divides by zero and dies with SIGFPE at load time, with no error message. Report the malformed metadata instead. * tests : exclude HY_V4 from WebGPU test-llama-archs tests (#28855) Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> * grammar : coalesce find + insert into a single insert and adjust move/copy mechanics (#26885) 1) Combine two consecutive lookups (find + insert) into a single insert-attempt/lookup routine so that we don't per form two O(log(n)) lookup operations in a row anymore -- we only need to do it once and then see if the insert succeeded. 2) Instead of copying every potential stack (expensive) and then moving it (cheap) to new_stacks when it's a final output state, we switch the order so that we move every potential stack (cheap), and then only copy it (expensive) to new stacks when it's a final output state. There are a LOT of intermediate states that get generated, and unless they become final output states, then all of these expensive intermediate copies are wasted. Before: lookup -> lookup/insert + copy -> optional move to output New: lookup/insert + move -> optional copy to output * ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (#28846) * ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA) * apply logic to NVIDIA as well --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * common : move llama_n_rs_seq to before llama_decode (#28749) This commit moves the llama_n_rs_seq function call to before the llama_decode call and returns directly if the check is true, removing the setting of res and the goto statement. The motivation for this change is to avoid the llama_decode call if it is not needed. * sycl : fix oneDNN scratchpad breaking the pool free order (#28704) * ci : remove gg_sum summary logic (#28857) Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * ci : trigger self-hosted CI on changes to ci/run.sh (#28859) The workflow's push/pull_request path filters did not include the ci/run.sh script that all of its jobs execute, so changes to it never re-triggered the self-hosted CI. Assisted-by: pi:llama.cpp/Qwen3.8-27B * ggml-cpu : disable PCH and fix CACHE_LINE_SIZE ambiguity to fix heap corruption (#28882) Disable the ggml-cpu precompiled header and remove the std::hardware_destructive_interference_size branch from CACHE_LINE_SIZE. The PCH force-includes ggml-impl.h before ops.h, which pulls in <new> via <array>/<vector> and defines __cpp_lib_hardware_interference_size. This makes the C++ kernels use CACHE_LINE_SIZE = 256 (hardware destructive interference size) while the C work-buffer sizing code in ggml-cpu.c always uses the fallback 64. The mismatch undersizes the rope work buffer by (CACHE_LINE_SIZE/4 - 16) * n_threads * 4 bytes, causing a heap-buffer-overflow that corrupts the heap and later crashes in ggml_compute_forward_rope_flt. Disabling the ggml-cpu PCH restores the natural include order so ops.h is processed before <new>, keeping CACHE_LINE_SIZE consistent. Removing the std::hardware_destructive_interference_size branch makes the value deterministic and include-order independent. ref: https://github.com/ggml-org/llama.cpp/issues/28858 Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * sycl: rfc: Use radix select for top_k (#28670) * sycl: GPU-resident TOP_K for large k, parallelised over the device The SYCL backend refused GGML_OP_TOP_K above k = 32 and let it fall back to the CPU, a backend round-trip per call. The limit was not conservatism: the scan-merge kernels keep (split_block + 1) * k candidate (value, index) pairs in SLM, so at k = 128 a work-group already needs 132 KB and cannot launch. qwen4exp's sparse-attention indexer asks for k = 2048 in 12 layers on every token, so this fired at every context length. Add a radix select for large k. The k-th largest is found by four most-significant-first passes over an order-preserving unsigned key: histogram the digit over the candidate set, walk the buckets from the top, and recurse into the one where the running count reaches what is still needed. SLM holds the histogram rather than candidates, so the footprint is independent of k. A final pass emits every column beating the pivot plus exactly as many pivot-equal columns as are still missing, so duplicate keys still yield exactly k distinct indices. Output order is not required and is not paid for: ggml-cpu/ops.cpp swaps its first two outputs to say so. The key folds -0.0 onto +0.0 so its equivalence classes match the reference comparator, under which the two tie. NaN has no defined order in the reference (its comparator is not a strict weak order there); here +NaN keys above +inf and -NaN below -inf, which at least makes the result deterministic. One work-group per row leaves the device idle whenever a graph has fewer rows than it has cores, which at batch size 1 means one work-group full stop: qwen4exp tops-k a tensor of shape [n_kv, n_tokens/n_stream, n_stream], so token generation gives nrows == 1, and the backend sampler reshapes logits to a single row as well. Measured, ne=[200000,1] and ne=[200000,16] cost 358.0 us and 363.4 us -- sixteen rows for 1.5% more wall-clock. So also spread a row over several groups when there are too few rows to cover the device. Per-pass state moves to global memory and each digit pass becomes its own launch, since a work-group barrier can no longer span the row. Groups accumulate in SLM and contribute 256 global atomics each, keeping global traffic per-group rather than per-element, and the last group of a row -- the one whose fetch_add returns G-1 -- performs that pass's scan, holding the launch count at one per digit plus one emit. The group count comes from the device and is floor-divided by nrows, so a row count that already covers the device is left whole and pays nothing. Below 64K columns the single-group kernel finishes inside the cost of the extra launches and stays in charge. Reading the row's prefix/mask/need through a device-scope atomic_ref costs more than the sweep it guards: those loads are uncached, so passes 2-4 ran at 49 us against 12 us for pass 1. One lane reads them into SLM and the group takes them from there -- 208 us -> 44.6 us at ne=[131072,1], k=2048. The block size now takes the device's max_work_group_size instead of a cap of 512. The cap was never a floor, so a device reporting 512 is unaffected; one allowing 1024 was being given half its width. Finally, put the scan-merge gate where the two paths actually cross. That kernel's cost climbs with k while the radix select's does not; measured over widths from 2 to 200K columns and row counts from 1 to 8192, radix is ahead everywhere from k = 8 up and behind at k <= 2, where scan-merge's smaller fixed cost wins. The short-row corner (ncols=2, nrows=65536, as in bailingmoe2 group selection) is exactly where radix loses at low k, and the gate keeps it on scan-merge. Op-level against the CPU-fallback path this replaces, and against the single-group radix select for the split: 4.98x at ne=[131072,1] k=2048, 6.65x at ne=[151936,1] k=40, 13.35x at k=20, 118x at ne=[65000,16] k=32. No measured shape regressed. End to end on 3x Arc Pro B60 with Qwen3.8-Flash-Next UD-IQ4_XS, llama-bench tg64, the parallelisation is worth 5.91 -> 6.05 t/s at d=131072 and a wash at shallower depths. Perplexity over wikitext-2 is unchanged within noise at both 512 and 81920 context. test-backend-ops: 525/525 TOP_K (previously every k > 32 case was refused), 880/880 MUL_MAT_ID. Perf coverage added for k > 32 at large widths and for the short-row corner, neither of which was exercised before. * move topk-select to topk-radix.{cpp|hpp} --------- Co-authored-by: cwriter <cwriter@localhost> * llama: add Maple 20B-A1B ternary MoE architecture (CPU) (#27000) * gguf-py: add Maple tensor constants Add MODEL_ARCH.MAPLE, its "maple" name, and the tensor list for the Maple 20B-A1B ternary MoE architecture: token embeddings, output, attention with Q/K RMS norms, and per-expert FFN tensors. * convert: add Maple HF->GGUF converter Register MapleForCausalLM in the HF architecture map and add the converter for the Maple 20B-A1B ternary MoE model: 24 layers, 256 experts with 8 active, sliding-window attention (SWA-512) interleaved with global attention at a 3:1 ratio, partial rotary factor 0.5, and per-expert weight stacking into merged 3D tensors. * llama: add Maple architecture (20B-A1B ternary MoE) Add the Maple 20B-A1B ternary MoE architecture: 24 layers, 256 experts with 8 active, sliding-window attention (SWA-512) interleaved with global attention at a 3:1 ratio, and ternary TQ1_0/TQ2_0 quantization support. - register LLM_ARCH_MAPLE between MAMBA2 and JAMBA - implement llama_model_maple: Q/K RMS norms after projection (GEMMA4 style), rope applied only on SWA layers (nope_on_global_attention), ISWA KV cache, and MoE FFN with swiglu gate clamp at +7 (DEEPSEEK4 style) - mark MAPLE as unsupported by the model saver (roundtrip skipped) * tests: mark Maple as MoE-mandatory Maple is always-MoE: the model throws when n_expert == 0, so the test harness must only run the MoE config for LLM_ARCH_MAPLE. * maple: apply review feedback (n_ff_exp_arr, get_arr, rope params) - load_arch_hparams: use n_ff_exp_arr + n_ff_exp() accessor (upstream changed these from a scalar member during the rebase) - sliding_window_pattern: get_arr, the pattern is mandatory for this arch - partial_rotary_factor: read only from rope_parameters (base.py mirrors the top-level key automatically) - document why TOKEN_EMBD/OUTPUT are forced to F16 (they are the two dense tensors in Maple, and the reference GGUFs ship them as F16) - add @ModelBase.example("deepgrove/maple-preview") * tests: add Maple to the SWA pattern array list get_arr for maple.attention.sliding_window_pattern requires an array, but the harness only emitted a per-layer array for the arches in its list, so test-llama-archs -a maple failed to load the model. Assisted-by: DeepSeek Harness * maple: move swiglu_clamp_exp to the converter The loader prefilled 7.0 and read the key optionally. The converter now writes it and the loader reads it as required, because llama-graph.cpp skips the clamp when the limit is 0 and an optional read would silently run unclamped. The test harness provides the key for the same reason. Also drops tensor_force_quant: base.py already forces FFN_GATE_INP to F32 and TOKEN_EMBD/OUTPUT to F16 for ternary file types. Assisted-by: DeepSeek Harness * convert: fix the LazyBase func signature in the Maple converter ty flagged the stack() closure: it takes no argument, while LazyBase is annotated with func: Callable[[Any], Any]. Pass the tensor list through args instead of closing over it, the same way kimi_k3 does, so the callable shape matches. Assisted-by: DeepSeek Harness * tests(s390x): add non-vxe build to tests (#28776) * tests: add non-vxe build to tests Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> ggml-cpu: add unused macro to fix ci Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> Revert "ggml-cpu: temporarily add #28775 patch until its merged" This reverts commit d4645257b6b7e65c47b1b46baec3eb46a3f40968. Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ggml-cpu: revert back to upstream/master Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * models : fix mimo2 swa pattern load (#28865) * models : fix incorrect uses of get_key_or_arr (#28868) * tests : add fusion baseline README and broaden fusion CI triggers (#28893) * tests : add README for updating the per-backend fusion baselines Assisted-by: pi:llama.cpp/Qwen3.8-27B * ci : trigger fusion on changes to test-llama-archs.cpp and src/models the dummy models and their architectures drive the fusion baselines, so a change to either can alter the per-fusion counters and should re-run the fusion job. Assisted-by: pi:llama.cpp/Qwen3.8-27B * tests : merge the fusion build commands in the README assisted-by: pi:llama.cpp/Qwen3.8-27B * pi : require explicit permission before posting PR/issue comments assisted-by: pi:llama.cpp/Qwen3.8-27B * ggml : bump version to 0.24.0 (ggml/1627) * sync : ggml * llama.cpp : bump version to 0.4.1 (#28900) * scripts: Add script to verify API/ABI compatibility (#28579) * cmake : remove precompiled headers (#28892) This commit removes the precompiled headers that I added in Commit 3bcfeb700 ("cmake : add PCH and unity build to improve build times (#28091)"). The motivation for this is that this looked good when developing this but has caused multiple issues that I had taken into consideration and we have decided to remove it and only keep the unity builds from the above commit. Refs: https://github.com/ggml-org/llama.cpp/pull/28882#issuecomment-5662272126 * release : added gfx1103 to ubuntu rocm build (#28423) * qwen4exp: enable rms_norm + mul fusion (#28896) * qwen4exp: enable rms_norm + mul fusion * use TENSOR_ALLOW_RESHAPE * ci : add ubuntu-cuda builds to release (#28186) * release : add ubuntu-cuda build job (12.8/13.3, x64+arm64) * Add GCC 14 for CUDA arm64 builds in CI * Eplicit bash * Install git for CCCL fetch * Install git before we clone/checkout * Match CI names for WIndows * Whitelist llama.cpp repo to git * Use $GITHUB_WORKSPACE * Also ship dependent libs on Ubuntu Need NCCL additionally as it's pre-built available on Linux * Avoid duplicate files in packaged cudart * Copy NCCL license * Install CURL to fetch NCCL license * Update .github/workflows/release.yml Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Remove NCCL until licensing has been confirmed --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * HIP: fattn-mma: use fp32 accumulation on MFMA devices (#28576) use fp32 accumulators in fattn-mma on CDNA * CI: hip-quality-check: ignore spill added in bfdc32183d57f1e35bacf35c47d6311e2028bbbc (#28909) the kernel spills 5 registers but is still faster than before the change * ci : reuse build tag name when used instead of safe one (#28911) * webui: stop re-probing disabled /tools endpoint on every message (#28646) When /tools returns 403 (server started without tools), the web UI refetched the tool list before every chat message, since the guard treated an empty tool list as "not yet fetched". Each retry returned 403 and could trip fail2ban. Skip the refetch once the store flags the endpoint as disabled, and detect that state via the response status code instead of string- matching the error message. The tools panel keeps probing on open so the UI recovers once the server is restarted with tools enabled. Fixes #28299 * cmake : use PROJECT_SOURCE_DIR instead of CMAKE_SOURCE_DIR (#28771) This commit updates cmake to use PROJECT_SOURCE_DIR instead of CMAKE_SOURCE_DIR for paths in function calls. The motivation for this is that when using add_subdirectory, CMAKE_SOURCE_DIR is fixed to the top-level projects source directory, that is the caller of add_subdirectory and not the llama.cpp root which means that common/common.h header will not be resolved. Refs: https://github.com/ggml-org/llama.cpp/pull/28091#issuecomment-5636106377 * cuda : enable i16 and i32 for DUP (#28897) * cuda : enable i16 and i32 for DUP * docs : update ops table for DUP on CUDA * ci : fix android release (#28936) * ci: Bump CUDA Windows x64 builds to 13.4.1 (#28930) * metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3) (#28599) * metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3) MiniCPM3 sets attention.key_length to 96 and does not set attention.value_length, which defaults to n_embd / n_head = 64. Metal had no (96, 64) instantiation, so -fa auto aborted on the missing kernel_flash_attn_ext_vec_f16_dk96_dv64. Instantiate the tile kernel at (96, 64) for every K/V type that already has (96, 96), and the vec kernel for the NE=4 configurations. Of the NE values the vec dispatch considers, only NE=4 works here, because NL = 32/NE has to divide both DK/4 = 24 and DV/4 = 16. * tests : avoid redundant FA vec slice coverage * ci: bump kleidiai runners from 22.04 to 24.04 (#28885) * ci: bump kleidiai runners from 22.04 to 24.04 Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ci: promote warnings to hard errors for ci Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * opencl: add generic ssm_scan (#28881) * opencl: add generic ssm_scan * opencl: fix whitespace * OpenVINO: optimize stateful decode and GPU MoE inference (#28638) * exclude GPU/NPU failing POOL_2D case * Fix pool case * ggml-openvino: fix stateful decode for Gemma-4 per-layer-type head sizes * ggml-openvino: fix MSVC narrowing error in permute * ggml-openvino: classify sliding-window layers structurally on interleaved-SWA models * ggml-openvino: add GGML_OPENVINO_REQUANT_KQUANT to select a 4-bit requant target * ggml-openvino: add GGML_OPENVINO_SPILL_DIR to spill weight buffers to disk * Stateful Performance: Added pass::KVStateSeqAxis to change KV layout * ggml-openvino: fix stateful decode past the sliding-window size Assisted-by: Claude Sonnet * ggml-openvino: refuse stateful decode that cannot resume from the KV state The stateful path seeds its KV state from ggml's cache when the decode position is ahead of what the state holds. That only works when ggml's cache is a plain prefix, where cell i holds position i. A sliding-window layer keeps just the last n_swa positions and drops the rest, so past the window cell i no longer holds position i and the seeded state is wrong. Slicing the state to the decode position also had no bounds check, so a position past the end surfaced as a bare ov::Exception from the ROI constructor (llama_decode ret = -3, with no reason given at default verbosity). Refuse both cases with a clear message instead, and refuse on the compile path too, where a new model starts with an empty state and so can only serve a sequence from its beginning. Reproducible with llama-bench -d, which restores a saved sequence state rather than recomputing the depth prefill. Assisted-by: Claude Opus 5 * ggml-openvino: use the per-layer KV head count for the stateful KV state The stateful path reinterprets ggml's KV buffer [1, 1, seq, n_heads_kv * head_size] as [1, seq, n_heads_kv, head_size]. The head size is already taken from the tensor's own combined dim, because gemma-4 varies it per layer type, but the head count still came from a model-level scalar that compute_llm_params() overwrites per attention node, so it ended up holding whatever the last layer said. gemma-4 varies the head count per layer too: 12B has 8 x 256 sliding layers and 1 x 512 full layers, 31B has 16 x 256 and 4 x 512. So 40 of 12B's 48 layers were split as 1 x 2048 instead of 8 x 256, and attention read the state with the wrong head split - both models decoded garbage on CPU and GPU. E2B is unaffected, its head count is 1 everywhere. Record the count per layer instead and look it up by the cache_k_l<N> leaf name. Key it by layer, not by layer type: the sliding/full classification comes from cache extents, which tie at a small -c, while the head count does not. The stateful state trim now derives its sequence axis per state for the same reason, since pass::KVStateSeqAxis matches per state on the head count. Assisted-by: Claude Opus 5 * ggml-openvino: apply the KV state relayout to any KV head count pass::KVStateSeqAxis was limited to states with a single KV head, where moving the sequence axis from dim 1 to dim 2 is a pure metadata change. The limit was also based on a measurement showing no gain for a multi-head model, but that was taken at depth 0, which is the one depth where this change does nothing. With several heads the pass does more than move metadata: it drops the reader side transpose of the whole accumulated state, which the graph otherwise redoes every token at a cost that grows with the context length, and replaces it with a transpose of the single new row. Measured on GPU, tg128, alternating arms: gemma-4-12B 6.27 -> 9.11 t/s at depth 8192 (stateless is 7.69, so stateful now wins at depth instead of losing), Llama-3.2-1B 47.8 -> 59.6 t/s. Both are within noise at depth 0, which is why the earlier check saw nothing. The state refill needs the rows copied rather than reinterpreted now: ggml stores [seq][n_heads_kv * head_size], and a relayout state with several heads is a different element order. Without that, a refill would seed wrong data - it is reachable today through llama-bench -d. Assisted-by: Claude Opus 5 * ggml-openvino : support ggml_rope_set_offset and simplify op support gating * add more cpy cases * reject BF16 cpy on NPU * Remove mul_mat_id fallback, gate large mul_mat_id only for mxfp4 * ggml-openvino: fuse the MoE expert block into MOECompressed on GPU * ggml-openvino: skip GPU MUL_MAT_ID for unbound expert tensors * ggml-openvino: requantize grouped 8-bit MoE experts on GPU * Enable special strided CPY for conv state writeback * openvino: support cacheless encoder models on NPU Packed QKV views used by mmBERT were rejected by the ROPE support check. This split Q/K RoPE onto CPU, prevented cacheless attention detection, and sent fragmented encoder graphs through the decoder-oriented NPUW path. Accept packed QKV RoPE views, detect cacheless attention from its mask, and run these models as a single full-sequence prefill without NPUW or a decode graph. Also provide static mask, output index, and mean-pooling shapes and inputs. * openvino: optimize norm and RoPE translation Replace the decomposed mean/variance normalization graph with an opset6 MVN operation. This preserves the GGML epsilon placement while allowing OpenVINO plugins to compile normalization as one operation with fewer intermediate tensors. Cache RoPE sine and cosine outputs in the graph-wide tensor map. Build the cache key from all RoPE parameters and the optional frequency-factor input so compatible Q/K and layer nodes share one subgraph without mixing different RoPE configurations. Expose NodeContext::put_shared() to publish translator-created outputs for graph-level reuse. * ggml-openvino : simplify op translators and enable IMROPE/NEOX RoPE fusion * remove unnecessary include and clean up PAD * fix mulmat bug * use ov::as_type_ptr instead of std::dynamic_pointer_cast * ggml-openvino: fix mixed-dtype ADD/SWIGLU_CLAMP, gate unsupported ROPE/SOFTPLUS cases - translate_add: upcast mismatched operand types (e.g. f16/f32 in fused ADD_ADD) to f32, add, then cast once to the output type. opset1::Add requires matching input types and downcasting first lost precision. - translate_glu_swiglu_clamp: same fix, f16 Swish/Clamp rounding was drifting past the test tolerance. - supports_op: reject ROPE with ne[3] > 1 (multi-sequence) since the cos/sin tables only cover one sequence, and SOFTPLUS on GPU since the OpenVINO GPU kernel overflows to inf for large inputs (CPU is fine). - ci/run.sh: serialize test-backend-ops on OpenVINO GPU; running two workers concurrently crashes the GPU plugin (CL_OUT_OF_RESOURCES). * openvino: share compiled models with per-context inference state; fix thread-safety * ggml-openvino: gate MoE expert-sum ReduceSum shortcut past 8 experts The ReduceSum shortcut for the MoE expert-plane-sum ADD chain drifts past the 1e-7 test tolerance for >8 experts (f32 accumulation order vs CPU reference), intermittently, like the existing Q4_K/Q5_K NMSE case. Expose is_moe_expert_sum_add() so supports_op can gate on expert count and fall back to CPU for just that reduction op. * ggml-openvino: gate degenerate m=1,n=1 MUL_MAT on GPU CI hit ERR=1.8e-3 (> 5e-4 tolerance) for a scalar-output f32 dot product (m=1,n=1,k=2048); didn't reproduce locally in 8 tries, so likely an internal fp16 accumulation path the GPU plugin picks for this tiny shape. m=1 output dim doesn't occur in real model weights, so gate it. * ggml-openvino: make SoftPlus decomposition opt-in native Assisted-by: Codex --------- Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com> Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com> Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com> Co-authored-by: ravi9 <ravi.panchumarthy@intel.com> * vulkan: support sparse Flash Attention (#28105) * vulkan: add sparse Flash Attention support for DSV4/GLM * tune implementation * add tests * avoid nondeterministic atomicAdd * add cm2 decode vector support * simplify logic and make variable names more consistent * add cm2 f16vec4 binding for decode vector * models : move build_arch_graph() after graph() template specialization (#28934) Move build_arch_graph()'s function definitions after the graph<true> and graph<false> template specializations have been explicitly defined. * cuda: support row-contiguous SUM_ROWS (#26308) * cuda: support row-contiguous SUM_ROWS * organize the code and add GGML_OP_MEAN to support row-contiguous tensors using the same shared kernel, and add a test to MEAN permute/slice * Keep original comments and add if/else branch * rpc : hash-cache only weights (#28789) * rpc : hash-cache only weights ggml_backend_rpc_buffer_set_tensor and ggml_backend_rpc_set_tensor_async hashed every transfer above HASH_THRESHOLD and let `rpc-server -c` serve it from its file cache. The cache is meant for weights, but the activations ggml_backend_sched copies between backends took the same path: with a two-node split of Qwen3.8-Flash-Next every prefill ubatch above 10 MB was hashed, written to the worker's cache directory (1.4 TB after a day) and later served from there. Use the hash path only for tensors in buffers marked GGML_BACKEND_BUFFER_USAGE_WEIGHTS. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * rpc : save a cache entry only for the tensor that missed the hash check With the client hashing weights only, the server still wrote every SET_TENSOR above HASH_THRESHOLD to the cache directory, so the compute data the scheduler sends kept filling the disk. Remember the hash of the last SET_TENSOR_HASH that missed and save only the SET_TENSOR that follows it with that hash - the weight the client is re-sending. * rpc : signal the cache decision in the SET_TENSOR payload Replace the server-side `pending_cache` state with a `cache_flag` byte in the SET_TENSOR message: the client sets it when SET_TENSOR_HASH reported a miss, the server saves a cache entry only when it is set. Bump RPC_PROTO_MAJOR_VERSION since the wire format changes. --------- Co-authored-by: Patrick Hoffmann <patrickhoffmann@MacBook-Pro-14-HOP.local> Co-authored-by: Claude Opus 5 <noreply@anthropic.com> * docs: Rule of thumb for AI review time [no ci] (#28945) * ci: build MUSA for only 1 arch (#28944) * ci: optimize * keep only the MUSA changes * opencl: choose the MoE expert matmul by batch size for speculative decoding/MTP (#27637) * opencl: gate the prebuilt q4_0 MoE GEMM on routing count * opencl: stop writing zeros into the padded MoE activation slots * opencl: rephrase claude's comments --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * HIP: Enable AllReduce for ROCm (#27825) * hex-cpy: use dma if src and dst are contiguous (#28906) * hexagon: add back missing contiguous fast-path and hvx_copy_uu for each run (#28886) * llama-bench: support --version to print build info (#28971) * ci : add self-hosted webgpu to hf-jobs (#28712) * add self-hosted vulkan and webgpu to hf-jobs * try t4-medium * cont : adjust cpu backend threads * try t4-small again * restore cm jobs --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * metal: fix NaN in mul_mm_id when activations exceed f16 range (#26223) * test-backend-ops: reproduce MUL_MAT_ID NaN for activations beyond f16 The Metal mul_mm_id path narrows src1 to `half` for the simdgroup MMA (`S1 = half` in every instantiation; ggml-metal.metal:10582 and :10595, mirrored at :10643/:10654 in the tensor-ops path). f16 saturates at 65504, so a model whose activations exceed that produces inf, and `simdgroup_multiply_accumulate` then turns the whole 8x8 accumulator tile into NaN. The mul_mv_id path used below `ne21_mm_id_min` (32) carries the same values in f32 and is correct, as is every CPU path. This was untestable before: `init_mul_mat_id_tensors` initializes uniform [-1, 1], so no existing case can drive an operand out of f16 range. `test_mul_mat_id` gains an `amax` parameter (default 1.0f, preserving the historical init exactly) that scales only the f32 activations, leaving the quantized weights in their normal range. Six cases: n=16 sits below the mul_mv_id -> mul_mm_id switch and is the control that must stay green; n=32 and n=64 are above it and fail on Metal today. Two shapes, because this is not model- or size-specific — q4_K at 128 experts / 4 active / 4096x2048 mirrors a real model, and q8_0 at 8 experts / 2 active / 512x256 shows the same failure at minimal size. Observed on Apple M2 Max, macOS, llama.cpp b10156: MUL_MAT_ID(type_a=q8_0,...,n=32,k=256,amax=100000.000000): [MUL_MAT_ID] NaN at index 0 (MTL0=nan CPU=583442.375000) FAIL The real model behind this is Mistral Small 4 (arch mistral4, 128 experts / 4 active), one of whose layers reaches ~1e5 activations: on Metal every prefill of >=32 tokens returns an entirely NaN vocabulary, while <32 tokens is correct. Note kernel_mul_mm (dense) has the identical conversion at :10273 and :10286 and is expected to fail the same way; it is not covered here. Found and written by Claude Opus 5 (via Claude Code). * metal: fix NaN in mul_mm_id when activations exceed f16 range kernel_mul_mm_id narrows src1 to `half` for the simdgroup MMA operands (`S1 = half` in every instantiation). f16 saturates at 65504, so a model whose activations exceed that produces inf on load, and simdgroup_multiply_accumulate then propagates NaN across the whole 8x8 accumulator tile. The result is an entirely NaN output — not a precision loss, a total loss. The mul_mv_id path taken below ne21_mm_id_min (32) keeps the same values in f32 and is correct, as is every CPU path, so the same model produces correct logits for short inputs and NaN for long ones. Fix: rescale src1 by a power of two so it fits, and undo the scale on the f32 accumulator at the store. A two-stage reduction computes max(|src1|) and writes the pair (1/scale, scale) into scratch chained off the destination buffer, in the same style as the existing tpe/ids id-mapping scratch. The matmul multiplies on load and on store. This is exact, not approximate, for two reasons: the dot product is linear, so one tensor-wide factor commutes through the accumulation; and the factor is a power of two, so both multiplications are exact in binary floating point. When max(|src1|) already fits — every model that works today — the factor is exactly 1.0 and the output is bit-identical to before. Accumulation was already f32 and is unchanged; only the operand narrowing was ever the problem. The reduction is two-stage (256 threadgroups into partials, then one threadgroup folding them) specifically so it stays bandwidth-bound. A single-threadgroup version was measured first and cost up to +451% median on prefill — the scan serialized against an otherwise idle GPU. It is also dispatched only on the mm path, so decode never pays for it. Measured on Apple M2 Max, `test-backend-ops perf -o MUL_MAT_ID -b MTL0`, 99 cases, versus the same build without this change: n=1/4/8 (mul_mv_id, decode) : -0.8% / -0.8% / -0.4% median (noise) n=32 (mul_mm_id, prefill) : +1.73% median n=64 : +1.30% median n=128 : +1.80% median n=256 : +3.98% median n=512 : +3.74% median, +7.20% worst overall : +1.14% median Correctness, same machine: - the six new test-backend-ops cases go from 4 FAIL / 2 OK to all OK, with the n=16 controls (mul_mv_id path) unchanged; - `test-backend-ops -b MTL0` full run: 0 failures, no regression; - Mistral-Small-4-119B (arch mistral4, 128 experts / 4 active) now generates correctly at the default n_ubatch of 512, in both UD-IQ3_S and UD-Q4_K_XL quantizations. Before this, every prefill of >= 32 tokens returned an all-NaN vocabulary and only n_ubatch <= 31 (forcing the mul_mv_id path) worked. Likely fixes #25722 (mistral4 empty output on Metal above ~300 tokens, FA on and off, generation degenerating to a single control token — the signature of argmax over an all-NaN distribution). #20668 may be the same defect attributed to a bad GGUF. Note kernel_mul_mm (dense) has the identical narrowing at the corresponding load sites and is expected to fail the same way; it is left alone here to keep this change reviewable. Also possible, and left for later: scaling per output column rather than per tensor, which would preserve more precision when a single token is the hot one. Found, diagnosed and fixed by Claude Opus 5 (via Claude Code). * metal : make requested edits - remove verbose comments - explain rationale as requested Generative AI disclosure: Claude made the edits as requested. * metal : stack mul_mm_id map0 with amax_part Implement @ggerganov suggestion to stack amax_part + map0. Mean 2.6% faster (worst -0.7%, best -4.1%). Win grows with batch size. Benchmarked on a hot M2 Max after reboot. Generative AI disclosure: Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * cont : fix var scope * cont : comment out tests temporarily Comment out tess to not break CI temporarily Assisted-by: Claude Fable 5.1 --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * vulkan: make MUL_MAT_ID BN/2 tail unconditional (#28923) Use BN/2 as the default for BNover2 and as the disabled fallback for BNover4, and remove the enable gate from the MUL_MAT_ID BN/2 branch. The BN/4 branch remains gated by enable_smaller_matrices, while the p.N path is unchanged. * chat : force `\n</think>` on reasoning budget end for qwen3-coder (#28869) * HIP: broaden MoE ncols_opt tile heuristic on RDNA3.5 architecture (#28935) It's found the MoE ncols_opt tile heuristic needs to be broadened to include the RDNA3.5 architecture. The code change is implemented in ggml/src/ggml-cuda/mmq.cu and just change the GGML_CUDA_CC_IS_RDNA3_0 to GGML_CUDA_CC_IS_RDNA3 in the condition. The dense dispatch logic remains unchanged. The Test machine configuration we used is AMD Radeon 8060S, gfx1151 (RDNA3.5), 20 CU, wave32 + AMD Ryzen AI MAX+ 388, 8C/16T, 23.79 GB RAM we complete the Correctness verification and performance evaluation as follows: test-backend-ops test -b ROCm0 -o MUL_MAT -p type_a=<q4_K|q5_K|q4_0|q5_0> test-backend-ops test -b ROCm0 -o MUL_MAT_ID -p type_a=<q4_K|q5_K|q4_0|q5_0> all pass: MUL_MAT 64/64, 29/29, 48/48, 14/14; MUL_MAT_ID 84/84, 3/3, 74/74, 3/3 Performance result on target machine: LFM2.5-8B-A1B-UD-Q4_K_M (Q4_K MoE) +16.198% [+12.704, +19.799] 8/8 Qwen1.5-MoE-A2.7B-Q2_K (Q2_K MoE) +6.189% [ +5.245, +7.141] 8/8 pooled (16 pairs) +11.081% [ +7.972, +14.279] 16/16 Token generation (tg128) is unchanged on the Q4_K MoE model and +2.188% [+0.905, +3.488] on the Q2_K one. * qwen4exp: add hc ops (#28901) * Change max context length for auto-fitting with unified KV (#28849) * rpc : invalidate cached compute graph when a referenced buffer is freed (#24292) The server caches the most recent compute graph per device so that GRAPH_RECOMPUTE can re-execute it without resending tensor data. The cached graph nodes hold direct pointers to backend buffers that were live at graph_compute() time. If any of those buffers is later released via FREE_BUFFER, the next GRAPH_RECOMPUTE re-executes the cached graph through the dangling pointers (use-after-free). The bug is reachable by an unauthenticated remote client. The dangling pointers point into chunks an attacker can reshape via subsequent ALLOC_BUFFER/SET_TENSOR commands, and the resulting read/write through the cached graph is sufficient to leak libc addresses and hijack the buffer iface vtable used by BUFFER_CLEAR, yielding remote code execution. Discard all cached graphs in free_buffer(). The existing null-check in graph_recompute() then rejects the request and the client falls back to GRAPH_COMPUTE on the next call. No protocol or API change. * spacemit : fix wrong transpose function for int16 data (#25161) The `sizeof(int16_t)` branch in `permute_transpose_impl` calls `rvv_transposed_s32_mn_to_nm` instead of `rvv_transposed_s16_mn_to_nm`. This is a copy-paste bug from the `sizeof(int32_t)` branch above it. The s32 function uses 32-bit segment load/stores (`vssseg8e32.v`) on 16-bit data, reading 2x bytes per element and producing completely wrong transposition results -- 14 out of 16 positions are corrupted for a 4x4 int16 matrix. The correct function `rvv_transposed_s16_mn_to_nm` already exists (line 390) and is used elsewhere in flash attention (line 1488). * CUDA/HIP: improve access patterns in im2col (#28013) * model : add support for HrmTextForCausalLM (DFM Mimir 1B) (#27625) * model : add support for HrmTextForCausalLM (DFM Mimir 1B) HRM-Text runs two transformer stacks (low, high) in an alternating cycle over the same token stream. The low-cycle state z_l starts from a learned [n_embd] tensor and is broadcast over positions. - conversion: new writer for the fused gqkv projection (order gate,q,k,v) remapped to llama.cpp q/k/v plus a separate sigmoid gate tensor - loader: block_count = lps * h_cycles * (l_cycles + 1) cache slots aliasing 2*lps physical blocks via struct copies - graph: looped build with sigmoid-gated attention, SwiGLU FFN and parameterless RMS norms; learned embedding_scale applied in build_inp_embd - saver: pointer-deduplicated layer loop (looped archs alias tensors) - tests: hrm_text fixture (lps 1, h 2, l 3) in test-llama-archs Limitations: causal attention only - the upstream prefix-LM mode is not implemented (the prefix_lm GGUF key round-trips unused). The KV cache holds one entry per pass: 128 layers for Mimir 1B, i.e. 4x a same-width 32-layer model - about 3072 MiB at ctx 4096 in F16 (halves with q8_0 KV + FA). Every token runs all 128 block passes, so decode cost is roughly 4x a dense model of equal width (2.65 t/s BF16, 8-thread desktop CPU). Verified against the HF reference: identical argmax at 334/334 positions across 20 prompts (BF16 GGUF vs FP32 golden). q8_0 requant: 95.8% top-1, all remaining misses inside the HF top-5 (accumulated error over 128 sequential blocks). AI usage disclosure: YES Used GLM-5.3 for the majority of code AI-generated under my direction, all gates verified locally. All in all I could say that I have written less than 20% of the code and most of the heavy lifting has been done by the model. As such, this should be considered experimental. * Update conversion/hrm_text.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update src/llama-arch.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * convert : add gguf_writer methods for hrm_text metadata replace raw add_uint32/add_bool calls with dedicated GGUFWriter methods, following the add_embedding_scale pattern Assisted-by: GLM-5.3 * convert : map regular hrm_text tensors via tensor_mapping delegate unfused checkpoints to the base tensor mapping; training-style attn. names are renamed to self_attn. so the patterns match Assisted-by: GLM-5.3 * model : format hrm-text build_* calls as in other models one argument group per line, matching sibling model files Assisted-by: GLM-5.3 * llama : move hrm z_l_init table entries out of the nemotron group place the name and tensor-info entries with the other global input tensors Assisted-by: GLM-5.3 * convert : slim down hrm_text comments Assisted-by: GLM-5.3 * convert : build hrm_text block tensor names from the {bid} template The tensor map holds concrete per-block names, so format the template with the computed layer index before handing it to super(). * llama : name hrm metadata keys in their own hrm. namespace The four keys are arch-independent, unlike the arch-substituted Keys.LLM entries, so group them under Keys.HRM (like Keys.Split) and rename the llm_kv entries to LLM_KV_HRM_*. Only our own GGUFs carry the old hrm_text.* keys; they are regenerated. * Update src/llama-model-saver.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * llama : keep hrm metadata keys arch-substituted Per review: the GGUF keys stay "{arch}.h_cycles" style, so the Python members drop the LLM_KV_HRM_ prefix and keep arch templates; C++ keeps the LLM_KV_HRM_* enums. GGUF output is unchanged - existing files and HF uploads stay valid. * Update gguf-py/gguf/constants.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update src/llama-arch.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update src/llama-arch.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * convert : rename hrm writer methods to add_hrm_* Generic names like add_h_cycles/add_prefix_lm are too broad on the shared GGUFWriter; prefix them with hrm_ like the metadata keys. * model : fix meta-split lookup for archs with aliased cache slots Cache tensors of archs that alias physical blocks across looped slots (hrm_text, nanbeige with num_loops > 1) can reference block indices without weight tensor names. Take the output projection from the layer array instead of asserting; all other lookups are unchanged. * model : replicate hrm_text tensors on meta devices instead of splitting The aliased cache slots rotate split states differently from their physical weights, so the meta-split execution invariants (set_rows requires the cache state to match the token indices) cannot hold for any device count. Replicate all hrm_text tensors on every meta device instead; single-device and non-meta paths are unchanged. Assisted-by: Claude Sonnet --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * models : allow Nemotron-H models to only define layer_norm_epsilon (#28989) * allows nemotron models to get by with just defining layer_norm_epsilon * made changes to load_arch_hparams instead * hexagon: accept the zeroed rope probe in supports_op (#28995) llama probes weight placement with a rope where all params are 0, so rejecting n_dims == 0 or freq_base == 0 puts rope_freqs on the CPU. That splits the decode graph at every full-attention layer (gemma-4-E2B: 5 splits instead of 2). Assisted-by: Claude Opus 5 * hexagon: Support for K-Quants Q4_K and Q6_K (#28994) implement q6k/q4k kernels Squashed from: feat: implement q6k kernel hex-q6k: improve unpack accuracy hex-q4_k: add support for Q4_K kernels Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com> * Enable CUDA graph for MTP draft (#28549) * Improve CUDA graph usage for MTP * Rename field * Address review feedback * ci: switch fast jobs back to github (#28959) * switch jobs to ubuntu-slim * ubuntu slim almost takes 15 minutes for check requirements so use something faster * TP: fix split state and granularity for fused QKV gemma4, qwen35 (#28965) * model: calculate split states for attn_qkv from n_head * n_embd_head_k required for gemma4 with --fuse-qkv, where n_embd is 5376 but Q is 8192. * model: handle fused full attention layers for qwen35/qwen35moe * model: add TODO: [TAG_SPLIT_QGATE_QWEN] * vulkan: work around NV bug with argsort_large.comp (#28975) * [SYCL] Fix function signature for `ggml_backend_sycl_split_buffer_type` (#28981) * vulkan: support qwen4exp hc ops (#28988) * vulkan: support qwen4exp hc ops * fix stale comment [no-ci] * vulkan: fix buffer_reference alignment in im2col shaders (#28996) Both im2col.comp and im2col_3d.comp declare D_ptr without an explicit buffer_reference_align, so glslang emits writes through it as Aligned 16. The shaders advance the pointer by D_SIZE, a per-variant define set to 4 for float and 2 for float16_t, so most write addresses are not 16-byte aligned. This triggers VUID-RuntimeSpirv-PhysicalStorageBuffer64-06315 under GPU-AV. Declaring buffer_reference_align = D_SIZE matches the alignment to the actual write stride and takes validation hits from 20 to 0 for both IM2COL and IM2COL_3D. Fixes #28960 * docs: remove JG as CODEOWNER for test-llama-archs (#29003) * opencl: fix various warnings (#28984) * opencl: fix warnings * opencl: fix warnings for non adreno * sycl: ssm_conv: fuse the SiLU epilogue into the ssm_conv kernel (#28929) * vulkan: skip unneeded MoE work in mul_mm coopmat1 path (#25483) * sycl : fix the B70 mem allocate error when >19.3GB (#28953) * gguf : align the data section relative to the GGUF start, not the file (#28993) * gguf : align the data section relative to the GGUF start, not the file gguf_init_from_file_ptr reads a GGUF from the current file position, but padded the data section from file offset 0, so a GGUF embedded at an offset that is not a multiple of the alignment loaded without error and returned wrong tensor data. Also adds llama_adapter_lora_init_from_file_ptr, and disables mmap with a warning when an embedded data section is not aligned, instead of asserting in ggml. Assisted-by: Claude Opus 5 * llama : load lora from path through the FILE* variant The test now checks that mmap is disabled only for an unaligned offset. Assisted-by: Claude Fable 5.1 * Update ggml/src/gguf.cpp Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * Update include/llama.h Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * llama : error on unaligned mmap of an embedded GGUF, drop test-load-file-ptr --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * chat : add message delimiters to the DeepSeek V3.2/V4 parser (#29008) * chat : add message delimiters to the DeepSeek V3.2/V4 parser Assisted-by: Claude Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * ci : add API/ABI check to make-release workflow [no ci] (#28947) * ci : add API/ABI check to make-release workflow [no ci] This commit adds an API/ABI compatibility check to the make-release workflow. The motivation for this to allow us to detect any potential breaking changes in API/ABI compatibility between releases and fail the the release if there are any. The workflow can be triggered manually as before and this check can be skipped if needed as it does take some time which might be useful when doing a dry-run and not specifically interested in the API/ABI check. By default this will check the current release against the latest release, but this can also be configured in the workflow, or in the script run on the command line, to check a different tag. * add check for minor version bumps [no ci] This commit also changes the build type to be RelWithDebInfo so that the reported information is more useful. * vulkan: split buffers and debug code into separate files, add shared headers (#28732) * ui: fix removed reasoning menu in single model mode on desktop (#27985) * ui: fix accidentally removed reasoning menu in single model mode on desktop * ui: formatting task run to fix storybook test * ui: mount the add menu reasoning submenu outside router mode only The models selector already owns the reasoning submenu in router mode, so the add menu only mounts it in single model mode. The first enabled item of the add menu is now the reasoning submenu, the accessibility story expects it. --------- Co-authored-by: Ben Babik <work@benjaminbabik.com> Co-authored-by: Pascal <admin@serveurperso.com> * openvino : Update OpenVINO to 2026.4;fix clangd,MSVC warnings; (#29009) * Update to openvino-2026.4 * Update OV docs * ggml-openvino : fix clangd and MSVC warnings * fix int to ptr cast, more internal linkage enforcement, and avoiding duplicate switch case --------- Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com> * model : extend Nemotron MTP support (#29018) * first fix * removed unnecessary declarations * model : skip gate_up_exps if TENSOR_SKIP is set (#29014) required for qwen35moe if MTP tensors are fused but not loaded * rpc : skip ACCEL devices (#29020) * ci : add missing evict-old-files (#29041) * vulkan: raise the hoisted row-id limit for mul_mat_id from 256 to 512 experts (#28501) * vulkan: raise the hoisted row-id limit for mul_mat_id to 512 experts The expert-count shader (count_experts.comp) sizes its shared arrays with BLOCK_SIZE, which is 256. Because of that, row-id hoisting is switched off for any model with more than 256 experts, and every mul_mat_id workgroup has to rescan the whole ids tensor on its own. Qwen3.8-Flash-Next has 512 experts and was quietly running on that slow path. This change sizes the arrays with a separate MAX_EXPERTS constant (512), clears them in a loop instead of one entry per thread, and raises the matching limit on the host side. On Strix Halo at batch 2048 the expert matmuls drop from 12.5 to 9.5 ms (iq3_s) and from 14.0 to 7.5 ms (iq4_nl) per op, and prompt processing gets about 19 % faster at 8k tokens. test-backend-ops MUL_MAT_ID passes (891/891) with new 512-expert test cases. Assisted-by: Claude Fable 5.1 * vulkan: raise the hoisted row-id limit for mul_mat_id to 1024 experts Follow-up to review feedback: 1024 matches LLAMA_MAX_EXPERTS instead of stopping at 512. The three shared arrays in count_experts.comp grow to 3 * 1024 * 4 = 12 KiB, which fits the 16 KiB that Vulkan guarantees for maxComputeSharedMemorySize. Adds mul_mat_id test cases at 1024 experts alongside the existing 512 ones. test-backend-ops MUL_MAT_ID passes on Vulkan (RADV, Strix Halo, Radeon 8060S): 889/889. * ci : bump android-actions/setup-android to 4.0.4 (#29065) * ci : disable GHA cache for copilot (#29068) * gguf-py: fix Q8_1 block size in GGML_QUANT_SIZES (2+2+32) (#29036) * gguf-py: fix Q8_1 block size in GGML_QUANT_SIZES * --whitespace --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * cmake : fix build when GGML_CPU=OFF and GGML_CUDA=ON (#29026) * fix: build fails when GGML_CPU=OFF and GGML_CUDA=ON * fix: eol in examples/convert-llama2c-to-ggml/CMakeLists.txt file * vocab : add ufakzeka pre-tokenizer (#29033) * vocab : add ufakzeka pre-tokenizer * vocab : move ufakzeka to the models list and regenerate the hash mapping * vulkan: add IQ3_S MMQ matmul kernels (#28822) * vulkan: add IQ3_S MMQ matmul kernels * Make block_a_to_shmem do 2-byte loads (110 bytes is divisible by 2) * Align the check, IQ3_S is also using K tile size * ggml : handle graph buffer reservation failure (#26070) * ggml-webgpu: fix supports_op condition for GET_ROWS (#28978) * fix get_rows vec4 handling * Add src strides checking to vec4_aligned of get_rows and the new test case. * ci: change ubuntu-latest to ubuntu-24.04 (#29079) * Model-Saver: Write the SWA pattern, 15 more architectures roundtrip (#29042) * llama: read the SWA pattern as a period or a per-layer array Add llama_model_base::load_swa_pattern(), which reads sliding_window_pattern either as one flag per layer or as a period expanded by set_swa_pattern(), and use it in every loader that reads the key as a period. These loaders silently ignored an array and applied their default period, although the converters of olmo2, gemma3n and exaone4 write arrays. The published GGUFs match the defaults, so their outputs do not change. The loaders that already accepted both forms lose their duplicated scalar-then-array block, and use their declared default period when the key is absent. * model-saver: write the SWA pattern and the MLA SWA geometry Write sliding_window_pattern as one flag per layer, nextn layers included, for every model using SWA. The array is never collapsed to a scalar, since the loaders read a scalar as a period. Also write the MLA key/value lengths and KV LoRA rank of the SWA layers, required by dots3note. This enables the saver for plamo3, gemma3, cohere2, cohere2moe, olmo2, exaone-moe, afmoe, mimo2, spark2_5, muse-glimmer, mellum, laguna, granite_swa, dots3note and maple, all passing the bit-exact roundtrip of test-llama-archs. * ggml : check for allocation failures to prevent crashes (#28149) * ggml : check for allocation failures to prevent crashes * wording * ggml-cpu: add F16 input to the FWHT (#27779) * ggml-cpu: add F16 input to the FWHT The CPU FWHT accepts F32 input only. This change makes the source type a template parameter. The CPU path now accepts F16 input and F32 input. The CPU MUL_MAT reference now converts an F16 src1 to F32. It does this when the caller sets the Hadamard hint. No backend has an F16 FWHT kernel yet. The test cases come with the backend changes that add one. * ggml-cpu: assert the F16 FWHT input path, and use the bulk converter Address review feedback. The F16 branch writes plain floats into wdata, which is only correct when vec_dot_type is F32. That invariant held because supports_op only accepts an F16 src1 for the Hadamard hint with F32 src0 and dst, but nothing enforced it. Assert it next to the existing src1 type check so widening supports_op cannot silently break the write. Replace the hand-rolled conversion loop with ggml_cpu_fp16_to_fp32. * opencl: add bin kernel `kernel_gemm_noshuffle_q6_k_f32_32b_trans_ila_a8_bin` (#28678) * opencl: add A8 Q6_K non-MoE binary kernel * opencl: fix layout compatibility * hexagon: HMX flash-attention head_dim padding (support DK=DV=72) (#26539) Allow HMX flash-attention to run with head_dim not a multiple of 64 (e.g. SigLIP head_dim=72), by operating on DK/DV rounded up to 64 with zero-filled tail lanes. * hexagon: im2col update (#29103) * ggml-hexagon: accept 1D and padded IM2COL ops * ggml-hexagon: make pure-DDR IM2COL kernel is_2D-aware * ggml-hexagon: extend IM2COL DMA patch-embed fast path to 1D * ggml-hexagon: add blocked-staging general IM2COL DMA kernel * hexagon: add ROLL op support (#29105) * opencl: add support for bin kernel `flash_attn_f32_f16_bin` (#29046) * opencl: add `flash_attn_f32_f16_bin` * opencl: guarded prefill fa * cuda : fix CUB argsort corruption caused by in-place keys (#28389) argsort_f32_i32_cuda_cub called the one-shot DeviceRadixSort::SortPairs API with d_keys_in == d_keys_out (temp_keys, temp_keys). CUB's internal double-buffer ping-pong requires distinct key buffers: with aliased buffers the sort partially overwrites its own input mid-pass and emits a corrupted permutation, surfacing as intermittent garbage indices (e.g. backend top_k over a 248k-column vocab on Maxwell/CUDA 12.5/CCCL 2.x, which then triggered out-of-bounds gathers in downstream get_rows). Use a distinct keys-out buffer for all six call sites (plain and segmented, ascending and descending, size-query and execute). --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Oliver Simons <osimons@nvidia.com> * metal : support qwen4exp hc ops (#29000) Add support for the new DSV4 HC op variants used by qwen4exp: - hc_pre with per-element sigmoid gate (gated variant) - hc_post with identity mixing (comb == nullptr) Assisted-by: pi:llama.cpp/Qwen3.8-27B * test-llama-archs : generate dummy test vocab (#29084) Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : fix FA support checks (#29122) * metal : add MoE and SSM_CONV fusion optimizations (#28948) * metal : add top-k MoE fusion Adds a Metal fusion for SOFT_MAX + ARGSORT + GET_ROWS with optional routing-weight normalization and scale, matching the top-k MoE fusion available in the CUDA and Vulkan backends. The fused kernel writes the selected expert ids and routing weights directly, eliding the separate softmax, argsort, get-rows, sum-rows, clamp, div and scale kernels. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : add MoE weighted reduction fusion Fuses MUL(experts, weights) plus the expert VIEW/ADD chain into one kernel that computes the weighted sum directly. The graph_optimize hook keeps the expert and weight buffers alive until the fused output so the allocator cannot reuse them while the kernel is still reading them. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * tests : expose MoE weighted reduction in fusion baseline Use 2 experts per token in the generated MoE test models so the Metal MoE weighted reduction fusion (MUL + ADD) is exercised by test-fusion. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : fuse RMS_NORM + SCALE Adds NORM/RMS_NORM + SCALE fusion to the Metal backend by reusing the norm+mul kernel with a scalar scale flag. Adds test coverage for both NORM+SCALE and RMS_NORM+SCALE and regenerates the fusion baseline. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : use function constant for RMS_NORM + SCALE Replaces the runtime use_scale karg with a Metal function constant. The norm+mul kernel is compiled with FC_norm_use_scale=false for MUL fusion and FC_norm_use_scale=true for SCALE fusion, so the fused kernel has no runtime branch. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : use function constant for top-k MoE with_norm Replaces the runtime with_norm karg with a Metal function constant. The top-k MoE kernel is compiled separately for the normalized and non-normalized routing variants, removing the runtime branch. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : rename moe_weighted_reduction suffix to moe_reduce Shortens the MoE weighted-reduction fusion identifiers, kernel, pipeline, matcher, args struct, and test op name from moe_weighted_reduction to moe_reduce. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : add MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion Adds dense mat-vec activation fusion for sigmoid/silu and bias+softplus. The mat-vec kernels apply the activation/bias epilogue via function constants, avoiding the separate unary/add passes. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : revert MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion The mat-vec activation fusion regressed decode throughput on Qwen3.6-35B-A3B by ~8% (tg32 81.5 vs 88.5 t/s). The regression is caused by loss of concurrency: the standalone unary kernels previously overlapped with other mat-vec work, while fusing the activation into the mat-vec kernel serializes it on the critical path. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : add SSM_CONV + UNARY (silu) fusion The SSM_CONV kernels apply silu directly via a function constant, eliding the separate unary pass. Regenerates the fusion baseline. Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : address fusion review comments - Fix declaration/table alignment - Rename top-k MoE kargs fields to val_clamp / val_scale - Move moe-reduce alloc-deps handling into a general fusion helper - Remove the public moe-reduce matcher API Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : fix unused parameter in top-k MoE fusion check Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : guard SSM_CONV fusion lookup behind use_fusion Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : track all fused outputs in graph reorder Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : keep top-k MoE logits alive until fused output Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : refactor alloc deps to pattern-driven approach Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : check fused kernel destination in concurrency tracking Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * meta : forward graph_optimize to underlying backends Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : use vector for fusion table Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * meta : keep graph_optimize unimplemented Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * parallel : fix non-deterministic prompt selection Assiste…

github-actions Bot added model Model specific testing Everything test related Vulkan Issues specific to the Vulkan backend ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Aug 31, 2026

LynxPDA mentioned this pull request Sep 10, 2026

vulkan: sparse prefill flash attention for qwen4exp top-k masks Nathanw1014/llama.cpp#10

Closed

0cc4m added 3 commits September 11, 2026 11:52

vulkan: add sparse Flash Attention support for DSV4/GLM

f4d1e06

tune implementation

9d624c7

add tests

5cfc32c

0cc4m force-pushed the 0cc4m/vulkan-fa-sparse branch from bad0e3e to 5cfc32c Compare September 11, 2026 11:24

0cc4m marked this pull request as ready for review September 11, 2026 11:34

0cc4m requested review from a team and ggerganov as code owners September 11, 2026 11:34

jeffbolznv reviewed Sep 12, 2026

View reviewed changes

0cc4m added 2 commits September 14, 2026 08:34

avoid nondeterministic atomicAdd

4191ad1

add cm2 decode vector support

8d9b181

jeffbolznv reviewed Sep 14, 2026

View reviewed changes

0cc4m added 2 commits September 14, 2026 16:27

simplify logic and make variable names more consistent

de18aa9

add cm2 f16vec4 binding for decode vector

201031a

jeffbolznv approved these changes Sep 14, 2026

View reviewed changes

0cc4m removed model Model specific ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Sep 15, 2026

0cc4m added the merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. label Sep 15, 2026

ggerganov merged commit fc82583 into master Sep 15, 2026
24 of 31 checks passed

ggerganov deleted the 0cc4m/vulkan-fa-sparse branch September 15, 2026 09:30

LynxPDA mentioned this pull request Sep 15, 2026

qwen4exp: grouped-union sparse prefill flash attention — fixes the pp depth cliff Nathanw1014/llama.cpp#11

Open

thomas9120 mentioned this pull request Sep 21, 2026

Upstream sync (#15) thomas9120/llama.cpp-halo-windows#16

Merged

Conversation

0cc4m commented Aug 31, 2026

Overview

Requirements

Uh oh!

LynxPDA commented Sep 11, 2026

Uh oh!

0cc4m commented Sep 11, 2026

Uh oh!

0cc4m commented Sep 11, 2026

Uh oh!

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

Uh oh!

Uh oh!

0cc4m commented Sep 11, 2026

Uh oh!

jeffbolznv left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

LynxPDA commented Sep 11, 2026 •

edited

Loading