Commit Graph
3044 Commits
Author SHA1 Message Date
Max Krasnyansky 960dffab05 hexagon: MUL_MAT and MUL_MAT_ID fusion and fixes (#28202)
* hex-mm: fuse QKV and FFN matmuls that land on HMX

* hex-mm: remove hardcoded ne[1] < 32K restriction

* hex-get-rows: explicitly reject repacked Q8_0 just in case somebody decided to add an override

* hex-mm: correct overhead sizing to make sure we dont exceed vtcm budget for large dims

* hex-mm: fuse MUL_MAT_ID into MUL_MAT_ID_NX (2x,3x,...) where possible

* hex-fusion: update opbatch and opqueue sizing to acount for new fusion and reduce overhead for trace buffer alloc

* hex-bufs: sort buffers while finalizing opbatch, helps avoid va space fragmentation

* hex-bufs: add simple va defrag to make sure we dont abort just because the va space is fragmented

* hex-mm: replaced more scalar divs with fastdiv and minor cleanup

* hex-mm: tighten up supported fusion checks to exactly match supported kernels
2026-09-02 09:15:21 +03:00
Laurent ZuijdwijkandMarshall ba8818cbf3 vulkan: handle larger batch sizes (>4) efficiently for IQ3_S mat-vec (#27449)
* vulkan: handle larger batch sizes (>4) efficiently for IQ3_S mat-vec when NUM_COLS > 4. 5x perf at n=8

Assisted-by: Claude Opus 5

* adds 2 cases per quant type at `k=16*256` to the `all_types` mat-vec sweep

---------

Co-authored-by: Marshall <assistant@llama.cpp>
2026-09-02 09:14:52 +03:00
Mads Marquart 56dd8150cc vulkan : only request VK_KHR_shader_bfloat16 extension if supported (#28155) 2026-09-02 09:13:25 +03:00
Alan Tseng 2637dfe373 ggml-cpu : conditionally add SpacemiT IME kernel sources (#27961)
When building with gcc < 15, CMakeLists.txt unconditionally adds
ime2_kernels.cpp, which fails to compile. FindSMTIME.cmake only defines
RISCV64_SPACEMIT_IME2 when the IME2 instructions are detected, and gcc 14
only has IME1, so ime2_kernels.cpp hits its #error.

This PR fixes it by using IN_LIST to add each kernel source according to
the spec that was actually detected.
2026-09-02 09:12:28 +03:00
Hongqiang Wang 43d87ff2dd opencl: fix out‐of‐bound reads in the Adreno image kernels (#27632)
* opencl: clamp the q4_K decode GEMV's fetch row on a padded x-grid

* opencl: enforce the tiling contract of the image KQ/KQV GEMMs

* opencl: decide the image KQ/KQV split at the dispatch, not from strides
2026-09-01 22:28:45 -07:00
Trivikram Reddy 69320fef12 hexagon: add missing FARF logs for cpy/get_rows/set_rows/gdn ops (#28217)
* hexagon: fix bug ne[2] printed in proc_op_req prep-src log

* hexagon: add shape/VTCM farf logs to cpy, get/set rows, gdn
2026-09-01 22:20:29 -07:00
Jhen-Jie Hong b96806d960 metal : add metallib build support for xcframework (#28163) 2026-09-02 07:45:56 +08:00
anujj 3466812d1f cuda: fuse MoE weighted expert reduction (#25952)
* cuda : fuse MoE weighted reduction (mul + view + add)

The MoE combine tail currently writes weighted expert outputs to
global memory before reducing them. That intermediate global-memory
traffic is the main cost. The production baseline generally runs two
physical fused kernels; this path runs one.

This change matches the full expert-weighting plus ordered-reduction
subgraph and replaces it with one weighted-reduction kernel.

Supported graphs:
- unscaled: experts * router_weights
- scaled:   (experts * expert_scale) * router_weights

k = 2..15 is handled by one runtime-k kernel.

Matching is structural: op sequence, shapes, strides, expert views,
and the left-to-right ADD chain. The fused kernel keeps that same
reduction order. Results are not claimed bit-identical; CUDA FP32
contraction can change rounding slightly.

Allocator integration uses add_alloc_dep from the graph-optimizer
API so experts, router weights, and optional expert scales stay live
until the fused destination is written. Memory ranges are rechecked
before the fused kernel runs.

Unrecognized or unsafe graphs are left alone and keep the existing
per-op path. Set GGML_CUDA_MOE_WEIGHTED_REDUCTION=0 to disable the
fusion.

test-backend-ops covers scaled/unscaled, aligned/unaligned, and
representative values across k=2..15, plus a k=16 case that must
stay on the per-op path.

* Pruned the test matrix from 15 to 6

* Addressed the aman and olivers review comments
2026-09-01 21:48:47 +02:00
Titaniumtown c845263f8b Revert "sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 12…" (#28184)
This reverts commit 1f3d318734.
2026-09-01 19:04:31 +03:00
Jingxin (Philip) Li 1f3d318734 sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 1280 (#28016) 2026-09-01 11:47:08 -04:00
Lukasz Stolcman 8887a48f05 metal : add fa-vec tuning for M2 Pro (#28122)
* metal: add fa-vec tuning for M2 Pro

* metal : update fa-vec tuning for M2 Pro with new dtypes
2026-09-01 21:24:44 +08:00
Jhen-Jie Hong be789c3448 metal : add fa-vec tunings for A18 Pro (MacBook Neo) (#28152) 2026-09-01 21:15:59 +08:00
Niklas WenzelandYiChen Lv fe2120bc9d metal : fix more leaks due to missing autoreleasepools (#27883)
* metal : fix more leaks due to missing autoreleasepools

* metal : rename variable

* metal : fix another missing pool warning

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>

---------

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
2026-09-01 13:50:47 +02:00
Georgi Gerganov d08c7872d6 metal : add fa-vec tuning for M2 Max (#28015)
Rows for M2 Max (30 GPU cores) collected with 'ggml-metal-tuning fa-vec
--dtype f16,q8_0', pasted into fa_vec_tuned_table.

ref: https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18205786

Assisted-by: pi:llama.cpp/Qwen3.8-27B
2026-09-01 13:37:40 +03:00
Neo Zhang 5eec3ad017 sycl : support limit max alloc memory within 2GB for host-pinned memory (#27559) 2026-09-01 13:35:47 +03:00
James Francis d5d993a093 metal: enable Metal 4.0 tensor API on M5+/A19+ (#27461)
* metal : request Metal 4.0 language version for the tensor API

* metal : load the tensor API kernels from a separate metallib

* tests : add external-metallib tensor API regression test

* metal : fix metallib build order for the tensor API kernels
2026-09-01 12:02:42 +03:00
Buğra Özgürsoy 458681e1d5 metal : add fa-vec tunings for M1 Ultra (#28088)
* metal : add fa-vec tunings for M1 Ultra

* metal : move M1 Ultra tunings after M1 Max section

* metal : remove duplicate blank line
2026-08-31 23:47:27 +02:00
ynankani e4b9af007b CUDA: XOR swizzle flash attn K,V smem fp16 tiles (#25635)
* CUDA: XOR swizzle flash attn  K,V smem fp16 tiles

Signed-off-by: ynankani <ynankani@nvidia.com>

* Fix use 64bit generic pointer instead of 32bit shared pointer

Signed-off-by: ynankani <ynankani@nvidia.com>

* fix shared memory race in FA on DGX Spark

* Handle corener case

Signed-off-by: ynankani <ynankani@nvidia.com>

* Add swizzle test cases and gate sync for swizzled path only

Signed-off-by: ynankani <ynankani@nvidia.com>

* gate CUDA PTX

Signed-off-by: ynankani <ynankani@nvidia.com>

* offset calculation specific for swizzle branch

Signed-off-by: ynankani <ynankani@nvidia.com>

* Reafctor code

Signed-off-by: ynankani <ynankani@nvidia.com>

* Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)

Signed-off-by: ynankani <ynankani@nvidia.com>

* rebase and update test case args

Signed-off-by: ynankani <ynankani@nvidia.com>

* Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0

Signed-off-by: ynankani <ynankani@nvidia.com>

---------

Signed-off-by: ynankani <ynankani@nvidia.com>
2026-08-31 22:18:01 +02:00
Georgi Gerganov ab0b3bd3c8 metal : add concat support for quantized types (#28116)
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
2026-08-31 23:16:04 +03:00
BartowskiandGeorgi Gerganov 85c55223ca AVX2: Speed up large batch size prompt processing of IQ models (#27402)
* Batched gemm for grid IQ quants

Style updates and a bit more performance

Clean up comments

Move code around

Vectorize IQ panel decode, lower threshold for speedup

IQ panel: single-source gather layout, gate bias, vectorize interleave

Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer

Move IQ panel code out of repack into iqp.cpp, clean up comments

Another comment sweep

* Add myself as iqp.* codeownder

* Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size

* Renaming and moving

* The other half of renaming and moving

* Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition

* Update ggml/src/ggml-cpu/iqp.h

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

* Add iqp_rows work buffer

* Revert "Add iqp_rows work buffer"

This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb.

* Add NUMA fallback

* Add 10 row batch tests for IQP coverage on all grid IQ types

* Swap assert for return false in support check

* Move IQP mul_mat_id test

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-08-31 14:33:50 -04:00
Georgi Gerganov 2a74817f93 metal : add top-k radix implementation (#28073)
Assisted-by: DeepSeek-v4-Flash-0731
2026-08-31 21:31:53 +03:00
Hongqiang Wang 010be9683a opencl: tune the quant paths for Intel Xe-LP GPUs to improve its TG and PP performance (#26438)
* opencl: Q4_K/Q5_K mul_mv N_DST 4->8 on Intel for 2x activation reuse

* opencl: Q4_K mul_mm 8x8 tile fot Intel

* opencl: Q5_K mul_mm 8x8 tile for Intel

* opencl: Q4_K mul_mv N_DST 8->16 for Intel
2026-08-31 08:56:22 -07:00
8e53fcefd2 webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation (#28045)
* webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation

* chore : improve code readability

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>

---------

Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
2026-08-31 16:04:38 +02:00
Jaden_Mach f8dbcd6189 ROCm: add radix TOP_K for long rows (#27466)
* ROCm: add radix TOP_K for long rows
2026-08-31 15:00:04 +02:00
Niklas Wenzel 5d4a3be26d metal : add fa-vec tunings for M1 (#28078) 2026-08-31 13:58:55 +02:00
ynankani 41ef91f7c8 CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were restricted to 1 token (#27621)
* CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were resticted to 1 token

Signed-off-by: ynankani <ynankani@nvidia.com>

* Address review comments

Signed-off-by: ynankani <ynankani@nvidia.com>

* Add SWIGLU_CLAMP case to multi-token moe fusion

Signed-off-by: ynankani <ynankani@nvidia.com>

---------

Signed-off-by: ynankani <ynankani@nvidia.com>
2026-08-31 19:22:28 +08:00
Neo Zhang a32af33de2 sycl : Enhance to get the free memory of Intel GPU (#27968)
* enhance get mem info by l0 an SYCL API

* remove debug code, format the code

* update SYCL.md for GGML_SYCL_GET_MEM_API
2026-08-31 13:33:02 +03:00
Simon Teixidor 2cdae802e4 vulkan: tune mat-vec rows for batched inference on Strix Halo (#27909)
* vulkan: RDNA3 static mat-vec rows above four columns

On RDNA3 above four columns a static 4 rows for all types benches faster than
the default.

* vulkan: RDNA3 static mat-vec-id rows

mul_mat_vec_id has no column dimension to switch on. On my Strix Halo machine,
a static 4 is faster here than the defaults across types and batch sizes.
2026-08-31 12:07:53 +03:00
fairydreamingandStanisław Szymczyk 557614e029 ggml : add MUL_MAT to the list of ops that may need additional memory (for WebGPU) (#28071)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-08-31 10:17:23 +02:00
Ruben Ortlam daef7b6874 vulkan: top_k radix select for k >= 1024 for Qwen 3.8 Flash Next (#28032)
* vulkan: add top-k radix sort shader for k >= 1024

* add Qwen 3.8 Flash Next top-k tests

* add top-k qsa fusion

* clean up code
2026-08-31 07:04:34 +02:00
Shenghan Yang 9723942adc hexagon: fix CPY fence bug (#28033) 2026-08-30 11:18:24 -07:00
codemonkey bd55e6aae8 metal : add remaining Q4_1/Q5_0/Q5_1 fa-vec tunings for M2 (#28017) 2026-08-30 20:00:10 +02:00
hmirinandGeorgi Gerganov a7cc83bbae rpc: avoid serializing buffers from other servers (#26500)
* rpc: avoid serializing buffers from other servers

Only include remote buffer pointers when the buffer belongs to the RPC dispatcher receiving the graph. Add a two-server regression test for cross-server tensor serialization.

Assisted-by: Codex

* cont : add ref

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-08-30 20:26:16 +03:00
Georgi Gerganov 6d1479c148 ggml : fix ggml_backend_buft_get_alloc_size() guard (#28038) 2026-08-30 20:25:15 +03:00
Aman Gupta 0190529ec4 ggml: add SWIGLU_CLAMP (#27930)
* ggml: add SWIGLU_CLAMP

* add vulkan shader
2026-08-30 23:00:02 +08:00
Pascal f1793c1c4e CUDA: use the fast mm_ids_helper path for any n_expert_used (#27978)
The optimized path grouped warp lanes by token and required
warp_size % n_expert_used == 0, with a single hardcoded exception
padding 6 up to 8. Every other count fell back to the generic path,
which walks the tokens one at a time with a warp reduction per token,
for each of the n_expert blocks.

The lane group only has to divide the warp, and the loop body already
guards the padded lanes with iex < n_expert_used, so the padding
generalizes to the next power of two. The 6 -> 8 case and every count
already dispatched keep the exact same padding as before.

n_expert_used = 10 now reaches the fast path. Measured on
Qwen3.8-Flash-Next (512 experts, 10 used) at 55k context on an
RTX PRO 6000, warm runs with the first one discarded:

  prompt processing   2334 -> 2600 t/s

Token generation is unaffected, since a single token leaves nothing to
walk. Other expert counts reach the fast path by adding their case to
the dispatch.
2026-08-30 16:06:32 +02:00
itterative 0b5be7e4a2 hip: tune rdna 3 mmq config (#26284) 2026-08-30 13:47:21 +03:00
LunalFresh e422148047 hip : optimize Q2_0 dot-product path for gfx1201 (#26753)
* hip/gfx1201: optimize q2_0 vec_dot_q2_0_q8_1 with native amdgcn perm

* Broadened HIP's Q2_0 perm optimization

* Remove redundant HIP perm availability guard

* Optimize HIP Q2_0 MMQ unpack with native perm

* cuda: label HIP preprocessor guard

* cuda: label HIP preprocessor guard

* Restore MMQ tile index handling
2026-08-30 13:18:36 +03:00
Georgi Gerganov 73f56d105b ggml : add ggml_backend_op_alloc_size_may_expand, use it in RPC (#27960)
some backends (Metal, SYCL, WebGPU) require additional memory for
fleeting data for certain ops, which is reflected in their
get_alloc_size implementations.

add ggml_backend_op_alloc_size_may_expand() to the backend utils,
listing these ops, and assert in ggml_backend_buft_get_alloc_size
that a backend expanding the alloc size of a compute op only does so
for ops listed in the helper.

use the helper in the RPC backend to decide whether to query the
remote server for the actual alloc size, instead of a hardcoded list.

Assisted-by: pi:llama.cpp/Qwen3.8-27B
2026-08-30 09:17:47 +03:00
Ryan C 742347b2e7 rpc: fix apple rdma error spew on teardown (#27908) 2026-08-30 09:16:26 +03:00
Nils Gladitz 093adb242e metal: add fa-vec tunings for M3 Ultra (#27999) 2026-08-30 09:06:29 +03:00
Daya Adianto b8b743c3c1 metal : Add fa-vec tuning for M3 Pro (#27963)
Related issue: #27668
2026-08-30 09:02:22 +03:00
Ryan C 2bf0415152 rpc : fix pre-rdma macOS versions (#27815) 2026-08-30 08:59:25 +03:00
9e54e687cb hexagon: support for device discovery and create sessions on demand (#27785)
* hex-devices: add support for lazy session allocation and cleanup dev interfaces

Co-authored-by: Marco Colombo <mcolombo@qti.qualcomm.com>

* hex-devices: support for runtime discovery of available NPU cores

Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com>
Co-authored-by: Ehsan Bateni <ebateni@qti.qualcomm.com>

* hex-devices: reject non-existing devices early during init

---------

Co-authored-by: Marco Colombo <mcolombo@qti.qualcomm.com>
Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com>
Co-authored-by: Ehsan Bateni <ebateni@qti.qualcomm.com>
2026-08-30 08:57:55 +03:00
Titaniumtown 370cb12e8b sycl: split long rows in TOP_K instead of one work-group per row (#27847) 2026-08-30 08:57:08 +03:00
QuintinShaw d882575cc8 metal : fix null-pipeline crash for F16 src1 mul_mat/mul_mat_id (#25648)
* metal : fail closed on mul_mat shapes with missing F16 kernels

* metal : abort on nil pipeline in encoder_set_pipeline

* metal : address review comments

* metal : share mul_mat mm dispatch with supports_op
2026-08-30 08:56:35 +03:00
Aman Gupta 57291f2644 ggml: allow passing alloc dependencies in graph_optimize (#27301)
* ggml: allow passing alloc dependencies in graph_optimize

* add alloc dep tests

* add TODO about using flat array
2026-08-30 11:34:20 +08:00
codemonkey c589f0ed10 metal : add fa-vec tunings for M2 (#27940) 2026-08-30 01:44:53 +02:00
Hongqiang WangandLi He c841aeeb8b opencl: use a better matmul path on two Adreno GPU generations (#27640)
* opencl: default the Adreno xmem F16xF32 GEMM on for X2E

kernel_mul_mm_f16_f32_l4_lm is the slowest matmul this backend has on Adreno: on
the X2-90 it runs the gpt-oss-20b attention projections at roughly a quarter of
what the tuned dense q4_0 GEMM reaches on the same device. That matters for any
model whose non-expert weights stay f16 -- the stock gpt-oss-20b release is
exactly that, and its prefill spends 40.8% of GPU time in that one kernel. The
xmem route already existed but was left opt-in, so nobody hit it.

Worth about 25% prefill on gpt-oss-20b on an Adreno X2-90. Gated to X2E: the
Adreno 840 measures neutral. Decode is untouched -- the dispatch gate needs
N >= 16. It is worth nothing on the q8attn variant, whose attention weights
already take the dp4a dense GEMM.

The env var was presence-tested before, so =0 previously enabled it; it is now
atoi()'d. MUL_MAT 963 OK / 0 FAIL on both arms.

* opencl: bypass the tiled f32 GEMM on the Adreno A7X

The A7X (E031.41) compiler executes kernel_mul_mm_f32_f32_l4_lm at roughly a
tenth of what the same silicon reaches in its own f16 and q4_K kernels. It
allocates 488 B/WI of private memory against 304 for the same source on the
following generation, i.e. the older register allocator spills in the K-loop.
Models with per-layer F32 projection pairs kept F32 by quantization policy land
on this kernel twice per layer, and it dominates their prefill on that part.

Route batched f32xf32 (ne11 > 8) around the tiled path on the A7X and let it
fall through to the per-row f32 kernel, which that compiler handles fine; small
batches keep the tiled path. Weights stay GPU-resident, so decode placement is
untouched -- declining the op in supports_op instead was measured first and
rejected, because the per-layer CPU round-trips cost more decode than the
prefill it gained.

Worth about 9% prefill on gemma-3n-E4B on an Adreno 740, with MUL_MAT counts
identical on and off. No other generation is affected. Override with
GGML_OPENCL_A7X_F32_LM_BYPASS=0.

* opencl: enable xmem GEMM for adreno by default

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-08-29 10:46:27 -07:00
Georgi Gerganov 3173a56471 metal : assert shared memory padding (#27951)
* metal : assert shared memory padding

* cont : add ref
2026-08-29 17:55:15 +03:00