3123 Commits
Author SHA1 Message Date
PascalandGitHub 3f5e94d7c2 webgpu: align tensor bindings to the type block size (#28382)
Walk the binding offset back until the distance to the tensor is a
whole number of blocks, so block quantized views get a valid element
offset in the shader.
2026-09-12 06:40:21 +02:00
eafe15a5e3 hexagon: support for multi-device model split (aka row-split) (#28589)
* hex-row-split: add support for multi-device row spliting

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>

* hex-mdev: add work splitting to fused kernels

* hex-mdev: use mdev_ prefix for all multi-device state

* hex-mdev: make device configuration more expressive to support device groups

* hex-mdev: fix mdev session init

* hex-mdev: fused nx (2x,3x) matmuls must update row counts for each w/o

* hex-mdev: fix MUL_MAT work partitioning bugs introduced by mdev

* hex-cont: fix crashes with new tests due to wrong striding

* hex-mdev: move fences after l2flushes

* hex-cont: fix work splitting for mnpu -- align chunks to cachelines

* hex-mdev: fix CPY tests with multi-dev

* hex-mmid: fix work partitioning with mnpu

* hex-mm: fix test failures with mdev

* hex-binary: fix work partitioning for mdev

* hex-argsort: fix mdev partitioning

* hex-mdev: fix work partitioning and general updates for all simple ops

* hex-fa: fix mdev work splitting issues

* hex-mdev: fixing more failing ops test

* hex-mdev: update the rest of the ops

* hex-mdev: refactor all mdev splitting logic to be contained within if (mdev_count > 1) {...}

* hex-mdev: fix macros

* hex-mdev: simplify session flush logic

* hex-sync: fix recursion in session flush

* hex-mdev: factor out fence buffer and allocator

* hex-fence: make fence allocation more robust with reserved slots for mdev

* hex-mdev: keep all mdev state in htp_mdev_group

* hex-mdev: further cleanup mdev group handling at the host

* hex-mdev: update group idx in the opbatch before serializing

* hex-batch: remove separate op_pending and use batch_req/rsp_seq

* hex-async: workaround another missing tensor_init in ggml-meta

* hex-fence: cleanup and robustify fences and error handling in multi-device scenarios

* hex-ar: improve ALLREDUCE error handling

* hex-async: robust error handling for op_cpy_fence

* hex-async: use seq0 from allreduce context to allocate fence_seq

* hex-mdev: fix remaining issues with fence and barrier clearing in CPY_FENCE

* hex-misc: realign macros and fix misplaces trace events

* hex-misc: align macros

* hex-mdev: fix unclone buffer re-entrancy

* hex-glu: fix mdev partitioning logic

* hex-mdev: make buffer uncloning/cleanup work with tensor-split scenarios

* hex-mdev: tighten up the can_split check in act-ops

* hex-mdev: factor out common bits of the partitioning logic

* hex-mm: minor realignment of the macros

* hex-bufs: fix incorrectly placed assert for MAX_BUFS

* hex-pad: tighten up gating checks for PAD

* hex-kparams: make sure all kernels properly use kparams->n_threads

* hex-docs: update user and developer docs with new features and detailed guide for ops development

* hex-scripts: update run script to properly parse dev groups

* hex-misc: formatting

* hex-sess: minor cleanup for session init

* hex-ar: fix vtcm size calc in allreduce kparams

* hex-scripts: fix flake8 warnings

* hex-rope: update ROPE to support mdev work split

* hex-ops: remove redunant checks and minor reformat

* hex-dev-guide: update dev-guide to avoid redundant null checks

* hex-async: improve event_wait, event_sync and fence implementations

* hex-async: remove synchronous flush from event_sync

* hex-async: symplify fence recovery protocol and make sync more robust

* hex-async: futher simplify error recovery for fences

* hex-err: return status instead of just -1

* hex-async: print all seq nums in hex

* hex-async: make sure fences flush dirty ranges

* hex-async: add dirty ranges merging to reduce fence flushes

* hex-async: properly sync before freeing the event

* hex-async: make sure fence owner session is not overriden

* hex-async: more fence write order more robust

* hex-async: make sure not to fuse ALLREDUCE+ADD if their dsts overlap

* hex-fusion: cleanup redundant checks

---------

Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com>
2026-09-11 20:46:51 -07:00
d3146f2b56 ggml-webgpu: Update to a recent version of Dawn (#28683)
* ggml-webgpu: Update to a recent version of Dawn

* No module scanning

* Accept review suggestion to update comment

Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>

---------

Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
2026-09-12 10:47:29 +09:00
Aman GuptaandGitHub 43f3dda623 ggml: skip 0-sized ids tensor when offloading selected experts (#28739) 2026-09-11 15:17:08 +02:00
Foad Abo DahoodandGitHub 5bda51bfbc metal : skip the empty half of the mul_mm_id token tile (#28301)
kernel_mul_mm_id splits its NR1 = 32 token tile into two 16-row halves and skips
the upper half when the expert did not fill it, on both the tensor and simdgroup
paths. The tB extents are corrected to (NK, NR1H) for the [NR1][NK] row-major tile.

The B tile is staged unconditionally, as on master: rows past nr1 restage a clamped
duplicate of a valid row, lie in the output-row dimension so they never contribute
to a valid row, and are dropped by the final store loop.

test-backend-ops: re-draw the expert ids between perf iterations of test_mul_mat_id
so MoE perf numbers are not warm-cache, and add token-tile boundary coverage using
n_used == n_mats, which routes every token to every expert so each expert receives
exactly n rows; n = 32, 33, 47, 48, 49 reach mul_mm_id and leave a last tile of 32,
1, 15, 16 and 17 rows.
2026-09-11 14:12:55 +03:00
Daniel BeveniusandGitHub 3bcfeb700f cmake : add PCH and unity build to improve build times (#28091)
* scripts : add initial profiling script (wip)

* src : add precompile headers (PCH) for models.h

* common : add common.h as PCH

* ggml : add PCH for ggml-impl.h

* mtmd : use PCH for models.h

* scripts : add script to build with Server/Tools/Tests

* server : add PCH for common.h

* docs: add profiling progress notes (wip)

* ggml : add exclude for GCC + SVE on ARM

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33393906061/job/99493756214?pr=28091

* ggml : attempt to fix use of std::hardware_destructive_inference_size

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33396221677/job/99501265689?pr=28091

* squash! ggml : attempt to fix use of std::hardware_destructive_inference_size

Add a version check for GCC 12 to conditionally apply the `-Winterference-size`
pragma.

* editorconfig : exclude profiling reports dir

This directory will not be included in the merge later and this commit
can be ignore at that point. Just fixing to keep CI happy.

* ggml : skip PCH for gcc on non-x86 architectures

* tests : add PCH for peg-parser/tests.h

There are 7 peg-parser tests that can share one PCH instead of then each
parsing the full tests.h.

* common : add PCH for chat.h

* docs : update linux build profiling full results

Just updating after a number of PCH additions. These are not exact
figures and will vary a bit from run to run, but they give a general idea
of the performance impact of PCH.

* cmake : introduce unity build for models

This commit introduces a unity build for the models to improve
compilation time.

The improvements were roughly the following:
```console
+------------------------+-----+------------+------------+------------+
| Build                  | TUs | Frontend   | Backend    | Total      |
+------------------------+-----+------------+------------+------------+
| Full,    master        | 396 |   811.0 s  |   692.2 s  | 1,503.2 s  |
| Full,    with PCH      | 405 |   380.0 s  |   664.7 s  | 1,044.7 s  |
| Full,    with PCH + UB | 264 |   357.7 s  |   635.7 s  |   993.4 s  |
+------------------------+-----+------------+------------+------------+

TU   = Translation Unit.
Full = includes Server, Tools, and Tests.
PCH  = precompiled headers.
UB   = unity build for models.
```

* docs : update linux profiling table with unitiy build results

* docs : update mac profiling results to include unity build [no ci]

* docs: remove profiling reports

* scripts : merge build profile scripts into one script

I was lazy before and just copied the first script to enable Tests,
Server, and Tools. This now merges them into a single script.

* Revert "editorconfig : exclude profiling reports dir" [no ci]

This reverts commit 2922a12118.

* src : rename ggml_view_2d_slice to gemma3n_view_2d_slice

This is to be consistent with the rename in gemma4.cpp which was
required to avoid a name clash.

* cmake : add build profile script for windows [no ci]

This commit adds a port of the scripts/build-profile.sh script to
windows powershell.

This was developed on Windows on ARM but should work on X64 as well but
needs to be tested there as well.
2026-09-11 13:01:29 +02:00
Georgi GerganovandGitHub a2878d30df metal : single-source fusion table + fusion debug rework (#28164)
* metal : rework fusion patterns into a single table

All fusable op patterns for the Metal backend are now declared once in a
fusion table (ggml-metal-fuse.cpp) and consumed by both the graph optimizer
(ggml_metal_fuse_max, packing) and the op encoders (ggml_metal_fuse_next,
compute). The two phases share the same pattern table plus ggml_can_fuse_subgraph_ext
for the structural checks, and differ only in the mode used for the pattern
check (STRUCTURAL at optimize time, since tensors are not allocated yet, and
FULL at compute time, including Metal buffer placement). This also protects the
snake activation (MUL + SIN + SQR + MUL + ADD) from being reordered during graph
optimization, which was previously unprotected.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : fix absolute output indices in fusion patterns

ggml_can_fuse_subgraph_ext expects the outputs array to contain absolute graph
node indices (it indexes cgraph->nodes[outputs[i]]), but the fusion table query
was passing a relative index (n_ops - 1). As a result the last node of every
pattern was not recognized as an output and was subjected to the elidable
use-count check, which failed for essentially all fusions. This silently
disabled the norm/MUL fusion and caused a ~5% token-generation regression.

Pass the absolute graph index of the last node instead.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : fuse gated_delta_net with cache cpy

Add GGML_METAL_FUSE_GDN_CACHE to the fusion table: when the gated_delta_net
kernel is followed by a cpy that scatters its recurrent state snapshots into
the KV cache, the kernel writes the snapshots straight into the cache buffer
and the trailing cpy is elided.

The gdn output has other consumers (the attn scores view), so unlike the
elision-chain patterns this is not a simple chain: a 'raw' flag on the fusion
pattern skips the generic chain/shape and ggml_can_fuse_subgraph_ext checks,
making the pattern-specific check callback the sole validator. Packing
(ggml_metal_fuse_max) now matches on the same view-transparent node sequence
that the compute phase uses, so the gdn + cache cpy group is packed along with
any intermediate views and stays adjacent through the reorder.

The fused cpy is a view consumer of the gdn (it writes the cache directly),
so its mem-range is skipped in the encoder; the skip is restricted to CPY
nodes consuming the previous fused node through a view so other fusions are
unaffected.

Add test_gated_delta_net_cache_fusion and register 5 cases.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : drop is_view_consumer mem-range skip

The is_view_consumer skip was carried over from the upstream gated_delta_net
cache-fusion draft, but it is not needed: keeping the elided cpy's mem-range in
the concurrency tracker only ever adds a (conservative) memory barrier at the
fusion point. It can never remove a barrier, so it cannot introduce a race. The
worst case is one spurious barrier per gdn+cache-cpy fusion, which is within
run-to-run noise on Qwen3.5-0.8B Q8_0.

Dropping the check keeps the mem-range loop uniform for all fused groups.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : rename gated_delta_net fused state output args

Rename the fused cache-write kernel argument to match the rest of the kargs:
state_out_stride -> nb_out (and widen it to uint64_t), and the local buffer id
bid_state_out -> bid_out.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : rename raw fusion flag to unsafe

raw did not convey that the flag opts a fusion pattern out of the generic
elision-chain safety net (ggml_can_fuse_subgraph_ext + chain/shape checks).
rename it to 'unsafe' to make explicit that the pattern's check callback is the
sole validator and must re-establish the safety guarantees itself.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : tidy fusion pattern checks and table

- const-correct ggml_metal_fuse_outputs buffer
- annotate unused check-callback parameters
- drop a redundant size_t cast
- align the ops/table initializers and add blank-line separation

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* metal : add generic fusion stats via ad-hoc proc-address API

Add a device-owned fusion context that lets a test tool count how many
times each fusion pattern fires and toggle fusion. It is exposed through
the ad-hoc ggml_backend_reg_get_proc_address mechanism with generic names
so the testing tool is backend-agnostic:

- ggml_backend_fusion_stats_init: start collecting fusion stats; when a
  context is created afterwards it registers the labels/counters and
  encodes single-threaded (n_cb == 0) so the counters are race-free
- ggml_backend_fusion_stats_reset / _get_stats / _set_enabled

The context lives on the metal device (not on the last backend context),
so counters accumulate across contexts and reads are always consistent.
The enable/disable toggle is initialized from GGML_METAL_FUSION_DISABLE
and can be overridden by the test through set_enabled. Labels are
synthesized from the fuse table via ggml_metal_fuse_label (e.g.
"GATED_DELTA_NET+CPY").

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : add fusion count regression test with per-backend baseline

test-fusion runs every dummy model generated by test-llama-archs on a
single backend (single-threaded encoding, n_cb == 0) with fusion enabled
and disabled, and for each mode (prefill / decode) reports the per-fusion
counters and the NMSE between the fused and unfused logits, plus the NMSE
against a CPU reference.

A fusion pattern that silently stops matching (or fires when it should
not) is caught as a regression by comparing the counters against a
committed per-backend TSV baseline:

- --record writes the golden baseline, --check (default) validates it
- the unfused run doubles as a control: its counters must be all-zero
- NMSE is skipped when it is NaN or the arch is already broken on the
  device (e.g. plamo2 on Metal), so the count check is the hard gate
- baseline counts depend only on graph structure, not weights (verified
  stable across weight seeds)
- the fusion stats API is resolved through the ad-hoc get_proc_address
  mechanism with generic names; a backend that does not export it makes
  the test fail with an error

The committed MTL0.tsv baseline covers 110 dummy archs (298 rows).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : rename fusion api helpers to match stats_init signature

Align the test with the ad-hoc fusion stats API: fusion_stats_init no
longer takes an enable bool (stats are turned on by calling it), so the
proc-address wrappers and typedefs are renamed to the api_* convention.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : rename backend to device in fusion test CLI

The fusion test operates on a compute device (e.g. MTL0), not a backend,
so rename the --backend argument to --device and the backend_name
variable to device_name. Keep "backend" where it refers to the ggml
backend interface (the ad-hoc proc-address mechanism).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : add --model and --help to fusion test

--model FILE runs the fusion regression test over a single model file
instead of enumerating a --models DIR. --models and --model are mutually
exclusive. Also add a --help/-h option that prints the usage.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : use backend base name for fusion baseline output

The fusion test is invoked with a specific device name (e.g. MTL0), but
its output - the recorded baseline and the header it writes - should be
named after the backend base name (e.g. MTL, via ggml_backend_reg_name),
since the counters depend on the backend, not on the specific device
index. Rename the committed baseline to MTL.tsv.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : run fusion test from ci instead of ctest

The fusion test needs Metal and generates a lot of dummy models, so it
does not belong in the generic ctest suite. Move it to ci/run.sh as
gg_run_test_fusion, gated on GG_BUILD_METAL like
gg_run_test_llama_archs_tensor_split: it generates the dummy models with
test-llama-archs -o and then validates the fusion counts against the
committed baseline. test-fusion.cpp is still built (llama_build) but no
longer registered as a ctest.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : align fusion baseline TSV columns

Pad the TSV fields to fixed widths so the columns line up regardless of
the variable arch and fusion-label lengths, and trim each field on parse
so the padded file is still accepted. Regenerate the committed MTL.tsv
baseline in the padded format (data unchanged, verified identical modulo
padding).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : widen label column and align fusion TSV header

Give the label column more room (28 chars) and fix the column header
widths so they match the data rows (moe/mode/label), keeping the header
aligned with the values. Regenerate the MTL.tsv baseline in the new
format (data unchanged).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : switch fusion baseline from TSV to CSV

Use comma-separated values like the rest of the project, keeping the
padded, aligned columns. Split on ',' and trim on parse. Rename the
committed baseline to MTL.csv (data unchanged, verified identical modulo
padding/separator). Update the ci/run.sh check path accordingly.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* cont : rebase + update MTL stats

* tests : avoid graph reallocations for some archs

* metal : tidy fusion debugging context and op init

- simplify the shared fusion debugging context comments
- shorten the ggml_metal_fusion struct comment
- align the ggml_metal_fuse struct fields and comments
- move the fusion parameter of ggml_metal_op_init right after dev

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : dedup fusion baseline into any mode

prefill and decode always produce the same per-graph fusion count, so
store a single row per label with mode = "any" and the per-graph count
instead of two rows. this halves the baseline size and keeps the check
stable.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* ci : move fusion model generation to a separate step

the dummy models generated by test-llama-archs are reused by other tests,
so generate them once in their own step instead of inside test_fusion.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : bump nmse thold

* models : fix plamo2 graph

* tests : remove "skip" logic from test-fusion

* tests : set qwen3tts dummy vocab to codec head size

the dummy qwen3tts model used a vocab of 4096 while the codec head is
3072, so the graph padded the output with -inf which made the NMSE in
test-fusion produce NaN. use the exact codec head size instead so the
padding is not generated at all.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* tests : regen fusion baseline

reflect the plamo2 graph fix, which changed its fusion pattern split
(RMS_NORM+MUL 11->10, RMS_NORM+MUL+ADD 3->4; same total).

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* ci : skip dummy model generation on OpenVINO

test-llama-archs does not build on the OpenVINO platform, so do not try
to generate the dummy models there.

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

* cont : minor

* tests : enable test-llama-archs on windows

* cont : disable on windows + workaround

* metal : naming nits

* test-fusion : add instructions to update baseline

* context : fix Kimi-K3 graph reserve

* fusion : update MTL

* cont : fix naming

* metal : rework fusion info storage

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* metal : align fusion info API

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* metal : use opaque fusion handle in ad-hoc API

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* ci : move fusion test to dedicated workflow

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp

* cont : run only on ggml changes

* cont : simplify

* fusion : remove multi-output stuff for now

* ci : fix typo
2026-09-11 12:41:54 +03:00
Foad Abo DahoodandGitHub aac810230f metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024 (#28692)
* metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024

Generalize the row split from #28086 to the six other kernels that use the
same lane-to-block mapping: iq1_s, iq1_m, iq2_xxs, iq2_xs, iq2_s and iq3_s.

Each of them assigns one 32-element chunk per thread, so when a row has
fewer than 32 chunks the rest of the simdgroup is idle. When nb32 < 32 and
nb32 divides 32, 32/nb32 threads now share each chunk and each takes a
slice of the rows, reusing the FC_mul_mv_split function constant and the
dispatch wrapper introduced for iq3_xxs.

The plain path is untouched: wide matrices keep one thread per chunk and
N_R0_<TYPE> = 4. Only the split path uses N_R0_<TYPE>_SPLIT = 8. The
K-quants have the same idle-thread issue but a different lane mapping, so
they are left for a separate change.

* metal : offset the src0 row pointer once in the iq mul_mv kernels

q2, dh, sc, qh and signs are all derived from xr, so the row slice
offset only has to be applied to xr.

* metal : fold iq mul_mv row split into offset0

Compute row0 and row1 before initializing the source pointers and apply
the row slice directly to offset0.

This keeps x and its derived pointers on the existing path while applying
the split row offset once.
2026-09-11 12:30:20 +03:00
16378d93f9 CUDA/HIP: Flash Attention tuning (gfx1201) (#28102)
* HIP: enable mma FA for head size 256 on RDNA4, tune configs

Assisted-by: Claude
Assisted-by: Codex

* HIP: prefer whole-tile FA grids over stream-k on AMD WMMA

Assisted-by: Claude
Assisted-by: Codex

* revise stream_k logic

* revise kernel selection logic

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-09-11 09:58:20 +02:00
Jeff BolzandGitHub 481c65f091 vulkan: fix data race and OOB access in argsort(large) (#28705)
argsort had a data race in the inner loop, which VVL caught. But I don't think
this was causing failures in practice.

argsort_large has OOB accesses which might explain the failures in CI, but I
couldn't reproduce it locally and I don't think it's a convincing explanation
of the failures.
2026-09-11 08:44:13 +03:00
shaofeiqiandGitHub df03399b88 opencl: add A8 Q4_0 mm binary kernel support (#28268) 2026-09-10 11:25:40 -07:00
Jeff BolzandGitHub 28ff095829 vulkan: use CPU writes in ggml_backend_vk_cpy_tensor_async if the context is idle (#28618) 2026-09-10 20:22:46 +03:00
Jeff BolzandGitHub 50182a53fa vulkan: use add_alloc_dep to enable topk_moe fusion for prefill (#28422) 2026-09-10 20:21:29 +03:00
Jeff BolzandGitHub 6788edb4f3 vulkan: small M matrix optimizations for qwen (#28457)
* vulkan: optimize m=1 mul_mat by swapping A/B

* vulkan: Improve small M perf

Allow split_k with small M.

Make small vs med tile selection (for coopmat2) depend on M, not just N.
2026-09-10 20:20:18 +03:00
Julian PscheidandGitHub 3ff67eb43d vulkan: fall back to shared-memory reduction for dmmv on PowerVR (#28341)
The Imagination proprietary Vulkan compiler returns VK_ERROR_UNKNOWN from
vkCreateComputePipelines for every dequant mul_mat_vec shader built with the
subgroup-only reduction that requires a subgroup size >= 16. That covers the
k-quants, the i-quants, TQ2_0, MXFP4 and NVFP4. ggml rethrows, so the first
generated token of any such model kills the process.

Reproduced on a Pixel 11 Pro (PowerVR C-Series CXTP-48-1536 MC1, driver
1.662.3024, subgroup size 128, min 32, max 128). The failure is independent of
subgroup size: 32, 64 and 128 all fail, as does dropping the full-subgroups
flag and the required-subgroup-size pNext. The legacy quants, which use the
plain subgroup reduction, compile and run fine.

The shared-memory reduction variant compiles and matches the CPU reference for
q2_K, q3_K, q4_K, q5_K and q6_K. The hybrid variant also compiles but costs
27% of token throughput (3.78 vs 5.20 t/s on Qwen3.5-2B-Q4_K_M).
2026-09-10 13:42:23 +03:00
72797e8919 vulkan : add command-buffer debug labels for GPU profilers (#28101)
* vulkan : add command-buffer debug labels for GPU profilers

Co-authored-by: gabby-zy <z2262718160@gmail.com>
Assisted-by: Claude Code

* vulkan : close the queue debug label with the label struct

---------

Co-authored-by: gabby-zy <z2262718160@gmail.com>
2026-09-10 10:06:07 +03:00
Aaron TeoandGitHub 4ea6d1bb6d ggml-cpu(s390x): add repack support for q4_0 (#28667)
ggml-cpu: clean comments

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2026-09-10 09:51:17 +03:00
Aaron TeoandGitHub f1b6fbf35c ggml-cpu(s390x): add Q1_0 vector intrinsic support (#28606)
* ggml-cpu: add `ggml_vec_dot_q1_0_q8_0` support

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: clean up variable naming for understanding

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* docs: update support for Q1_0

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2026-09-10 09:50:28 +03:00
Ruben OrtlamandGitHub 91f6a6cf36 vulkan: use spec constant for matrix matrix multiplication A-type (#25773)
* vulkan: use spec constant for mul mat type_a

vulkan: use map for mul_mm shapes

cleanup

fix indentation

fix cm2 and shmem init

fix cm2 spec constants

fix cm2 bindings

consolidate shmem tables and reduce size by type spec constant

fix compiler warning

fix missing Q2_0 type

fix unused warning when integer dot glslc support is missing

use minimal shmem size 8 instead of 1 to workaround cm2 compiler bug

fix missing Q2_0 type in cm2 matmul

fix types

* remove LUT quants from unified shader

* clean up

* restore coopmat2 q4_k/q5_k optimization

* split out q4_k/q5_k cm2 shader to fix Ampere regression

* revert iq shmem table renames

* simplify cm2 code with single uint8_t buffer

* fix fp4 extension use switch being overwritten by generic shader

* clean up

* adapt TQ1_0 changes

* adapt #27471 f16 Intel tuning changes
2026-09-09 18:53:39 +02:00
6d9c82ea2b hexagon: rope updates (#28628)
* hexagon: vectorize RoPE theta cache on v75

* hexagon: vectorize MROPE/IMROPE theta pick

* hexagon: tighten NEOX RoPE rotate and aligned tail copy

* hex-rope: use inplace rope for all scenarios

* hex-rope: remove ctx->spad usage and legacy timers

* hex-rope: add kernel params and enforce vtcm reqs at the host

* hex-rope: cleanup unused params and tighten the mode checks

* hex-rope: add  missing ops header

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-09-09 08:40:24 -07:00
Masato NakasakaandGitHub 22397c31a0 vulkan: Convert FILL to distribute workgroups in 2D to avoid exceeding maxComputeWorkGroupCount (#28592)
* divide workload to 2D

This is to workaround FILL exceeding maxComputeWorkGroupCount for Intel GPUs on Qwen 3.8 flash next

* minor change

* Fixed comment
2026-09-09 16:54:15 +02:00
d4abd573f6 CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#28552)
Recreated from #24546

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>

* CUDA: pick MMQ tile size against ncols_opt set on the host side

Assisted-by: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_011SYPfRhKoUpU3gMsGxq6go

---------

Co-authored-by: ravel7524 <58877666+ravel7524@users.noreply.github.com>
Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
2026-09-09 13:25:54 +02:00
5a4d0fecae CUDA: replace GGML_FA_ALL_QUANTS with GGML_FA_QUANTS, more control over what is compiled (#28079)
* CUDA: add configurable FA quant combinations

Assisted-by: Codex

* remove all flags but , add runtime fallback with warning for uncompiled combination

* Update docs/build.md

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>

* apply code review comments

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-09-09 12:50:08 +02:00
WakeUpMortyandGitHub df750f76bb vulkan: add dedicated iq4_xs mat-vec shader (#28426)
* vulkan: add dedicated iq4_xs mat-vec shader

Dedicated mul_mat_vec_iq4_xs for the dmmv path, replacing the generic fallback. ~+6-17% token generation on RDNA4 depending on model.

Assisted-by: Pi agent with Qwen3.8 27B

* vulkan iq4_xs: remove dead n_it unroll branch

Remove the n_it <= 8 experimental branch that attempted to fully unroll
the block loop. Since n_it is a runtime value, [[unroll]] is ignored by
the compiler, making both branches equivalent. Kept the simple loop
matching mul_mat_vec_iq3_s.comp.
2026-09-09 09:26:36 +03:00
Jiang, FishandGitHub 1945e09203 vulkan: add f16 B-type matmul pipelines and warp tile size tuning for Intel coopmat1 (#27471)
* vulkan: add f16 B-type matmul pipelines and warp tile size tuning for Intel coopmat1

* simplify mmp selection in mul_mat_id per review comment

* vulkan: enable f16 B-type pipeline for dense matmul on all vendors (MoE stays Intel-only)

* vulkan: add missing ocp_fp4 branches, restrict required_subgroup_size to Intel, fix mmp selection

* refine mmp selection in mul_mat_id

* add f16B pipeline selection just like q8_1

* update f16B pipeline selection in dense function
2026-09-09 09:03:32 +03:00
304665fe7a Add IQ type handling for MoE (#28476)
Co-authored-by: cwriter <cwriter@localhost>
2026-09-09 09:25:41 +08:00
Sarah WuandGitHub 9113cc1880 ggml : fix msvc+clang ggml_vld1q_u32 (#28284) 2026-09-08 17:40:26 +03:00
uvosandGitHub d4389a4dd9 Revert "ggml-cuda : restore prop.integrated on HIP builds (#24233)" (#28604)
This reverts commit c7d8722922.
2026-09-08 16:19:53 +02:00
Foad Abo DahoodandGitHub 88ada91c18 metal : fix idle threads in mul_mv_iq3_xxs for ne00 < 1024 (#28086)
* metal : fix half-idle simdgroup in kernel_mul_mv_iq3_xxs_f32 for ne00 < 1024

* metal : keep N_R0_IQ3_XXS = 4, dispatch a separate 8-row split kernel for ne00/32 < 32

The plain kernel is unchanged from master (4 rows per simdgroup, one thread per
chunk). The row-split mapping now lives in a separate kernel_mul_mv_iq3_xxs_f32_split
instantiation with N_R0_IQ3_XXS_SPLIT = 8, and the host selects it only when
ne00/32 < 32 and divides 32, so wide matrices keep the master kernel bit for bit.

* metal : select the iq3_xxs row split with a function constant instead of a separate kernel
2026-09-08 15:54:42 +03:00
Pepper GrayandGitHub ca86fb222e llama : add missing headers (#28566)
* fix compile-error: add missing header

Bug: #28557
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28559
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28560
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28561
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28562
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

* fix compile-error: add missing header

Bug: #28564
Signed-off-by: Pepper Gray <hello@peppergray.xyz>

---------

Signed-off-by: Pepper Gray <hello@peppergray.xyz>
2026-09-08 12:59:53 +02:00
Ankit KhandelwalandGitHub 64e9bceb2c vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL (#27220)
* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL

* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL

- implement fusion in unary.comp behind UNARY_MUL_FUSION ifdef,
  specialized pipelines per op instead of runtime branching
- fuse adjacent nodes only, ordering handled by graph_optimize
- drop runtime consumer scan and pending_unary_mul deferral

* vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL

1. GELU: gelu_mul_f32/f16 pipelines registered, CREATE_UNARY_MUL(gelu), GELU in dispatch + fuse gate + perf fusion name
2. Renamed/moved: gate is now ggml_vk_can_fuse_unary_mul(cgraph, unary_idx, mul_idx), placed with the other can-fuse helpers
3. norepeat both variants: each op gets plain (spec {0}) + _norepeat (spec {1}) pipelines from the same SPIR-V, selected via ggml_are_same_shape(src0, src1); the shape gate now allows broadcast (other dims equal-or-1)
4. graph_optimize: lambda deleted; standard "// UNARY + MUL: pull the consuming MUL forward" block added alongside the SSM_CONV/ROPE/MUL_MAT reorderings, with the same "other src must be weights or already processed" readiness check

* vulkan : align unary_mul fusion with binary kernel layout, relax gelu test tolerance

- schedule the fused kernel like mul.comp (256 threads x 2 unrolled
  iterations), recovering a 10-18% prompt-processing regression
- allow 5e-7 f32 error for gelu_mul: the shader evaluates gelu with an
  exp-based tanh identity while the CPU reference uses tanhf (~1 ulp)

* vulkan : use ggml_can_repeat in UNARY+MUL fusion shape check

The fused kernel indexes src1 via per-dim fastmod (generic_binary_head.glsl),
which is exact whenever the other operand tiles into the unary result -- not
just when its dims are equal or 1. Replace the hand-rolled loop with
ggml_can_repeat(other, unary) so the check matches the kernel's actual
capability and reuses the standard helper. Argument order matters: reversed,
it would wrongly admit graphs where the unary result is mul->src[1] and the
other operand is larger, producing truncated output.

Also add a rep_ne0 layout to the fused unary+mul backend tests covering a
non-1 repeat factor along dim 0.

* vulkan : fuse UNARY+MUL pairs separated by zero-compute nodes

gemma4's per-layer embedding gating builds gelu -> view_2d_slice -> mul,
where the intervening view is a zero-compute node aliasing an input that
was computed much earlier. Strict adjacency requirements meant neither
CUDA nor the vulkan unary+mul fusion handled this pattern.

Extend ggml_vk_graph_optimize to detect a UNARY whose consuming MUL is
separated only by unscheduled zero-compute nodes (GGML_OP_NONE, VIEW,
RESHAPE, TRANSPOSE, PERMUTE) and schedule those nodes ahead of the pair,
making it adjacent so the existing fusion applies. The reorder is guarded
by ggml_vk_can_fuse_unary_mul, a source-availability check for every
interleaved node, and the protected fusion patterns (topk_moe*, snake);
if fusion is later rejected the reordered graph still executes correctly,
just unfused.

Add a view_mid layout to the fused unary+mul backend tests replicating
the gemma4 pattern.

* vulkan : support OP-on-B in UNARY+MUL fusion

Some models apply the unary activation to the smaller MUL operand, e.g.
qwen3next/qwen35moe shared-expert gating builds ffn_shexp * sigmoid(gate)
with a [1,n_tokens] gate tensor. This shape was correctly rejected before:
the fused kernel derives its iteration extent from the unary tensor and
would leave most of the destination unwritten, and the generic same-shape
requirement in ggml_can_fuse blocked the pair outright.

Add UNARY_MUL_B_FUSION shader variants computing dst = src0 * OP(src1):
the OP operand rides the existing per-dim fastmod indexing, while the
iteration extent now comes from mul. Route {UNARY, MUL} pairs through a
local can-fuse variant that drops the generic same-shape rule and instead
requires the unary result to tile into mul->src[0] (ggml_can_repeat);
pairs with the unary as src0 keep the previous direction check, and
equal-shape pairs keep using the original pipelines.

Add a "gate" layout to the fused unary+mul backend tests covering the
shared-expert gate shape for gelu/sigmoid/silu/softplus in f32 and f16.

* vulkan : fold unary+mul view-hoisting into graph_optimize dep checks

Replace the dedicated UNARY + EMPTY* + MUL scanning block with two small
extensions to the existing scheduling logic:

- a consuming MUL may now join its in-set UNARY across a gap of unused
  zero-compute nodes (NONE/VIEW/RESHAPE/TRANSPOSE/PERMUTE), instead of
  requiring strict adjacency
- while doing so, such zero-compute blockers are ignored for this pair

Fusion validity is still decided later by ggml_vk_can_fuse at dispatch
time, so a rejected pair simply executes adjacent-but-unfused. Note the
relaxation must stay scoped to this pattern: exempting zero-compute
blockers globally reproduces silent output corruption on gemma3n.

* vulkan : select unary_mul OP-on-B via specialization constant

Replace the UNARY_MUL_B_FUSION compile-time shader variants with an
op_on_b specialization constant on the existing unary_mul SPIR-V,
mirroring how the norepeat flag is handled. The four {op}_mul_b_{f32,f16}
shader artifacts are gone - the OP-on-B pipelines reuse the base SPIR-V
with two-entry {norepeat, op_on_b} spec lists - and the duplicated store
expression is collapsed into a single runtime branch that the driver
prunes per specialization.

The constant is declared only under UNARY_MUL_FUSION so every other
binary pipeline keeps its single-entry specialization list.

* vulkan : replace unary_mul pipeline switches with a lookup table

Collapse the four nested selection switches in ggml_vk_unary_mul into a
single indexed lookup against a pipeline_unary_mul[4][2][2][2] table
([unary op][f16][norepeat][op_on_b]), whose trailing dims mirror the
{norepeat, op_on_b} spec constant list. The op axis uses a small shared
index helper that also replaces the switch in ggml_vk_can_fuse_unary_mul,
making it the only place that maps ops to the table.

Pipeline names are unchanged. Adding another supported op now requires
one macro invocation line and one helper case instead of edits in four
separate switches.

* vulkan : use ggml_can_fuse_subgraph for unary_mul pairs

Replace the hand-rolled pair validation in ggml_vk_can_fuse_unary_mul_pair
(bounds, op match, compute flags, single-use elision) with the shared
ggml_can_fuse_subgraph helper; backend-specific shape/type rules remain in
ggml_vk_can_fuse_unary_mul. Unlike ggml_can_fuse, the subgraph helper has
no same-shape requirement, so it covers both operand slots including
OP-on-B gates, and additionally rejects intermediates flagged as graph
outputs and validates view-source confinement.

The outputs parameter takes absolute node indices into the cgraph.

* Fix Whitespace

* vulkan : drop redundant unary_mul gap check in graph_optimize

The zero-compute nodes separating a UNARY from its consuming MUL are
already scheduled ahead of the pair by pass 2 of an earlier
optimization window, so the scoped gap tolerance added for this pattern
is unreachable in practice - disabling it leaves gemma-3n dispatch
counts unchanged (841 GELU_MUL per pass). Remove the flag, the empty
blocker exemption, and the now-unused gap helper, restoring the strict
adjacency requirement of the UNARY -> MUL pull-forward.

Keep the relaxation scoped out entirely: generalizing "zero-compute
nodes never block" beyond this pattern previously reproduced silent
output corruption on gemma3n.

* vulkan: fix whitespace (tab in indent)

* vulkan: fix whitespace (extra blank line)

* vulkan : move op_on_b spec constant to unary.comp

op_on_b is only used by the fused unary*mul path. Keep
generic_binary_head.glsl generic by defining it in unary.comp
instead. Same constant_id=1 and guard, no functional change.

* vulkan : make RMS_NORM/UNARY fusion gap-tolerant for views

Strict j==c+1 blocked RMS_NORM->MUL and UNARY->MUL when a
VIEW sits between (e.g. rms_norm -> view -> mul). Allow
c==back() with an empty-or-scheduled gap, matching the
review suggestion to check src linkage instead of adjacency.
Scoped to the two blessed pairs; safe because gaps can only
contain zero-compute nodes.

* vulkan : trim comments in UNARY+MUL fusion

Assisted-by: Muse Spark
2026-09-08 09:35:02 +02:00
miyanandGitHub f014bfef8b Fix Vulkan-Hpp handle usage on 32-bit targets. (#22892)
On 32-bit platforms, Vulkan non-dispatchable handles such as VkBuffer are
represented as uint64_t, and Vulkan-Hpp disables implicit conversions for
type safety. This exposes two issues in ggml-vulkan:

1. vk::Buffer is streamed directly into std::ostream in debug/memory logs.
2. vk::Buffer is cast to VkBuffer before being passed to Vulkan-Hpp
   CommandBuffer::copyBuffer APIs.

Fix these by add the operator<< for vk::Buffer, and
by passing vk::Buffer directly to Vulkan-Hpp copyBuffer calls.
2026-09-08 09:34:12 +02:00
lhezandGitHub 7d701b5929 opencl: properly handle non-contiguous inputs to conv2d (#28503)
* opencl: fix conv2d non-contiguous strides

* opencl: format
2026-09-08 09:26:34 +03:00
Georgi GerganovandGitHub 5a6caa05fc ggml : update ggml_prec specification (#26675)
* ggml : update ggml_prec specification

[no ci]

* cont : add GGML_PREC_BF16

* cont : rework API

* cont : use new API

* cont : swap arg order

* cont : support for MUL_MAT_ID

* cont : fix accidental remove of "break;"

* cont : return bools, add doc TAG_GGML_PREC, clean-up

* cont : add search tag

* cont : ws
2026-09-08 09:06:24 +03:00
Todor BoinovskiandGitHub 050dde50c9 hexagon: add RELU and LEAKY_RELU ops (#28585)
* hexagon: add RELU op

* hexagon: add LEAKY_RELU op too
2026-09-07 17:04:25 -07:00
Piotr Wilkin (ilintar)andGitHub e71b80510c Revert "CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)" (#28551)
This reverts commit 0c963452ea.

Assisted-by: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Q7rfnkjzgnfhvJsdeXhdoH
2026-09-07 16:28:19 +02:00
PascalandGitHub c0b1871bc7 webgpu: format the GET_ROWS case block (#28542)
Brace on its own line and body indented one level, matching the
surrounding cases, so the webgpu clang-format check passes.
2026-09-07 15:55:14 +02:00
TitaniumtownandGitHub dbeb37548e sycl: add a batched L2_NORM kernel (#28222)
* sycl: add a batched L2_NORM kernel

* sycl: batch consecutive L2_NORM siblings in the graph dispatch

Measured on Intel Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, f16 KV,
npp=128 ntg=128 npl=2, GGML_SYCL profiler:

    L2_NORM dispatches       12480 -> 6240
    L2_NORM device time      68.77 -> 39.14 ms   (-43%)
    total device time        6782 -> 6748 ms     (-0.5%)
    wall decode t/s          flat

* tests: add L2_NORM_BATCH coverage
2026-09-07 15:24:14 +02:00
7a333e7240 vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST) (#26578)
* vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST)

CUDA has these ops from the DeepSeek-V4 merge and Metal gained them in
PR 26459. Vulkan was the last major backend running the unfused primitive
chain. On DeepSeek-V4-Flash the unfused Sinkhorn comb chain alone takes
about 32% of decode op time on gfx1151 (Strix Halo), spread over roughly
16k dispatches per token.

dsv4_hc_comb runs the full 20-iteration Sinkhorn in registers. A token's
4x4 comb matrix lives in 16 consecutive subgroup lanes, with idst in bits
0-1 and isrc in bits 2-3 to match the CPU reference layout, so
subgroupShuffleXor by 1|2 reduces rows and by 4|8 reduces columns. One
dispatch replaces about 137 strictly ordered node executions per site.
The shuffle masks never cross a 16-lane boundary, so a subgroup of size
64 packs 4 independent tokens.

dsv4_hc_pre and dsv4_hc_post handle the elementwise stream collapse and
fan-out, with per-token coefficients staged in shared memory.

GGML_VK_DISABLE_DSV4_HC disables all three ops. The _COMB, _PRE and
_POST variants gate each op independently so a single kernel can be
bisected against the unfused graph.

Adds eval cases at the production n_iter=20 across batch sizes that
cross subgroup and workgroup boundaries.

* vulkan: dsv4 hc review fixes

Drop the per-op env-var disables and device flags, the stride divisibility
check (ggml guarantees it) and the workgroup-count fallback in supports_op.
Trim the comb shader comments to the lane layout.

---------

Co-authored-by: Kevin Hopper <no-reply@maestro.press>
2026-09-07 15:24:03 +02:00
0c963452ea CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)
* adjust ncols_picker for routed MoE in mul_mat_q_case function

* Adding CDNA, RDNA2 and RDNA4

* fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support

* feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support

* refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files

* HIP: mmq: enable typical moe ncols on RDNA4

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
2026-09-07 15:22:42 +02:00
AuroraRASandGitHub 4735997382 ggml: add gfx90c HIP support (#26454)
* ggml: add gfx90c HIP support

* ggml: make gfx90c HIP support compliant with specifications
2026-09-07 15:21:42 +02:00
73ab7599b5 CUDA: branchless Q4_K/Q5_K unpack to speed up mmvq, L2 prefetch on DGX Spark (#26705)
* Update Q4_K and Q5_K to use branchless computation, which stops the scale unpack being re-executed for every column in mmvq, improving perf at batch sizes > 1

* Gating the change off from DGX Spark due to no gain

* Adding prefetch gated to Spark, making branchless change in Q4_K and Q5_K general and modifying switch points based on latest perf data

* Guard the mmvq L2 prefetch against MUSA as well as HIP

* Define the mmvq L2 prefetch only under the Spark guard

* Update switch point for Q4_K to accommodate more models

* Remove stale comments

* Add block_size to ggml_cuda_type_traits and create a separate mmvq_should_prefetch function

* Rename block_size to bs for cleaner indentation

* Fix build error on non-Spark CUDA arch with appropriate conditional around new function added

---------

Co-authored-by: praneshgo <227579474+praneshgo@users.noreply.github.com>
2026-09-07 19:36:58 +08:00
0cae43063c vulkan: support type-aligned GET_ROWS (#28253)
* vulkan: fall back to CPU for GET_ROWS with misaligned offsets

The Vulkan GET_ROWS shader asserts when a tensor's backing-buffer offset
plus view_offs is misaligned w.r.t. minStorageBufferOffsetAlignment
(see init_pushconst_tensor_offsets). Previously this caused a hard crash
on models using ggml_view + ggml_get_rows (e.g. Qwen3-TTS, Qwen3-VL).

Return false from supports_op() in the misaligned case so the scheduler
falls back to CPU, matching the existing pattern for PAD_REFLECT_1D and
other unsupported op/shape combinations.

Repro: llama-tts -m Qwen3-TTS-*.gguf -mm mmproj-*.gguf -ngl 99
Crash: GGML_ASSERT(dst->op != GGML_OP_GET_ROWS || (a_offset == 0 && ...)) failed

* vulkan: trim comment for GET_ROWS misalign fallback

* vulkan: fix file corruption in gated_linear_attn struct

* vulkan: properly handle misaligned offsets in GET_ROWS quantized path

- get_rows_quant.comp was missing get_aoffset()/get_boffset()/get_doffset()
  calls that are already present in get_rows.comp, causing GGML_ASSERT crashes
  when GET_ROWS operates on views with non-zero view_offs, as produced by
  KV cache slices in Qwen3-TTS and Qwen3-VL.
- Remove the defensive misalignment GGML_ASSERT in init_pushconst_tensor_offsets
  for the binary push-constants specialization, since both get_rows.comp and
  get_rows_quant.comp now correctly apply per-tensor base offsets.
- Remove the workaround CPU fallback in supports_op() for GET_ROWS, since the
  Vulkan backend now handles misaligned offsets natively (no more bailout).
- Add backend test coverage with view_src0=true (ggml_view_4d into a padded
  tensor) for F32, F16, Q4_0, Q4_K, Q8_0, and I32 types, exercising both the
  non-quantized (get_rows.comp) and quantized (get_rows_quant.comp) paths
  with non-zero view_offs that reproduce the original Qwen3-TTS crash.

* tests: trim redundant comments in test_get_rows vs0 region

* tests: trim redundant comments in test_get_rows vs0 region (follow-up)

* vulkan: bind tensor base for binary ops, pass full view_offs via push constants

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, MUL, etc.),
bind the view_src base and pass the full view_offs divided by type_size via
push constant misalign_offsets. This avoids truncation when misalign_bytes is
not a multiple of quantized block size.

ggml_vk_tensor_subbuffer gains a use_view_offs parameter. When false, the
binding points to vk_tensor_offset (base) and size includes view_offs.
init_pushconst_tensor_offsets<binary> computes a/b/d_offset directly from
tensor->view_offs, which is always row-aligned and therefore exact.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).

* vulkan: bind aligned offset for binary ops, pass adjusted misalign via push constants

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.

ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types).

ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (target - adjusted_misalign) instead of the view_src base, keeping the
offset small enough for 16-bit/8-bit push constant fields.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).

* vulkan: bind aligned offset for binary ops, fix UMA offset mismatch

For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.

Added ggml_vk_tensor_physical_offset to unify physical offset lookup across
UMA and non-UMA devices. On UMA, resolves via ggml_vk_host_get(tensor->data);
otherwise uses vk_tensor_offset(t) + t->view_offs. Both get_misalign_bytes and
the new ggml_vk_get_adjusted_misalign helper build on top of this function,
so buffer bindings and push constant offsets are always consistent regardless
of device memory model.

ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types) while
remaining small enough for 16-bit/8-bit push constant fields
(adjusted_misalign < lcm(align, type_size)).

ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (physical_offset - adjusted_misalign) on both UMA and discrete GPUs,
fixing a bug where the UMA host_get path previously skipped the adjusted
misalign binding and returned the target offset directly.

Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).

All 223 GET_ROWS tests pass on Vulkan (NVIDIA GeForce RTX 5060 Ti).

* finish misalignment fix

* supports_op changes for openvino/webgpu

---------

Co-authored-by: AiChiTuDouPian <15327701848@qq.com>
2026-09-07 12:22:10 +02:00
Siavash NorouziandGitHub b74f590eaf ggml-cuda: fix divergent barrier in f16 flash attention (#27870)
* ggml-cuda: fix divergent barrier in f16 flash attention

* ggml-cuda: avoid duplicate metadata pointer setup
2026-09-07 09:23:21 +03:00
Aman GuptaandGitHub 992cb503cd ggml: allow backend inputs to not create another split (#28387) 2026-09-07 09:10:40 +03:00
Jeff BolzandGitHub 9ac8c408a3 vulkan: rms_norm fusion opportunities (#28024)
Support RMS_NORM + MUL + ADD (+ MUL) and RMS_NORM + VIEW + SET_ROWS.
Extend ROPE + VIEW + SET_ROWS to support IMROPE.

Worth around 4% in gemma4 on my system.
2026-09-07 09:08:28 +03:00
AnjielonandGitHub 8fe90e1fbf vulkan: add TQ1_0 support (mm, mat-vec, mat-vec-id, dequant, get_rows) (#27765)
* vulkan: add TQ1_0 support (mm, mat-vec, dequant, get_rows)

* vulkan: pack TQ1_0 powers of 3 into a 32-bit constant

Replaces the constant array with a packed 32-bit value (7 bits per entry,
max 81 < 128) extracted with shift/mask, as suggested in review — avoids a
constant array that may not be kept in registers.

test-backend-ops on gfx1151: tq1_0 MUL_MAT 11/11, MUL_MAT_ID 6/6,
GET_ROWS 4/4, unchanged.

* vulkan: address review - shared TQ1_0 decode helpers, fix standalone dequant shader

Review feedback from jeffbolznv, all points:

- Move the packed-pow3 decode into shared helpers in types.glsl
  (tq1_0_byte_of / tq1_0_digit_of / tq1_0_trit) and use them from
  dequant_funcs.glsl, mul_mm_funcs.glsl, dequant_funcs_cm2.glsl and
  dequant_tq1_0.comp instead of repeating the logic. The cm2 path also
  drops its constant array for the packed-constant extraction.
- Translate all remaining comments to English.
- dequant_tq1_0.comp: use dequant_head.glsl. The shader previously declared
  its own single-field push constant while the pipeline is created with the
  5-field layout, so p.ne read the wrong field - confirmed broken, as
  suspected in review.
- Fix wg_denoms for the standalone dequant pipeline: one invocation decodes
  4 elements with local_size 256, so a workgroup covers 256*4 elements, not
  256*16. With the old value the dispatcher launched a quarter of the
  required workgroups.

Verified by temporarily forcing the dequant + f16 matmul path for TQ1_0
(hack not committed): test-backend-ops MUL_MAT passes through the rewritten
standalone shader, and the standard MUL_MAT / MUL_MAT_ID / GET_ROWS
tq1_0 cases still pass on Vulkan (AMD gfx1151).

* vulkan: address review — English comments, shared tq1_0_trit, trim TQ1_0 test cases

- mul_mat_vec_tq1_0.comp: drop leftover non-English comment and the local
  POW3_PACKED constant; all decode sites now call tq1_0_trit() from types.glsl
- types.glsl / dequant_funcs_cm2.glsl: ASCII-only, drop stale reviewer note
- test-backend-ops: remove the oversized MUL_MAT_ID case (432 MiB A tensor,
  ~172 GFLOP reference); move the two remaining ones next to the other
  backend-specific mul_mat_id one-offs and document why they are needed

* metal: decline TQ1_0 for GET_ROWS and mat-mul in supports_op

The new TQ1_0 cases in test-backend-ops exposed that the Metal backend
claimed support for GET_ROWS/MUL_MAT/MUL_MAT_ID with TQ1_0 sources while
having no such kernels (ggml_metal_library_compile_pipeline aborted on the
missing kernel_get_rows_tq1_0). Decline the type so the ops fall back to
the CPU, matching the existing NVFP4 handling on the same lines.

Assisted-by: Claude Fable 5

* vulkan: trim the TQ1_0 comments

Addresses @0cc4m's review: keep only what the code does not already say.

Removed the block-format recaps (the layout is right there in the struct) and
the step-by-step decode walkthrough. Kept the two facts a reader cannot infer:
the 8-bit truncation is part of the format, not an optimisation, and the powers
of 3 are packed into one uint so they do not end up in a constant array that
may miss the registers.

No functional change.

* vulkan: address review — trim comments, fold Metal check, drop unused _v

Per @0cc4m's review:

- dequant_funcs.glsl, dequant_funcs_cm2.glsl: drop the "see types.glsl"
  pointers — they apply to every quant and say nothing specific.
- dequant_tq1_0.comp: drop the wg_denoms note. It is a precondition, not
  information.
- mul_mm_funcs.glsl: same pointer removed.
- types.glsl: the comment on tq1_0_trit is down to the one fact the code
  cannot show — the 8-bit truncation is part of the format, matching the C
  reference, not an optimisation.
- dequant_funcs_cm2.glsl: removed dequantFuncTQ1_0_v and its define. You were
  right that it is optional: it wrapped four scalar decodes and vectorised
  nothing, and mul_mm_cm2.comp already guards the path with
  `#if defined(dequantFuncA_v)` (DATA_A_F32 omits it the same way).
- ggml-metal-device.m: folded TQ1_0 into the existing NVFP4 check instead of a
  separate block, and dropped both comments.
- test-backend-ops.cpp: the two mul_mat_id cases stay — they cover the
  block-stride loop and the per-expert base offset that k == 256 alone never
  reaches — but the comment is now one line instead of five.

Kept: the one-line labels on the three block regions in mul_mat_vec_tq1_0.comp
and on tq1_0_byte_of(). Those state the 5-trits-per-byte packing, which the
loop bounds do not show. Happy to remove them too if you prefer.

Re-verified on AMD gfx1151 (Vulkan), test-backend-ops, 2/2 backends passed:
MUL_MAT 9 TQ1_0 cases, MUL_MAT_ID 5, GET_ROWS 4 — all OK, no failures.
The coopmat2 path is unchanged apart from the removed _v define.
2026-09-07 06:35:30 +02:00
lhezandGitHub d03efa5d53 opencl: properly choose weights pack for q4_K, q5_K mul_mat (#28402) 2026-09-06 08:33:08 -07:00
Aman GuptaandGitHub 73a43d1f69 cuda: fixes races in mmid and mmf (#28475) 2026-09-06 19:45:01 +08:00
Niklas WenzelandGitHub 971595d669 metal : add remaining fa-vec tunings for M2 Max (#28458) 2026-09-06 07:37:01 +02:00