Repository navigation
Sync upstream master (2026-09-28, ggml-org 6c7a87f7e) - #12
Merged
Merged
Conversation
Assisted-by: OpenCode
* hexagon: fix accuracy issue in Q8_0 N=1 MUL_MAT * hex-quant: fix register spills * hex-mm: use dma for all dyn.quant paths Co-authored-by: Aparna M P <aparmp@qti.qualcomm.com> * hex-mm: remove obsolete run_quant_task * hex-mm: update tracing to properly wrap the events * hex-mm: use act for activation data in all paths * hex-mm: use act_ instead of src1_ to avoid confusion in fused kernels * hex-mm: remove/reroute the rest of the non-DMA act (aka src1) logic * hex-dma64: yet another pass at cleaning up the dma_addr_t casts * Update ggml/src/ggml-hexagon/htp/matmul-ops.h Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update ggml/src/ggml-hexagon/htp/matmul-ops.c Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update ggml/src/ggml-hexagon/htp/matmul-ops.c Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Update ggml/src/ggml-hexagon/htp/matmul-ops.c Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com> Co-authored-by: Aparna M P <aparmp@qti.qualcomm.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* ggml : bump version to 0.25.2 (ggml/1642) * ggml : fix ubsan error in `ggml_graph_nbytes` (ggml/1644) * ggml : bump version to 0.25.3 (ggml/1645) * sync : ggml
* metal : cache sparse FA indices in shared memory Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * metal : simplify shared memory size calculation Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * pi : update general * metal : unroll sparse index load Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
- return early when the graph has no nodes - drop the redundant reset of capture_compute: the decrement at the top of the function already transitions the counter from 0 to -1, so a capture happens exactly once - hint at METAL_CAPTURE_ENABLED=1 in the capture error message - pass capture_compute == 0 (not the raw counter) as use_capture to ggml_metal_op_init, so GPU debug-group markers are only emitted on the captured compute Assisted-by: pi:llama.cpp/Qwen3.8-27B
* hexagon: use DMA for contiguous dim1 CONCAT Assisted-by: OpenCode * hexagon: update CONCAT DMA for DMA64 Assisted-by: OpenCode
…ml-org#29294) * llama : fix tensor split for fused qkv with uneven K/V head sizes Assisted-by: Qwen3.8-27B * fix v granularity * convert: fix mtp conversion * convert: add support for mtp flags * fix loader
- ggml-org#28068 builds the GDN q/k l2norm as ggml_scale(ggml_rms_norm(x, eps/n), 1/sqrt(n)). This adds 2 SCALE nodes per GDN layer, 96 extra kernel launches per ubatch on Qwen3.8-27B (48 GDN layers). - The extra kernels take no measurable GPU time, but each launch has a host/driver cost. It is small with plain batch processing and about 10x larger with draft-mtp speculative decoding. - rms_norm_f32 gets a do_scale flag, the same pattern as do_multiply/do_add, so the fused path shares the kernel, the reduction and the launcher. It computes scale * (rsqrt(mean + eps) * x), which matches the unfused rms_norm + scale bit for bit, so ggml-org#28068 numerics are kept. - Fusion only fires when SCALE has no bias and the rms_norm output has a single consumer (ggml_can_fuse). - Metal (ggml-org#28948) and SYCL (ggml-org#28931) already fuse the same pattern. Measured on 2x GTX 1080 Ti (sm_61, PCIe 3.0 x16 + x4), i7-13700KF, Windows 11, driver 582.66, CUDA 12.9. Qwen3.8-27B-UD-Q4_K_XL, -ngl 99 -ts 53,47 -ot token_embd=CPU, master fee39dd. llama-bench -ub 128,512 -p 512,2048 -n 128 -r 5, tok/s: build pp512@128 pp2048@128 pp2048@512 tg128 master 367.7 419.1 385.4 12.90 master + fix 372.6 420.6 388.6 12.98 +1.3% +0.4% +0.8% +0.6% llama-server cold prefill, -c 56000 -ub 128 -b 2048, draft-mtp n-max 3 p-min 0.5, mean of 2 rounds x 3 reps: build pp 8000 pp 20000 master 356.5 322.0 master + fix 371.4 (+4.2%) 337.4 (+4.8%) - Launches per ubatch go from 1032.9 + 841.7 back to 978.9 + 799.7 (CUDA0 + CUDA1), the b10828 count. The GPU op sum is unchanged. - test-backend-ops RMS_NORM_SCALE, NORM_SCALE, RMS_NORM_MUL_ADD, RMS_NORM_MUL_ROPE, RMS_NORM, RMS_NORM_BACK, NORM, L2_NORM and SCALE all pass on both GPUs. - Perplexity is identical to the unfused build: 3.2030 +/- 0.0559 at -c 2048, 16 chunks. - Draft acceptance counts per request match the unfused build. Assisted-by: Claude Opus 5.5
…g#29193) * musa: use 16-byte copies for MUSA like sm_70+ ggml_cuda_get_max_cpy_bytes() derives the copy width from __CUDA_ARCH__. mcc never defines it, so MUSA fell into the generic branch and returned 8 bytes instead of the 16 bytes that every sm_70+ target gets. The value sizes the per-thread copy unit of the FlashAttention K/V staging code (fattn-common, fattn-vec, fattn-tile, fattn-mma-f16 shared-memory loads) and of mmq-vec-dot, so every MUSA FlashAttention kernel moved half as many bytes per instruction. On an MTT S5000 (mp_31, MUSA SDK 5.2.0) with Qwen3.8-27B-UD-Q4_K_M, -ngl 999, -p 512 -n 64, -fa on: 751.15 -> 794.73 t/s prefill and 15.59 -> 15.69 t/s decode. -fa off is unchanged (1050.05 -> 1052.86 t/s prefill), FLASH_ATTN_EXT is unchanged (3984 ok / 0 fail / 1323 unsupported) and perplexity is unchanged. * musa: enable the CUB paths on MUSA GGML_CUDA_USE_CUB and USE_CUB are selected by "CUDART_VERSION >= 11070", which the MUSA SDK never satisfies: CUDART_VERSION is not defined anywhere under /usr/local/musa/include, so the condition is always false and every CUB-based path stayed compiled out on MUSA even though the SDK ships CUB and the kernels build for mp_31. Select them from GGML_USE_MUSA as well. The device-wide algorithms are usable too: cub::DeviceSegmentedSort compiles and produces correct results on mp_31. This lifts the ne[0] <= 1024 limit that ggml_backend_cuda_device_supports_op applied to ARGSORT and TOP_K on MUSA. On an MTT S5000 (S5000, mcc 5.2.0): ARGSORT 48 ok / 52 not supported -> 100 ok / 0 (CUDA parity), TOP_K 0 ok / 354 not supported -> 527 ok / 0. The other 20 per-op suites are unchanged, the Qwen3-0.6B f16 (14.4679) and Qwen3.8-27B iq4_nl (5.1724) perplexities are unchanged, and the 0.6B graph keeps the same nodes and splits (18 CPU + 18 MUSA0, SET_ROWS 1008) as before. * musa: take the upstream code path where the toolkit supports it Several guards were written for an older MUSA toolkit. Verified against MUSA SDK 5.2.0 and on an MTT S5000 (mp_31): - device init: query cudaDevAttrCooperativeLaunch instead of hardcoding false. The device reports cooperativeLaunch=1 and musaLaunchCooperativeKernel works (verified with a kernel whose result was checked). - device init: keep prop.warpSize instead of overriding it with 32. The device reports 32 anyway, so this only removes the divergence. - CUDA_SET_SHARED_MEMORY_LIMIT and the FA shared-memory raise: musaFuncSetAttribute returns success and sharedMemPerBlockOptin is 192 KiB, so the kernels can use more than the default 48 KiB. - vendors/musa.h: add the cudaDeviceGetAttribute and cudaDevAttrCooperativeLaunch mappings the device-init change needs. Measured on one S5000 with Qwen3.8-27B Q4_K_M (-ngl 999, -r 3): pp512 968.27 -> 957.09 t/s, tg64 10.09 -> 10.23 t/s, FLASH_ATTN_EXT sweep identical (3975/3982 both), perplexity identical (80.2841 +/- 7.26772 both). * musa: drop compile-time guards that MUSA's runtime gates already cover mcc never defines __CUDA_ARCH__, so the arch-gated fallbacks in this group were already taken on MUSA and the GGML_USE_MUSA guards on top of them only kept the upstream text from being compiled: - wkv.cu: the "#pragma unroll" suppression has no effect on the generated code that is not already covered by the surrounding guards - common.cuh: the MUSA-only __builtin_unreachable() in no_device_code() is not needed to silence the compiler - ssm-scan.cu: the SSD (Mamba-2 prefill) block and its dispatch are gated at runtime by GGML_CUDA_CC_IS_NVIDIA(cc) and turing_mma_available(cc), which are both false for PH1 (cc 0x100310), so compiling them changes nothing - common.cuh: warp_reduce_max(half2) is guarded the same way as warp_reduce_sum(half2) (FP16_AVAILABLE); the MUSA-only guard left the function with no return statement. It has no caller today. MTT S5000 (mp_31, MUSA SDK 5.2.0), MUSA_ARCHITECTURES=31: build rc=0. Against an unmodified build of the same tree on the same card, FLASH_ATTN_EXT (3984 ok / 0 fail / 1323 unsupported), SSM_SCAN (15/0), RWKV_WKV6 (6/0), GATED_DELTA_NET (38/0) and MUL_MAT (1299/0/385 unsupported) are identical, and perplexity with -fa on is bit-identical (5.1639 +/- 0.36673, 4 chunks). * musa: do not use MMQ on PH1 test-backend-ops on an MTT S5000 (mp_31, MUSA SDK 5.2.0) fails 260 cases and every one of them goes through the MMQ path: - MUL_MAT with a batched src1 (any bs/nr != [1,1]): 109 cases across all quantized types, e.g. 12 of 13 cases at n=16, while the plain [1,1] layout passes - every quantized MUL_MAT_ID: 147 cases, while the f16/f32 variants of the same shapes pass - MUL_MAT with more than ~512 tokens: 4 cases (n=509..4096); the small-n cases pass The cuBLAS/dequant path is correct for all of them and the MMVQ path used for small batches is unaffected, so quantized matmuls now take that path on PH1 instead of returning wrong values. 27B perplexity with default flags goes from nan to finite, and the full suite reports 0 failures out of 22237 cases. The MMQ defect itself (fastdiv, __umulhi, uint3 kernel parameters and __CUDA_ARCH__-based MMA availability were all checked and are correct on this part) is not addressed here. * musa: keep the block barrier of the fused TOPK_MOE kernel reachable topk_moe_cuda returns early for the rows past the end of the graph, but one block covers TOPK_MOE_ROWS_PER_BLOCK (8) rows, so the last block is only partially filled whenever n_rows is not a multiple of 8. On MUSA a warp that has already returned blocks the block wide __syncthreads() below, which makes the kernel hang and the launch time out. CUDA tolerates the exited warps, which is why the CUDA numbers never showed it. For MUSA, clamp the row index of those warps to the last row so that every warp of the block reaches the barrier; they recompute the last row and write the same values. The CUDA code path is unchanged. On an MTT S5000 (mp_31) the fused TOPK_MOE cases change from a launch timeout with no completed case to 418 ok / 0 not supported / 0 failed, i.e. the CUDA result, and the other 101 per op suites are unchanged (0 failed, no count changes). * musa: enable GATED_DELTA_NET The op was turned off for every MUSA target because mcc could not build the kernel at the time. The current toolkit builds it: with mp_31 and MUSA SDK 5.2.0 the file compiles with zero errors and all 36 test-backend-ops GATED_DELTA_NET cases pass against the CPU reference. 27B perplexity is unchanged. While the op is refused, the scheduler has no choice but to run it on the CPU: 48 GATED_DELTA_NET nodes per forward pass. On an MTT S5000 (Qwen3.8-27B Q4_K_M, -ngl 999, one container, -r 3): pp512 (FA off) 964.51 -> 2119.26 t/s tg64 (FA off) 10.15 -> 15.50 t/s * musa: name the stream capture query API for the graph aware kernels argsort.cu and mean.cu call cudaStreamCaptureStatus, cudaStreamIsCapturing and cudaStreamCaptureStatusNone inside their USE_CUDA_GRAPH blocks, but the MUSA compatibility headers do not alias those names, so building with the experimental GGML_MUSA_GRAPHS option fails with 7 errors in those two files. Map the three names to their musa* counterparts, under the same guard that enables the graph code, so the default build is untouched. The option stays off by default: on an MTT S5000 the captured path measured slower (pp512 693 vs 772 t/s, tg128 15.20 vs 15.39 t/s over two sessions) and the borderline MUL_MAT cases are not reproducible between runs. * musa: build the CI and docs for PH1 (MTT S5000) The MUSA CI job and the documented default still targeted the first generation (MTT S80, MUSA_ARCHITECTURES=21) while the current MUSA SDK targets PH1 (MTT S5000, 31). Move the job, ci/run.sh's default and the build docs to 31, and run the job in the PH1 MUSA SDK devel image: registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64 That image needs two things the previous one did not: python3-venv for the ccache-buckets step, which builds a virtual environment for the Hugging Face CLI, and no time prefix on the build command, because container jobs run their steps with sh and the image ships no time binary.
* fix conflict * fix format issue * rm unused code
…at ggml_nbytes (ggml-org#29283) * rpc : include nb in the get_alloc_size cache key and floor the result at ggml_nbytes * cont : remove redundant comment * cont : add TODO --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
) * Rebase and update based on ggml-org#26675 Signed-off-by: ynankani <ynankani@nvidia.com> * CI failure fix(launh_bounds overload on HIP) and cleanup Signed-off-by: ynankani <ynankani@nvidia.com> * Address review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Use ggml tensor instead of name in act policy map Signed-off-by: ynankani <ynankani@nvidia.com> * Address review comments and cleanup Signed-off-by: ynankani <ynankani@nvidia.com> * Address review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Rename changes Signed-off-by: ynankani <ynankani@nvidia.com> * Update ggml/src/ggml-cuda/mmq.cu Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * MXFP4 dispatch changes for higher src prec Signed-off-by: ynankani <ynankani@nvidia.com> * Refactor and address review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Updates based on review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Apply batched suggestions from code review Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * Address review comments Signed-off-by: ynankani <ynankani@nvidia.com> * Apply patch from review Signed-off-by: ynankani <ynankani@nvidia.com> --------- Signed-off-by: ynankani <ynankani@nvidia.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* metal : split fa kernels into per-dtype libraries Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp * cont : minor fix comment
* metal: FWHT kernels for block widths above 512 The Metal FWHT covers widths 64 to 512, one row per simdgroup with N/32 values per lane. Wider blocks need more registers per lane than that layout allows. kernel_fwht_tg runs one row per threadgroup with 256 threads, so each thread keeps N/256 values. Butterflies below the simdgroup width still shuffle, those up to the threadgroup width go through threadgroup memory, and the rest stay in registers. Same butterfly and sign convention as the simdgroup kernel. Widths 64 to 512 keep the simdgroup kernel. 1024 through 8192 use the new one, for both F32 and F16 sources. The wide kernels allocate float[N] of threadgroup memory, 32 KB at 8192, so the size check takes the device limit and reports those widths as unsupported where they would not fit. Without that a device with less threadgroup memory would accept the op and then abort on a nil pipeline. test-backend-ops on M5 Pro: MUL_MAT_HADAMARD 26/26, MUL_MAT 1265/1265. * cont : add TODOs --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
…l-org#29417) * templateprocessing must win over tokenizer config * remove obsolete override
…perativeMatrix API support (ggml-org#29373) (ggml-org#29409) * vulkan : fix build issue of legacy glslc version by adding GGML_VULKAN_COOPMAT_GLSLC_SUPPORT macro check for Intel FA shader compiling * vulkan : add preprocess condition to filter out unsupported FA 2 phases kernels before creation. * vulkan : move lock_guard for Intel FA shader pointer creation under CM1 compiling preprocessor
…a8_bin`, `kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin` (ggml-org#29401) * opencl: add A8 Q5_K non-MoE non dp4a + dp4a binary kernel * opencl: fix s transpose - s only transposed for bin kernels --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
which resulted in different greedy transcripts for 4.5% of English and 6.5% of Japanese test utterances. In Japanese, some differences changed entire words. This change: * uses `log(x + 2^-24)` instead of clamping to the log floor * uses a symmetric Hann window, equivalent to `torch.hann_window(periodic=False)` * adds the normalization epsilon to the standard deviation instead of inside the square root Only the `lfm2a` preprocessor opts into these behaviors. Other audio preprocessors are unchanged. Tested on top of 84e76d8 using `llama-server` with CUDA and `temperature=0`, compared against http://github.com/Liquid4All/liquid-audio fp32. Test set: * 200 LibriSpeech `test-clean` utterances (EN) * 200 Common Voice `ja` test utterances (JP) * identical 16 kHz audio passed to both implementations | Greedy transcript identical to `liquid-audio` | Without fix | With fix | | --------------------------------------------- | ----------: | ----------: | | EN F16 | 191/200 | 200/200 | | JP F32 | 187/200 | 200/200 | | JP F16 | 187/200 | 199/200 | The remaining JP F16 difference is a comma and matches the reference implementation's own bf16 output. Mel relative L2 error versus `liquid-audio`: * EN: 3.2% -> ~2e-6 median * JP: 3.9% -> ~2e-6 median
) The original function was broken on Windows for some unicode paths Paths without a trailing separator now create the last directory too, matching the function name. All current callers already include a trailing separator, so this change does not affect them. Signed-off-by: Adrien Gallouët <angt@huggingface.co>
…l-org#29449) * hex-scripts: fix table alignment * hex-scripts: find sw div calls using binary inspection tool
* Added tiled mul_mat. For each mul_mat_one_chunk, quants are unpacked into (max) 256x256 tiles of int8, one routine per quent. Then microkernel computes 16x16 tiles before writing out 256x256 float reults to main memory. Tests/benches in tests/test-tiled-mulmat.cpp. 3-6x speed improvement for large matmul, break even at 4096x64 * 64x4096, 80% performance (net loss) for GEMV. Error rates trivial (order of 1-e04 max, 1-e05 rmse). * Fixes for ARM/windows builds * more windows fixes, ggml-cpu.h isn't visible in MSVC for some reason * unified iqp + tiled on the Q5_K, IQ4_XS set for benchmarking, updated benchmark * Fixed accidental removal of llama_build_and_test(test-backend-ops.cpp) * First integration of iqp code Co-authored-by Bartowski <3266127+bartowski1182@users.noreply.github.com> * Cleaning up declaration of iq unpacking helpers to align with the bit unpackers * Removed iqp path * Fix cross-platform warnings * Disabling benchmarks unless explicitly enabled * Fix backend_init for DLL-based builds, add self and bartowski to CODEOWNERS for tiled * Put benchmarks behind a flag * kernel fix for AVX2, iq quants * Fix for asan, leaking memory in test-tiled-mulmat and avoid stack use after return * guarding env flags with std::call_once * Simplified repacking for VNNI to a single call per macrotile * No threadlocals anymore, aligned wdata access * Doing aligned reads since we ensure alignment with padding in wdata * Eliminated per-thread gather of Q8_K rows in mul_mat_id, we now gather/repack in a single pass. Repack method now takes pointer array to support both dense/normal and mmid paths. Interface with ggml-cpu.c simplified as a result * Unified/simplified dispatch and support checks. Put details on wdata needed inside the kernel.h body, simplified interactions with ggml-cpu.c. * Cleanup includes and whitespace, update src1_repack to return false if we don't need a special repack, so the common case is handled by driver * Better detection of win32 and additional whitespace fixes * Gating fuzz tests behind a parameter and some extra prints to try and fix slow CI hosts * Optimized AVX2 kernel * Changed interleave format and added ability to interleave in-place after dequant * Repacks now happen in-place, 16x64 microtiles are independent of each other * Only repack rows in groups of 16 as they're needed. Save work in low n_rows cases and optimize L1 usage in other cases * Use long panels for memory-bound regime (M <= 16), reintroduce IQP path for benchmarks * Fix unused warnings and cleanup. Improved IQ dequantization speed. * Removed separate process benchmarks * Revert "Removed separate process benchmarks" This reverts commit 0688cf4. * AVX2 optimizations and guards for tests on windows * Removed temp perf harness * Remove perf-mulmat from build * Removed IQP path, simplified tests to not use sub processes * Cleaning up alignment of wdata * Whitespace fixes and aligning L2 workspace to clean 512kb boundaries * Update ggml/src/ggml-cpu/tiled/tiled-kernel.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Cleanup merge-duplicated declaration of test-backend-ops target * Undo accidental line deletion in ggml.c --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
…29514) * ci : enable GGML_SCHED_DEBUG_REALLOC=1 for ctest workflows * cont : metal paravirtual device is not compatible
* RPC: use RDMA completion queue to not spin * add TODO for apple RDMA
…rg#29516) Signed-off-by: Adrien Gallouët <angt@huggingface.co>
…zes (ggml-org#28907) * HIP: Enable fattn-mma kernel on cdna for dkq > 256 for large batch sizes * CI: hip-quality-check: ignore spills for very large mfma mma kernels
The SYCL FWHT covers 64 to 512 via the standard butterfly network, plus 384/640/768/1280 via the Kronecker/Paley construction added separately in Hadamard hint can produce (1024, 2048, 4096, 8192); those still fall through to the default case and run as a dense GEMM against the materialized rotation tensor, correct but O(n^2) instead of O(n log n). fwht_kernel_wide runs one row per work-group instead of per sub-group, so each work-item keeps N/NT values rather than N/WARP_SIZE. Butterflies below the sub-group width still shuffle; those up to the work-group width go through work-group local memory; the rest stay in registers. Same butterfly and sign convention as the existing narrow kernel. ggml's SYCL backend registration (dpct::dev_mgr) unconditionally requires a GPU-labeled platform to exist and throws before any op-level test can run, so test-backend-ops could not be exercised on this box (a GPU-less pod) even via the CPU device. Verified instead with a standalone harness: the same kernel body run through a real SYCL CPU device (Intel oneAPI DPC++ 2026.1, OpenCL CPU backend), checked against an independent recursive-doubling Hadamard reference, cross-validated by first running the existing unmodified narrow kernel through the identical harness and confirming it passes (rules out a reference-convention bug before trusting a pass on the new code). Random-input results for all four widths, single- and multi-row: N=1024 NT=256 rows=1 max_abs_err=1.7e-07 max_rel_err=4.9e-04 PASS N=2048 NT=256 rows=1 max_abs_err=1.9e-07 max_rel_err=2.0e-04 PASS N=4096 NT=256 rows=1 max_abs_err=2.0e-07 max_rel_err=1.4e-04 PASS N=8192 NT=256 rows=1 max_abs_err=2.5e-07 max_rel_err=3.8e-03 PASS N=1024 NT=256 rows=7 max_abs_err=2.4e-07 max_rel_err=1.0e-03 PASS N=2048 NT=256 rows=5 max_abs_err=3.0e-07 max_rel_err=9.4e-04 PASS N=4096 NT=256 rows=3 max_abs_err=2.7e-07 max_rel_err=1.7e-03 PASS N=8192 NT=256 rows=2 max_abs_err=2.5e-07 max_rel_err=1.9e-03 PASS This covers the kernel algorithm itself; it does not exercise the ggml dispatch/supports_op integration end to end, which needs a real GPU (or a SYCL GPU plugin) to get past backend registration. test-backend-ops build is verified: fwht.cpp recompiles with zero warnings as part of ggml-sycl.
* add support for dict builtin * add tests
* bump ty to 0.0.84 * fix assertion bug caught by ty
Recent PLaMo-3 models use YaRN, while some earlier PLaMo-3 models do not. The recent PLaMo-3 store their YaRN settings as flat config keys (rope_scaling_factor, initial_context_length) and build the dict at runtime in Plamo3Config.rope_parameters. The current converter misses these settings and writes plain RoPE metadata to GGUF. Mirror the runtime settings into rope_parameters so the corresponding rope.scaling.* is written to GGUF.
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
- register --rpc unconditionally and call llama_supports_rpc() only from its handler - print server "initialization ..." log after args are parsed Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
…(ie. Qwen3 and Qwen3-VL) (ggml-org#28876) * server : allow splitting RANK pooling for causal LLM rerankers Rerank models fall into two categories: bidirectional cross-encoders (BERT, etc.) that require all tokens in a single physical batch, and causal LLMs repurposed as rerankers (Qwen3, Qwen3-VL) that can use chunked prefill like any other decoder. Previously the server rejected all RANK-pooling inputs larger than n_ubatch, and the graph builder hardcoded QWEN3/QWEN3VL arch checks to determine last-token pooling. This broke long-document and multimodal reranking for causal models. Fix: expose llama_get_causal_attn(ctx) so the server can check the effective runtime attention type (reflecting any --attention override or set_causal_attn call). Also expose llama_model_is_causal(model) for querying the static architectural property from GGUF metadata. can_split() now permits chunked prefill for RANK pooling when the context is causal. The graph builder's inline arch check is replaced with the same cparams.causal_attn predicate, removing the duplication. Assisted-by: Opencode/Qwen3.8-27B * remove unused llama_model_is_causal, fix whitespace Assisted-by: opencode --------- Co-authored-by: timothywang21 <timothywang21@users.noreply.github.com>
…che (ggml-org#28956) * vulkan: read the batch stride of an in place src0 from nb[2] A dim01 contiguous tensor can still be a view whose batches are strided by more than ne[1] rows, the first rows of a KV cache for example. Both the mat-vec and the matrix paths read such a tensor in place but passed ne00*ne01 as the batch stride, so every head past the first read the wrong rows. The same applies to src1. The stride now comes from nb[2] whenever the tensor is used in place; the value is unchanged for a contiguous tensor. test-backend-ops gets an m_v parameter on test_mul_mat, the number of rows of a in memory, and two cases at the shapes of a decoder self attention over a cache. * vulkan: size the in place A and B ranges by their strided extent The matrix path bound src0 and src1 to the shader with a range of elements times type size, which ends before the batches of a strided view. Pipelines with bounded access read zero past that range, so the same view that the mat-vec path already handles gave wrong results on Intel and on NVIDIA without coopmat2. The range now comes from ggml_nbytes when the tensor is read in place. * vulkan: address review from jeffbolznv Bind the in place A and B of the matrix path with ggml_vk_subbuffer, which spans to the end of the buffer, so a strided view is in range without computing its extent. mul_mat_id reads the batch stride of an in place src0 and src1 with the same helper as mul_mat. test_mul_mat_id gets an m_v parameter, the number of rows of as in memory, and a case whose experts are strided by more rows than it uses. * vulkan: read the batch stride of an in place src0 in mul_mat_vec_id The single token path of mul_mat_id passed ne00*ne01 as the batch stride of A, so a strided expert view read the wrong rows. The stride now comes from ggml_vk_batch_stride like the other three paths, and src1 follows the same rule. test_mul_mat_id gets a single token case over the strided view. * vulkan: address review from jeffbolznv The batch stride of an in place tensor is taken from nb[2] as nb[2] / type_size * block_size, which holds when nb[2] is padded and not a multiple of nb[1]. A test_mul_mat case with a padded batch stride covers it. * vulkan: keep the A and B ranges exact in mul_mm The quantized A loads of mul_mm carry no row bound and rely on the descriptor range to read zeros past the last row of a partial tile. Binding A and B up to the end of the buffer let those tiles read the leftovers of a previous node and hung the NVFP4 mul_mm on NVIDIA without coopmat2. The range is the strided extent of a tensor read in place and the staged size otherwise.
* tests : init ggml for test-recurrent-state-rollback * cont : same for test-save-load-state * cont : add to test-state-restore-fragmented + add TODOs
* can reproduce the issue vlad sees * fix fma issue * drop volatile * fix volatile runtime task * add arm flag if needed * fix hsum compile error * fix syntax in quants * strengthen sve probing * make the syntax fixes one liners * remove debug code * formatting * remove macro for float * drive down gcc instruction count * support armec * fix CI comments address CI comments fix cross compile issue remove warning fix style and fix fma probing fix style * add documentation * update documentation
…gml-org#28751) * context : do not re-reserve the scheduler when toggling causal_attn `llama_context::set_causal_attn()` marks the scheduler to do a full re-reserve on every change of the flag. For vision inputs, this flag is flipped twice around each non-causal image chunk for Gemma models, resulting in two expensive `sched_reserve()` passes per image. This is especially slow for multi-image or video inputs. The cost of a re-reserve scales with context and ubatch configurations, so larger settings pay more per image (see table below). The re-reserve is unnecessary in this case because `causal_attn` only changes the values written to KQ mask, not tensor shapes or any other buffer sizes. Note: `causal_attn` is a graph reuse key (`llm_graph_params` via `cparams`), so a new graph is built regardless of `sched_need_reserve`, so this doesn't change the graph rebuilding behaviour. llama-server with gemma-4-26B-A4B Q4_0 + BF16 mmproj, 130-token images, cache_prompt=false, prompt_ms median of 3 (before -> after): | images | config | H200 before -> after | RTX 4090 before -> after | |-|-|-|-| | 1 | `-c 8192 -ub 512` | 134 -> 105 ms (1.27×) | 201 -> 119 ms (1.69×) | | 24 | `-c 8192 -ub 512` | 2278 -> 1562 ms (1.46×) | 3559 -> 1748 ms (2.04×) | | 24 | `-c 32768 -ub 2048` | 5379 -> 1584 ms (3.40×) | 13377 -> 1759 ms (7.61×) | Generated output remains identical before and after. * qwen4exp : make the indexer bias shape independent of causal_attn The block/cell bias path was selected on cparams.causal_attn, so the causal and non-causal graphs differed in tensor shapes and ops. With the re-reserve removed (previous commit), a runtime flip resulted in reallocating the compute buffers, which would fail under GGML_SCHED_NO_REALLOC. This commit selects the block path from the mask shape only, independent of causal_attn. causal_attn is instead passed to set_input_qsa. causal_attn is fixed per graph as it's part of the reuse key. Causal values are unchanged. Non-causal values now follow the reference rule, where every visible block competes on score and only unpooled cells are always selected. * context : state the causal_attn shape rule in the comment * cont : add TODOs --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* metal: support left and circular padding in GGML_OP_PAD Align Metal with CPU, CUDA and Vulkan: shift the source coordinates by the left paddings, wrap them around with the same wrap_around when circular, and read the source through nb00, which also fixes a right padding of a permuted source. A test case covers it. Drop the f32_4 kernel: its selection is disabled as slower, and it fails two pad cases once enabled. * metal: use a function constant for the circular pad variant Address review from ggerganov: replace the bool template with FC_PAD, as FC_upscale_aa does, so the pad kernel is compiled once and specialized per pipeline.
…ims on x86 (ggml-org#29423) * ggml-cpu: enable tiled flash attention for non-vector-multiple head dims on x86 * add AVX2 support for masked loading and storing in simd_gemm_ukernel_tail * ggml-cpu: fix FA softcap handling for padded KV tiles
* tests : use llama_context_ptr in test-recurrent-state-rollback Replace raw llama_context pointers with llama_context_ptr and drop the manual llama_free calls and cleanup lambda. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL * tests : run test-recurrent-state-rollback over all dummy models Add a --models DIR mode that mirrors test-save-load-state: iterate every dummy model, report PASS/FAIL/SKIP in a table and fail only when a model fails. Register a single ctest entry with ARGS --models instead of the four per-model registrations. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL * cont : fix typo * metal : allow fusing 0-element nodes to keep graph packing shape-independent The fusion packing in ggml_metal_fusion_max excluded 0-element tensors and the topk_moe/moe_reduce checks rejected n_tokens == 0, so graphs decoding batches with no outputs packed differently from the worst-case reserved graph. The Metal optimizer then reordered the nodes differently and ggml_gallocr_needs_realloc failed on the layout mismatch, forcing an unexpected graph re-reserve (caught by GGML_SCHED_DEBUG_REALLOC). Treat empty tensors like their non-empty counterparts: match them in the pattern sequence and only reject genuinely malformed shapes. Fused kernels dispatch zero threadgroups for empty graphs, which is a legal no-op. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL * tests : run test_multi_seq_split_replay as a separate test test_multi_seq_split_replay was invoked at the end of test_rollback, so its result was folded into the rollback status and it only ran when the rollback part passed. Give it its own test_status return, run both tests independently over both cache fills via a shared run_tests helper, and report them as separate rollback / split replay columns in the --models table with per-test summaries. The exit code fails when either test fails. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL * tests : loosen the split replay nmse bound to 1e-4 test-generate-models seeds its weights from std::random_device, and some generated lfm2 models drift up to ~1.7e-5 nmse on the split replay due to rounding noise, tripping the previous 1e-5 bound. Raise the bound to 1e-4 so the random generations stop flaking. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL * tests : reuse run_tests_for_model in single-model mode The single-model path duplicated the model init and the non-recurrent check from run_tests_for_model; route it through the shared helper instead. Model load failures now return FAIL rather than SKIP so that --model with a broken file still exits non-zero, and the helper loads with model_only like the --models loop does since the tests create their own contexts. Assisted-by: pi:llama.cpp/MiMo-V2.6-Flash-RL
…sor (ggml-org#29471) * Fix: Handle unaligned writes in ggml_backend_webgpu_buffer_set_tensor * Clang formatting
Supersedes ggml-org#29158 Signed-off-by: Adrien Gallouët <angt@huggingface.co>
70 upstream commits (8212c78..6c7a87f). No carried patch landed upstream. Conflicts, all between ggml-org#28702 (fused Q4_K gate/up + SwiGLU MMQ) and upstream ggml-org#24364 (llama_prec_policy + W4A4 MMQ path): - ggml/src/ggml-cuda/mmq.cu: take upstream's prec_src1 plumbing (GGML_CUDA_MMQ_PREC env, ggml_cuda_mmq_get_prec_src1, extra switch_type arg) and keep our impl/public-wrapper split and x_gate dispatch. The fused path pins prec_src1 to GGML_PREC_Q8 instead of reading op_params[3] of the GLU node, which is not a MUL_MAT. - ggml/src/ggml-cuda/mmq.cuh: keep both; our gate/up launcher block, then upstream's prec_src1 template on launch_mul_mat_q; our DECL_MMQ_GATE_UP_SWIGLU_CASE next to upstream's DECL_MMQ_CASE_W4A4. - ggml/src/ggml-cuda/template-instances/generate_cu_files.py: keep both appends (Q4_K gate/up instance, W4A4 instances for MXFP4/NVFP4); generated files already match.
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.
Merges 70 upstream commits from ggml-org/llama.cpp:
8212c7802..6c7a87f7e.Merge with "Create a merge commit" only. Squash or rebase merges drop the upstream ancestry.
Dropped patches
None. No carried patch has landed upstream.
Conflicts
All three are between ggml-org#28702 (fused Q4_K gate/up + SwiGLU MMQ) and upstream ggml-org#24364 (
llama_prec_policy+ W4A4 MMQ path):ggml/src/ggml-cuda/mmq.cuprec_src1plumbing (GGML_CUDA_MMQ_PREC,ggml_cuda_mmq_get_prec_src1, extraswitch_typearg) and keep our impl/public-wrapper split andx_gatedispatch. The fused path pinsprec_src1 = GGML_PREC_Q8instead of reading op_params[3] of the GLU node, which is not a MUL_MAT.ggml/src/ggml-cuda/mmq.cuhprec_src1template onlaunch_mul_mat_q, withDECL_MMQ_GATE_UP_SWIGLU_CASEnext toDECL_MMQ_CASE_W4A4.template-instances/generate_cu_files.pyUpstream gave every device-side MMQ config helper a defaulted
prec_src1 = GGML_PREC_Q8and removed the unused hostget_nthreads/get_occupancy. The fused Q4_K kernel's calls keep their meaning, andmmq_argsis unchanged.Other upstream changes worth noting for the benchmark: a new RMS_NORM+SCALE CUDA fusion, and GATED_DELTA_NET enabled on MUSA (does not affect CUDA).