Halo/upstream sync 2026 09 05 - #20
Open
LaurentZuijdwijk wants to merge 111 commits into
Open
Conversation
…l-org#28032) * vulkan: add top-k radix sort shader for k >= 1024 * add Qwen 3.8 Flash Next top-k tests * add top-k qsa fusion * clean up code
… (for WebGPU) (ggml-org#28071) Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
…rg#27909) * vulkan: RDNA3 static mat-vec rows above four columns On RDNA3 above four columns a static 4 rows for all types benches faster than the default. * vulkan: RDNA3 static mat-vec-id rows mul_mat_vec_id has no column dimension to switch on. On my Strix Halo machine, a static 4 is faster here than the defaults across types and batch sizes.
…27310) * dflash : fuse the encoder into the KV injection decode The encoder is a single fc + norm, but running it as a separate llama_encode forced a device-to-host round trip of its output before the injection decode could re-upload it, plus a second graph build per round. Fold the encoder into the decoder's embd branch and feed the target features directly to one llama_decode. Assisted-by: Claude Fable * nit * Apply batched suggestions from code review Co-authored-by: Ruixiang Wang <wangruixiang07@outlook.com> * Fix missing references from renaming --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> Co-authored-by: Ruixiang Wang <wangruixiang07@outlook.com>
* enhance get mem info by l0 an SYCL API * remove debug code, format the code * update SYCL.md for GGML_SYCL_GET_MEM_API
…outer fusion were restricted to 1 token (ggml-org#27621) * CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were resticted to 1 token Signed-off-by: ynankani <ynankani@nvidia.com> * Address review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Add SWIGLU_CLAMP case to multi-token moe fusion Signed-off-by: ynankani <ynankani@nvidia.com> --------- Signed-off-by: ynankani <ynankani@nvidia.com>
* ROCm: add radix TOP_K for long rows
…backend_tensor_get() implementation (ggml-org#28045) * webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation * chore : improve code readability Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
) * ui: copy the displayed text of grouped agentic responses Agentic sessions render as a single entry anchored on the first assistant turn, whose content is typically just the first tool call, so the copy button wrote an empty string to the clipboard. Derive the text sections of the whole session and copy them joined, matching the visible response. Plain messages keep the previous behavior. * const
…nd PP performance (ggml-org#26438) * opencl: Q4_K/Q5_K mul_mv N_DST 4->8 on Intel for 2x activation reuse * opencl: Q4_K mul_mm 8x8 tile fot Intel * opencl: Q5_K mul_mm 8x8 tile for Intel * opencl: Q4_K mul_mv N_DST 8->16 for Intel
* kv cache : batch state restore scatter reads per contiguous run When restoring state into non-contiguous destination cells (e.g. a prompt-cache snapshot into a fragmented ring), state_read_data issued one small copy per KV cell - ~1.4M copies of a few KiB each for a 40k+ token restore, taking 25-63 s on the CUDA backend. The snapshot stores cell rows in cell order, so a maximal run of consecutive destination indices maps to one contiguous block and can be restored with a single copy. Precompute the runs once and use them in all three scatter loops (K, V, transposed V). Byte-identical. The on-device reader copies with a byte cursor when the read and write chunking differs, so the batched reads are safe for it as well. Batching makes equal tensor counts with a different split reachable (save ranges [2,1] vs restore runs [1,2]); the next commit teaches the reader's 1:1 path to fall back to the byte cursor in that case. Verified in a production setup: 1,363,616 copies / 25-63 s -> 224 copies / 221-424 ms for the same restores (42,603 cells, 4 runs). Assisted-by: Claude Code (unsloth/qwen3.8-27b) * context : fall back to the byte cursor when read and write chunking differ the on-device reader copies saved state back with a 1:1 copy by tensor index whenever the write and read sides recorded the same number of tensors, guarded by a per-tensor size assert. equal tensor counts do not imply equal chunking: a state restore may batch its reads per contiguous run of destination cells while the save used per-range reads, so both sides can record two tensors that split the same data differently, and the assert aborts in all builds. compare the per-tensor sizes and only take the 1:1 path when the chunking actually matches, otherwise fall through to the existing byte-cursor copy. both sides enumerate the same logical data in the same order, so the cursor copy is well-defined across tensor boundaries. Assisted-by: Claude Code (unsloth/qwen3.8-27b) * tests : cover state restore scatter reads on host and on-device paths decode the same prefix on two sequences, interleaving the seq 0 cells between the seq 1 cells, so the seq 1 cells are isolated from each other in the kv cache (three cells, two saved ranges). save the seq 1 state, free the interleaved seq 0 cells, and restore: the destination is then non-contiguous (two runs), and the restore-side chunking has the same tensor count as the save-side with a different split, so the scatter path is batched per contiguous run and the on-device reader's byte-cursor fallback is exercised. the restored state is saved again on the host and compared byte for byte with the first save: the blob is serialized in sequence cell order, so the two saves are identical if and only if the scatter restore wrote exactly the same KV content. this documents the byte-identical guarantee of the run-batched scatter reads. one test per io backend: the host (CPU) path and the on-device path. Assisted-by: Claude Code (unsloth/qwen3.8-27b)
Assisted-by: DeepSeek-v4-Flash-0731
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 4255429. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* CUDA: XOR swizzle flash attn K,V smem fp16 tiles Signed-off-by: ynankani <ynankani@nvidia.com> * Fix use 64bit generic pointer instead of 32bit shared pointer Signed-off-by: ynankani <ynankani@nvidia.com> * fix shared memory race in FA on DGX Spark * Handle corener case Signed-off-by: ynankani <ynankani@nvidia.com> * Add swizzle test cases and gate sync for swizzled path only Signed-off-by: ynankani <ynankani@nvidia.com> * gate CUDA PTX Signed-off-by: ynankani <ynankani@nvidia.com> * offset calculation specific for swizzle branch Signed-off-by: ynankani <ynankani@nvidia.com> * Reafctor code Signed-off-by: ynankani <ynankani@nvidia.com> * Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset) Signed-off-by: ynankani <ynankani@nvidia.com> * rebase and update test case args Signed-off-by: ynankani <ynankani@nvidia.com> * Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0 Signed-off-by: ynankani <ynankani@nvidia.com> --------- Signed-off-by: ynankani <ynankani@nvidia.com>
* metal : add fa-vec tunings for M1 Ultra * metal : move M1 Ultra tunings after M1 Max section * metal : remove duplicate blank line
* qwen4exp: sum the indexer heads by slices The head reduction went through a transpose and a sum_rows over ne[1], which left sum_rows with ne0 = 4, one block per row for a four element reduction, and the transpose copied the whole block by token surface twice on the way in. The heads are adjacent on ne[1], so each one is a strided view and the sum is a short chain of adds. RTX PRO 6000, Qwen3.8-Flash-Next UD-Q4_K_XL, fa on, 55k context, warm runs on top of ggml-org#28011: prompt processing 2170 -> 2366 t/s Generation is unaffected. The removed work scales with n_blocks by n_tokens, so the gain grows with context and with ubatch size. * qwen4exp: drop the redundant cont on the indexer query rope returns a freshly allocated, contiguous tensor, so the reshape that feeds the matmul does not need a copy. ggml_reshape_3d asserts contiguity, so a layout that would need the cont cannot slip through silently. Greedy output is unchanged token for token. Address review from @ggerganov
MTP speculative decoding needs the target state to move back by the number of rejected draft tokens. Without rollback support the context is classified as SEQ_RM_TYPE_FULL and the server serializes the whole recurrent state to host memory on every round, which costs more than the drafting saves. The recurrent cache already holds n_rs_seq + 1 snapshot planes and the delta net writes its SSM state into them, but build_conv_state_at wrote a single plane, so a rollback restored a convolution history that was never captured. It now writes one snapshot per slot, each ending one token earlier, for the delta net QKV convolution and for the PLE convolution alike. Measured on Qwen3.8-Flash-Next UD-Q4_K_XL with the standalone MTP draft, n-max 3 and a single slot: decoding reaches 183 tok/s on code and 144 tok/s on prose. The same branch before this change, where the server falls back to checkpointing the state to host memory, reaches 123 and 83 tok/s, for 108 tok/s without a draft.
Signed-off-by: Jonathan Clohessy <Jonathan.Clohessy@arm.com>
* metal : request Metal 4.0 language version for the tensor API * metal : load the tensor API kernels from a separate metallib * tests : add external-metallib tensor API regression test * metal : fix metallib build order for the tensor API kernels
* tests : fix log verbosity for test-llama-archs * cont : naming * cont : add note
…add tests (ggml-org#27941) * qwen4exp: follow up fixes * -kvu NaN collapse fix Assisted-by: Claude * indexer cache ext.x/ext.y restore fix Assisted-by: Claude * kv-cells: rename seq_set to seq_get_all seq_get is already taken by the single-id getter, so the suggested name cannot be overloaded on return type alone. Assisted-by: Claude * memory-hybrid-idx: implement set_input_qsa on the memory class The context held the whole implementation, where the pattern elsewhere is a thin context forwarding to the memory class, as llama_kv_cache_context does for set_input_kq_mask. The body reads no context state, so it moves unchanged and the context keeps a forwarder. Also shortens the seq_get_all comment as suggested. * tests: check that a sequence state survives a save/restore round-trip Saves seq 0, erases it, restores the blob and saves again, requiring the two blobs to match. Compares blobs rather than generated text, which cannot see a field dropped on the way back in. Note this passes on master for qwen4exp, so it does not demonstrate the ext.x/ext.y drop this PR fixes; reaching that needs 2D mrope content. * tests: give the synthetic qwen4exp a PLE so the state test bites has_cell_ext() is n_pos_per_embd() > 1 || ple_n_heads > 0, and the indexer cache sets rope_type = NONE, so without a PLE it serializes no cell ext at all and the round-trip test cannot see a dropped ext.x/ext.y. With one, removing the ext_set restore in state_read_meta fails the test: 198 of 335692 bytes differ, first at offset 282092. Loading such a model needed two fixes: - the row count of per_layer_token_embd came from require_weight(), which a model synthesised from metadata alone has no file to answer. Derive it from the head ranges and prefer the file's padded count where there is one. - the PLE conv history is a row of the recurrent cache, so a PLE on a full attention layer dereferenced a null p_l. Reject it at load time instead. The meta mirror is skipped for qwen4exp. It returned NaN logits before this fixture carried a PLE, which the nmse check passes since a NaN comparison is false, and aborts with one. -sm tensor on real devices works. Assisted-by: Claude * llama: disable -sm tensor for qwen4exp test-llama-archs skipped the tensor split for this arch from inside the test, so the arch still advertised support it does not have. Declare it in llm_arch_supports_sm_tensor instead and drop the test-side exception; the existing llm_arch_supports_sm_tensor branch then does the skipping. Assisted-by: Claude
Rows for M2 Max (30 GPU cores) collected with 'ggml-metal-tuning fa-vec --dtype f16,q8_0', pasted into fa_vec_tuned_table. ref: ggml-org#27668 (comment) Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : fix more leaks due to missing autoreleasepools * metal : rename variable * metal : fix another missing pool warning Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
…8159) * load hparams.n_layer_nextn before n_layer() calls * remove duplicate loads
* move more jobs to ccache-buckets * add venv deps * also jq
* opencl: add Adreno xmem SDPA path Assisted-by: Codex * Removed the Adreno-specific queue profiling override * Clean up formatting * 修复数值误差优化gqa/mask attn Assisted-by: Codex * add env GGML_OPENCL_XMEM_SDPA Assisted-by: OpenAI Codex --------- Co-authored-by: happyyzy <happyyzy@users.noreply.github.com>
…g#27633) * opencl: add extended elementwise unary ops (sgn, step, elu, hardswish, hardsigmoid, floor, ceil, round, trunc) Adds nine GGML_UNARY_OP_* elementwise ops that were falling back to CPU on the OpenCL backend, following the same variant shape as the existing ABS op: f32, f32_4 (vec4), f16, f16_4 (vec4), and stride-addressed f32_nc / f16_nc for non-contiguous inputs. New kernels/unary_ext.cl (macro-generated), a shared ggml_cl_unary_ext dispatch helper mirroring ggml_cl_abs, the supports_op cases, and the compute-forward cases. Values are computed in float (the f16 variants read/write half and convert), so the conditional ops (step, elu) match the CPU reference; the vec4 forms use select() for the branch. Validated with test-backend-ops on Adreno 840 and 850 (E17): all nine ops pass every case including the vec4 and non-contiguous variants (8/8 or 14/14). * opencl: dispatch a contiguous f32 copy over the whole device kernel_cpy_f32_f32 maps one workgroup to each (i01,i02,i03) row and strides the row across that workgroup's lanes, and the host launches ne01*MIN(64,ne00) work items. A tensor with few long rows therefore runs on a single workgroup. The mamba2 and gated-delta-net recurrent state cache is one row of 524288 floats, copied once per layer per graph, and lands on 64 work items. When both sides are contiguous the copy is a linear move, so dispatch it over the whole device: one work item per float4. Gated on ggml_is_contiguous for both tensors and equal element counts, so copies already spread over many rows keep the existing path. The kernel is created optionally, so a driver that rejects it falls back rather than aborting. vload4/vstore4 rather than a float4 cast: they require only the scalar type's alignment, and these buffers carry an arbitrary 4-byte view offset. CPY, DUP and CONT are 217/217 on Adreno 840 and 740 with the path enabled and disabled. GGML_OPENCL_CPY_FLAT=0 forces the old kernel. * opencl: support all easy-copy types in CONCAT CONCAT was F32-only. Extend it to every "easy-copy" type -- any non-quantized type with a block size of 1 and an element size of 1, 2, 4 or 8 bytes, i.e. f16/bf16/i8/i16/i32/i64 as well as f32. The kernels are keyed by element SIZE rather than by type, which is what CUDA already does for the same op: one kernel per byte width (b1/b2/b4/b8) plus the packed b4 fast path, instead of one per ggml type. supports_op gates on the same property, so a new type of a supported width is picked up with no further work. Validated with test-backend-ops on Adreno 840 / A8X and X2-90 / X2E.
…25789) (ggml-org#28089) * ui : fix MCP image attachments not displayed in tool block (ggml-org#25789) Fixes regression from ggml-org#25450 where ChatMessageAgenticContent passed message.extra instead of section.toolResultExtras to tool blocks, leaving tool images invisible. Also fixes TOOL_RESULT_JSON_OPEN_REGEX which misclassified "[Attachment saved: ...]" as JSON. Fixes ggml-org#25789 Assisted-by: Muse Spark * Addressed PR comments: 1.- Removed ·?? mesage?extra· as it has no case left to cover 2.- Added ·[\· to cover the case of ·[[1, 2], [3, 4]]· case suggested in the PR comment 3.- Added unit test for covering up this regex case * ui : fix MCP image attachments not displayed in tool block (ggml-org#25789) - Addressed lint error on regex (redundant \)
* addition of m3 in fa_vec_tuned_table * adding q4_0,q4_1,q5_0,q5_1 in ggml-metal-tuning * Fix formatting in ggml-metal-tuning.cpp
…-org#27631) define two new environment variables to better understand how much memory is being allocated, and when. This has been invaluable in inproving the --fit algorithm, and is likely to be useful when debugging other memory-related issues. `-lv 4` will be required to enable the following: GGML_SYCL_MEMTRACE=1 will show per-site memory usage, updated whenever it increases by more than 64MiB. GGML_SYCL_MEMTRACE=2 will show every allocation and deallocation. To change the default 64MiB threshold for reporting memory usage increases, use GGML_SYCL_MEMTRACE_STEP. A sample log line: [SYCL-MEMTRACE] device memory query (dev): total 59493 MiB, free 4494, in use 54998; allocated 0 (buffers 0 + scratch 0), peak 0 MiB
…T support (ggml-org#28016) (ggml-org#28254) * Reapply "sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 12…" (ggml-org#28184) This reverts commit c845263. * tests : fix unused variable M in test-backend-ops * tests: fix trailing space error and isolate kronecker tests for sycl backend only
Conflicts and how they were resolved: - ggml-vulkan.cpp: upstream landed its own radix top_k (ggml-org#28032), which supersedes ours (#17) and adds the qwen4 QSA fusion the rest of the file now needs. Took upstream's top_k wholesale and dropped topk_radix.comp. Ours stays available in ca062e9 for the benchmark that decides this. - qwen4exp.cpp: kept upstream's padded row-count validation, then our --ngram-on-disk branch. get_weight() is nullable now, so a model with no tensor to read from stays on the lazy path. - test-backend-ops.cpp: additive collision, kept both sides. - qwen4exp.cpp MTP loop: hparams.n_ff_exp is a function upstream now. Assisted-by: Claude Opus 5
github-actions
Bot
added
documentation
dzannotti
approved these changes
Sep 5, 2026
Three resolutions that were verified locally but left out of the merge commit itself: - ggml-vulkan.cpp: take upstream's ggml_vk_topk whole. Our fork had split it into _small/_large plus a dispatcher, so keeping the renamed half left the GGML_OP_TOP_K call site with no ggml_vk_topk to call. - vulkan-shaders-gen.cpp: drop the entry for topk_radix.comp, deleted with the merge. - qwen4exp.cpp: hparams.n_ff_exp is a function upstream now. The trunk loop was updated by the merge, the fork's MTP loop was not. Assisted-by: Claude Opus 5
O_DIRECT is a Linux extension, so the macOS build did not compile. Darwin turns the page cache off per descriptor with F_NOCACHE instead. Assisted-by: Claude Opus 5
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
20 participants
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.
Merge of upstream