Commit Graph

4700 Commits

Author SHA1 Message Date
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
Daniel Bevenius f049fff95a
release : v1.9.1 (#3892) 2026-06-19 06:12:37 +02:00
Daniel Bevenius 200b119790
ci : add GGML_NATIVE=OFF and GGML_BMI2=OFF to windows-blas (#3891)
* ci : add GGML_NATIVE=OFF and build all cpu-variants

This commit adds -DGGML_BACKEND_DL=ON, -DGGML_NATIVE=OFF, and
-DGGML_CPU_ALL_VARIANTS=ON to the releases.

The motivation for this is that currently the Windows BLAS build
uses the native CPU instructions and if target systems do not support
these instructions, the build will fail like the linked issue reports.

Resolves: https://github.com/ggml-org/whisper.cpp/issues/3889

* ci : update ubuntu-cpu release job for all variants [no ci]

This commit enables the ubuntu-cpu job to include all cpu variants and
ensures that the ggml backend libraries are built into the bin directory
similar to how llama.cpp does it.

The following is a build on my fork with this change:
https://github.com/danbev/whisper.cpp/releases/tag/untagged-fc3c71f0bf0f7bf19d19
2026-06-18 14:49:08 +02:00
Daniel Bevenius 86c40c3bd6
release : v1.9.0 (#3886) 2026-06-17 11:36:57 +02:00
KITAITI Makoto 0d14756929
ruby : add support for Parakeet (#3885)
* Add Whisper::Parakeet::Params

* Add tests for Parakeet::Params

* Remove unused variabel

* Add callbacks to Parakeet::Params

* Group callback and user_data params

* Undefine local macros

* Define GetParakeetParams

* Remove unused variable

* Use ITERATE_CALLBACK_PARAMS

* Use ITERATE_CALLBACK_PARAMS instead of ITERATE_USER_DATA_PARAMS

* Fix memsize

* Remove unnecessary macros

* Simplify params registration

* Define Parakeet

* Add hook methods to Parakeet::Params

* Fix typo

* Check callback container in GetParakeetParams

* Reduce if

* Free parakeet_full_params

* Implement Parakeet::Context#initialize

* Add TestParakeetContext

* Add Parakeet::Segment

* Prevent double-free

* Add Parakeet::Context#transcribe

* Add Parakeet::Context#each_segment

* Define Parakeet::Segment attributes

* Define Parakeet::Segment#deconstruct_keys

* Add tests for Parakeet::Segment#deconstruct_keys

* Run Parakeet::Context#transcribe without GVL

* Make it to abort for Parakeet

* Add Parakeet.log_set

* Define Parakeet::Token

* Define Parakeet::Segment#each_token

* Implement some hooks of Parakeet::Params

* Convert int to VALUE

* Implement hooks for Parakeet

* Implement Parakeet::Context#full

* Add tests for Parakeet::Context#full

* Add Parakeet to RBS

* Fix ruby_whisper_parakeet_params_free

* Free ruby_whisper_parakeet_context

* Add tests for hooks

* Add Parakeet section to README

* Add more attributes of Parakeet::Context

* Add tests for Parakeet::Context's attributes

* Update RBS

* Register parakeet-tdt-0.6b-v3

* Narrow scope of log constants

* Extract activate and deactivate of log_queue

* Make start_log_callback_thread private

* Don't call start_log_callback_thread unncecessarilly

* Early return from log_queue_enqueue when not active

* Gropu log_queue members

* is_active -> is_open

* Fix English

* Share parakeet full body function

* ruby_whisper_parakeet_abort_callback_user_data -> ruby_whisper_abort_callback_user_data

* NULL check for callback containers

* Fix Parakeet.log_set

* Omit Parakeet tests on CI

* Extract Whisper::LogSettable

* Join log callback thread in a log queue function

* Revert Join log callback thread in a log queue function

* Extract output methods to modules

* Move Parakeet init functions into init_parakeet()

* Add output methods to Parakeet classes

* Add Parakeet's output methods to RBS

* Use Whisper::Output in RBS

* Add LogSettable to RBS

* Fix module Token -> class Token

* Add Parakeet::Model

* Add test for Parakeet::Model

* Add Parakeet::Model to RBS

* Move position of Parakeet::Model in RBS

* Parakeet -> TestBase::Parakeet

* Add Parakeet::Context#model in RBS

* Add Whisper::Output

* Fix nil check

* Define ruby_whisper_parakeet_model_memsize

* Fix order of declaration in ruby_whisper_parakeet_model_get_xxx

* Define Parakeet.system_info_str

* Add test for Parakeet.system_info_str

* Add signature of Parakeet.system_info_str

* Define Parakeet::VERSION

* Add test for Parakeet::VERSION

* Add signature of Parakeet::VERSION

* Add Parakeet::Context::Params

* Make Parakeet::Context.new accept Context::Params

* Add test for Parakeet::Context.new with Context::Params

* Update RBS

* Remove params from Parakeet::Params which are moved from whisper_parakeet_full_params

* Remove tests for removed params

* Make Parakeet tests follow original behavior changes

* Add Parakeet model shortcuts

* Alloc token data in factory instead of alloc func

* Fix variable name

* Update RBS

* Refactor log settable module

* Use log settable for Whisper

* Address deadlock

* Make test follow change of log queue implementation

* Refactor to make abort callback use the same way to parakeet's way

* Remove redundant structs

* Fix test name

* Fix README

* Add missing parallel transcription

* Fix test for parakeet info

* Remove removed params

* Wait for logs dequeued

* Fix instance variable name

* Load etc feature

* Remove unnecessary comment

* Remove unnecessary thread safety check

* Remove outdated comment

* Skip downloading model if cache exists

* Change Hugging Face URI for Parakeet models

* Bump required Ruby version to 3.3

* Fix English
2026-06-17 06:42:09 +02:00
Daniel Bevenius 9efddafb91
parakeet : add support for NVIDIA Parakeet (#3735)
* parakeet : add support for NVIDIA Parakeet


Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-06-16 20:44:10 +02:00
Daniel Bevenius 3805e602d3
ci : only trigger release jobs for tags (#3883)
* ci : only trigger release jobs for tags

This commit removes the building of the release jobs on pushed to
master.

The motivation for this is that it can be confusing at the momement when
releasing that the push to master also triggers the release jobs but
the actual release will be skipped. With this change the release job is
only run when a tag is pushed which should result in a single Release
github actions job and make it easier to follow.

* ci : add GGML_NATIVE=OFF for ubuntu-22-gcc
2026-06-16 14:33:42 +02:00
Daniel Bevenius 48f628a848
release : v1.8.7 (#3881) 2026-06-16 12:28:23 +02:00
Rum Nguyen db5a84bd79
cli : add --version flag (#3878)
Adds a `--version` option to whisper-cli that prints the library version
via `whisper_version()` and exits, plus a corresponding entry in the help
output. Mirrors the existing `-h`/`--help` handling.

Closes #608
2026-06-16 08:58:09 +02:00
Georgi Gerganov 0ec0845110 talk-llama : sync llama.cpp 2026-06-15 10:33:53 +03:00
Georgi Gerganov 0a3fa9ca17 sync : ggml 2026-06-15 10:33:53 +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
Daniel Bevenius df7638d822
ci : pin github actions to commit sha's (#3865) 2026-06-09 12:51:00 +02:00