Commit Graph

2811 Commits

Author SHA1 Message Date
Jassieluo 02be871218 sycl : clamp softmax input to avoid underflow (llama/24941) 2026-06-26 16:03:57 +03:00
leonardHONG 96996579f0 CUDA: batch out_prod broadcast (dps2>1) path with cublasSgemmBatched (llama/24426) 2026-06-26 16:03:57 +03:00
shaofeiqi 7ce3616ae8 opencl: flush profiling batch at shutdown for incomplete batches (llama/25016) 2026-06-26 16:03:57 +03:00
Oliver Simons 9cad795918 CUDA: Various fixes to `cpy.cu` (llama/25000)
* Add failing test-case to test-backend-ops

Extracted from https://github.com/ggml-org/llama.cpp/issues/24072

* Minimize repro with help of AI

N = 8 * (65535 - 1) + 1 = 524273

* Port and adjust workaround from 0ba798341e

Fall-back should share code, also relax y-z constraint to be inclusive

* Add test-case + fallback also for y dim

* Fix x-guards which is 2^{31}-1, so inlusive of INT_MAX

* Fix overflow problems for transposed copy kernel
2026-06-26 16:03:57 +03:00
fairydreaming 82686b518f ggml : address integer overflows in binary ops CUDA implementation (llama/24706)
* ggml : address integer overflows in binary ops CUDA implementation

* ggml : add size_t casts to avoid integer overflows

* ggml : add more asserts checking integer overflows in binary ops CUDA implementation

---------

Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-06-26 16:03:57 +03:00
David Spruill ffd5b0b5c7 sycl : support --split-mode tensor (llama/24152)
* Sycl tp stage1 (llama/1)

* SYCL: tensor parallelism (--split-mode tensor) for dual-GPU

Adds the comm_init/comm_free/comm_allreduce_tensor trio that the
meta-backend queries via get_proc_address to enable backend-specific
all-reduce, mirroring the pattern used by ggml-cuda.cu.

For N=2 (the common dual-GPU case) implements a degenerate ring
all-reduce with two size-branched paths:

  * Small (nelem < 32768): FP32 direct memcpy + per-device ADD kernel
    chained via depends_on(memcpy_event). 4 SYCL submissions/call.

  * Large (nelem >= 32768): BF16-compressed. Each device compresses
    FP32 -> BF16 in a local outbox, cross-device memcpys to the peer's
    inbox (HALF the PCIe bytes), then decompresses + adds into the
    local FP32 partial. 6 SYCL submissions/call but PCIe bytes halved
    -- wins for any tensor where PCIe dominates kernel time.

Threshold and BF16 path pattern mirror the CUDA NCCL allreduce.

Storage: ONE persistent uint8_t buffer per device, 4 * nelem bytes
(matches both path layouts: FP32 nelem floats; BF16 outbox+inbox =
2 * nelem uint16_t each). Single alloc+free per device keeps the
SYCL pool's strict-LIFO invariant trivial.

Initial impl handles N=2 FP32 contiguous tensors. Other cases return
false, causing the meta-backend to use its generic butterfly fallback.

Per-call sync is intentionally omitted. SYCL in-order queue semantics
ensure that the meta-backend's next compute on the same per-device
queue waits for our final ADD, and the next allreduce's first op on
the same persistent buffer waits via the same queue. Only comm_free
does an explicit final wait.

OneCCL is NOT used: OneCCL 2021.17 hardcodes single-device-per-process
in communicator_impl.hpp:47 (condition devices.size() == 1), which is
incompatible with llama.cpp's single-process multi-GPU model.

Measured on dual Intel Arc Pro B70 (NEO 26.05.x, oneAPI 2025.3 +
DPC++ nightly):

  Llama-3.3-70B Q4_K_M, -sm tensor -fa 1 -ctk f16 -ctv f16:
    pp512 = 377.08 t/s  (vs 313.65 layer mode = +20.2%)
    tg128 = 17.40 t/s   (vs   9.74 layer mode = +78.6%)

  Qwen3-Coder-Next-80B-A3B Q3_K_M (MoE):
    pp512 = 216.56 t/s  (vs 156.58 meta-backend butterfly = +38.3%)
    tg128 = 17.60 t/s   (vs  14.31 meta-backend butterfly = +23.0%)

  Qwen3-4B Q4_K_M:
    pp64  = 984.51 t/s, tg16 = 49.29 t/s

Llama-3.3-70B in SYCL TP now comfortably beats production layer mode
on both prefill and decode. Coder-Next-80B-A3B (MoE) also wins on
both — the BF16 path is what unlocks the many-medium-allreduces
prefill pattern.

Build/CMake: no changes. No new dependencies. ~210 lines added across
ggml-sycl.h and ggml-sycl.cpp.

* Fix comments

* documentation update to address PR feedback

* Bring over my device-to-device memcpy chagnes

* move the dev2dev_memcpy calls to the upstream 7-parameter variety

* Fix a typo and remove a trailing whitespace
2026-06-26 16:03:57 +03:00
Neo Zhang c4ba5e2e13 sycl : fix the failed UT cases of conv_3d (llama/24900) 2026-06-26 16:03:57 +03:00
lhez fe37d2e489 opencl: support non-contig rows in norm (llama/24965) 2026-06-26 16:03:57 +03:00
Max Krasnyansky e0a0f1256f hexagon: MUL_MAT and MUL_MAT_ID rework : 32x32 tiled weight repack, kernel-params, cached graphs (llama/24954)
* hex-mm: new weight layout and fusion updates

* hvx-mm: unroll the new tiled vec_dots to optimize hvx register util

* hex-mm: optimize dyn.quant format for q8_0 and q8_1 to reduce overhead in vec_dots.

* hvx-mm: parallel quantizer per block for large rows

* hvx-mm: simplify and futher optimize dyn.quant and vec_dots

* hvx-mm: keep intermediate per tile accumulators in fp16

* hmx-mm: optimize weight dequant by aligning the repacked tiles with the DMA

* hmx-mm: remove qweight scratch and just use vtcm_weight

* hmx-mm: remove all unused and obsolete code

* hmx-mm: the new tiled repack format is here to stay -- rename all x4x2 to _tiled

* hmx-mm: improve activation processing with dma prefetch

* hex-mm: fix hmx/hvx fallback logic and MUL_MAT_ID allocation (unbreaks OLMoE)

* hex-mm: align the weight tiles with dma just like we did in hmx-mm

* hex-mm: factor out common mm bits into htp/matmul-ops.h

* hex-mm: start moving mm kernel selection to the host

* hex-mm: move all of the matmul param compute into the host

* hmx-mm: restore pipelined mode

* hmx-mm: unroll the dequant functions to optimize register usage

* hmx-mm: further improve activation process

* hex-mm: use vtcm_seq_alloc for all vtcm allocations and define more common functions

* hex-mm: improve mm optimizer to acount for number of activation threads

* hex-mm: fix matmul-id kernel params selection (unbreaks OLMoE and LFM)

* hexagon: remove support for arch < v73 since HMX is now required for most use-cases

* hex-mm: cleanup naming for consistency

* hex-mm: make sure matmul fusion accounts for vtcm allocation

* hex-mm: minor cleanup for kernel_params definition

* hex-mm: replace hardcoded limits with proper checks for vtcm requirements

* hex-mm: add support for non-tiled mm as a fallback option and factor out hvx kernels into separate header

* hex-mm: remove unused functions

* hex-mm: add shorthand for MM_SELECT in run-tool script

* hvx-mm: factor out hvx/hmx microkernels and unify matmul entry and dispatch

* hex-mm: further cleanup matmul fallback path

* hex-mm: refactor matmul entry point and dispatch a bit further

* hexagon: update cmake build to enable hmx for everything

* hex-ops: optimize kernel_param updates and include summary in the logs

* hex-mm: add support for GGML_HEXAGON_MM_SELECT

* hex-mm: add hex-common header

* hex-mm: pass correct number of tasks to workpool

* hex-mm: add proper checks for no-work in dyn.quant tasks

* hex-mm: convert all quantizers into a macro

* hex-mm: fix hvx-flat fallback to pass all MUL_MAT tests

* hex-mm: vectorize q8_1 quantizer

* hex-mm: improve fused ffn mm stride handling

* hex-mm: consistent use of n_threads and pipeline in kernel_params

* hexagon: minor formatting

* hex-mm: update MUL_MAT_ID kernel_param handling to make sure host/npu are in sync

* hvx-mm: go back to accumulating in fp32 in tiled hvx kernels, more accurate and same perf

* hvx-mm: unroll the loops and remove masking that is not needed for tiled accums

* hmx-mm: optimize activation processing (slit loops, some unrolling, etc)

* hmx-mm: minor optimization for output processing

* hex-mm: consistent use of uint32_t and size_t in mm kernels

* hex-mm: remove legacy restrictions for rows to be multiple of 256

* hexagon: replace sprintf with snprintf

* hex-mm: relax hardcoded nrows checks and rely on VTCM size requirements

* hexagon: minor alignment fix

* hexagon: fix trailing spaces

* hex-mm: relax padding from 256 to 128 (leftovers)

* hex-mm: remove redundant checks for weight align to 128

we always use 2D dma for the weights and align them properly

* hmx-mm: MUL_MAT_ID better work distribution between hvx threads and hmx tracing

* hex-mm: specialize per-token mmid activation handling

* hex-profile: update python scripts to handle kernel-params section in the logging output

* hex-mm: move n_prefetch (aka dma_depth) into kernel params and remove unused fields

* hex-trace: use easier to parse format, simply and fix post-proc scripts

* hmx-mm: relax 32 row limit for output processing which helps utilization

* hmx-mm: use start-chunk idx for tracing info

* hmx-mm: parameterize activation dma pipeline

* hexagon: add support for simple graph caching to avoid recomputing kernel-params

* hex-mm: remove left-over repack functions

* hex-mm: tighten n_prefetch asserts

* hex-mm: remove duplicate round/align_up helper

* hexagon: cleanup common header used in host/npu

* hexagon: update early wakeup threshold

* hmx-mm: define cost constants and update solver to assume that repacked ne[1] is padded to 32

* hmx-mm: make precompute_matmul a bit more readable (split into smaller functions, etc)

* hex-mm: remove n_threads constraint

* hex-mm: minor formatting updates

* hex-mm: remove obsolete profiling logs

* hex-mm: restore hardcode gate to refuse lm-head to avoid repacking that tensor
2026-06-26 16:03:57 +03:00
Wagner Bruna a201c20505 vulkan: allow reducing the graph submission batches to avoid timeouts (llama/24872) 2026-06-26 16:03:57 +03:00
liminfei-amd f081d84a4e vulkan: fail the build when a shader fails to compile (llama/24450)
* vulkan-shaders-gen: fail the build when a shader fails to compile

vulkan-shaders-gen did not detect shader-compile subprocess failures, so a
broken libggml-vulkan could be produced while the build reported success and
the breakage only surfaced at run time. execute_command() discarded the child
exit code (POSIX waitpid passed nullptr for status; the Windows branch never
called GetExitCodeProcess) and string_to_spv decided success only from whether
stderr was empty, so a non-zero exit with empty stderr, or a subprocess that
failed to launch, was treated as success.

Return the child exit code from execute_command() (WEXITSTATUS on POSIX,
GetExitCodeProcess on Windows), treat a non-zero exit or non-empty stderr or a
launch exception as a failure, and record it in an atomic flag. main() checks
the flag after process_shaders() and returns EXIT_FAILURE before writing the
output files, so the build stops instead of emitting a broken backend.

Fixes #24393

Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com>

* vulkan-shaders-gen: simplify compile_failed access and drop unreachable return

Address review feedback on #24450:
- Access the std::atomic<bool> compile_failed directly (= / implicit bool)
  instead of .store()/.load(); the flag stays atomic because the worker
  threads in process_shaders() set it concurrently.
- Remove the unreachable trailing return -1 in execute_command(): on POSIX the
  child _exit()s after execvp and the parent returns (fork()<0 throws); on
  Windows the block returns the exit code.

Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com>

---------

Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com>
2026-06-26 16:03:57 +03:00
Jeff Bolz 792da0ecdb vulkan: Apply bias before softmax in FA, to avoid overflow (llama/24909) 2026-06-26 16:03:57 +03:00
Jeff Bolz 3fc2c57a63 vulkan: support all backend tests for SQR/SQRT/SIN/COS/CLAMP/LEAKY_RELU/NORM (llama/24582)
* vulkan: make SQR/SQRT/SIN/COS/CLAMP/LEAKY_RELU use unary.comp

* vulkan: make NORM support noncontig

* add noncontiguous row test cases for norm/l2_norm, handle this in the CPU backend and l2_norm.comp

* fix supports_op for cuda and webgpu
2026-06-26 16:03:57 +03:00
Jeff Bolz 1fdbfb19eb vulkan: Support GET_ROWS_BACK (llama/24883) 2026-06-26 16:03:57 +03:00
Jeff Bolz c25eb26c32 vulkan: support CONV_3D (llama/24612)
* vulkan: support CONV_3D

This is a pretty direct port of conv2d_mm.comp to CONV_3D, done by codex
and cleaned up by me.

* disable slower perf tests
2026-06-26 16:03:57 +03:00
Jeff Bolz cc0d00f610 vulkan: make mul_mm ALIGNED a spec constant (llama/24689)
This trims down some of the shader variant explosion and reduces binary size.
2026-06-26 16:03:57 +03:00
Wyatt Caldwell e3c18388c5 vulkan: link ggml-cpu when GGML_VULKAN_CHECK_RESULTS / RUN_TESTS are enabled (llama/24444)
The result-checking and test debug paths in ggml-vulkan.cpp call ggml_graph_compute_with_ctx() to compute a CPU reference graph, but that symbol is defined in ggml-cpu, which ggml-vulkan does not link. Enabling -DGGML_VULKAN_CHECK_RESULTS=ON (or -DGGML_VULKAN_RUN_TESTS=ON) therefore fails to link with an unresolved external (e.g. LNK2019 on MSVC, undefined reference on GCC/Clang). This regressed after ggml-cpu was split into its own library. Link ggml-cpu under those two options so the debug builds link again.

Signed-off-by: Wyatt Caldwell <218154709+Detensable@users.noreply.github.com>
2026-06-26 16:03:57 +03:00
Masashi Yoshimura 2955e84ec3 ggml-webgpu: improve MTP inference by using mat-vec path for small batches (llama/24811)
* ggml-webgpu: improve small batches decoding

* Add barrier to the NUM_COLS loop in mul-mat-vec
2026-06-26 16:03:57 +03:00
Shawn Gu 0badff27cb opencl: q8_0 gemv precision improvement (llama/24923) 2026-06-26 16:03:57 +03:00
Neo Zhang 1b3f8434f9 support bf16 on bin_bcast OP and unary OPs (llama/24838)
* support bf16 on bin_bcast OP and unary OPs

* support the older Intel compiler than 2026.0
2026-06-26 16:03:57 +03:00
Guanhuai Zhang 556d900cfb fix(hexagon): use padded stride for ssm-conv weights (llama/24470) 2026-06-26 16:03:57 +03:00
Adrien Gallouët 420518d88e ggml : optimize AMX (llama/24806)
Flatten the partition over n_batch * M so every thread participates in
the quantization

    | CPU                             | Model                         | Test   |   t/s OLD |   t/s NEW |   Speedup |
    |:--------------------------------|:------------------------------|:-------|----------:|----------:|----------:|
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_NL - 4.5 bpw  | pp512  |    730.71 |    779.86 |      1.07 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_NL - 4.5 bpw  | tg128  |     87.88 |     86.79 |      0.99 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_XS - 4.25 bpw | pp512  |    725.09 |   1023.31 |      1.41 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_XS - 4.25 bpw | tg128  |     83.64 |     83.62 |      1.00 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_0              | pp512  |    820.51 |    924.05 |      1.13 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_0              | tg128  |     90.59 |     92.46 |      1.02 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_1              | pp512  |    776.88 |    872.79 |      1.12 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_1              | tg128  |     89.39 |     90.94 |      1.02 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_M            | pp512  |    719.28 |   1009.27 |      1.40 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_M            | tg128  |     80.62 |     80.86 |      1.00 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_S            | pp512  |    732.29 |   1077.29 |      1.47 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_S            | tg128  |     86.42 |     83.53 |      0.97 |

Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-06-26 16:03:57 +03:00
Masashi Yoshimura 75b90fb0ea ggml-webgpu: add adapter toggles for F16 on Vulkan + NVIDIA 2026-06-26 16:03:57 +03:00
Xuan-Son Nguyen adbcfa0d46 mtmd, arg: fix utf8 handling on windows (llama/24779)
* mtmd, arg: fix utf8 handling on windows

* also fix ggml_fopen

* fix build fail

* also fix CLI
2026-06-26 16:03:57 +03:00
Georgi Gerganov f92382f1c2 ggml : bump version to 0.15.2 (ggml/1548) 2026-06-19 12:53:43 +03:00
shalinib-ibm 587f602ea1 ggml-cpu: support K tails in power10 Q8/Q4 MMA matmul (llama/24753)
* ggml-cpu: support K tails in Power10 MMA Q8/Q4 matmul

This patch removes the requirement that K be divisible by kc in the tinyBlas_Q0_PPC tiled matmul path. Process the final K panel using its actual depth and pass the reduced panel size through packing and kernel execution.  This allows more workloads to use the MMA kernel and reduces fallback to mnpack.

* Apply suggestion from @taronaeo

Co-authored-by: Aaron Teo <taronaeo@gmail.com>

---------

Co-authored-by: Aaron Teo <taronaeo@gmail.com>
2026-06-19 12:53:43 +03:00
Pascal 50fd261d0d Ggml/cuda col2im 1d (llama/24417)
* cuda: add GGML_OP_COL2IM_1D, follow-up to the CPU op

* cuda: col2im_1d use fast_div_modulo for the index decomposition

* cuda: col2im_1d tighten supports_op, type match and contiguous dst
2026-06-19 12:53:43 +03:00
Max Krasnyansky 5cbc2153c2 hexagon: support for op-trace (fine-grain tracing of HVX/HMX/DMA events) (llama/24592)
* hex-optrace: add support for optrace and instrument matmul and flash-atten code

* hex-trace: improve trace event and prefetto generator

* hex-trace: add new script dedicated to handling traces, specifically perfetto traces

* hex-trace: add --head/--tail options to profile and trace tools

* hex-trace: fix whitespaces

* hex-trace: fix flake8 warnings

* hex-trace: fix flake8 warnings

* hmx-fa: restore q_tiles clearing

* hex-profile: remove circular dep in includes

* hex-trace: simplify trace sizing check

* hex-profile: sort events in the summary by name
2026-06-19 12:53:43 +03:00
Neo Zhang 4c4f6eaca7 rename GGML_SYCL_SUPPORT_LEVEL_ZERO (llama/24719)
* rename GGML_SYCL_SUPPORT_LEVEL_ZERO to GGML_SYCL_SUPPORT_LEVEL_ZERO_API, and GGML_SYCL_ENABLE_LEVEL_ZERO to  GGML_SYCL_USE_LEVEL_ZERO_API

* fix code format

* fix error when rebase
2026-06-19 12:53:43 +03:00
Neo Zhang 69a9798a82 sycl : support MUL_MAT and OUT_PROD with Q1_0 (llama/24721) 2026-06-19 12:53:43 +03:00
Neo Zhang 312033789f support OPs: conv_2d, conv_2d_dw, conv2d_transpose (llama/24600)
* fix conflict

* fix format issue, rename

* rm debug code

* correct the file name
2026-06-19 12:53:43 +03:00
Georgi Gerganov c39dd2db8e metal : check for BF16 support in concat kernel (llama/24747) 2026-06-19 12:53:43 +03:00
shalinib-ibm 6aff09ac48 ggml-cpu: Conditionally enable power11 backend based on compiler support (llama/24687)
* ggml: Conditionally enable power11 backend based on compiler support

Guard POWER11 backend creation behind a compiler flag check for -mcpu=power11. This avoids build failures on current GCC/Clang toolchains while preserving forward compatibility once POWER11 support becomes available.

* Update CMakeLists.txt

ggml-cpu: Use -mcpu=power10 for P10 and P11
2026-06-19 12:53:43 +03:00
Georgi Gerganov d911569230 metal : implement rope_back operator (llama/24725)
Reuse existing rope kernels with a function constant to toggle forward/backward
rotation, avoiding duplicate kernel code.

Assisted-by: pi:llama.cpp/Qwen3.6-27B
2026-06-19 12:53:43 +03:00
Georgi Gerganov ad19c98240 metal : add f16 and bf16 support for concat operator (llama/24724)
* metal : add f16 and bf16 support for concat operator

Extend the Metal backend concat operator to support f16 and bf16 tensor
types in addition to the existing f32 and i32 support.

- Template kernel_concat on type T with specializations for float, half,
  bfloat, and int
- Add type-specific pipeline getter ggml_metal_library_get_pipeline_concat()
- Update device support check to allow f16 unconditionally and bf16 when
  device supports bfloat16
- Update dispatch to select the correct kernel specialization by type

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

* metal : extend concat operator to support f16, bf16, i8, i16 and i64

Assisted-by: pi:llama.cpp/Qwen3.6-27B
2026-06-19 12:53:43 +03:00
Neo Zhang 6772827b26 add dev2dev memcpy by SYCL API (llama/24476)
* add dev2dev memcpy by SYCL API

* mv GGML_SYCL_DEV2DEV_MEMCPY to runntime table

* update the detect method for p2p comm

* fix the erro created during fix confilct

---------

Co-authored-by: Neo Zhang <NA>
2026-06-19 12:53:43 +03:00
Neo Zhang 01b1c3ada7 Add conv_3d (llama/24691)
* add conv_3d

* optimize

* update ops.md

* restore test script

* rm unused code

* rm copyright notes
2026-06-19 12:53:43 +03:00
Winston Ma ae09736e83 vulkan: record actual memory properties during buffer creation (llama/24326) 2026-06-19 12:53:43 +03:00
Ruben Ortlam 201f69ce6a Revert "cuda: reset cuda context after reading memory size (llama/23935)" (llama/24715)
This reverts commit 0f7fada56bea32c9e0db3874b801d887ff6c4529.
2026-06-19 12:53:43 +03:00
kononnable 055bf5c73c ci: fix vulkan docker images (llama/24595)
* Update vulkan-shaders-gen.cpp

* Update vulkan-shaders-gen.cpp

add comment describing code change intention

* Update vulkan-shaders-gen.cpp

fix potential UB
2026-06-19 12:53:43 +03:00
lhez 9f785839a3 opencl: optimize mul_mat_f16_f32_l4 for decode (llama/24504) 2026-06-19 12:53:43 +03:00
Zijun Yu 5fa14e9931 openvino: OV 2026.2, context-shift, Q5_1 support, gemma4 dense/embedding, and -fa off (llama/24503)
* Add interface is_model_splitted() to check the c-graph is splited or not

* Infer and propagate dynamic-dimension indices for all tensors in the GGML graph in api compute_model_outputs()

* Only do this for fallback sub graph

* Move dynamic dims compute in graph missmatch

* ggml-openvino: fix tensor data handling for PERMUTE/VIEW ops in split models

* ggml-openvino:add comments

* ggml-openvino: override VIEW op_case to 0 for split model inputs

* openvino backend: Handle unsupported VIEW shape-mismatch in OpenVINO backend

* Enable additional mul_mat tests and add tensor data saving function (llama/81)

* ggml-openvino: fix CONT/TRANSPOSE mapping and improve dynamic-dimension handling

* OpenVINO: add NORM/TANH support and rework SOFT_MAX translation

* ggml-openvino: extend VIEW handling

* Enable -fa off (llama/118)

* Enable --context-shift

* Fix llm param compute error for normal softmax not the softmax in attention

* OpenVINO backend: fix error for attention size compute in llm param

* use tensor->extra in infer_request i/o

* OpenVINO backend: refacter the compute_llm_params() func add get_attention_pattern_case to easy extand

* OpenVINO backend: clean unused code

* 1to1 match op update (llama/146)

* added translate_1to1_match_1_input function and updated gelu and tanh translations

* Remove unused translation function calls

---------

Co-authored-by: Mustafa Cavus <mustafacavus@intel.com>

* initial gemma4 support

* removed hardcoded names for kv cache slicing

* OpenVINO backend: Add new attention pattern for llm parameters compute

* flash attn Q shape static conversion

* Remove slice in permute translation when n_seq is 1

* return optional in extract_layer_from_name

* OpenVINO backend: refactor VIEW related operation (llama/148)

* OpenVINO backend: refactor VIEW related operation

* Enable VIEW handling in following ops

* OpenVINO backend does not support GGML_OP_NORM & GGML_OP_L2_NORM with VIEW input accuracy issue from OpenVINO

* OpenVINO backend: Add ops l2_norm & pad

* OpenVINO backend does not support CPY with non-contiguous data or mismatched types

* add op SSM_CONV GATED_DELTA_NET

* OpenVINO backend: fix error for bf16 in OV gpu plugin

* reverted static Q input shape for attention layer

* OpenVINO backend: remove hardcode name inp_tokens, which ignore some leaf case

* Disable remote tensor due to bug in ov gpu

* Disable n_token > 1 GATED_DELTA_NET on gpu

* OpenVINO backend: fix the view op dynamic handling issue in gemma4 & enable view + get_row

* OpenVINO backend: clean code

* OpenVINO backend: enable view + norm/rms_norm

* OpenVINO backend: concat op

* OpenVINO backend: argsort op

* OpenVINO backend: enable unary + view & GGML_UNARY_OP_SOFTPLUS

* Fix issue for test-backend-ops in TOPK_MOE, which compare VIEW ops result, VIEW node in OpenVINO no need compare, the whole graph result is correct

* OpenVINO backend: enable sum_rows

* OpenVINO backend: enable clamp

* OpenVINO backend: enable DIV

* OpenVINO backend: enable GGML_OP_MUL_MAT_ID

* OpenVINO backend: disable MUL_MAT_ID_FUSION case with large mem needed

* OpenVINO backend: Disable GGML_OP_ARGSORT, cause test_backend-ops failed

* OpenVINO backend: fix issue in mul_mat_id

* OpenVINO backend: Disable DIV with broadcast on GPU

* OpenVINO backend: update DIV

* use ov internal op GatedDeltaNet

* OpenVINO backend: enable llama erch test qwen3next

* OpenVINO backend: enable RMS_NORM + VIEW & remove op_case 2 for rope

* OpenVINO backend: fix error

* suggested changes, need review

* suggested changes, need review

* OpenVINO backend: clean unused code & fix build warning

* OpenVINO backend: enable minicpm3 for arch test

* Disable GDN op (llama/177)

* disable gated_delta_net

* update stateful_kv_size correctly in mismatch case

* OpenVINO backend: enable arch test for qwen3vl

* OpenVINO backend: enable cohere2 for arch test

* OpenVINO backend: enable t5 for arch test

* OpenVINO backend: enable jamba for arch test

* OpenVINO backend: remove warning for tmp

* OpenVINO backend: enable kimi-linear for arch test

* Remove unused

* Fix gpt-oss accuracy issue

* OpenVINO backend: enable arctic for arch test

* OpenVINO backend: enable grok for arch test

* Gemma4 initial npu support (llama/179)

* Initiall gemma4 npu support

* temp. fix for gemma4 accuracy bug on npu

* Remove hardcoded names for npu-fold handling

* revert static n tokens for cont translation as it is not needed

* removed unused variable

* ggml-openvino: add GGML_OPENVINO_ENABLE_CACHE env var to control decoder cache. Add environment variable GGML_OPENVINO_ENABLE_CACHE (default: YES). When set to NO, the decoder_cache is bypassed and models are rebuilt from the cgraph on every inference call in both dynamic and static compute paths. This is useful for debugging and verifying correctness without caching interference.

* Revert "Gemma4 initial npu support (#179)"

This reverts commit 0d29a9c4a52dc2c8aa52990f1a3854cfb01768ad.

* OpenVINO backend: disable debug log print

* Update TBB discovery. Delegated to OpenVINOs own config.

* OpenVINO backend: GGML_OPENVINO_ENABLE_CACHE YES -> 1

* OpenVINO backend: fallback FLASH_ATTN_EXT in gemma3n to CPU backend

* Add raw ov infer profiling metric

* Add OV raw infer time metric to static compute path

Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com>

* Modify precision of static profiling

* update to OV 2026.2, add OV windows CI

* fix editorconfig-checks

* Initiall gemma4 npu support

* temp. fix for gemma4 accuracy bug on npu

* Remove hardcoded names for npu-fold handling

* revert static n tokens for cont translation as it is not needed

* removed unused variable

* test-llama-archs fix

* Fix gemma4 flash_attn fallback

* support im2col

* fix code style

* disable add_rope_sin_cos optimization

* stateless boradcast and rope optimizations

* Enable manual gqa attn by default for stateless gpu

* manual gqa: fixed static batch

* gemma4 llama-bench ctx update fix

* Update OV win CI

* stateful rope fusion temp. fix

* OpenVINO backend: Conslolidate supported ops

* Exclude unsupported GGML_OP_SUB cases

* Exclude unsupported TOPK_MOE cases

* OpenVINO Backend: MUL_MAT enhancements

* Update OV CI

* support f16 mask input for npu

* Make GGML_OPENVINO_* env vars usage uniform

Standardize all GGML_OPENVINO_* env flags:
positive integers >0 to enable. Unset, empty, =0, or non-numeric values to disable.
This fixes cases where text values or empty strings enabled features.

* OpenVINO backend: Enhance envvar handling

* more cleanup

* move ggml_openvino_env_flag to appropriate place

* OpenVINO backend: add REPEAT translator, Q5_1 weights, and GLU view-input fix

* ggml-openvino: fix -Werror=cast-qual in extract_q5_1_data

* Update openvino.Dockerfile

Use BuildKit cache mounts for faster Docker rebuilds.
Use apt instead of dpkg, remove unused .ddeb downloads, add DLLAMA_BUILD_TESTS=OFF.

* ggml-openvino: centralize env var access via *getenv_str/getenv_int helpers

Replace getenv and legacy flags with _str and _int helpers.Minor cleanup, doc updates.

* OpenVINO backend: Enable GGML_OP_ADD_ID

* Uptade openvino backend clamg-format

* clang-format

* Update OPENVINO.md (llama/211)

* OpenVINO backend: fix accuracy issue for op CONCAT with i64 precision

* Remove strict concurrency for gpu-openvino-low-perf

* Update openvino CI keynames; add ccache-clear

* Apply suggestions from code review

Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>

* Fix formatting

---------

Co-authored-by: Xuejun Zhai <Xuejun.Zhai@intel.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: Mustafa Cavus <mustafacavus@intel.com>
Co-authored-by: Xuejun <XuejunZhai@intel.com>
Co-authored-by: Wang Yang <yang4.wang@intel.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com>
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
2026-06-19 12:53:43 +03:00
Neo Zhang da66f048b1 sycl : Enable to support fp16 by OPs: SQR, SQRT, LOG, SIN, COS, CLAMP (llama/24692) 2026-06-19 12:53:43 +03:00
Alexey Kopytko fddcda58a3 SYCL: fix use-after-free bug with async memcpy in MoE prefill (llama/24676)
* SYCL: fix a bug with async memcpy

* make mmid_row_mapping_host persistent

* comment on stream->wait

* Apply suggestion from @sanmai

* Apply suggestion from @sanmai

* Apply suggestion from @sanmai
2026-06-19 12:53:43 +03:00
Francois Dugast dd1a6ca897 sycl: Add optional USM system allocations (llama/22526)
This introduces an optional feature to allocate large GPU buffers (≥ 1GB)
using USM system allocations if supported by the device. It allows using
buffers from the system allocator then letting the system manage memory
migrations between host and device as necessary.

This feature is disabled by default and requires the GGML_SYCL_USM_SYSTEM
environment variable to enable. If USM system allocations are not supported
by the device or the system, we fallback to regular allocations.

This feature can allow VRAM overcommit. For example, the test below fails
on B580 due to lack of memory for allocation, but it passes when enabling
USM system allocations:

  ./examples/sycl/test.sh -m Qwen3.5-27B-Q3_K_M.gguf -lv 4

Signed-off-by: Francois Dugast <francois.dugast@intel.com>
2026-06-19 12:53:43 +03:00
Winston Ma 694579182f vulkan: prefer host-visible memory buffers on UMA devices (llama/22930)
* implement UMA host-visible memory

* update based on 0cc4m's suggestion
2026-06-19 12:53:43 +03:00
Jeff Bolz 93c02083bd vulkan: Support gated_delta_net with S_v=16 (llama/24581) 2026-06-19 12:53:43 +03:00
Frosty40 fa204c82f4 sycl: support reordered Q4_K/Q5_K/Q6_K MoE MUL_MAT_ID (llama/24452)
* sycl: support reordered Q4_K and Q5_K MoE MUL_MAT_ID

Extend reordered-weight handling to fused MoE MUL_MAT_ID for Q4_K and Q5_K expert tensors and add Q5_K reordered DMMV coverage. Unsupported 3D reorder cases now fall back instead of aborting.

* sycl: extend MoE reorder to Q6_K mul_mat_id
2026-06-19 12:53:43 +03:00
Neo Zhang 79f88a1104 Support OP EXPM1, support all UT cases of FLOOR, TRUNC, ROUND (llama/24363)
* support OP EXPM1, support all UT cases of FLOOR, TRUNC, ROUND

* fix conflict

* rebase, support new UT case of repeat, concat
2026-06-19 12:53:43 +03:00
Pascal c8f370a460 vulkan: add col2im_1d op (llama/24425)
* vulkan: add GGML_OP_COL2IM_1D, follow-up to the CPU op

* vulkan: col2im_1d bounded gather loop instead of full-K scan with modulo

* vulkan: col2im_1d address review from @jeffbolznv

* vulkan: col2im_1d return nullptr for unsupported types, address review from @0cc4m
2026-06-19 12:53:43 +03:00
Jeff Bolz d77b2f704c vulkan: support more CONCAT types (llama/24579) 2026-06-19 12:53:43 +03:00
Andrei bd3912b0a8 wasm : fix fallback symbol collision (llama/24639) 2026-06-19 12:53:43 +03:00
Katostrofik 1fd857e2f9 SYCL: use native subgroup size for K-quant DMMV (llama/21700) 2026-06-19 12:53:43 +03:00
someoneinjd 96051f04a7 sycl: fix soft_max_f32 max reduction (llama/24451) 2026-06-19 12:53:43 +03:00
Neo Zhang e958dcead1 sycl : fix reorder function; add fp32/fp16 in build script (llama/24578) 2026-06-19 12:53:43 +03:00
Neo Zhang d20057908a sycl : enhance set_rows to support q1_0, mxfp4, nvfp4 (llama/24564) 2026-06-19 12:53:43 +03:00
Neo Zhang 3cb087c42a add to support pool_1d, move pool_1d/2d code to pool.cpp/hpp (llama/24584)
* add to support pool_1d, move pool_1d/2d code to pool.cpp/hpp

* update ops.md
2026-06-19 12:53:43 +03:00
Alexey Kopytko 5832e734d4 Remove per-allocation Level Zero runtime checks (llama/23399)
* [SYCL] Centralize Level Zero detection in ggml_sycl_init

* use the same wording

* get back the warning

* [SYCL] Remove per-allocation getenv() for GGML_SYCL_ENABLE_LEVEL_ZERO

* bring back the comment

* move it up to make sure devices call the shots

* move the env detection early

* replace g_ggml_sycl_enable_level_zero with a direct call to .ext_oneapi_level_zero

* update the comment

* switch back to g_ggml_sycl_enable_level_zero with a sentinel

* remove the check

* Reduce the diff

* reword, move lower

* move things aroudn

* remove forward declaration if favor of a full replace

* pre-cache results of zeDeviceGetProperties

* put ggml_sycl_get_env back

* replace get_sycl_env with ggml_sycl_get_env

* add whitespace back

* Apply suggestion from @sanmai
2026-06-19 12:53:43 +03:00
Georgi Gerganov 1b2d6d2c23 metal : add repeat bf16 (llama/24638) 2026-06-19 12:53:43 +03:00
leonardHONG 7349e5ae11 CUDA: only support F32/F16 for GGML_OP_REPEAT (llama/24533) 2026-06-19 12:53:43 +03:00
Masashi Yoshimura 3e0b917514 ggml-webgpu: improve i-quants mul_mat performance and speed up prefill (llama/24530)
* Improve prefill speeds for i-quants

* Fix #if defined() usage in preprocessor guards.
2026-06-19 12:53:43 +03:00
Jeff Bolz 1216e0957b vulkan: support non-contig unary/glu ops (llama/24215)
* vulkan: support non-contig unary/glu ops

Change unary/glu ops to pass in all strides and use fastdiv for the index
calculation. Put all unary ops in one file, similar to glu, to share the
code. codex went ahead and added expm1 without me asking, but I had to
make it do a real precision analysis rather than just making stuff up.

unary.comp initially couldn't use generic_unary_head because there wasn't
space for xielu's additional constants. Fixing this required packing the
fastdiv 'L' values.

* attempt to workaround compiler bug

* resolve conflict from #23991

* use expm1
2026-06-19 12:53:43 +03:00
Ruben Ortlam dc195118ef vulkan: add pipeline barriers for memcpy read operations (llama/23770)
* vulkan: add pipeline barriers for memcpy read/write operations

* remove unnecessary host write pipeline barriers
2026-06-19 12:53:43 +03:00
Georgi Gerganov f35f47b5d2 ggml : bump version to 0.15.1 (ggml/1541) 2026-06-15 10:33:53 +03:00
ZihaoMu 882736f886 ggml: support concat for scalar types at cuda backend (llama/24011)
* cuda: support concat for scalar types

* Update concat.cu

* fix metal ci issue
2026-06-15 10:33:53 +03:00
shaofeiqi 2dcfd49d59 opencl: add q5_0/q5_1 gemm and gemv kernels for Adreno (llama/24319)
* opencl: add q5_0 adreno support

* opencl: add q5_1 adreno support

* opencl: cosmetic fix

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-06-15 10:33:53 +03:00
Jeff Bolz afd559279c vulkan: ifdef eMesaHoneykrisp (build fix) (llama/24479)
Fixes build/CI after #24306.
2026-06-15 10:33:53 +03:00
Georgi Gerganov b04008fcec ggml : bump version to 0.15.0 (ggml/1539) 2026-06-15 10:33:53 +03:00
Winston Ma 6870cfd616 vulkan: add fast path for contiguous buffer transfers (llama/23973) 2026-06-15 10:33:53 +03:00
Kevin Liu a512e4c5c3 vulkan: use medium matmul tile on Asahi Linux (llama/24306)
* vulkan: use medium matmul tile on Asahi Linux

* vulkan: switch Apple detection to Honeykrisp driver id
2026-06-15 10:33:53 +03:00
Gaurav Garg 1a1900f90c Remove padding and multiple D2D copies for MTP (llama/24086)
* Make ggml_gated_delta_net take only the initial recurrent state (D, 1, n_seqs) and passes the snapshot count K as an op parameter instead of inferring it from state->ne[1].

Remove the padding hack and copy all emitted snapshots into the recurrent cache with a single strided ggml_cpy

* Make GDN changes in all backends. Address review comments.

* Fix CI build errors
2026-06-15 10:33:53 +03:00
Oliver Simons ef85b26d9f CUDA: Fix ssm_scan_f32 data-races (llama/24360)
* Add missing syncthreads before resuing cub_temp_storage

__syncthreads() is required before being allowed to resue TempStorage
smem:
https://nvidia.github.io/cccl/unstable/cub/api/classcub_1_1BlockLoad.html#_CPPv4I0EN3cub9BlockLoad4LoadEv20RandomAccessIteratorRA14ItemsPerThread_1Ti

* Add one more missing __syncthreads

Could also double-buffer, but alternative is to simply ensure all
threads have read smem* before writing to it again in the next loop
iteration

* Remove unused smem from ssm_scan_f32
2026-06-15 10:33:53 +03:00
Jeff Bolz dc794303d8 vulkan: reduce iq1 shared memory usage for mul_mm (llama/24287) 2026-06-15 10:33:53 +03:00
Ruben Ortlam 686bc802d1 vulkan: add `v_dot2_f32_f16` support in matrix-matrix multiplication and Flash Attention (llama/24123)
* vulkan: add support for valve fp16 dot2 extension

* use macro for dot2 path choice

* properly check for the feature

* add dot_product abstraction to reduce preprocessor branching
2026-06-15 10:33:53 +03:00
Pascal 28c7ed3db7 ggml : add GGML_OP_COL2IM_1D (llama/24206)
* cpu: add GGML_OP_COL2IM_1D

Add the overlap-add (scatter-add) step of a 1D transposed convolution.
A ConvTranspose1d factorizes as a GEMM followed by col2im: a weight
pre-permuted to [IC, K*OC] is contracted against the [IC, T_in] input
with mul_mat to produce a column matrix [K*OC, T_in], and col2im_1d
scatters those columns back into the [T_out, OC] signal, with
T_out = (T_in - 1)*s0 + K - 2*p0.

Keeping the contraction as a plain mul_mat leaves the heavy work on the
optimized (and quantizable) matmul kernels, so col2im_1d only does the
cheap overlap-add.

CPU uses a gather formulation parallelized over output channels,
supporting F32, F16 and BF16 with an F32 accumulator.

* tests: add backend coverage for GGML_OP_COL2IM_1D

Add test_col2im_1d next to the conv_transpose_1d cases, covering F32,
F16 and BF16 across eight geometries: the canonical kernel = 2*stride
DAC upsampling shape, overlap, no overlap, cropping (p0 = 1 and
p0 = stride/2), kernel < stride with zeroed gaps, kernel not a
multiple of stride, and a single column unfold.

Perf mode gets three real vocoder stage shapes reporting memory
bandwidth. max_nmse_err relaxes to 5e-4 for F16 and BF16.

* cpu: harden GGML_OP_COL2IM_1D

ggml_col2im_1d validates s0, oc, p0 and input contiguity at graph
build time, before the oc division, protecting every backend at once.
The kernel asserts the contiguity its flat indexing assumes and its
doc states the full output length including the crop term.

The kernel parallelizes over the time axis: the split stays balanced
down to OC = 1, where the previous channel split was single threaded.
Values are bit identical on the three real vocoder chains, two out of
three improve.

* tests: extend the GGML_OP_COL2IM_1D grid

The eval grid grows to eleven geometries: OC = 1 (mono output stage),
K = 1 with stride > 1 (sparse scatter, every gap position zeroed) and
a crop down to T_out = 2 where all the gather bounds act at once.

* tests: add col2im_1d equivalence test

tests/test-col2im-1d.cpp proves mul_mat + col2im_1d matches the
native ggml_conv_transpose_1d on the CPU backend, F32 bit exact, F16
and BF16 through casts of the column matrix. test-backend-ops cannot
cover this for a CPU only op since the CPU backend is its own
reference there.

* rpc: bump protocol patch version for GGML_OP_COL2IM_1D

GGML_OP_COUNT goes from 96 to 97 with the new op, which trips the
static_assert in ggml-rpc.h. Bump RPC_PROTO_PATCH_VERSION since the
op is appended and no existing op code shifts.
2026-06-15 10:33:53 +03:00
Yash Raj Pandey 2d68a3066f ggml-cpu : fix rms_norm_back wrong output under in-place aliasing (llama/24305)
* ggml-cpu : fix rms_norm_back wrong output under in-place aliasing

* cont : clean-up comment

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-06-15 10:33:53 +03:00
ravel7524 72894aa250 Remove case for GGML_TYPE_Q4_K in mvvq.cu (llama/23528) 2026-06-15 10:33:53 +03:00
Reese Levine e69e5138fe ggml-webgpu: Add clang-format job (llama/24308)
* Add clang-format job

* try local formatting
2026-06-15 10:33:53 +03:00
Masashi Yoshimura aa42b48312 ggml-webgpu: Improve prefill speeds for k-quants + refactor matmul for Q4/Q5/Q8 and k-quants (llama/24225)
* ggml-webgpu: Improve prefill speeds + refactor matmul for quants

* Fixes for editroconfig checker
2026-06-15 10:33:53 +03:00
Nikhil Jain 15e5d401d1 Handle buffer overlap / buffer aliasing for concat operator (llama/24000)
* Only run webgpu CI on my fork

* Add webgpu only workflow

* handle buffer overlap case for concat operator

* restore build-webgpu.yml

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* Run clang-format

* Update ggml/src/ggml-webgpu/wgsl-shaders/concat.wgsl

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>
2026-06-15 10:33:53 +03:00
Nikhil Jain 490e50056c Implement 2D workgroups for scale, binary, and unary ops (llama/24044)
* Only run webgpu CI on my fork

* Add webgpu only workflow

* Implement 2d workgroups for more operations

* fix

* Fix type

* Move back to global_invocation_id
2026-06-15 10:33:53 +03:00
Jeff Bolz fbf720dc9f vulkan: Use cm2 decode_vector for mul_mat_id B matrix loads (llama/23991)
This allows vec4 loads of the B elements. Also increase BK to 64 when this is
enabled. Neither of these alone is consistently faster, but together these give
a nice speedup.

In ggml-vulkan.cpp, we need to make sure the B matrix alignment and stride are
multiples of 4.
2026-06-15 10:33:53 +03:00
Ruben Ortlam 782f1226c8 cuda: reset cuda context after reading memory size (llama/23935)
* cuda: reset device in get_memory function if no backend is active

* also count device and host buffers

* exclude hip and musa from counting and device reset

* use device mutex instead of atomic

* undo backend_free function move
2026-06-15 10:33:53 +03:00
Georgi Gerganov b31466b4a1 ggml : bump version to 0.14.0 (ggml/1533) 2026-06-08 14:36:36 +03:00
Harkirat Gill 4669631d20 HIP: add gfx1152 and gfx1153 to RDNA3.5 (llama/24129) 2026-06-08 14:36:36 +03:00
Xuan-Son Nguyen 2c139c2e5e metal : fix im2col 1D case (audio models) (llama/24220) 2026-06-08 14:36:36 +03:00
Ruben Ortlam 1777deff4c vulkan: check coopmat2 features before reporting support (llama/24186) 2026-06-08 14:36:36 +03:00
lhez a87e950a06 opencl: improve get_rows, cpy, concat and q6_k flat gemv (llama/24160)
* opencl: allow multiple workgroups for large rows

* opencl: improve small cpy

* opencl: packed concat for small input

* opencl: tweak flat q6_K gemv, increase N_DST and remap threads
2026-06-08 14:36:36 +03:00
Ruben Ortlam 5a1feed8ca vulkan: add fwht support for Intel with shmem reduction (llama/23964)
* vulkan: add fwht support for Intel with shmem reduction

* don't use N as workgroup size

* disable subgroup shuffle on MoltenVK AMD

* disable fwht shader on Intel Windows due to driver bug
2026-06-08 14:36:36 +03:00
Charles Xu facb02c4c3 kleidiai : dynamic chunck-based scheduling for hybrid execution (llama/23819) 2026-06-08 14:36:36 +03:00
Oliver Simons 4fa1e0687e CUDA: enroll mul_mat_vec_q_moe into pdl (llama/24087)
* Enroll mul_mat_vec_q_moe into PDL, boosting MTP performance on BW

Data collected on a B4500:

Before
```
(llama.cpp) ➜  llama.cpp git:(master) ✗ python mtp-bench.py
  code_python        pred= 192 draft= 150 acc= 116 rate=0.773 tok/s=202.8
  code_cpp           pred= 192 draft= 147 acc= 117 rate=0.796 tok/s=212.8
  explain_concept    pred= 192 draft= 161 acc= 110 rate=0.683 tok/s=196.4
  summarize          pred= 192 draft= 138 acc= 122 rate=0.884 tok/s=226.6
  qa_factual         pred= 192 draft= 138 acc= 121 rate=0.877 tok/s=225.1
  translation        pred= 192 draft= 158 acc= 112 rate=0.709 tok/s=201.5
  creative_short     pred= 192 draft= 160 acc= 110 rate=0.688 tok/s=197.2
  stepwise_math      pred= 192 draft= 150 acc= 115 rate=0.767 tok/s=209.2
  long_code_review   pred= 192 draft= 148 acc= 116 rate=0.784 tok/s=208.9
```
After
```
(llama.cpp) ➜  llama.cpp git:(master) ✗ python mtp-bench.py
  code_python        pred= 192 draft= 150 acc= 116 rate=0.773 tok/s=211.9
  code_cpp           pred= 192 draft= 147 acc= 117 rate=0.796 tok/s=224.6
  explain_concept    pred= 192 draft= 161 acc= 110 rate=0.683 tok/s=207.8
  summarize          pred= 192 draft= 138 acc= 122 rate=0.884 tok/s=240.2
  qa_factual         pred= 192 draft= 138 acc= 121 rate=0.877 tok/s=238.5
  translation        pred= 192 draft= 158 acc= 112 rate=0.709 tok/s=213.4
  creative_short     pred= 192 draft= 160 acc= 110 rate=0.688 tok/s=208.8
  stepwise_math      pred= 192 draft= 150 acc= 115 rate=0.767 tok/s=221.7
  long_code_review   pred= 192 draft= 148 acc= 116 rate=0.784 tok/s=220.7
```

Server launched with:
```
➜  llama.cpp git:(osimons/enroll_mul_mat_vec_q_moe_into_PDL) ✗ ./build-x64-linux-gcc-reldbg/bin/llama-server \
    -m /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-MTP-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf -dio \
    --spec-type draft-mtp \
    --spec-draft-n-max 2 \
    -ngl all \
    -fa on \
    --host 0.0.0.0 \
    --port 8080 -np 1 --chat-template-kwargs "{\"preserve_thinking\": true}"
```

* LC to overlap with following kernels
2026-06-08 14:36:36 +03:00
Mason Milburn 4ecede8c8b sycl : port multi-column MMVQ from CUDA backend (llama/21845)
mmvq:

Port the ncols_dst optimization from ggml-cuda/mmvq.cu to SYCL.
Read weights once per dispatch instead of once per column.
Covers all standard quant types + reorder paths for Q4_0, Q8_0,
Q3_K, Q4_K, Q5_K, Q6_K. IQ types (except IQ4_XS) excluded due to
incompatible vec_dot signatures.

ggml-sycl:

The weight reorder was only bootstrapped on single-token mat-vec
(ne[1] == 1). Speculative / MTP verify issues only multi-column mat-vec,
so it never triggered the reorder and ran on the slower non-reorder
kernel. Bootstrap it on small multi-column batches (ne[1] <= 8) too.
2026-06-08 14:36:36 +03:00
Kartik Sirohi 991b5a8b4a ggml: vectorize ggml_vec_dot_q4_1_q8_1 with WASM SIMD128 (llama/22209)
* ggml: vectorize ggml_vec_dot_q4_1_q8_1 with WASM SIMD128

Optimize the inner loop of ggml_vec_dot_q4_1_q8_1_generic using
WASM SIMD128 intrinsics, gated behind #ifdef __wasm_simd128__ so
non-wasm builds are completely unaffected.

Approach:
- single wasm_v128_load covers all 32 packed 4-bit weights
- nibbles unpacked via AND/SHR into two u8x16 registers
- widened to i16 before multiply (WASM SIMD has no i8*i8 instruction)
- 4x wasm_i32x4_dot_i16x8 calls accumulate all 32 element pairs
- horizontal reduce via 4x wasm_i32x4_extract_lane

Benchmark (node v25, emcc -O3 -msimd128, 64 blocks x QK8_1=32,
200k iterations):

| impl   | ns/call | speedup |
|--------|---------|---------|
| scalar |   880.7 |   1.00x |
| simd   |   257.8 |   3.42x |

Correctness verified against scalar reference across 10 random seeds
with exact output match.

* ggml: move q4_1_q8_1 WASM SIMD implementation to wasm backend

Relocate the SIMD128 implementation of ggml_vec_dot_q4_1_q8_1 to ggml/src/ggml-cpu/arch/wasm/quants.c to follow architecture-specific layout. Restore the generic implementation in ggml/src/ggml-cpu/quants.c.
Move for loop in the else block.

* ggml: use generic q4_1_q8_1 fallback in wasm backend
2026-06-08 14:36:36 +03:00
Georgi Gerganov 9d6e561f69 metal : reduce rset heartbeat from 500ms -> 5ms (llama/24074) 2026-06-08 14:36:36 +03:00
Reese Levine e9dbd0c18a ggml-webgpu: FlashAttention refactor + standardize quantization support (llama/23834)
* Start work on flash_attn refactor

* Refactor

* Split k/v quantization

* Refactor and abstract quantization logic for flash_attn and mul_mat

* Add quantization support to tile path

* formatting

* Move to functions, add a check
2026-06-08 14:36:36 +03:00
rehan-10xengineer a1a3186887 ggml-cpu: extend RVV quantization vec dot to higher VLENs (llama/22754)
* ggml-cpu: add rvv 512b,1024b impls for iq4_xs

* ggml-cpu: refactor; add rvv 512b, 1024b impls for q6_K, i-quants

* ggml-cpu: refactor; add 512 and 1024 implementations of tq3_s, iq3_xxs, iq2_s, iq2_xs, iq2_xxs

improve iq2_xs impl for rvv 256

Co-authored-by: Rehan Qasim <rehan.qasim@10xengineers.ai>

---------

Co-authored-by: taimur-10x <taimur.ahmad@10xengineers.ai>
Co-authored-by: Rehan Qasim <rehan.qasim@10xengineers.ai>
2026-06-08 14:36:36 +03:00
Andreas Kieslinger 00a9728de3 Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030)
* Removes __restrict__ from PDL kernel headers due to incompatibility with
PDL. Adds preprocessor directives based on arch in kernel body to add
__restrict__ to retain performance on older architectures.

* Simplifies new __restrict__ usage via macro

* Add hopper to PDL __restrict__ fix.

Co-authored-by: Oliver Simons <osimons@nvidia.com>

---------

Co-authored-by: Oliver Simons <osimons@nvidia.com>
2026-06-08 14:36:36 +03:00
Charles Xu 750fa4ca35 ggml-cpu: use runtime SVE width in FWHT (llama/24059) 2026-06-08 14:36:36 +03:00
Aman Gupta d5a49ebec8 cuda: reserve space for quantize kv-cache at startup (llama/23907)
* cuda: reserve space for quantize kv-cache at startup

* address review comments

* remove forward decl

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

* remove assert in ggml-cuda.cu

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-06-08 14:36:36 +03:00
lhez f110ff540c opencl: use flat variants of q4_K and q6_K gemv for very large M (llama/24006) 2026-06-08 14:36:36 +03:00
Max Krasnyansky d31cb20b25 hexagon: profiler output fix and script updates (llama/24042)
* hex-ops: fix profiler output (ie remove the redundant NONEs)

* hex-prof: update profiling script to support tot.usec column
2026-06-08 14:36:36 +03:00
Max Krasnyansky 8d61a9edf0 hexagon: MUL_MAT, MUL_MAT_ID, FLASH_ATTN and GDN cleanup and optimizations for latest models (llama/23989)
* hex-mm: initial support for F32 * F32 -> F32 matmuls

* hex-rms-norm: fix src1 stride use in fused rms_norm_mul

* hex-ops: clear spad pointers in the ops that clober it

This fixes an odd case where fused rms-norm-mul was failing but only in qwen3.5-2B and only at searth op-bath sizes.

* hmx-mm: add support for F32 * F32 -> F32 matmul_2d on HMX

Decided to use Q4_0 * F32 -> F32 matmul for this.
Q4_0 gets dequantized and tiled into F16, and here we quantize and tile F32 into F16.
Super simple and pretty efficient.

* hmx-mm: route f16 2D matmuls through the same kernel used for all other types

* hmx-mm: re-introduce pipelined vs non-pipelined mode that we used to have but is much more generic way

This update futher improves matmul performance and at the same time removes most of the redudant logic
we had in different paths.

* hmx-fa: slighlty improved pipeline simimar to matmul updates

* hmx-mm: initial version of MAT_MUL_ID support for HMX

* hmx-mm: fixed mxfp4 handling for MUL_MAT_ID

* hex-gdn: optimize GATED_DELTA_NET

DMA prefetch/double-buff, vectorize everything with HVX, in other words -- the usual :)

* hmx-mm: missed one more case where we can use fastmod

* hexagon: update DCVS settings for a slight perf bump

* hmx-fa: use fastdiv in hmx-flash-attn

* hmx-fa: precompute slope values to avoid disrupting the inner loop

* hvx-utils/fa: new HVX helpers for powf and logf and using those to speed up FA alibi

* hex-ops: fixed a bug in fusion logic that was messing up the order of the src tensors when some srcs are empty

* hex-fa: correctly fallback to HVX if we have sinks or the dims are not quite right
2026-06-08 14:36:36 +03:00
Todor Boinovski 754247f28b hexagon: add gelu_quick (llama/24007) 2026-06-08 14:36:36 +03:00
Anav Prasad 79223704a1 clean up unused variables warnings (llama/23975) 2026-06-08 14:36:36 +03:00
lhez 9a0265d13b opencl: fix compiler warnings for non-adreno path (llama/23922)
* opencl: fix compiler warnings for non-adreno path

* opencl: fix const cast warning
2026-06-08 14:36:36 +03:00
Masashi Yoshimura db2a39507c revert to using global_invocation_id for cpy shader (llama/23955) 2026-06-08 14:36:36 +03:00
shaofeiqi e728bae159 opencl: add basic support for q5_0 and q5_1 (llama/23548)
* opencl: add general q5_0 support

* opencl: add general q5_1 support

* opencl: support non-uniform workgrp size

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-06-08 14:36:36 +03:00
Shrivas Shankar 050b8567a0 metal: template GLU kernels to support f16/f32 (llama/23882)
Drops the hardcoded f32 GLU kernels in favor of a single template. We now load/store in the native tensor type (half or float) to save memory bandwidth, but keep the actual ALU compute in float to avoid exploding math in geglu/swiglu. Also opened up the dispatch gate to allow f16 inputs.
2026-06-08 14:36:36 +03:00
Jeff Bolz 71d80aa49e vulkan: don't hold the device mutex while compiling pipelines (llama/23641)
* vulkan: don't hold the device mutex while compiling pipelines

We need to hold a lock while we traverse all pipelines and lazily initialize
them, but we don't need to hold it while the pipeline is being compiled. And
it doesn't need to be the same lock as the device mutex. We call load_shaders
each time a pipeline is needed, so we only need to compile that one pipeline
(and, for example, don't want to end up compiling a pipeline that another
thread should be compiling).

* remove 'needed'
2026-06-08 14:36:36 +03:00
Winston Ma c471bcce1b vulkan: reduce host memory lock contention (llama/23376)
* vulkan: reduces lock contention

* replace unique_lock with lock_guard
2026-06-08 14:36:36 +03:00
Johannes Gäßler e815b264eb TP: quantized KV cache support (llama/23792)
* TP: quantized KV cache support

* fix partial view

* remove overly strict assert
2026-06-08 14:36:36 +03:00
Matt Corallo 982533fc0c vulkan: Block-load Q3_K/Q6_K block data and subtract on 32b ints (llama/23056)
Q2_K/Q3_K/Q6_K do much better when using MMVQ on Intel BMG even
though they're only 2-byte aligned, and Q3_K still wins on
NVIDIA as well.

mesa isn't all that great at coalescing back-to-back loads from
alternating arrays, so we force it instead. Further, we can do
subtraction directly on a full int32_t rather than an i8vec4
with bit twiddling because the high bit is always free to start.

On Intel BMG on mesa, the switch to MMVQ provides an immediate
~57% perf increase in tg128 for unsloth/Qwen3.5-9B-GGUF:Q3_K and
~78% perf increase in tg128 for unsloth/Qwen3.5-9B-GGUF:Q6_K.

The futher switch to block loads leads to a ~24% perf increase in
tg128 for unsloth/Qwen3.5-9B-GGUF:Q3_K and a ~48% perf increase in
tg128 for unsloth/Qwen3.5-9B-GGUF:Q6_K.

Finally, Xe2 wins on MMVQ even for small k, so we take the NVIDIA
override for K quants on Xe2 as well.
2026-06-08 14:36:36 +03:00
Winston Ma aea93ada61 vulkan: Removed unused functions (llama/23175) 2026-06-08 14:36:36 +03:00
Neo Zhang ec0c661950 Support Q4_1, Q5_0, Q5_1 in Flash-attention (llama/23812)
* support Q4_1, Q5_0, Q5_1

* update ut case
2026-06-08 14:36:36 +03:00
Neo Zhang 20323e48c4 Add more types in GET_ROWS OP (llama/23710)
* add to support Q1_0, NVFP4, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ1_S, IQ1_M, IQ3_S, IQ4_NL, IQ4_XS, I32, MXFP4, Q2_K, Q3_K, Q5_K, and Q6_K in GET_ROWS OP

* correct the link
2026-06-08 14:36:36 +03:00
Neo Zhang 687fbcb149 sycl : Optimize Q3_K mul_mat by reorder (llama/23725) 2026-06-08 14:36:36 +03:00
lhez 1c0d1f0f7c opencl: support bf16 by converting to f16 (llama/23839) 2026-06-08 14:36:36 +03:00
Georgi Gerganov bf74b557d2 metal : restore im2col implementation for large kernels (llama/23901) 2026-06-08 14:36:36 +03:00
Jinyang He 64b0d6b7fc ggml : add some lsx support (llama/23798)
* loongarch : optimize LSX fp16 load/store with native intrinsics

Use __lsx_vfcvtl_s_h and __lsx_vfcvt_h_s instead of scalar loops in
__lsx_f16x4_load and __lsx_f16x4_store.

* loongarch : add LSX implementation for q8_0 dot product

* loongarch : add LSX implementation for q6_K dot product

* loongarch : add LSX implementation for iq4_xs dot product

* Improve reduce ops when sun int16 pairs to int32
2026-06-08 14:36:36 +03:00
Ruben Ortlam 4317ddbe2b vulkan: add Flash Attention support for BFloat16 KV cache (llama/23420)
* vulkan: add flash attention bf16 kv support

* vulkan: bf16 FA coopmat1 support

* vulkan: bf16 FA coopmat2 support

* fix FA bf16 f32 fallback

* fix FA bf16 coopmat1 shader

* fix FA bf16 coopmat2 shader

* code cleanup

* cleanup comment change

* address feedback

* add O_TYPE for cm2 FA

* use O_TYPE for gqaStore function

* reduce BFLOAT16 ifdefs
2026-06-08 14:36:36 +03:00
Reese Levine 9147a9676b ggml-webgpu: Check earlier for WebGPU required features (llama/23879) 2026-06-08 14:36:36 +03:00
Reese Levine acd91d2c38 ggml-webgpu: add q4_0/q8_0 SET_ROWS (llama/23760)
* Add q8_0 and q4_0 set_rows

* Add fast(er) quantization set_rows path

* formatting/naming

* a little more naming

* Remove unused constant

* Don't override other override

* Avoid bitcast

* Narrow relaxation
2026-06-08 14:36:36 +03:00
Oliver Simons f7aad4ed7e CUDA: Check PTX version on host side to guard PDL dispatch (llama/23530)
* CUDA: Check PTX version on host side to guard PDL dispatch

Checking on `__CUDA_ARCH_LIST__` alone is insufficient for JIT, as this
variable doesn't differentiate between compiling for say sm_90, sm_90a
or sm_90f (so forward-jittable PTX vs. arch/family-specific PTX).

Thus, one can have a bug when compiling with
`DCMAKE_CUDA_ARCHITECTURES="89;90a"`, where current code would wrongly
dispatch to PDL on sm_90/sm_120 in forward-JIT mode.

This PR fixes this issue by checking `cudaFuncAttributes::ptxVersion` of
the incoming kernel at runtime. A check on ptxVersion alone is
sufficient, as device-codes will always be >= ptxVersion (and any
violation of this would be a severe bug in CUDA/nvcc), see:
 https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/#gpu-code-code-code

* Implement MurmurHash3 mixer for better hash distribution

Magic constants were taken from boost:
2698b43803/include/boost/container_hash/detail/hash_mix.hpp (L19-L65)

* Update ggml/src/ggml-cuda/common.cuh

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

* Address review comments, make seed non-zero

* Apply code-formatting

* Replace std::size_t -> size_t for consistency

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-06-08 14:36:36 +03:00
fairydreaming c50e951afd model : support for DeepseekV32ForCausalLM with generic DeepSeek Sparse Attention (DSA) implementation (llama/23346)
* llama : support DeepSeek V3.2 model family (with DSA lightning indexer)

* convert : handle DeepseekV32ForCausalLM architecture

* ggml : support for f16 GGML_OP_FILL

* memory : separate hparams argument in llama_kv_cache constructor

* memory : add llama_kv_cache_dsa memory (KV cache + lightning indexer cache)

* llama : support for LLM_ARCH_DEEPSEEK32

* model : llama_model_deepseek32 implementation

* model : merge two scale operations into one in DSA lightning indexer implementation

* chore : remove unused code

* model : support NVFP4 in DeepSeek V3.2

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

* memory : refactoring TODO

Co-authored-by: ggerganov <ggerganov@users.noreply.github.com>

---------

Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
Co-authored-by: ggerganov <ggerganov@users.noreply.github.com>
2026-06-08 14:36:36 +03:00
Georgi Gerganov 92fc3f2a58 ggml : bump version to 0.13.1 (ggml/1523) 2026-05-29 09:47:30 +03:00
Andreas Kieslinger e90501e179 cuda : disables launch_fattn PDL enrollment due to compiler bug (llama/23825) 2026-05-29 09:47:30 +03:00
Matt Corallo f1b687da28 meta : Add missing `buffer` set in allreduce fallback !COMPUTE clear (llama/23480)
Without this at least the vulkan backend will skip the `* 0` for
!COMPUTE tensors, causing corrupt output.
2026-05-29 09:47:30 +03:00
Max Krasnyansky 442be1789d hexagon: basic/generic op fusion support and RMS_NORM+MUL fusion (llama/23835)
Updating infra to enable op fusion and using RMS_NORM+MUL as the use-case.
2026-05-29 09:47:30 +03:00
lhez 94922ce12c opencl: move backend info printing into its own function (llama/23702)
* opencl: move backend info print into its own function

* opencl: move new log line

* opencl: fix for non adreno path
2026-05-29 09:47:30 +03:00
fl0rianr e1faa7cb4d ggml: auto apply iGPU flag CUDA/HIP if integrated device (llama/23007) 2026-05-29 09:47:30 +03:00
redfox 4e8af441e5 mmvq Optim: add MMVQ_PARAMETERS_TURING(mmvq_parameter_table_id) for … (#23729)
* mmvq Optim:  add MMVQ_PARAMETERS_TURING(mmvq_parameter_table_id) for SM75 TURING

* avoid a mismatch for JIT compilation of Turing device code for Ampere or newer

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

---------

Co-authored-by: Copilot <copilot@github.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-05-29 09:47:30 +03:00
Jaden_Mach 04795e6272 CUDA: route batch>=4 quantized matmul to MMQ on AMD MFMA hardware (llama/23227)
* CUDA: per-quant MMVQ/MMQ batch threshold on AMD MFMA hardware

The dispatcher uses a single global threshold (MMVQ_MAX_BATCH_SIZE = 8)
to choose between mul_mat_vec_q (per-row GEMV) and mul_mat_q (MFMA-tiled
GEMM) for quantized matmul. On AMD CDNA, the optimal crossover differs
substantially by quant family because the per-row GEMV cost is dominated
by dequantisation, not the dot-product itself: K-quants pay a heavier
super-block decode and so MMQ wins sooner; legacy and IQ quants have
lean decode and stay ahead until the batch fully populates an MFMA tile.

This patch introduces ggml_cuda_should_use_mmvq(type, cc, ne11) -> bool,
mirroring the existing ggml_cuda_should_use_mmq, and gates per-quant
thresholds on amd_mfma_available(cc):

  Q3_K, Q4_K, Q5_K  : MMVQ <= 3   (MMQ wins from batch=4: +5% .. +76%)
  Q2_K, Q6_K        : MMVQ <= 5   (MMQ wins from batch=6: +8% .. +35%)
  others            : MMVQ <= 8   (legacy & IQ regress under MMQ; unchanged)

Non-AMD-MFMA paths (NVIDIA, RDNA, CDNA1 without MFMA) are byte-identical
to master. GGML_CUDA_FORCE_MMVQ=1 restores the original global threshold
for A/B testing.

Measured on MI250X (gfx90a, ROCm 7.2.1) with Llama-3.2-3B-Instruct,
llama-bench pp512 across all 20 supported quants, ubatch 1..8, 10 reps.
Full table in PR description.

  Selected pp512 throughput (tok/s, ub=8):
    Q4_K_S:  559 -> 940  (+68%)
    Q5_K_S:  503 -> 884  (+76%)
    Q3_K_S:  629 -> 879  (+40%)
    Q2_K  :  615 -> 809  (+32%)
    Q6_K  :  582 -> 776  (+33%)

  Selected pp512 throughput (tok/s, ub=4):
    Q4_K_S:  444 -> 480  (+ 8%)
    Q4_0  :  682 -> 685  (+ 0%)   (no regression - retains MMVQ)
    IQ4_XS:  706 -> 698  (- 1%)   (no regression - retains MMVQ)

* CUDA: address review — inline MMVQ batch table, drop env hatch & doc block

* tune kernel selection logic for CDNA1

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-05-29 09:47:30 +03:00
Max Krasnyansky 1b241b879c hexagon: minor refresh for HMX FA and MM (llama/23796)
* hex-fa: clean up qf32/fp32 handling and stride handling

* hex-fa: fix corner case fp NAN issues that were cause bad output from gemma4 on v79

* hex-fa: vectorize leftover handling

* hex-fa: avoid HVX fallback during token gen HMX has more FP16 compute capacity

* hmx-mm: remove dead code

* hmx-mm: use fastdiv in x4x2 dequant

* hmx-mm: sandwich dequant and scatter to improve perf

* hmx-mm: fixed rebase conflicts

* hmx-mm: further improve weight dequant by doing early type dispatch and precomputing fastdiv

* hmx-mm: an even earlier dispatch for per-type dequant

* hmx-mm: dequant linear types like q4_0 and q4_1 without the LUTs

This is a bit faster than LUT.

* hex-cmake: one more tweak for lto

---------

Co-authored-by: Trivikram Reddy <tamarnat@qti.qualcomm.com>
2026-05-29 09:47:30 +03:00
Jeff Bolz b896e91f18 vulkan: fast path for walsh-hadamard transform (llama/23687)
* vulkan: fast path for walsh-hadamard transform

* disable for intel due to segfault
2026-05-29 09:47:30 +03:00
Winston Ma 816c3029bc vulkan: fix wrong index variable in inner loop (llama/23665) 2026-05-29 09:47:30 +03:00
Winston Ma 5db94bac04 vulkan: Fix memory logger unsafe iterator access (llama/23667) 2026-05-29 09:47:30 +03:00
fairydreaming 60e420ff6a cuda : fix KQ mask offset integer overflow in fattn MMA kernel (llama/23610)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-05-29 09:47:30 +03:00
Martin Klacer 8e40325876 ggml: fixed Arm SVE usage bug in vec.h, vec.cpp (llama/22841)
* Updated vec.h/vec.cpp code to accumulate to F32 rather than F16

Change-Id: I0cb789347f2bf60ffaf9047319f727e788c825f8

Signed-off-by: Martin Klacer <martin.klacer@arm.com>
Co-authored-by: Milos Puzovic <Milos.Puzovic@arm.com>
2026-05-29 09:47:30 +03:00
ymcki d284e1c3aa Hexagon: OP_GATED_DELTA_NET K>1 support (llama/23531)
* K>1 state snapshot support

* removed picky indent multiple of 4 fixes
2026-05-29 09:47:30 +03:00
ymcki 7e843a80e1 opencl: OP_GATED_DELTA_NET (llama/23312)
* OP_GATED_DELTA_NET impl

* add back lanes_per_column declaration

* removed has_subgroup_arithmetic and has_subgroup_clustered_reduce

* removed trailing spaces and fixes indentation. Hard coded subgroup size for Adreno and Intel. Return not supported when K>1 state snapshot

* support for K>1 state snapshot

* removed picky indent multiple of 4 fixes

* removed return that won\'t be executed
2026-05-29 09:47:30 +03:00
Reese Levine 8c8f213dac ggml-webgpu: remove legacy constants (llama/23672) 2026-05-29 09:47:30 +03:00
Max Krasnyansky 3bbe93378c hexagon: add support for Q4_1 in MUL_MAT and MUL_MAT_ID (llama/23647)
* hex-mm: add support for Q4_1 matmul/matvec, hvx-only for now

* hmx-mm: add support for Q4_1

* hex-mm: use Q8_1 dynamic quantization to avoid having to compute sums in the vec_dot

* hexagon: fix repack scratch buffer overflow

* hex-mm: fix Q4_1 repack buffer sizing

* hexagon: flip the build order for mm and fa (seems to help LTO)

* hex-mm: add vec_dot 4x1s and minor HMX cleanup after adding Q4_1

* hex-mm: fix fp16 vec_dot fallback to 2x1 and another issue that could cause incorrect output

* hexagon: resurrect early-wake and add support for polling for op-batch completions

With Q4_1 ggml-hexagon now claims pretty much the entire graphs which gives the CPU more time to chilax.
This is a good thing! But it does add extra latency for the pure benchmark runs.
Early wakeup helps recover the latency a bit in the normals runs and op-batch polling is just for benchmarking.

---------

Co-authored-by: Todor Boinovski <todorb@qti.qualcomm.com>
2026-05-29 09:47:30 +03:00
Masashi Yoshimura a52bd385d6 ggml-webgpu: Fix how to dispatch WG to some ops (llama/23750) 2026-05-29 09:47:30 +03:00
Matt Corallo 8bce478ee8 vulkan: Switch MUL_MAT_VEC to 4 K per iteration for F16/32 (llama/22887)
* vulkan: Switch MUL_MAT_VEC to 4 K per iteration for F16/32

Against mesa git, this shows a 4.8% performance improvement for
tg128 on Qwen3.5-9B:BF16 on Intel BMG.

Note that this breaks some tests until the last commit which fixes
OOB A reads.

* vulkan: Use aligned loads in mul_mat_vec when available

Against mesa git, this shows a 3.3% performance improvement for
tg128 on Qwen3.5-9B:BF16 on Intel BMG.

* Make explicit that `num_rows` is <= `NUM_ROWS` in mul_mat_vec

Mesa's UUB logic can't see through conditionals, limiting its
ability to understand the bounds on the `num_rows` field in the
cleanup run. Making it explicit that `num_rows` is, indeed, always
<= `NUM_ROWS` helps mesa make slightly better codegen.

Against mesa git, this currently shows a 1% performance improvement
in tg128 on Qwen3.5-9B:BF16 on Intel BMG.

* vulkan: Fix OOB A reads in MUL_MAT_VEC for odd sizes

There was a TODO to fix the OOB reads from the A matrix which we do
here.

It is within performance noise (+<0.1%) in tg128 for
Qwen3.5-9B:BF16 on Intel BMG.
2026-05-29 09:47:30 +03:00
Jeff Bolz 1b590bbb9a vulkan: use GL_NV_cooperative_matrix_decode_vector for faster matmul (llama/23541) 2026-05-29 09:47:30 +03:00
l8bloom c5cde8c717 vulkan: add REPEAT op support for f16 to f16. (llama/23298)
* feat: extend repeat op for vulkan

* feat: add repeat_f16 vulkan pipeline

* fix: ensure same dst and src types

* fix: use type_size instead of data types

* fix: use int16 and int32 for repeat shader op

* chore: rename repeat_f* to repeat_i*

* chore: rename repeat vulkan pipelines
2026-05-29 09:47:30 +03:00
Oliver Simons 98c6722fec CUDA: restrict PDL to CTK >= 12.3 due to MSVC issues (llama/23742) 2026-05-29 09:47:30 +03:00
Winston Ma 80e87ec453 vulkan: avoid preferring transfer queue on AMD UMA devices (llama/22455) 2026-05-29 09:47:30 +03:00
Vladislav 6a249cd640 ggml-zendnn : fixed naming of matmul function (llama/20964)
* ggml-zendnn: fixed naming of matmul function

* ggml-zendnn: fixed naming of mul_mat_id function

* ggml-zendnn: fixed print in  mul_mat_id

---------

Co-authored-by: plotnikov.v10 <plotnikov.v10@wb.ru>
2026-05-29 09:47:30 +03:00
Jeff Bolz a0efd13f0f vulkan: optimize conv2d and implement coopmat1 support (llama/22620)
* vulkan: add CONV_SHAPE_64x128 for medium-K conv2d

* vulkan: skip conv2d bounds checks when shapes align with tile sizes

* vulkan: use WG_SIZE=128 for CONV_SHAPE_64x32 conv2d

* vulkan: stage cm2 conv2d accumulator through shmem before global store

* vulkan: add coopmat1 conv2d path

* fallback when using too much shared memory. clean up comments

* Require 16x16x16 and subgroup size 32 or 64

* check whether shared memory is sufficient before overwriting conv2d params with coopmat1 values
2026-05-29 09:47:30 +03:00
Max Krasnyansky f8df28d331 hexagon: add support for CONCAT op (llama/23648)
* hexagon: add support for CONCAT with optimized concat_2d_transposed

qwen3.5 models are quite heavy on the CONCAT with large and transposed src1.

* hex-concat: use fastdiv in generic version

* hex-concat: make checks for transposed a bit more readable

* hex-concat: reoder dma ops for better pipelining

* hex-cont/cpy: optimize CPY and CONT ops

The primary change is to avoid scalar divs in the inner loops.
We were calling hvx_copy_uu(... type_size) where type_size is non a constexpr.
This causes runtime divs by that value which is normally just 4 or 2 (f32/f16).

* hex-get-rows: optimize GET_ROWS for large rows

We now use DMA for larger rows and also split them into chunks to improve perf for Qwen3.5 and other models
that do lots of GET_ROWS with huge (2MB+ rows).

Also bump the DMA queue depth now that we can take advantage of it.

* hex-concat: unroll the inner loops of concat_2d

* hex-concat: more updates to concat_2d to improve perf a bit further

* hex-cpy: fixed n_rows per thread checks in the copy ops

* hmx-fa: fix alignment issues while computing dma sizes

* hex-set-rows: add early returns for idle threads

* hvx-rope: minor optimization to replace loops with fastdiv logic

* hex-rope: replace scalar tail processing with HVX

* hex-rope: optimize rope cache init with HVX

Add hvx-utils sin/cos helpers that use an aprox method (similar to rsqrt, inverse, etc)
Use the helpers to optimize ROPE.
2026-05-29 09:47:30 +03:00
Alexey Kopytko 049f0af339 SYCL: implement ggml_sycl_pool_vmm (llama/22862)
* SYCL: implement ggml_sycl_pool_vmm

* Add an option to bypass VMM with GGML_SYCL_DISABLE_VMM

* Clean up debugging logging

* document GGML_SYCL_DISABLE_VMM

* Multi-stream MoE optimization

* Revert "Multi-stream MoE optimization"

This reverts commit 938929c3f13a562ec67c59e87cc5d38595444cce.

* Update common.hpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Flip GGML_SYCL_DISABLE_VMM to GGML_SYCL_ENABLE_VMM

* add logging for GGML_SYCL_ENABLE_VMM when extension is not available (SYCL_EXT_ONEAPI_VIRTUAL_MEM macro)

* Apply suggestions from code review

Co-authored-by: Alexey Kopytko <alexey@kopytko.com>

* Apply suggestion from @sanmai

* Apply suggestion from @sanmai

---------

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
2026-05-29 09:47:30 +03:00
Masashi Yoshimura 00a5110b19 ggml-webgpu: Add MMVQ path for Q4/Q8/Q2_K/Q4_K and clean up legacy MUL_MAT pipeline (llama/23594)
* ggml-webgpu: Add MMVQ path for Q4/Q8/Q2_K/Q4_K

* Fix to editorconfig checking pass

* Remove mul-mat-legacy pipeline

* Fix to use vendor name as is and add dot_product/vendor to shader_lib_ctx
2026-05-29 09:47:30 +03:00
Nikhil Jain bc77933c2d Check batch_compute_passes before sending passes when not doing GPU profiling (llama/23457)
* Only run webgpu CI on my fork

* Add webgpu only workflow

* refactor batch_compute_passes to a per-thread variable, and submit individual passes when it is set to false and no GPU profiling is enabled

* restore build.yml
2026-05-29 09:47:30 +03:00
Johannes Gäßler 2307712d32 CUDA: missing PDL sync for FWHT, better fallback (llama/23690) 2026-05-29 09:47:30 +03:00
forforever73 1c477d4056 metal : add apple device id (llama/23566)
Co-authored-by: lvyichen <lvyichen@stepfun.com>
2026-05-29 09:47:30 +03:00
Aman Gupta 205ee5a189 CUDA: add fast walsh-hadamard transform (llama/23615)
* CUDA: add fast walsh-hadamard transform

* review: add unrolls + change size_t -> int

* warp size 64

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2026-05-29 09:47:30 +03:00
Georgi Gerganov 1cf8e3a903
ggml : bump version to 0.13.0 (ggml/1510) 2026-05-25 12:44:04 +03:00
Johannes Gäßler bcff515150
TP: fix ggml context size calculation (llama/22616)
* TP: fix ggml context size calculation, memory leak

* move split state cache back into the context

* revert to constant ggml context size for cgraphs

* increase headroom for statically allocated tensors

* remove obsolete include
2026-05-25 12:44:04 +03:00
Gilad S 2979e5f95f
ggml: `gguf_init_from_callback` and `gguf_init_from_buffer` (llama/22341)
* ggml: implement `gguf_init_from_buffer`

* test: `gguf_init_from_buffer`

* fix: memory breakdown for a model loaded with `no_alloc` from a file is consistent with being loaded from a buffer

* fix: use `GGML_UNUSED`

Co-authored-by: Copilot <copilot@github.com>

* fix: remove `total_size` from `gguf_reader`

* fix: file offset calculation, rename `offset` to `data_offset`

Co-authored-by: Copilot <copilot@github.com>

* refactor: extract model loader bug fixes to another PR

* feat: add `gguf_init_from_callback`

* fix: always require a max expected size

* fix: change `gguf_reader_callback_t`'s `output` type to `void *`, change `max_expected_size` and offsets to `uint64_t`

* fix: harden against offset overflow in buffer read

* fix: remove seek behavior from the callback

* feat: `max_chunk_read == 0` means `SIZE_MAX`

* fix: seeking in a gguf file with no tensors

---------

Co-authored-by: Copilot <copilot@github.com>
2026-05-25 12:44:04 +03:00
Georgi Gerganov 946d6813b9 ggml : bump version to 0.12.1 (ggml/1508) 2026-05-25 12:26:07 +03:00
Jeff Bolz a369b3949c ggml : Parallelize quant LUT init (llama/23595)
- Use OpenMP to parallelize iq2xs_init_impl and iq3xs_init_impl.
- Move the OpenMP detection from ggml-cpu to ggml-base.
- Update OpenMP dependencies in ggml-config.cmake.in.
2026-05-25 12:26:07 +03:00
Johannes Gäßler 3306af62b1 TP: fix entirely zero-sized slices per device (llama/23525) 2026-05-25 12:26:07 +03:00
shaofeiqi 1435988ab3 opencl: batch profiling to improve speed and prevent memory leaks (llama/23495) 2026-05-25 12:26:07 +03:00
Yiwei Shao b84d03487c hexagon: apply repl optimization in flash attn softmax as #22993 (llama/23455) 2026-05-25 12:26:07 +03:00
dskwe 511f8602b1 ggml : Check the right iface method before using the fallback 2d get (llama/23514) 2026-05-25 12:26:07 +03:00
Jeff Bolz 6b85d73b33 vulkan: fix windows find_package of SPIRV-Headers (llama/23215)
* vulkan: fix windows find_package of SPIRV-Headers

* not windows-only
2026-05-25 12:26:07 +03:00
Shawn Gu aefffa1fa5 opencl: generalize Adreno MoE kernels on M (llama/23449) 2026-05-25 12:26:07 +03:00
Alexey Kopytko b0c9f90059 SYCL: improve MoE prefill throughput (llama/23142)
- change `k_copy_src1_to_contiguous` so that uses a precomputed contiguous mapping where all rows "owned" by an expert are in one slice with a know starts and ends
- switch the `O(n_as * n_routed_rows)` contraption to a counting sort-based procedure with `O(n_as + n_routed_rows)` complexity
2026-05-25 12:26:07 +03:00
Alexey Kopytko 21c65a78a3 sycl : Level Zero detection in ggml_sycl_init (llama/23097)
* [SYCL] Centralize Level Zero detection in ggml_sycl_init

* use the same wording

* get back the warning
2026-05-25 12:26:07 +03:00
karavayev 0416feecfc SYCL : gated_delta_net K>1 (llama/23174)
* sycl_gated_delta_net K>1

* editor_config
2026-05-25 12:26:07 +03:00
Katostrofik 6fb7f1af2c SYCL: add BF16 to DMMV kernel path (~4x tg speedup on Intel Arc) (llama/21580)
* SYCL: add BF16 to DMMV kernel path for ~4x token generation speedup

BF16 models had no dedicated token generation kernel — they fell through
to the generic full-GEMM path, resulting in ~14% memory bandwidth
utilization on Intel Arc GPUs. This adds BF16 support to the DMMV
(dequantize mul-mat-vec) path, matching the existing F16 implementation.

Fixes #20478

* SYCL: fix BF16 DMMV out-of-bounds when ncols % 64 != 0

The qk=1 kernel (used for F16 and BF16) iterates with stride
2*GGML_SYCL_DMMV_X (= 64 on Intel targets where WARP_SIZE=16). When
ncols is a multiple of DMMV_X (32) but not of 2*DMMV_X (64), the last
warp iteration accesses elements at col >= ncols, producing NaN for the
final row and wrong values for interior rows.

Fix: tighten can_use_dequantize_mul_mat_vec to require ne[0] %
(2*DMMV_X) == 0 for F16/BF16 types, and update the ASSERT in the BF16
launcher to match. Quantized types use block-structured kernels with
different access patterns and keep the existing DMMV_X check.

Verified: test-backend-ops MUL_MAT passes 913/913 on Intel Arc Pro B70.
Previously failing: m=128/129 n=1 k=1056 cases (NaN and ERR > 0.0005).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-25 12:26:07 +03:00
Sachin Sharma 2d629533a5 ggml-zendnn : add Q8_0 quantization support (llama/23414)
* ggml-zendnn : add Q8_0 quantization support

* ggml-zendnn : sync with latest ZenDNN

* ggml-zendnn : address review comments for Q8_0
2026-05-25 12:26:07 +03:00
Johannes Gäßler ec183556c6 CUDA: fix PDL CC check for JIT compilation (llama/23471) 2026-05-25 12:26:07 +03:00
Pascal 8402c36039 vulkan: fuse snake activation (mul, sin, sqr, mul, add) (llama/22855)
* vulkan: fuse snake activation (mul, sin, sqr, mul, add)

Add snake.comp shader with F32 / F16 / BF16 pipelines and
ggml_vk_snake_dispatch_fused. The matcher recognizes the naive 5 op
decomposition emitted by audio decoders (BigVGAN, Vocos) for snake
activation y = x + sin(a*x)^2 * inv_b and rewrites it to a single
elementwise kernel.

test_snake_fuse from the CUDA PR now also compares CPU naive vs
Vulkan fused across F32 / F16 / BF16.

* vulkan: address jeffbolznv review for fused snake activation

Rename T / C to ne0 / ne1 in the shader and push constants to match
the standard naming convention used across the Vulkan backend.

Tighten ggml_vk_can_fuse_snake: require x and dst to be contiguous
(the shader uses idx = i0 + i1 * ne0) and require a / inv_b to be
tightly packed on the broadcast dim (the shader reads data_a[i1]).

* vulkan: tighten snake fusion type checks for all operands (address jeffbolznv review)

* vulkan: reject snake fusion when ne[2] or ne[3] > 1 (address jeffbolznv review)

* vulkan: address 0cc4m review for fused snake activation

snake.comp is renamed to follow the ggml DATA_A_* / A_TYPE convention.
A_TYPE now applies to the activation tensor data_a instead of the
broadcast multiplier, and the bindings become data_a (A_TYPE), data_b
(float), data_c (float) and data_d (D_TYPE). A header at the top of
the shader maps each buffer to its role in y = x + sin(b * x)^2 * c.

On the C++ side, ggml_vk_can_fuse_snake reuses the existing snake_pattern
constant instead of duplicating the op list, sin_node is extracted as a
named local alongside the other chain nodes, and the broadcast operands
a and inv_b are now required to be GGML_TYPE_F32 to match the hardcoded
float bindings on data_b and data_c (the previous a->type == x->type
would silently reject any future BF16 or F16 chain once the supports_op
gate for SIN / SQR is lifted). ggml_vk_snake_dispatch_fused gets an
explicit GGML_TYPE_F32 case and GGML_ABORT on default in place of the
silent f32 fallback, and a stale comment about data_a[i1] / data_inv_b[i1]
is refreshed to match the new binding names.
2026-05-25 12:26:07 +03:00
Chen Yuan c436f1419f fix(flash-attn): replace f32 with kv_type and q_type (llama/23372) 2026-05-25 12:26:07 +03:00
Georgi Gerganov 158d93c836 metal : optimize concat kernel and fix set kernel threads (llama/23411)
* metal : fix GGML_OP_SET kernel threads

* tests : extend test_cpy to support different src/dst shapes

Extend test_cpy to support different source and destination tensor shapes
for CPY operations (reshaping), where the total number of elements must match.

- Renamed ne -> ne_src, added ne_dst parameter (default: use src shape)
- Added 50 new reshaping test cases covering 1D<->2D<->3D<->4D conversions
- Tests exercise 1024 boundary, small shapes, and large dimensionality changes
- Fixed dangling reference bug (storing & to temporary std::array)
- Updated all existing test calls with permute/transpose args for compatibility

Assisted-by: llama.cpp:local pi

* metal : optimize concat kernel with row batching for small widths

When ne0 < 256, batch multiple rows into a single threadgroup to improve
occupancy. This avoids underutilizing the GPU when processing narrow tensors.

- Dispatch nth = min(256, ne0) threads per group
- Calculate nrptg (rows per threadgroup) to fill up to 256 threads
- Update kernel index calculation to handle the row batching
- Add boundary check for i1 >= ne1

Assisted-by: llama.cpp:local pi

* tests : clean-up

* tests : refactor CPY shape tests to use dimension permutations

Replace 75 hardcoded test cases with a loop over permutations of
{3, 5, 7, 32} (total elements: 3360). Each src permutation is tested
against canonical sorted and reverse dst, skipping identical shapes.
Covers F32, F16, and Q4_0 (when both src and dst ne0 == 32).

Assisted-by: llama.cpp:local pi
2026-05-25 12:26:07 +03:00
Matt Corallo 03da9f17f4 ggml : Check the right iface method before using the fallback 2d get (llama/23306)
Probably no backends implement only one of 2d get/set, but this
might be annoying for some future backend developer trying to add
2d get/set.
2026-05-25 12:26:07 +03:00
Todor Boinovski 6d1d66de40 hexagon: ssm-conv fix for large prompts (llama/23307)
* hexagon: remove gathers and better handling of vtcm in ssm-conv

* hexagon: relax ssm-conv gating requirements

* hexagon: add new prefill ssm-conv backend test

* hexagon: remove trailing white space

* hex-rope: uninline rope_cache_init, otherwise it breaks after rebaseing with SSM_CONV changes

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-05-25 12:26:07 +03:00
lhez 896718eacf opencl: refactor backend initilization (llama/23318)
* opencl: refactor initialization

* opencl: refactor GPU identification

* opencl: rename for consistency

* opencl: cache global mem size in dev_ctx

* opencl: adjust log level

* opencl: load argsort and flash_attn kernels in supports_op

* argsort kernel must be built for supports_op for querying the max
  workgroups
* flash_attn kernel has many variants, only load them when needed
2026-05-25 12:26:07 +03:00
Daniele ad717a6de0 vulkan: optimize operations in the IM2COL shader (llama/22685)
* vulkan: optimize operations in the IM2COL shader

* Add comments and improve the code formatting
2026-05-25 12:26:07 +03:00
Max Krasnyansky b93a5ba605 hexagon: HMX quantized matmul rework (llama/23368)
* hmx-mm: update debug logging in hmx-mm

* hmx-mm: update dequant logic to use HVX_vector_x2/4

* hmx-mm: remove non-pipelined version of the quantize matmul

It seems that we don't reall need non-pipelined version

* hmx-mm: use activation depth mode and update naming

Co-authored-by: Kim-Chyan Gan <kgan@qti.qualcomm.com>

* hex-mm: minor hmx matmul naming updates

* hmx-mm: remove unused vars

* snapdragon: scripts bump default ubatch-size to 1K

* hexagon: combine HMX and power and clock settings into a single set_power call

* hmx-mm: remove leftover of the scale repl helper

* hexagon: fix editconf error

---------

Co-authored-by: Kim-Chyan Gan <kgan@qti.qualcomm.com>
2026-05-25 12:26:07 +03:00
Andreas Kieslinger 3fa19558f2 Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522)
* Adds initial PDL setup.

* Adds PDL barriers based on simple heuristic: place "sync" before first input pointer access, and "launch" after last write, e.g. to tensors like dst.

* Further optimization pass of the first half of kernels

* Optimized PDL barriers for the second batch of kernels

* Further refinements after rebase.

* Moves pdl logic to separate function, removes some whitespace

* Strips post-hoc PDL logic

* Adds stream capture PDL setup. Enrolls quantize_q8_1 to leverage pdl to
overlap execution with previous kernels

* Enrolls mul_mat_vec_q, rms_norm_f32 and k_bin_bcast (partly) into PDL

* Enrolls mmvf, rope, set-rows and topk kernels for gpt-oss into PDL

* Introduce ggml_cuda_kernel_launch, to abstract away cudaLaunchKernelEx,
to enable hip/musa compatibility

* Enrolls cpy_scalar_contiguous, k_get_rows_float and rms_norm_f32

* Enrolls flash_attn_combine_results

* Fix: Drops needless and broken check of CUDA arch for PDL. PDL either
works or is without effect.

* Enrolls flash-attention kernels to pdl

* Fix: inlines ggml_cuda_kernel_launch, and uses perfect forwarding for
kernels args. This fixes PDL.

* Perf: Enrolls k_bin_bcast variadic template invocation into PDL, via
and template alias and template expansion

* Enrolls all remaining kernels for qwen3-coder-next into PDL

* Remove all PDL LC calls to create a baseline

* Added LC according to internal guidance and tested kernel performance.

* Enrols missing qwen3-5 kernels passively into PDL.

* Kernel optimizations (LC signals) for qwen3.5

* Enrolls ssm-scan kernels into PDL

* Adds GGML_CUDA_PDL command line option to toggle PDL.

* Fix: Ada and lower compilation by guarding PDL calls correctly

* Cleanup: Removes commented out GGML_CUDA_PDL_LC

* Cleanup: Removes experimental comments

* Adds 90-virtual to build script so that Hopper GPUs can leverage PDL.

* Adds stricter checks to enable PDL, adds env-check to disable it, and removes now superfluous compile option to enable PDL.

* Fix: Correct PDL en/disablement based on device-side arch check. Host
side check is UB. Required moving from macros to inlined functions

* Fix: default-disable PDL. Enable by setting GGML_CUDA_ENABLE_PDL=1

* Enable PDL by default for Hopper+ devices

* Enrolls softcap_f32 and two flash_attn kernels into PDL.

* Improves flash attn PDL barrier placement

* Fix: Perf regression on ada; excludes ada and below from PDL launches

* Improves some sync barrier placements

* Drops superfluous constructor

* Adds #endif guard comments

* Reverts experimental change to top-k-moe.cu, which moved expensive allocations
in front of the PDL barrier. It did not have a meaningful impact.

* Exchanges GGML_CUDA_DISABLE_PDL with GGML_CUDA_PDL. IFF GGML_CUDA_PDL=0
PDL is disabled

* Revert "Drops superfluous constructor". Adds const to remaining
arguments

This reverts commit 12b1d250da0089ae02a9bb71bbb3fd6d70f6f2f1.

* Cleanup: Removes and fixes some comments and whitespace

* Clarifies comment of sync-barrier position

* Relocates and refactors PDL launch functions and accessories

* Adds error checking to the regular kernel launch path

* Drops "auto" in favor of "ggml_cuda_kernel_params"

* Adds "const" to ggml_cuda_kernel_launch_params

* [Whitespace] Adds final newline to common.cuh to make editorconfig CI job happy
2026-05-25 12:26:07 +03:00
Georgi Gerganov c58fc465df metal : optimize pad + cpy (llama/23354)
* metal : optimize pad

* metal : optinmize cpy

* cont : better row packing in threadgroup
2026-05-25 12:26:07 +03:00
ravel7524 0a0a34287e ggml-cuda: tune RDNA3 Q6_K MMVQ nwarps (llama/23349) 2026-05-25 12:26:07 +03:00
shaofeiqi 37f17208c2 opencl: add MoE support for q4_k, q5_k, q6_k on Adreno (llama/23303)
* opencl: add q4_k moe support

* opencl: add q5_k moe support

* opencl: add q6_k moe support

* opencl: adjust format

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-05-25 12:26:07 +03:00
Aparna M P aca63e7638 hexagon: add MROPE and IMROPE support in HTP rope op (llama/23317) 2026-05-25 12:26:07 +03:00
Aparna M P 459ff0707b hexagon: enable support for NORM op (llama/23319) 2026-05-25 12:26:07 +03:00
Reese Levine 6090f39f36 ggml-webgpu : extend GDN for K>1 (llama/23299) 2026-05-25 12:26:07 +03:00
Intel AI Get-to Market Customer Success and Solutions 28edd0cb36 sycl: add GGML_SYCL_USE_ASYNC_MEM_OP env toggle (llama/22153)
* sycl: add GGML_SYCL_USE_ASYNC_MEM_OP env toggle

Signed-off-by: Chun Tao <chun.tao@intel.com>

* Use async mem ops for correctness when SYCL graphs are explicitly on.

Signed-off-by: Tao, Chun <chun.tao@intel.com>

---------

Signed-off-by: Chun Tao <chun.tao@intel.com>
Signed-off-by: Tao, Chun <chun.tao@intel.com>
Co-authored-by: Chun Tao <chun.tao@intel.com>
2026-05-25 12:26:07 +03:00
Radoslav Gerganov 3477fdb2e3 rpc : keep last_graph_uid in the device context (llama/23273)
With the introduction of MTP we can have multiple compute contexts for
the same RPC device. In this case last_graph_uid is not updated properly
when contexts are being switched. This patch fixes this by moving
last_graph_uid to the device context, making sure it is always updated.

closes: #23242
2026-05-25 12:26:07 +03:00
Pranav Dhinakar eb558f23cb hexagon: add support for TRI op (llama/22822)
* Hexagon: TRI HVX Kernel addition to ggml hexagon HTP ops and context

* addressed PR review comments for TRI op

* hexagon: clang format

* hex-unary: remove merge conflict markers

* hex-ggml: remove duplicate op cases (merge conflict)

* hex-ggml: fix editor config errors

---------

Co-authored-by: Todor Boinovski <todorb@qti.qualcomm.com>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-05-25 12:26:07 +03:00
Pranav Dhinakar 0a11c9fe83 ggml-hexagon: add PAD op HVX kernel (llama/23078)
* ggml-hexagon: add PAD op HVX kernel

Implements GGML_OP_PAD on the Hexagon HTP backend using HVX vectorized
kernels. Supports zero-padding and circular padding across all 4 tensor
dimensions.

* hex-ggml: remove duplicate op cases (merge conflict)

* hex-pad: fix editorconfig checks and macro alignment

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-05-25 12:26:07 +03:00
Intel AI Get-to Market Customer Success and Solutions c65b082c94 sycl: scalar SWAR byte-subtract in Q6_K MMVQ dot product (llama/22156)
Signed-off-by: Chun Tao <chun.tao@intel.com>
Co-authored-by: Chun Tao <chun.tao@intel.com>
2026-05-25 12:26:07 +03:00
Intel AI Get-to Market Customer Success and Solutions 619262ad24 sycl: route small f32 matmuls to oneMKL, bypass oneDNN (llama/22150)
Signed-off-by: Chun Tao <chun.tao@intel.com>
Co-authored-by: Chun Tao <chun.tao@intel.com>
2026-05-25 12:26:07 +03:00
Gabe Goodhart 4fb3ccabd3 feat: Support d_conv=15 for ssm-conv.cu (llama/23017)
Branch: ModalityConditionalAdapters
AI-usage: none
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
2026-05-25 12:26:07 +03:00
Oliver Simons 53736a3f0e CUDA: Continue directly including cuda/iterator (llama/23102)
Cont of #22936, forgot to update one site
2026-05-25 12:26:07 +03:00
Jan Ekström 9e96e0eaf1 ggml-vulkan/CMakeLists: add a check for SPIRV-Headers (llama/22009)
* ci/run: set explicit SPIR-V Headers search path for macOS vulkan CI

For whatever reason, the files are under additional sub-path
`vulkan/` under the cmake directory, which does not match either
current LunarG macOS Vulkan SDK structure (`lib/cmake/SPIRV-Headers`),
nor what gets installed when you run the cmake build+install for
SPIRV-Headers itself on at least Linux (`share/cmake/SPIRV-Headers`).

This allows for SPIRV-Headers to be found, as currently the CI
runner's setup does not seem to include the relevant path in
list of search locations.

* ggml-vulkan/CMakeLists: add a check for SPIRV-Headers

This is installed by the project if it is built and installed.
Receiving an error during the configuration step is generally
preferred to receiving an error in the middle of a build.
2026-05-25 12:26:07 +03:00
Pascal 50482cbd22 vulkan: add cpy bf16 -> f32 pipelines (llama/22677) 2026-05-25 12:26:07 +03:00
Jeff Bolz e417ce7aeb vulkan: Support unaligned tensors for ROPE (llama/22637) 2026-05-25 12:26:07 +03:00