* metal: fix memory leak if model is freed without any GPU operations
* metal: run dummy work only if residency sets are used
* metal: wrap function in #if defined
* metal: measure system-wide wired memory in test
* metal: always build regression test
Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
---------
Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
ggml_cuda_should_use_mmq() selects MMQ purely from the quantization
type. The current MMQ configurations are designed and maintained against
a minimum of 48 KiB per-block shared memory, the limit provided by
NVIDIA Pascal GPUs and later. On devices that report less, no supported
MMQ tile fits and mul_mat_q_switch_J() aborts when every tile size
exceeds the device's per-block shared memory budget.
Disable MMQ when smpbo < 48 KiB so the caller falls back to the BLAS
path instead of hitting GGML_ABORT. Some current MUSA QY1 devices
report only 28 KiB and are covered by this guard.
Reproduced on a Moore Threads MTT S70 (arch mp_21, 28 KiB shared memory
per block) with an RWKV-7 0.1B Q8_0 model:
$ llama-bench -m rwkv7-g1d-0.1b-Q8_0.gguf -p 128 -n 0
J_best=0
ggml/src/ggml-cuda/template-instances/../mmq.cuh:1521: fatal error
(core dumped)
Only prefill (batch > 1) is affected; token generation is fine. After
the fix the same device falls back to the BLAS path:
Q8_0 pp128 1470.7 t/s, tg8 55.3 t/s (was: abort)
FP16 unchanged
Q4_K_M unchanged
This matches a -DGGML_CUDA_FORCE_CUBLAS=ON build (pp128 1464.2 t/s),
which confirms the fallback path is the one being taken.
This is not MUSA-specific: any device with less than 48 KiB per-block
shared memory is affected.
Co-authored-by: KakaruHayate <KakaruHayate@users.noreply.github.com>
* Add overlap glu variant to support all archs, fix recurrent-state-rollback test
* format
* Fix all arch overlapped ranges
* format
* diagnose bus error on apple ci
* More testing
* more testing
* more targeted testing
* Fix bug in alignment for > 4gb buffer offsets
* Fix bug in view offsets
* Try avoiding multi_buffers
* not fixed yet, more logging :(
* Handle edge case in set_rows
* Try looking at view source
* Skip deepseek32 for now and clean up trace infrastructure
* simplify skipping
* last cleanup
* actually final cleanup
* update handling of overlap
* format
* try skipping other failing model
The Adreno KQ/KQV image1d kernels (ggml_cl_mul_mat_kq_kqv_adreno) ignore
dim 3 entirely: the sub-buffer covers only nb02*ne02 bytes and the kernel
receives no ne03/ne13/nb03/nb13 arguments. With the unified KV cache,
multi-sequence batches (e.g. llama-perplexity with its default -b 2048,
n_seq=4, or a multi-slot llama-server) present KQ/KQV as 4D tensors with
ne3 = n_stream, so every stream past the first reads the first stream's
K/V and produces garbage. Flash attention masks the bug where it is
enabled; devices where FA is declined (e.g. Adreno 740) hit it with
default settings.
Route ne03/ne13 > 1 to the general path, which handles dim 3, and honor
view_offs when creating the sub-buffers (currently always 0 for tensors
reaching this function, but the function would silently misread any
future view).
Llama-3.2-1B-Instruct Q4_0, wiki.test.raw, 8 chunks, -ngl 99:
- Adreno 740, default: PPL 1817.64 -> 15.61
- Adreno 740, -fa 0: PPL 1941.64 -> 15.61
- Adreno 840, -fa 0: PPL 1943.90 -> 15.50
- single-stream (-b 512) results unchanged (15.6090)
- test-backend-ops -o MUL_MAT on 740: identical before/after (909 OK,
12 pre-existing q6_K failures)
* vulkan: add iq4_nl support back to FA
I was originally concerned about wasting shared memory on the LUT, but it's small
and unlikely to matter in practice.
Also support q1_0 for non-coopmat2.
Fixes#23681
* remove q1_0 FA support
* ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration
* cuda: added SSD CICD fixes for CUDA / HIP / MUSA / MSVC.
* ggml-cuda: review comments fixed.
* ggml-cuda: Fuse M matrix materialization into pre_matmul kernel and enabled test.
* ggml-cuda: test updates and fixes
* ggml-cuda: test updates to remove hardcoding of tensor initialise data limits.
* ggml-cuda: ssd minor review comment fixed.
* ggml-cuda: ssd minor CICD fixed.
* CUDA SSD: Fixes correctness by promoting s0_stride_seq to int64_t, improves memory coalescing in ssm_ssd_prepare_dt_kernel, and boosts efficiency by merging B_weighted and C_scaled; also addresses prior review comments.
* cuda: fix sdata read-write race in prepare_dt fallback scan loop
* sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path
The scale was uploaded with an async memcpy sourced from a stack local. On the
in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is
large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host
stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale.
Output then collapses to a single repeated token and the KV cache is poisoned
for the rest of the session.
Short contexts win the race by accident, and test-backend-ops caps
FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous
device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the
symptom only by keeping the frame alive across the copy at the cost of a host
sync on every FA call.
Fix: cache one device scalar per (device, value) -- the scale is constant per
model -- and upload it synchronously once. The single-device fast path (no
per-call host sync) is then safe: every device-side hazard already serializes
on the in-order queue. The multi-GPU conservative wait is kept unchanged.
Also:
- GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that
routes very long sequences to the native FA kernel.
- test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B
geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind
spot. Note the race itself needs a live multi-op pipeline to reproduce;
single-op runs pass even on broken builds.
Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output
byte-identical at temp 0 to the native FA path through 32k-deep prefill, with
prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth).
Assisted-by: Claude Fable 5
* sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it
Review feedback on #25880:
- read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via
ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it)
- document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md
runtime table
Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536
shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels
(long-sequence fp16 accumulation drift, present before this PR) and would fail
CI for an unrelated reason.
Assisted-by: Claude Fable 5
* sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled
Assisted-by: Claude Fable 5
* sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly
Assisted-by: Claude Fable 5
* Update ggml/src/ggml-sycl/fattn-onednn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* sycl: write the SDPA scale from a kernel instead of caching it
The per-(device, value) scale cache was a function-local static
unordered_map with no synchronization, so concurrent backend instances
could access and rehash it at the same time.
Write the scalar with a single_task instead. The value is captured into
the command, so no host memory has to outlive the call -- which is what
the use-after-return fix needed in the first place. That removes the
shared container, the leaked device allocation and the string key, and
it also closes the remaining async-memcpy-from-a-stack-local on the
first flash-attention call.
Ordering does not rely on timing: the queue is created with
sycl::property::queue::in_order and the dnnl stream wraps that same
queue, so the write completes before the SDPA reads the scalar. The
multi-GPU wait_and_throw() branch is unchanged.
Also drop the <cstdlib> include, which is unused.
Assisted-by: Claude Opus 5
---------
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* hex-l2: use dirty ranges for flushing
* hex-l2: simplify range based flush logic
* hex-l2: optimize dirty range scans
* hex-hvx: support for reduce_max_i32
* hex-mm: optimize fused MUL_MAT+ADD to use vtcm for bias when it fits
* hex-mmid: optimize mmid row-mapping generation
* hex-mmid: optimize mmid row-mapping generation
* hex-mmid: optimize mmid row-mapping generation (round2)
* hmx-mm: optimize output proc by tiling (col-chunking)
* hex-fa: start the next q dmas a bit earlier
* hex-fa: prefetch Q even earlier
* hvx-fa: optimize softmax to keep things in hvx registers
* hex-fa: hoist const register init in softmax loop
* hmx-fa: kick off next-qkv DMAs before o-proc
* hmx-fa: hoist various checks out of the inner loop
* hmx-fa: adjust the cost model to better balance softmax work across hvx threads
* hmx-fa: overlap diag rescale build with last HMX task
* hmx-fa: optimize idx update in output proc
* hmx-fa: unroll the softmax loops for improved perf
* hmx-fa: overlap qk-dot with softmax, double-buffer p and s tiles
* hex-trace: double the default number of trace entries
* hex-trace: add trace events for opbatch and buffer mgmt
* hex-trace: overhaul tracing to simplify runtime event handling and support opbatch stats
* hex-trace: replace ascii timeline diagram with pipeline bubbles detector
* hex-trace: handle missing start/stop events
* hex-dma: always log stop/start trace events even for dummy dmas
* hex-scripts: fix flake warnings
* Squash history before conflict-resolution during rebase on master
WIP commit
Add 32-byte loads, restore per-block amax
Use nvfp4x4 intrinsic when available
Fuse per-channel amax and quantization kernels
Do pointer arithmetic only once on x
Remove unnecessary ternary in the load
We assert on host side that ne00 is 64-aligned
Add back scale-search, but optimize it with intrinsics
Code cleanup
Make scale in MMQ-epilogue NVFP4-specific/restrictive for now
Remove unneeded include, add comment
Fix trailing whitespace
Guard __builtin_align__(32) struct to NVIDIA
Seems like HIP doesn't have this available, see https://github.com/ggml-org/llama.cpp/actions/runs/29438651734/job/87431623001
* compiler massaging to avoid unnecessary LDCs
* kvalues_mxfp4 -> kvalues_nvfp4 in quantize_mmq_nvfp4
* Always pass in src1_scale.ptr
* Extract ggml_cuda_is_aligned helper
* hex-geglu: optimized all-in-one geglu microkernel
* hex-geglu: enable non-contiguous src and strided DMA
* hex-act: enable non-contiguous srs and strided DMA for rest of ACT ops
* hex-act: generalize GLU per-thread functions via DEFINE_GLU_PER_THREAD macro
* hexagon: move UNARY_SILU and UNARY_GELU to unary-ops
* hex-act: replace the generic ops_context scratchpad usage with a local htp_vtcm_layout computation per act op.
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* ggml: enable PowerPC backend variants on AIX
Allow the PowerPC CPU backend variants to be built on AIX by extending the platform check in the CMake configuration. This reuses the existing PowerPC backend implementations without changing their behavior.
Also fix a missing semicolon in the PowerPC Q0 matmul implementation.
* Fix missing semicolon in sgemm.cpp
* webgpu : add CONV_2D_DW (depthwise conv2d) kernel
Implement GGML_OP_CONV_2D_DW for the WebGPU backend,
ported from the Vulkan backend's conv2d_dw.comp.
Assisted-by: Claude Opus-4.8
* Remove unnecessary comments in webgpu support
* update supported ops tables, triggered by adding webgpu CONV_2D_DW
* cuda: add k-quant support to GET_ROWS
Device-side embedding lookups require GET_ROWS to handle the k-quants
used by common GGUF recipes (Q4_K_M stores token_embd as q6_K). Without
it the backend rejects the op and the scheduler falls back to the host,
copying the full embedding matrix back on every token in single-device
graphs.
Factor the super-block dequantizers out of the dequantize_block kernels
in convert.cu into shared device functions in dequantize.cuh and reuse
them from a new k_get_rows_kq kernel : one thread block dequantizes one
(dst row, super-block) pair with the existing thread layouts, 32 threads
for q4_K and 64 for the other k-quants.
Covers q2_K to q6_K in get_rows_cuda and supports_op. i-quants are left
as a TODO.
* cuda: add i-quant support to GET_ROWS
Extends the shared super-block dequantizers to the nine i-quants and
reuses them from k_get_rows_kq with the 32-thread layout of the matching
convert.cu kernels. supports_op gates the k-quant and i-quant path on
ne0 being a multiple of QK_K, which iq4_nl does not guarantee on its
own (QK4_NL sub-blocks). mxfp4 is left as a TODO.
* cuda: add mxfp4 support to GET_ROWS
Moves the mxfp4 dequantizer into the shared super-block helpers and
reuses it from k_get_rows_kq with the 32-thread layout of the matching
convert.cu kernel. mxfp4 joins the ne0 % QK_K gate in supports_op since
its 32-value sub-blocks do not guarantee QK_K-aligned rows on their own.
This closes GET_ROWS type coverage on CUDA: every quantized GGML type
now takes the direct device path.
* cuda: gate the GET_ROWS row size only for 32-value sub-block types
Address review from @pwilkin: the i-quant commit replaced the return
shared by the whole supported type cascade, so f16/f32/bf16/i32 and the
legacy quants also inherited the ne0 % QK_K == 0 gate and any row size
that is not a multiple of 256 fell back to the scheduler. Split the
cascade: unconditional support is restored everywhere, the gate stays
only on iq4_nl and mxfp4 whose 32-value sub-blocks do not guarantee the
QK_K super-blocks the kernel iterates on.
* Refactor vk_queue to use per-instance mutexes and unique handles
* integrates VK_KHR_internally_synchronized_queues, abstracting the queue submission into a polymorphic interface that completely bypasses host-side mutex locking when driver-side synchronization is supported
* fix compilation error
* fix duplicate pNext chain for VkPhysicalDeviceInternallySynchronizedQueuesFeaturesKHR
* add fallback defines for VK_KHR_internally_synchronized_queues
* add null checks for queues in vk_device_struct destructor
* use unique_ptr for outer queues to enforce exclusive ownership and optimize lifetime
* use static constexpr for eInternallySynchronizedKHR
* add lock guard to ggml_vk_create_aliased_queue for thread safety
* initialize sync_query_features.internallySynchronizedQueues to VK_FALSE
* reuse sync_query_features for internallySynchronizedQueues and simplify chaining
* refactor internallySynchronizedQueues detection
* fix internallySynchronizedQueues query guard
* use eInternallySynchronizedKHR constant
* fix self-referential alias for eInternallySynchronizedKHR
* use macro for eInternallySynchronizedKHR fallback
* fix internallySynchronizedQueues query timing in ggml-vulkan.cpp to prevent device creation mismatch
* reset sync_query_features.pNext before reusing in device creation chain, also removed the redundant second probe call
* refactor internally synchronized queues detection to use chained feature query and avoid redundant API calls
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* rename sync_enable_features to internally_synchronized_queues_features
* queue_flags is still computed before has_internally_synchronized_queues is set
* fix trailing whitespace
* replace eInternallySynchronizedKHR macro with static constexpr
* preserve source queue semantics in single-queue aliased transfer queue
* vulkan: fix cmd_pool access via pointer for compute_queue unique_ptr
* vulkan: lock queue during debug label emission when not internally synchronized
---------
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
k_get_rows_float did a scalar one-element-per-thread copy and recomputed the
row-invariant work (index load, fast_div_modulo, src/dst row pointers) for
every element. Hoist that out of the per-element loop, and add a vectorized
path (k_get_rows_float_vec) that copies one int4 (16 B) per thread for the
contiguous same-type (no-cast) case.
The vectorized path is gated at compile time (is_same<src0_t, dst_t>) and at
runtime on 16-byte alignment of the base pointers and all row strides and on
ne00 % VEC == 0. Vectorizing divides the block count by VEC, so a small
single-row gather can drop below the device CU count and regress; an
occupancy gate keeps those on the block-rich scalar path.
On Strix Halo (gfx1151) the DeltaNet recurrent-state gather (ne00=524288)
drops 18.6us -> 13.0us (rocprofv3 HW timestamps), faster than the Vulkan
backend, with no regression on the small conv-state gather; total get_rows
-27%. test-backend-ops GET_ROWS passes (47/47).
Assisted-by: Claude Opus 4.8
* vulkan: Support Q2_0
The backend perf tests for mat-vec-mul weren't very good at first (worse than
q2_k), doubling the rows per workgroup made a big difference.
* reorder
* resolve merge conflict, adjust err threshold for f16->q2_0 set_rows
Daniel Bevenius
Georgi Gerganov
Pasha Khosravi
Niklas Wenzel
Jonathan Clohessy
Kakaru
Titaniumtown
Aman Gupta
Geramy Loveless
Reese Levine
Hongqiang Wang
Jeff Bolz
Bhavik Sharda
meatposes
Nick Lafleur
Beinsezii
shalinib-ibm
yzyyzyhhh
Todor Boinovski
Yongmin Yoo 유용민
Johannes Gäßler
kumaal
Max Krasnyansky
adgup-qti
Ilia Ilmer
Oliver Simons
m1el
Pascal
Wei Wang
Kamalesh VS
Winston Ma
Markus Ebner
Piotr Wilkin (ilintar)
lhez
Aaron Teo