* 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>
* 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
* 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>
* 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>
* 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.
- #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 #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 (#28948) and SYCL (#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 fee39dd92.
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
- 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
* 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
* 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>
* vulkan: add int8 coopmat quantized matmul shader
* apply scales inline
* use scalar sums
* probe and directly access coopmat values instead of going through shmem
* add q8_0 support
* add BK_STEP to shader, default to 2
* use larger workgroups
* double buffering
* preload scales
* coopmat load first, then wmma
* use float for scales
* add faster RDNA int->float conversion
* workgroup scheduling for cache proximity
* clean up
* use wave32
* restructure for vgpr use
* skip computation for inactive tiles
* only force subgroup size 32 on AMD RDNA
* use BK_STEP 4
* fix compilation
* move quant-specific prefetch function out of main file
* add q4_1, q5_0, q5_1 support
* restructure mmq cm1 functions
* enable mul_mat_id support
* fix segfault
* fix mul_mat_id bug
* support iq4_nl and mxfp4
* remove elem row/col fast path, invalid for RDNA4
* use shmem arrays for LUTs
* use 4-byte loads where possible
* add q3_k, q4_k, q5_k, q6_k and nvfp4 support
* fix l warptile
* improve performance
* improve performance
* improvements
* dedup b scales
* merge shmem arrays
* undo uint8_t, gate to RDNA3/4
* add RDNA4 architecture, use for hardcoded coopmat elem thread access, set BK_STEP back to 4
* improve offset application
* clean up
* fix iq4_nl and nvfp4 performance
* rdna4 tuning
* use BK_STEP 2 on MUL_MAT_ID
* adapt to upstream changes
* fix shmem support function, clean up comments
* fix warptile logic
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
* vulkan: add IQ4_XS support to the coopmat1 integer matmul shader (#28440)
Adds IQ4_XS to mul_mmq_cm1: dedicated block_a_load/block_a_to_shmem that
expand both nibbles of each packed32 word through cm1_kvalues, LOAD_VEC_A 8
and an IQ4_XS-sized a_panel_bytes estimate for the L2-friendly scheduling.
Assisted-by: OpenAI Codex
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* avoid compiling f16 acc shader variants
---------
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
* model : fold Ling 3.0 VL into the BailingMoeV3 architecture
Assisted-by: Scout
* model : keep shared NORM rope list intact when gating bailingmoe3 on mrope sections
---------
Co-authored-by: aetherbird <aetherbird@users.noreply.github.com>
* test-save-load-state : print a per-model results table in --models mode
in --models mode the output was very heavy: every model printed its
token dumps, per-test headers and PASS lines. instead, silence all
logging except the table itself (common_log_set_verbosity_thold(0)
leaves only LOG / LOG_LEVEL_OUTPUT) and print one row per model with
one column per test, colored PASS/FAIL/SKIP cells, row by row.
- run_save_load_tests_for_model returns a test_suite with a dynamic
std::vector<test_status> and continues past failures: tests 3-5 are
SKIPped when the baseline (test 1) fails, model init failure skips all
- per-test token dumps, test headers and PASS lines are demoted to
LOGV(LOG_LEVEL_INFO, ...) so they still show in single-model mode
- the table header/rows derive their columns from test_names; the
model name is printed and flushed before the suite runs so the model
currently in flight is always visible
- single-model output and exit codes are unchanged
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* test-save-load-state : print example usage on -h
add a print_usage callback passed to common_params_parse, so -h/--help
also shows example commands for the tool-specific --models option and
the -lv verbosity level
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* test-save-load-state : remove comments
ref: https://github.com/ggml-org/llama.cpp/pull/29316
Assisted-by: pi:llama.cpp/Qwen3.8-27B
make-release-desc.sh now emits "Changelog since [vX.Y.Z](<repo>/releases/tag/vX.Y.Z)"
instead of a plain version string, so the release notes link back to the previous release.
The repo URL is derived from the origin remote (SSH or HTTPS); if it cannot be
resolved (local run without origin), the title falls back to plain text.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
Resolve the target arch with get_model_architecture so vision targets
(e.g. Lfm2VlForConditionalGeneration) map to their text model for the vocab.
Fix double rope reorder for LFM2/LFM2.5 DSpark drafters
* tests: add backend option to test-llama-archs
* Update tests/test-llama-archs.cpp
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* remove extra space
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
Restore get_cache_directory() as fs::path as string() can be lossy on Windows
Partially reverts #29125
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
ggml_conv_1d_dw builds its im2col as f32 when the kernel is bf16, then
multiplies the two, so a depthwise convolution over bf16 weights asks
for kernel_mul_mv_f32_bf16, which was never instantiated. The base, the
_4 and the _short families are filled in next to their bf16 neighbours,
inside the same runtime guard, so a device without bf16 support is
unaffected.
* CUDA: enable sparse-fa for dsv4 prefill (again)
* CUDA: unroll the query loop of the sparse mask scan
The query loop of flash_attn_mask_to_sparse_indices has a runtime trip
count, which keeps the unrolled scan over the values of a lane from
issuing its loads together. Template the kernel on ncols1 so the loop
is bounded at compile time: batch one decodes compile to straight line
code and the scan drops from 46 to 17 us at 49k columns on sparse
decode shapes.
* CUDA: pick the out of bounds check of the sparse mask scan in host code
The query loop of the ncols1 == 8 scan keeps a runtime bound and an
early exit, so it does not unroll past its first iteration. Template the
kernel on whether the last group of queries is partial, decided on the
host from n_queries, and hoist the column bound out of the loop: the
loop becomes straight line code and the batched sparse op at 49k
context drops from 586 to 244 us.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
