* cuda: add F16 input to the FWHT
The CUDA FWHT accepts F32 input only. This makes the source type a template
parameter, so the kernel reads an F16 source directly instead of requiring a
converted copy. The F32 path is unchanged.
supports_op accepts an F16 src1 against an F32 src0 for the Hadamard hint.
Every other F16 src1 against a non-F16 src0 is still refused.
ggml_cuda_op_mul_mat_use_fwht is the single predicate both supports_op and
the dispatch call now share, checking contiguity and same-shape(src1, dst)
in addition to the type/hint conditions above. Without a shared predicate,
supports_op could admit an op that ggml_cuda_op_fwht then rejects only after
the unconditional same-shape assert has already fired; that gap predates
this change (it applies to the existing F32 path too) but this PR is what
touches supports_op, so it closes it here.
test-backend-ops on an A10 (lambdalabs): MUL_MAT 1297/1297, including all
24 Hadamard cases (18 existing F32, 6 new F16).
* cuda: use ggml_cuda_cast in the FWHT load, drop the comment
* llama : add discard for deferred state writes
* llama : add tensor zeroing helper for backends without tensor memset
* llama : clear K/V data after failed sequence restore
* llama : clear recurrent state data after failed sequence restore
* llama : simplify discard and restore cleanup
* llama : report error when abnormal cell count is found in state_read_meta
* llama : clear attention state on hybrid restore failure
* tests : cover failed state restore cleanup
* llama : clear MLA state on dsa restore failure
* tests : update test for rebased test suite
* llama : clarify comment in llama_memory_recurrent::state_read
* 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 0688cf43d5.
* 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>
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>
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
* 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>
Anav Prasad