Commit Graph
2976 Commits
Author SHA1 Message Date
Georgi Gerganov 0df017d6dd metal : fix glu dispatch with ne00 = 1 (#28306)
* metal : fix glu dispatch with ne00 = 1

* tests : disable ill-defined tests
2026-09-03 13:25:41 +03:00
0ba6499c3b CUDA: Allow concurrent streams per split for multi-GPU (#28198)
* CUDA: Allow CUDA optimization per split for multi-GPU.

Previous guard caused multi-GPU to skip the graph optimization.  The
graph is already split per device and the optimization doesnt run
over the whole model but once per split, and thus should be allowed.
However, the CUDA event ggml_cuda_concurrent_event belongs to
whichever GPU was "current" when created. If the pass ran while
GPU 0 was current, it would stick and during event creation for the
second GPU it would land on GPU 0.

The fix: set the device explicitly ggml_cuda_set_device(cuda_ctx->device);
Default behaviour remains unchanged, only active for GGML_CUDA_GRAPH_OPT=1.
Explicit device setting pattern re-used from ggml_backend_cuda_graph_compute.

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Aman Gupta <amangupta052@gmail.com>

---------

Co-authored-by: tannerbruhn <tannerbruhn@users.noreply.github.com>
Co-authored-by: Aman Gupta <amangupta052@gmail.com>
2026-09-03 18:03:03 +08:00
Nathan Wilson c7bda030e7 vulkan: fix FA dequant path engagement (#28190)
Skip the nb[3] check when ne[3] == 1, the shader never reads it for a
single stream. Cache views carry the full-buffer stride there, so the old
check reduced to n_kv == kv_size and the path only engaged with the
cache full.
2026-09-03 10:40:34 +02:00
Neo Zhang 0df974d777 sycl : enhance the api to support peer-to-peer copy (#27550) 2026-09-03 10:41:07 +03:00
EurekaticandRaulAbejonDelgado 4aa6ffba25 sycl: reduce redundant work in Q4_K multi-column MMVQ (#27062)
* sycl: Q4_K Weight unpack optimization and reuse between destination Columns

* sycl: Q4_K small N (N=2..4) + two output rows by subgroup reuse of activation between two rows.

* sycl: gate Q4_K two-row reuse for small N=2

* sycl: Fix on magic number now uses Q4_K_MMVQ_ROW_PAIR_MIN_NROWS=6272 for it, added tests for coverage around Q4_K_MMVQ_ROW_PAIR_MIN_NROWS with perf support to test Q4_K MUL_MAT, applied the same  reuse pattern to the activation as the weights.

Assisted-by: GPT-5.6 Sol

---------

Co-authored-by: RaulAbejonDelgado <raul.abejon.delgado@gmail.com>
2026-09-03 14:59:06 +08:00
Xuan-Son Nguyen 159b741427 finetune: fix no KV cache (#27199)
* training: fix no KV cache

* apply @ ggerganov
 suggestion
2026-09-02 23:53:32 +02:00
cqderek f027c4f1b0 ggml-hexagon: add F16 support for unary ops (#28228)
Extend the HTP backend's F16 unary op coverage to include ABS on top
of the existing NORM/RMS_NORM/L2_NORM/SCALE/CLAMP/SQR/SQRT set.

- Add hvx_abs_f16_{aa,au,ua,uu} + dispatcher in hvx-arith.h, mirroring
  the sqr_f16 kernel structure and using the existing hvx_vec_abs_f16()
  sign-bit-clear helper
- Add abs_f16() row-wise dispatch and DEFINE_UNARY_TASK_F16(unary_abs, ...)
  in unary-ops.c, wired into execute_op_unary()'s op_type/task_func
  switches
- Register HTP_OP_UNARY_ABS in htp_op_is_unary() (unary-ops.h) so that
  ggml_hexagon_precompute_unary_params() fills kernel_params (n_threads,
  VTCM layout) for ABS nodes -- required for the F16 path to function
- Narrow the F16 GGML_OP_UNARY gate in ggml_hexagon_supported_unary()
  (ggml-hexagon.cpp) to allow GGML_UNARY_OP_ABS specifically, instead of
  rejecting all GGML_OP_UNARY ops for F16
- Merge the separate execute_op_unary_f32()/execute_op_unary_f16()
  functions into a single execute_op_unary(), branching on an is_f16
  flag for the parts that actually differ by type (elem_size, the
  early F16 op-support check, and which task_func table to use) while
  keeping the F32-only tiled/RMS_NORM_MUL paths intact -- per review
  feedback to avoid duplicating the shared VTCM/DMA plumbing

Verified on-device (QRD8850, Hexagon v81) via test-backend-ops -o ABS:
8/8 passing (F16 + F32, HTP0, no CPU fallback). Regression-checked
SQR/CLAMP/SQRT (F16+F32) and NORM/RMS_NORM/L2_NORM/SCALE (F32; their F16
paths have no CPU reference kernel in test-backend-ops and cannot be
correctness-tested there independent of this change).
2026-09-02 12:59:36 -07:00
Isaac 9cc33944f9 metal : add fa-vec tunings for M3 (#28236) 2026-09-02 20:13:12 +02:00
8c0b9cd04a metal : fix memory query under low-memory conditions (#27701)
* metal: Fix memory query under low-memory conditions

* Simply variable name

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

* Write it even shorter

Co-authored-by: Niklas Wenzel <dev@nikwen.de>

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Niklas Wenzel <dev@nikwen.de>
2026-09-02 20:09:46 +02:00
Adrien Gallouët 3d3d7c8181 ggml-cuda : remove unused vars (#28235)
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-09-02 18:54:11 +02:00
Aman Gupta 8e93a9773b CUDA + ggml: add sparse-fa for DSV4/GLM (#27970) 2026-09-02 17:27:37 +03:00
Aman Chadha(IVIXMMI)andAcmmi b81c99b479 ggml: avoid KleidiAI buffer type init on dispatch (#27891)
Co-authored-by: Acmmi <acmmi@Acmmis-MacBook-Air.local>
2026-09-02 09:16:15 +03:00
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