Commit Graph
5080 Commits
Author SHA1 Message Date
James Francis 5032008bc8 metal: enable Metal 4.0 tensor API on M5+/A19+ (llama/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-04 13:39:38 +03:00
Buğra Özgürsoy 8e54c659b5 metal : add fa-vec tunings for M1 Ultra (llama/28088)
* metal : add fa-vec tunings for M1 Ultra

* metal : move M1 Ultra tunings after M1 Max section

* metal : remove duplicate blank line
2026-09-04 13:39:37 +03:00
ynankani f22bb2ea4d CUDA: XOR swizzle flash attn K,V smem fp16 tiles (llama/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-09-04 13:39:37 +03:00
Georgi Gerganov dbc40efce8 metal : add concat support for quantized types (llama/28116)
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
2026-09-04 13:39:37 +03:00
BartowskiandGeorgi Gerganov 2f608ab4f8 AVX2: Speed up large batch size prompt processing of IQ models (llama/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-09-04 13:39:37 +03:00
Georgi Gerganov c6934d0fcf metal : add top-k radix implementation (llama/28073)
Assisted-by: DeepSeek-v4-Flash-0731
2026-09-04 13:39:37 +03:00
Hongqiang Wang 088c603e29 opencl: tune the quant paths for Intel Xe-LP GPUs to improve its TG and PP performance (llama/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-09-04 13:39:37 +03:00
c648b9a4d0 webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation (llama/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-09-04 13:39:37 +03:00
Jaden_Mach c1be45b890 ROCm: add radix TOP_K for long rows (llama/27466)
* ROCm: add radix TOP_K for long rows
2026-09-04 13:39:36 +03:00
Niklas Wenzel 7614a4c139 metal : add fa-vec tunings for M1 (llama/28078) 2026-09-04 13:39:36 +03:00
ynankani b0f4bc02ed CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were restricted to 1 token (llama/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-09-04 13:39:36 +03:00
Neo Zhang 76a51e82d7 sycl : Enhance to get the free memory of Intel GPU (llama/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-09-04 13:39:36 +03:00
Simon Teixidor 6ce7b89952 vulkan: tune mat-vec rows for batched inference on Strix Halo (llama/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-09-04 13:39:36 +03:00
fairydreamingandStanisław Szymczyk 96dddd87f2 ggml : add MUL_MAT to the list of ops that may need additional memory (for WebGPU) (llama/28071)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-09-04 13:39:36 +03:00
Ruben Ortlam db00b0196b vulkan: top_k radix select for k >= 1024 for Qwen 3.8 Flash Next (llama/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-09-04 13:39:36 +03:00
Shenghan Yang 01ebd225a7 hexagon: fix CPY fence bug (llama/28033) 2026-09-04 13:39:35 +03:00
codemonkey e5c96ca45d metal : add remaining Q4_1/Q5_0/Q5_1 fa-vec tunings for M2 (llama/28017) 2026-09-04 13:39:35 +03:00
hmirinandGeorgi Gerganov 4089fa628a rpc: avoid serializing buffers from other servers (llama/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-09-04 13:39:35 +03:00
Georgi Gerganov 749683d30f ggml : fix ggml_backend_buft_get_alloc_size() guard (llama/28038) 2026-09-04 13:39:35 +03:00
Aman Gupta e9583f075a ggml: add SWIGLU_CLAMP (llama/27930)
* ggml: add SWIGLU_CLAMP

* add vulkan shader
2026-09-04 13:39:35 +03:00
Pascal e900a732c8 CUDA: use the fast mm_ids_helper path for any n_expert_used (llama/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-09-04 13:39:35 +03:00
itterative 35d9e2237e hip: tune rdna 3 mmq config (llama/26284) 2026-09-04 13:39:34 +03:00
LunalFresh e5c9e3e3e9 hip : optimize Q2_0 dot-product path for gfx1201 (llama/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-09-04 13:39:34 +03:00
Georgi Gerganov 4b2243a6c2 ggml : add ggml_backend_op_alloc_size_may_expand, use it in RPC (llama/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-09-04 13:39:34 +03:00
Ryan C 43acf3d6e8 rpc: fix apple rdma error spew on teardown (llama/27908) 2026-09-04 13:39:34 +03:00
Nils Gladitz 1e0f382573 metal: add fa-vec tunings for M3 Ultra (llama/27999) 2026-09-04 13:39:34 +03:00
Daya Adianto b66593ef1b metal : Add fa-vec tuning for M3 Pro (llama/27963)
Related issue: #27668
2026-09-04 13:39:34 +03:00
Ryan C 5e494599a4 rpc : fix pre-rdma macOS versions (llama/27815) 2026-09-04 13:39:34 +03:00
3ad8b9b217 hexagon: support for device discovery and create sessions on demand (llama/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-09-04 13:39:33 +03:00
Titaniumtown 3d4e0e9858 sycl: split long rows in TOP_K instead of one work-group per row (llama/27847) 2026-09-04 13:39:33 +03:00
QuintinShaw c68f20555a metal : fix null-pipeline crash for F16 src1 mul_mat/mul_mat_id (llama/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-09-04 13:39:33 +03:00
Aman Gupta c969c68b54 ggml: allow passing alloc dependencies in graph_optimize (llama/27301)
* ggml: allow passing alloc dependencies in graph_optimize

* add alloc dep tests

* add TODO about using flat array
2026-09-04 13:39:33 +03:00
codemonkey b33bbc5c4e metal : add fa-vec tunings for M2 (llama/27940) 2026-09-04 13:39:33 +03:00
Hongqiang WangandLi He 2a11026cfe opencl: use a better matmul path on two Adreno GPU generations (llama/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-09-04 13:39:33 +03:00
Georgi Gerganov 285f1ffd99 metal : assert shared memory padding (llama/27951)
* metal : assert shared memory padding

* cont : add ref
2026-09-04 13:39:33 +03:00
Niklas Wenzel 308fa4f8a2 metal : add remaining fa-vec tunings for M4 Pro (llama/27915) 2026-09-04 13:39:32 +03:00
Nick Farrell 325c8d16c1 sycl: make --fit respect --fit-target better (llama/27629)
improve the --fit algorithm to take into account the actual peak
required VRAM for a given context size on a SYCL backend.

This includes both properly accounting for how much VRAM is required
when the allocated context is fully used (which makes the reported
context drop below what it did before, but stop it OOMing) as well
as preventing some overly-conservative calculations which meant too much
VRAM was being reserved.

Tested on a Arc b70 with unsloth's qwen3.8 (Q4_K_XL), able to get 262144 context,
fully usable, with q8_0 KV and MTP and 4k ubatch size using --fit-target 1
2026-09-04 13:39:32 +03:00
Jeff Bolz e644752070 vulkan: combine duplicated fastdiv functions, rename the one optimizing small divs (llama/27526)
* vulkan: combine duplicated fastdiv functions, rename the one optimizing small divs

* remove one more fastdiv
2026-09-04 13:39:32 +03:00
Jhen-Jie Hong 590fe18902 metal : add fa-vec tunings for M1 Max (llama/27932) 2026-09-04 13:39:32 +03:00
Jeff Bolz d501a0a3e8 vulkan: Change mul_mat_id to pad K rather than N (llama/27925)
The N padding is needed for mul_mat, but not mul_mat_id. For mul_mat_id,
we indirect the row index through a shared memory lookup table which avoids
any OOB row coordinate. But that callback doesn't bounds check K, so we
actually need K padding instead.
2026-09-04 13:39:32 +03:00
Eric A StaleeandJeff Bolz 4c38040fd3 vulkan: fix missing view-alias dependencies in ggml_vk_graph_optimize (llama/27812)
* vulkan: fix missing view-alias dependencies in ggml_vk_graph_optimize

is_src_of doesn't treat two views of one tensor as dependent, so the optimizer reorders nodes across aliased reads and writes.

Result: silently wrong tokens under greedy decoding, different output on every server start, and invalid speculative-decoding acceptance, with nothing logged.

Hits Qwen3.8's recurrent state (and any model with view-aliased state) on AMD and NVIDIA Vulkan.  CUDA is clean.

Compare view_src bases on both sides.

Fixes #27805

* vulkan: don't treat view/no-op nodes as aliasing dependencies

Nodes whose op is NONE, RESHAPE, TRANSPOSE, VIEW or PERMUTE execute nothing, so aliasing through them is not a real dependency. The previous base comparison matched them anyway, which only costs the optimizer reordering freedom.

Co-authored-by: Jeff Bolz <jbolz@nvidia.com>

* vulkan: make the lambda parameter const and capture is_empty in is_src_of

Code will not compile without these changes.
is_src_of has an empty capture list, so is_empty was not visible inside it, and is_empty took a non-const pointer, while is_src_of receives const ones. Other call sites pass non-const pointers, which still convert as usual.

---------

Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
2026-09-04 13:39:32 +03:00
Tekin ErtekinandGeorgi Gerganov caea96f6ce ggml : fix conv_transpose_2d for multiple batches (llama/26132)
* ggml : fix conv_transpose_2d for multiple batches

ggml_compute_forward_conv_transpose_2d_impl only computed the first
batch (ne[3] of the destination); every batch after the first was left
as zero. Both the src1 permutation and the main compute loop now iterate
over the batch dimension, and the work buffer size in ggml_graph_plan is
scaled by the src1 batch count so the extra permuted batches fit. A
multi-batch test case is added to test-backend-ops.

Fixes ggml-org/ggml#1448

* metal : fix conv_transpose_2d for multiple batches

The kernel only computed batch 0 of the input (src1->ne[3]); every
output batch after the first was left as zero, so multi-batch
conv_transpose_2d results diverged from the CPU reference.

The grid now covers all batches (OW x OH x OC x N), the kernel decodes
the batch from the grid z coordinate and offsets both the input and
destination indices accordingly. nb3 is passed in the kernel args.

Assisted-by: pi:llama.cpp/Qwen3.8-27B

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-09-04 13:39:32 +03:00
ravel7524andJeff Bolz ba99c09866 Vulkan: add hoisting support for row IDs and expert count in shaders (llama/26686)
* vulkan: add hoisting support for row IDs and expert count in shaders

* use hoisted row ids in coopmat2

* vulkan: address review feedback on count_experts
- use vk_op_count_experts_push_constants instead of a raw uint vector
- apply the fastdiv trick to the ne00 div/mod in count_experts
- compute the per-expert offsets with subgroupExclusiveAdd when the
  device supports it, keeping the serial path as fallback
- document the data_d layout and the hoisted_row_id_words bound
- drop a leftover debug print in ggml_vk_matmul_id

* vulkan: use init_pushconst_fastdiv for count_experts push constants

* vulkan: refine comments for row ID hoisting and data layout in count_experts shader

* Whitespace

---------

Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
2026-09-04 13:39:31 +03:00
StrongtutandStrongtut 0a150873ef metal : add fa-vec tunings for M4 (llama/27875)
This adds fa_vec_tuned_table records for Apple M4 to ggml-metal-tuning.cpp.

Includes F16, Q4_0, Q4_1, Q5_0, Q5_1, and Q8_0. (M4, 10 GPU Cores)

Co-authored-by: Strongtut <8432058+Strongtut@users.noreply.github.com>
2026-09-04 13:39:31 +03:00
7f78e1b469 OpenVINO: Update OV to 2026.3.1, whisper.cpp support, Qwen3.5 on NPU, and new ops (llama/27843)
* OpenVINO Backend: Fuse IM2COL + MatMul convolution into OpenVINO convolution

* ci:ggml-ov: Skip recurrent state rollback tests

* ci:ggml-ov: Skip recurrent state rollback tests

* Update OPENVINO.md

* ggml-openvino : add env-var gated op support debugging

* Fix ggml_rope_set_offset case

* OpenVINO backend: Support Whisper.cpp

* Fix code style

* openvino : enable qwen35 on NPU

Static shapes:
- get_graph_input_shape() left the s_copy / s_copy-leaf inputs dynamic
  ([1,1,1,-1]) even in static mode, which propagated a dynamic slot dim through
  GET_ROWS into the conv/GDN state, the state reshapes and the GDN output.
- With -np 1 the s_copy defrag remainder gathers zero rows; short-circuit that
  CPY to the untouched cache instead of emitting a degenerate Slice/Concat, and
  skip binding its zero-byte ggml tensor as an output (the dynamic path already
  did the latter, the static path wrote the full cache over a 0-byte buffer).

Token-count independence:
- In static mode the compiled model's token count is the prefill chunk size or
  1, not the captured cgraph's. Offsets derived from the captured count were
  therefore wrong. Anchor the GDN state slice at the end of the packed
  [attn | state] output and drop the rs_src_begin runtime inputs, and make
  VIEWs over the GDN output / conv_input pass through so the consumer does the
  slicing.
- CONT could not identify its token axis when the graph was captured with a
  single token (every trailing dim has the same stride and size 1) and baked
  the captured shape into the prefill model.

Chunked prefill:
- The last chunk is padded with fabricated tokens. Attention masks them, but
  the recurrent path folded them into cache_r/cache_s permanently. Add a
  chunk_valid_len runtime input, use it to zero g and beta for padded steps
  (making the recurrence an exact identity) and to end the conv snapshot window
  at the last valid token, and disable the recurrent-cache reset after the
  first chunk so earlier chunks are not wiped.
- get_is_prefill() and the chunk loop bound read inp_pos->ne[0] directly, but
  IMROPE stacks 4 position planes, so every decode step was run through the
  padded prefill model and the loop ran extra out-of-bounds chunks.

cache_rs_reset_idx/len now stay runtime Parameters in static mode, since
can_reuse_statically() does not invalidate the cached model on ComputeParams
changes. Add GGML_OPENVINO_FORCE_STATIC to exercise the static path on CPU.

* Update to OpenVINO 2026.3.1

* ggml-openvino: forward NPU compilation mode parameters

Add GGML_OPENVINO_NPU_COMPILE_CONFIG to the backend's cached environment so callers can configure the NPU compiler without using the generic property escape hatch.

When the value is non-empty, pass it to OpenVINO as NPU_COMPILATION_MODE_PARAMS. This enables settings such as optimization-level=3 for NPU compilation while preserving the existing behavior when the variable is unset and leaving CPU and GPU configuration unchanged.

Document the variable, its NPU-only scope, and the optimization-level=3 example in the OpenVINO backend runtime configuration table.

* ggml-openvino : support RELU, POOL_2D, QUICK_GEGLU, and ROLL ops

* reorder op table

* exclude GPU/NPU failing POOL_2D case

* move op type detection to compute_op_case

* Relax rope supported cases

* Fix pool case

* Update openvino doc, gpu driver in ov docker

* openvino: remove unused static remote context branch

* openvino: parallelize static model build

* Apply editorconfig

---------

Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
2026-09-04 13:39:31 +03:00
Ozymandias_EBON fa4d244c93 sycl: use TILE for quantized KV decode on BMG (llama/26689)
Route quantized KV decode to TILE on Xe2 (BMG) only, keep VEC on other archs until validated there.
2026-09-04 13:39:31 +03:00
Titaniumtown 97d0da26a2 sycl: bind the f16 KV cache in place for the oneDNN SDPA path (llama/27468)
Measured at a live KV length of 34816 (32768 depth plus one 2048 ubatch),
on Qwen3.8 27B Q4_K_S:

  per tensor         4 * 34816 * 256 * 2 B  =  71.3 MB
  staged per call    K and V, so 2x         = 142.6 MB
  traffic per call   read once, write once  = 285.2 MB
  traffic per ubatch 285.2 MB * 16 calls    =   4.56 GB

One ubatch is one ggml_cgraph submission (llama_context::process_ubatch ->
graph_compute), so that 4.56 GB is the cost of a single 2048-token prefill
chunk, and it scales with the live KV length: the first ubatch of the same run,
at seq = 2048, moves 0.27 GB.

Reproduce the two measured inputs with:

  GGML_SCHED_DEBUG=2 llama-bench -m MODEL -p 8 -n 0 -r 1 -ngl 0 \
      -fa on -ctk f16 -ctv f16 -v > nd.txt 2>&1
  grep -E 'n_layer|n_head_kv|n_embd_head_k' nd.txt
  awk '/node #  0 /{g++} g==1 && /\(FLASH_ATTN\)/{n++} END{print n+0}' nd.txt
2026-09-04 13:39:31 +03:00
Georgi Gerganov 530e3f4834 metal : add fa-vec tunings for M3 Max, M5 and M5 Pro (llama/27863)
* metal : add fa-vec tunings for M5

This is a followup contribution to efeda76b948f59ee52ea20db640bc4cf3dfe8ac1 as requested in https://github.com/ggml-org/llama.cpp/discussions/27668 to add support for additional Apple GPUs. I generated this output using the provided instructions:

```sh
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp

cmake -B build -DGGML_METAL=ON
cmake --build build --target ggml-metal-tuning -j

./build/bin/ggml-metal-tuning fa-vec --dtype f16,q8_0 > fa_vec_rows.txt 2> fa_vec_sweep.log
```

This ran on a machine with Apple M5.

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* metal : add fa-vec tunings for M5 Pro

This adds fa_vec_tuned_table records for Apple M5 Pro to ggml-metal-tuning.cpp.

Contributed by SerayaEryn in https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18157544 (F16, Q4_0, Q8_0; M5 Pro, 20 GPU cores).

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* metal : add fa-vec tunings for M3 Max

This adds fa_vec_tuned_table records for Apple M3 Max to ggml-metal-tuning.cpp.

Contributed by TeeAaTeeUu in https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18175220 (F16, Q8_0; M3 Max, MacBook Pro 64GB, low power mode).

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* cont : whitespaces
2026-09-04 13:39:31 +03:00
Brad Smith ff38b98e5a metal : add fa-vec tunings for M4 Pro (llama/27824)
This is a followup contribution to efeda76b948f59ee52ea20db640bc4cf3dfe8ac1 as requested in https://github.com/ggml-org/llama.cpp/discussions/27668 to add support for additional Apple GPUs. I generated this output using the provided instructions:

```sh
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp

cmake -B build -DGGML_METAL=ON
cmake --build build --target ggml-metal-tuning -j

./build/bin/ggml-metal-tuning fa-vec --dtype f16,q8_0 > fa_vec_rows.txt 2> fa_vec_sweep.log
```

This ran on a MacBook Pro (14-inch, Nov 2024) with Apple M4 Pro. The `ggml-metal-tuning` command completed successfully in 1h 13m 1s with no other notable load on the system.
2026-09-04 13:39:31 +03:00
cqderek b6571e4a55 ggml-hexagon: add HTP unary ops for ABS and LOG (llama/27786)
Add HVX-accelerated implementations for GGML_OP_LOG and
GGML_UNARY_OP_ABS on the HTP backend.

- Register HTP_OP_UNARY_ABS and HTP_OP_UNARY_LOG in op_remap_to_htp()
- Add ABS and LOG to ggml_backend_hexagon_device_supports_op()
- Implement hvx_abs_f32_aa() in hvx-arith.h using hvx_vec_abs_f32()
- Implement hvx_log_f32_aa() in hvx-log.h using hvx_vec_log_f32()
- Add abs_f32() and log_f32() row-wise dispatch in unary-ops.c
- Define tiled and non-tiled task functions via DEFINE_UNARY_TASK and
  DEFINE_UNARY_TILED_TASK macros
- Route HTP_OP_UNARY_ABS and HTP_OP_UNARY_LOG through execute_op()
  in main.c
2026-09-04 13:39:30 +03:00