Commit Graph

1437 Commits

Author SHA1 Message Date
Reese Levine 1361f679cc
GGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (llama/16018)
* Add paramater buffer pool, batching of submissions, refactor command building/submission

* Add header for linux builds

* Free staged parameter buffers at once

* Format with clang-format

* Fix thread-safe implementation

* Use device implicit synchronization

* Update workflow to use custom release

* Remove testing branch workflow

* some f32 tests passing

* Disable set_rows until it's implemented

* f32 add all tests passing

* Begin work on set_rows

* Work on set rows

* Add error buffers for reporting unsupported SET_ROWS indices

* Remove extra comments

* Add templated addition, clean up code

* Get addition and multiplication working

* Implement rms_norm

* Add get_rows implementation

* Add new get_rows files

* Refactor use of wg size entry

* Fix compilation

* Try manually unrolled q4_0 quant

* Revert "Try manually unrolled q4_0 quant"

This reverts commit 77f8b96515f7e640ae4b0e44f066321fbc4a6166.

* Move to constant max wg size

* Check for tensor size in supports_op

* Vectorize f32 and change default workgroup size

* Move f32 get_rows from < 4 to % 4 != 0

* fix linter errors

* Add in-place tests

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
2025-09-20 13:46:37 +03:00
Georgi Gerganov eb2c01f92e
metal : refactor + optimize v2 (llama/15995) 2025-09-20 13:46:10 +03:00
Johannes Gäßler d452f0cf8c
CUDA: fix FA occupancy, optimize tile kernel (llama/15982) 2025-09-20 13:45:30 +03:00
Eve e96b285011
vulkan: automatically remove unsupported devices (llama/15976)
* remove unsupported vulkan devices

* make this happen during selection instead

* pass by reference
2025-09-20 13:45:30 +03:00
Chenguang Li e32c3b0fd3
CANN: Optimize ggml_cann_set_device (llama/15935)
* CANN: Fix ggml_cann_set_device to avoid redundant device switches

- Added a check to skip aclrtSetDevice if the current device is already set.
- Prevents unnecessary context switches while keeping thread/device consistency.

* CANN: add device default id
2025-09-20 13:45:30 +03:00
Daniel Bevenius 5c524bb879
ggml : fix padding in timestep embedding kernels (llama/15932)
* ggml : remove adding extra dim timestep embedding

This commit updates the ggml_timestep_embedding function to no longer
add an extra dimension when the specified dimension is odd.

The motivation for this change is that this introduces an unnecessary
dimension when the dimension is odd, which caused an issue in the
kernels which were not expecting this extra dimension and it resulted in
uninitialized memory for the second to last dimension.

* ggml-cuda : fix padding in timestep embedding kernel

This commit removes the zeroing out of the last dimension now that we
are not adding the extra padding dimension.

* ggml-metal : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel

* ggml-opencl : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-sycl : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-vulkan : fix padding in timestep embedding kernel

This commit fixes the zero padding for odd dimensions in
the timestep embedding kernel.

* ggml-cpu : fix padding in timestep embedding function

This commit removes the zeroing out of the last dimension now that we
are not adding the extra padding dimension.
2025-09-20 13:45:30 +03:00
Jake Karnes f72ec185fb
CUDA: fix im2col_3d to respect non-contiguous inputs (views) (llama/15956)
* fix im2col_3d to respect non-contiguous inputs (views)

The CUDA 3D im2col kernel computed source addresses assuming compact layout (products of dims), ignoring nb[] strides.

This patch switches im2col_3d source indexing to use true strides derived from src1->nb[] (in elements), mirroring the approach used in the 2D CUDA im2col path. Destination indexing is unchanged.

* use ggml_element_size() for src strides

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2025-09-20 13:45:30 +03:00
yael-works a642b533a4
SYCL: Add COUNT_EQUAL operator support (llama/15991)
* SYCL: Add COUNT_EQUAL operator support (rebased on master)

* SYCL: remove duplicate op_count_equal definition

* tests: remove test_count_equal_typed and use test_count_equal for all cases

* tests: keep only I32 case for COUNT_EQUAL as suggested

* tests: keep only I32 case for COUNT_EQUAL as requested
2025-09-20 13:45:30 +03:00
Aman Gupta 10bd5d3626
CUDA: some micro-optimizations in mmf.cuh for mul_mat_id (llama/15926) 2025-09-20 13:45:30 +03:00
Georgi Gerganov 82a8c141ea
metal : remove memory pools (llama/15966)
* metal : remove mem pool usage

ggml-ci

* metal : remove mem pool implementation

ggml-ci

* metal : take into account the actual allocated memory of the tensor

ggml-ci

* cont : use ggml_backend_buft_get_alloc_size

ggml-ci

* cont : improve, comments

ggml-ci

* cont : add functions for the extra tensor sizes

* metal : add comments

ggml-ci

* metal : implement .get_alloc_size for the rest of the buffer types

ggml-ci

* metal : remove ggml_metal_heap

ggml-ci
2025-09-20 13:45:29 +03:00
Ruben Ortlam c36358cb3c
Vulkan: Clean up mul_mm shader (llama/15987)
* vulkan: move mul_mm dequantization steps into a separate file and functions

* improve mul_mm vector load code

* fix debug mode issues and warnings
2025-09-20 13:45:29 +03:00
Georgi Gerganov 2d3f15607f
metal : fix kernel requirements (llama/15983)
* metal : fix kernel requirements

ggml-ci

* cont : fix supports_op

* cont : fix supports_op for ARGMAX
2025-09-20 13:45:29 +03:00
Aaron Teo 7dca05ca77
ggml-zdnn: rm user mapped buffers (llama/15965)
* ggml-zdnn: rm user mapped buffers

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-zdnn: rm dead code

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-zdnn: attempt to fix missing extra data buffer free

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2025-09-20 13:45:29 +03:00
Jeff Bolz 1789ed3f2c
vulkan: fix failing dequant shaders (llama/15862)
* vulkan: fix failing dequant shaders

* add missing const
2025-09-20 13:45:29 +03:00
Jeff Bolz a3defb0a3b
vulkan: initialize vulkan-hpp to allow using extension function pointers (llama/15705)
Use this to query register count for shader compiles on NVIDIA. Currently
this is only for performance debug, but it could eventually be used in some
heuristics like split_k.
2025-09-20 13:45:29 +03:00
Georgi Gerganov 2caf15d68a
metal : refactor kernel loading (llama/15964)
* metal : refactor bin kernels loading

ggml-ci

* metal : refactor rms kernel loading

ggml-ci

* ci : try to add memory leaks check

ggml-ci

* ci : try to enable memory leak detection for Mac

* cont : seems to be working
2025-09-20 13:45:29 +03:00
Georgi Gerganov 0d36ba9e1a
metal : allow ops to run concurrently (llama/15929)
* metal : run graphs ops concurrently

ggml-ci

* cont : add flags for debugging and disabling concurrency

ggml-ci

* cont : refactor and handle fusing

ggml-ci

* cont : simplify - no need to use GPU address

ggml-ci

* cont : prepare mem ranges for reuse + add ggml-metal-common.cpp

ggml-ci

* cont : avoid redundant keywords in cpp [no ci]

* metal : reorder graph for better concurrency

ggml-ci

* metal : fix race on mem pool buffers

ggml-ci

* cont : add env GGML_METAL_GRAPH_OPTIMIZE_DISABLE

ggml-ci

* cont : refactor, optimize, add comments

ggml-ci

* cont : refactor ggml-metal.m

ggml-ci

* minor : update logs [no ci]
2025-09-20 13:45:29 +03:00
Georgi Gerganov 20a930ec94
metal : fix memory leaks (llama/15962)
ggml-ci
2025-09-20 13:45:28 +03:00
Aaron Teo e902731ccc
ggml-zdnn: fix #15414, activate FP16 and BF16 acceleration and incorrect zTensor free (llama/15839) 2025-09-20 13:45:28 +03:00
Ruben Ortlam 424c85f22a
Vulkan iGPU device selection overhaul and PCI ID API support (llama/15947)
* vulkan: implement ggml igpu device type, implement pci id support

* fix compiler warning

* prevent printf overflow warning
2025-09-20 13:45:28 +03:00
Mathieu Baudier 5a752bab84
vulkan: Make device memory check more portable (llama/15939) 2025-09-20 13:45:28 +03:00
Neo Zhang Jianyu cd764eaf2b
Revert "sycl: add usage of enqueue_functions extension (llama/14244)" (llama/15910)
* Revert "sycl: add usage of enqueue_functions extension (#14244)"

This reverts commit 8308f98c7fb778e54bf75538f5234d8bd20915e9.

* fix missed revert code, format the code
2025-09-20 13:45:28 +03:00
Diego Devesa 555dcb3e01
ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type (llama/15797)
* ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type

ggml-backend : add device id to device props

llama : only use iGPU devices if there are no GPU devices

llama : do not use multiple devices from different backends with the same device id
2025-09-20 13:45:28 +03:00
Johannes Gäßler f0768eb575
CUDA: larger SRAM reads for tile FA, AMD FP16 dot (llama/15927)
* CUDA: larger SRAM reads for tile FA, AMD FP16 dot

* fix logic for availability of v_dot2_f32_f16
2025-09-20 13:45:28 +03:00
Daniel Bevenius 020eb19eb3
ggml-cpu : add check for ARM MATMUL_INT8/i8mm support (llama/15922)
This commit adds a check for GGML_MACHINE_SUPPORTS_i8mm when enabling
MATMUL_INT8 features, ensuring that i8mm intrinsics are only used when
the target hardware actually supports them.

The motivation for this is to fix ggml CI build failures where the
feature detection correctly identifies that i8mm is not supported,
adding the +noi8mm flag, but MATMUL_INT8 preprocessor definitions are
still enabled, causing the compiler to attempt to use vmmlaq_s32
intrinsics without i8mm support.

Refs: https://github.com/ggml-org/ggml/actions/runs/17525174120/job/49909199499
2025-09-20 13:45:28 +03:00
Charles Xu b079d9c8b0
kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed (llama/15614)
* kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed

* removes the Whisper-specific check for GET_ROWS support
2025-09-20 13:45:27 +03:00
hipudding dadf73665a
CANN: Disable acl_graph for prefill stage (llama/15933)
Since the prefill length is not fixed, graphs constructed for the
prefill stage cannot be reused. For this reason, ACL graph
execution is disabled by default during prefill.
2025-09-20 13:45:27 +03:00
Oliver Simons f5ef0e25e2
CUDA: Add `fastdiv` to `k_bin_bcast*`, giving 1-3% E2E performance (llama/15872)
* Add fastdiv and fastmodulo to k_bin_bcast kernel

* Address review comments

* `prod_` instead of `prod` suffix

* Add test case for `k_bin_bcast_unravel` in CUDA backend
2025-09-20 13:45:27 +03:00
Daniel Bevenius 3617008c37
ggml-cpu : fix padding in ggml_timestep_embedding (llama/15917)
This commit fixes the zero padding for odd dimensions in
ggml_compute_forward_timestep_embedding_f32.
The motivation for this is that currently if an odd dimension is used,
the padding check incorrectly uses the dimension value for indexing.
For example, with dim=15:

Elements 0-6 are set to cosine values
Elements 7-13 are set to sine values
Element 14 is left uninitialized (contains garbage)
Element 15 is correctly set to zero

This fix changes embed_data[dim] to embed_data[2 * half] so that
element 14 (the first unused element) is properly set to zero as well
as the last element.

Resolves: https://github.com/ggml-org/ggml/issues/1324
2025-09-20 13:45:27 +03:00
Georgi Gerganov 7eae055e61
metal : make the backend async (llama/15906) 2025-09-20 13:44:27 +03:00
Chenguang Li 4d453b14a9
CANN: Add ROPE sin/cos cache for reuse (llama/15912)
* CANN: Add ROPE sin/cos cache for reuse

Introduce sin/cos caching mechanism in ROPE to avoid redundant
computation across layers. The cache is built on the first layer
per device and reused by subsequent layers if parameters match.

- Added sin_cache / cos_cache pointers and position_length tracking
- Introduced cache validity flags and properties:
  (ext_factor, theta_scale, freq_scale, attn_factor, is_neox)
- Accelerates ROPE by eliminating repeated sin/cos generation

This change reduces overhead in multi-layer scenarios while
preserving correctness by verifying parameter consistency.

Co-authored-by: hipudding <huafengchun@gmail.com>

* fix typo

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
Co-authored-by: hipudding <huafengchun@gmail.com>
2025-09-20 13:42:53 +03:00
Chenguang Li 9b773acac0
CANN: implement LRU cache for ACL graphs (llama/15814)
* CANN: implement LRU cache for ACL graphs in CANN backend

- Introduce ggml_cann_graph_lru_cache to store multiple ggml_cann_graph objects.
- Graphs are loaded on demand and evicted using LRU policy when capacity is exceeded.
- Updated push, move_to_front, and clear methods to manage cached graphs efficiently.
- Ensures reuse of graphs, reducing graph reconstruction overhead in CANN backend.

* fix typo

* The LRU cache capacity can be configured via an env variable

Signed-off-by: noemotiovon <757486878@qq.com>

* refactory acl graph

* refactory && fix review comments

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:53 +03:00
Ruben Ortlam 7abe187860
vulkan: throw the oom error instead of no memory type found (llama/15905) 2025-09-20 13:42:53 +03:00
Jeff Bolz d0e98656c3
vulkan: Fix OOB accesses in soft_max_back (llama/15861) 2025-09-20 13:42:52 +03:00
Johannes Gäßler e35d1375ee
HIP: use v_dot2_f32_f16 instruction for FA (llama/15884) 2025-09-20 13:42:52 +03:00
lksj92hs 7fbbb67b47
Workaround for subgroup arithmetic failing on MoltenVK with AMD GPUs (issue 15846) (llama/15886) 2025-09-20 13:42:52 +03:00
Aman Gupta 621764b1a5
CUDA: Add mul_mat_id support for the mmf kernel (llama/15767)
* CUDA: Add mul_mat_id support the mmf

Add support for mul_mat_id for bs < 16

* Review: use warp_size, fix should_use_mmf condition

* Launch one block per expert, stride along n_expert_used

* templatize mul_mat_id

* Pad shmem to 16 bytes, add helper function mul_mat_f_switch_ids

* Reduce compile times by dividing mmf into f16, bf16 and f32 variants

* Divide mmf by ncols_dst

* Add missing files

* Fix MUSA/HIP builds
2025-09-20 13:42:52 +03:00
Johannes Gäßler 260982232c
CUDA: fix GET_ROWS for large tensors (llama/15882) 2025-09-20 13:42:52 +03:00
Jeff Bolz c29cd54818
vulkan: sort graph to allow more parallel execution (llama/15850)
* vulkan: sort graph to allow more parallel execution

Add a backend proc to allow the backend to modify the graph. The
vulkan implementation looks at which nodes depend on each other
and greedily reorders them to group together nodes that don't
depend on each other. It only reorders the nodes, doesn't change
the contents of any of them.

With #15489, this reduces the number of synchronizations needed.

* call optimize_graph per-split
2025-09-20 13:42:52 +03:00
Aman Gupta 70ee808f3d
CUDA: generate_cu_files.py - add missing mxfp4 (llama/15880) 2025-09-20 13:42:52 +03:00
Georgi Gerganov ae6cc6a386
cuda : fix supports_op condition for get_rows when number of blocks is too large (llama/15868)
* cuda : fix supports_op condition for get_rows when src1->ne2 > 1

ggml-ci

* ggml : add comment about ggml_get_rows

ggml-ci

* cuda : add FIXME [no ci]

* cuda : update support condition

ggml-ci
2025-09-20 13:42:52 +03:00
Georgi Gerganov e9cb59e970
metal : refactor + optimize (llama/15857) 2025-09-20 13:42:51 +03:00
Xuan-Son Nguyen 40bcd1a469
ggml: allow casting between f32 and i32 (llama/15783)
* ggml: allow casting between f32 and i32

* fix cuda

* add vulkan

* fix CPU non-cont

* add non-cont test case

* add note

* extend test number range

* correct note

* add cont version for vulkan
2025-09-20 13:42:51 +03:00
Sigbjørn Skjæret 0175a1df8d
CUDA: non-contiguous src0 not supported for PAD (llama/15869) 2025-09-20 13:42:51 +03:00
Chenguang Li d9c0ead2ab
CANN: Stream sync between devices for acl_graph (llama/15809)
* CANN: Switch to stream synchronization

Switch to stream synchronization because events are not effective.

Co-authored-by: hipudding <huafengchun@gmail.com>

* CANN: add Comments

---------

Co-authored-by: hipudding <huafengchun@gmail.com>
2025-09-20 13:42:51 +03:00
Jeff Bolz dfa7722e2e
vulkan: support im2col_3d (llama/15795) 2025-09-20 13:42:51 +03:00
Aaron Teo db4f504b69
ggml-cpu: clean up s390x SIMD (llama/15855)
* ggml-cpu: clean up s390x simd

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
(cherry picked from commit 0da4b6aa07d96b758812d17b2c82267632fa4ba5)
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: fix hsum data types

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2025-09-20 13:42:51 +03:00
Jeff Bolz 9523fd8de6
vulkan: Support pad_ext (llama/15794) 2025-09-20 13:42:51 +03:00
Jeff Bolz 647e2d7de5
vulkan: Use larger loads in scalar/coopmat1 matmul (llama/15729)
I think glslang will translate an access like x[i][1].z to
OpAccessChain ... x, i, 1, 2
OpLoad float16_t ...

rather than loading all of x[i] in a single OpLoad. Change the
code to explicitly load the vector/matrix.
2025-09-20 13:42:51 +03:00
Daniel Bevenius cda7d4e5ac
ggml WebGPU: remove userdata from request adapter callback (llama/15527)
* ggml WebGPU: remove userdata from request adapter callback

This commit removes the `userdata` parameter from the WebGPU request
adapter callback in `ggml-webgpu.cpp`. Instead, the lambda function
captures the `webgpu_context` directly.

The motivation for this change is to simplify the code and improve
readability.

* inline the callback lambda into the RequestAdapter call

This commit removes the callback lambda variable and inlines it directly
into the RequestAdapter call.
2025-09-20 13:42:50 +03:00
Johannes Gäßler cd70d89628
CUDA: faster tile FA (Pascal/AMD), headsize 256 (llama/15769) 2025-09-20 13:42:50 +03:00
Charles Xu be2676bb1c
kleidiai: generalize compute_forward_kv_cache to compute_forward_fp16 (llama/15817) 2025-09-20 13:42:50 +03:00
Johannes Gäßler 69400f16f1
ggml-cpu: document use of "free" memory [no ci] (llama/15834) 2025-09-20 13:42:50 +03:00
Aaron Teo f499271c4e
ggml-cpu: drop support for nnpa intrinsics (llama/15821) 2025-09-20 13:42:50 +03:00
Johannes Gäßler 6ff468cfaa
CUDA: fastdiv, launch bounds for mmvq + q8_1 quant (llama/15802)
* CUDA: fastdiv, launch bounds for mmvq + q8_1 quant
2025-09-20 13:42:50 +03:00
Daniel Bevenius 4d6e1144b1
ggml : introduce semantic versioning (ggml/1336)
* ggml : introduce semantic versioning

This commit introduces semantic versioning for the GGML library.

The motivation for this is that the current versioning, using build
numbers, makes it difficult to track changes and releases for projects
that use ggml.

The release steps are the following:
1. Sync the changes from llama.cpp using sync-llama-am.sh and after the
   PR has been approved and merged move to step 2.
2. Run scripts/release.sh and specify the type of release, major, minor,
   or patch. This script will handle incrementing the version
   (major|minor|patch), create a new commit with the version change,
   create a tag for the version, and prepare for the next development
   iteration.
3. Inspect the commits/tag and push to master. This will trigger the
   github release workflow which is triggered for new tags which will
   then publish a new release on github.

Example usage:
```console
$ ./scripts/release.sh major --dry-run
[dry-run] - No changes will be made

Step 1: Reading current version...
Current version: 0.9.0-dev
New release version: 1.0.0

Step 2: Updating version in CMakeLists.txt...
  [dry-run] Would update GGML_VERSION_MAJOR to 1
  [dry-run] Would update GGML_VERSION_MINOR to 0
  [dry-run] Would update GGML_VERSION_PATCH to 0
  [dry-run] Would remove -dev suffix

Step 3: Committing version bump...
  [dry-run] Would commit: 'ggml : bump version to 1.0.0'

Step 4: Creating git tag...
  [dry-run] Would create tag: v1.0.0 with message 'Release version 1.0.0'

Step 5: Preparing for next development cycle...
  [dry-run] Would update GGML_VERSION_MINOR to 1
  [dry-run] Would add -dev suffix back

Step 6: Committing development version...
  [dry-run] Would commit: 'ggml : prepare for development of 1.1.0-dev'

[dry-run] Summary (no changes were made):
  • Would have released version: 1.0.0
  • Would have created tag: v1.0.0
  • Would have set next development version: 1.1.0-dev
```

Refs: https://github.com/ggml-org/ggml/issues/1333

* ggml: create branch for release candidate and check master

* ggml : sign the git tag
2025-09-20 13:42:50 +03:00
Gregor Jasny c80f78cc7b
CUDA : conditionally add cuda architectures (ggml/1341) 2025-09-20 13:42:50 +03:00
Gabe Goodhart ffe560cbb1
metal : Add template specialization for mul_mm_id w/ ne20 == 10 (llama/15799)
Branch: GGMLMetalNE20

Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
2025-09-20 13:42:49 +03:00
Chenguang Li 3780a3c917
CANN: Refactor ND to NZ workspace to be per-device (llama/15763)
* CANN:Refactor ND to NZ workspace to be per-device in Ascend backend

- Replaced the previous single global ND→NZ workspace with a per-device
  cache using unordered_map keyed by device ID.
- Functions `release_nz_workspace`, `relloc_nz_workspace`, and
  `get_nz_workspace` now manage workspace independently for each device,
  preventing memory conflicts in multi-device / pipeline parallel scenarios.
- This change fixes potential precision issues caused by workspace
  overwrites when multiple devices perform ND→NZ conversions concurrently.

Co-authored-by: hipudding <huafengchun@gmail.com>

* refactor

Signed-off-by: noemotiovon <757486878@qq.com>

* rename

Signed-off-by: noemotiovon <757486878@qq.com>

* fix review comments

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
Co-authored-by: hipudding <huafengchun@gmail.com>
2025-09-20 13:42:49 +03:00
leejet 2228462b19
ggml: add ops for WAN video model (cuda && cpu) (llama/15669)
* add conv3d support

* add ggml_pad_ext for cpu & cuda backend

* cuda/cpu: add im2col_3d support

* cuda: make im2col a little faster

* fix cuda pad/scale/im2col3d

* make im2col_3d faster

* gguf: support loading tensors which n_dims > GGML_MAX_DIMS

* fix cuda get_rows

* avoid ggml_conv_3d conflict

* correct GGML_OP_COUNT assertion

* avoid build failure

* avoid build failure on MacOS

* cuda: remove unnecessary MIN define

* fix cpu im2col_3d

* adjust the code style

* cuda: use simpler loop in get_rows

* add test_im2col_3d to test-backend-ops

* test-backend-ops.cpp: remove trailing whitespace

* cpu: im2col_3d support non continuous src

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

* fix test_im2col_3d

* remove unused variables

* cuda: get_rows: dfloat2 -> float2

* add test_pad_ext to test-backend-ops.cpp

* add gguf_init_from_file_ext impl

* Revert "gguf: support loading tensors which n_dims > GGML_MAX_DIMS"

This reverts commit d8377a0a37f314bd3713fe043b4333ad661610c1.

* Revert "add gguf_init_from_file_ext impl"

This reverts commit d9f1d13208c68ef83b3538201ac7f31614fb1994.

* update ggml_backend_vk_device_supports_op

* fix ggml_backend_vk_device_supports_op

* update other backend supports op for ggml_pad_ext

* metal/opencl/sycl/vulkan: fix GGML_OP_PAD check in supports_op

---------

Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
2025-09-20 13:42:49 +03:00
hipudding 96efb472b4
CANN: Fix precision issue on 310I DUO multi-devices (llama/15784) 2025-09-20 13:42:49 +03:00
rmatif 1569daf524
opencl: add hs=40 to FA (llama/15758) 2025-09-20 13:42:49 +03:00
Chenguang Li 5c860e94c6
CANN: fix acl_rstd allocation size in ggml_cann_rms_norm (llama/15760)
Fixes #15330

Adjust the allocation size of acl_rstd. The parameter `dims` is set to 3 according to the CANN documentation.

Co-authored-by: Yuchuan <yuchuan-cao@users.noreply.github.com>
2025-09-20 13:42:49 +03:00
Ruben Ortlam 719a05c665
vulkan: fix mmv subgroup16 selection (llama/15775) 2025-09-20 13:42:49 +03:00
Jeff Bolz 4a702a867c
vulkan: don't use std::string in load_shaders, to improve compile time (llama/15724)
* vulkan: don't use std::string in load_shaders, to improve compile time

* keep the string version for those calls that use it
2025-09-20 13:42:49 +03:00
Daniel Bevenius 4144ae10e9
vulkan : update ggml_vk_instance_validation_ext_available (llama/15666)
* vulkan : update ggml_vk_instance_validation_ext_available

This commit updates ggml_vk_instance_validation_ext_available() to
check for VK_EXT_validation_features instead of
VK_KHR_portability_enumeration.

Based on how the returned boolean is used later in the code (to enable
both the validation layer and the VK_EXT_validation_features extension),
it appears the function may have been intended to check for the
validation layer features extension.

* remove try/catch

This was a left over from a previous iteration where I was explicitly
quering for a specific validation layer first, which would throw.

* update warning message about validation layers
2025-09-20 13:42:48 +03:00
Shin-myoung-serp 85c7aa3750
ggml vulkan: add hardsigmoid and hardswish operations (llama/15762) 2025-09-20 13:42:48 +03:00
Oliver Simons 9eef377330
CUDA: Optimize `rms_norm_f32` kernel and its fused variants, giving 1-6% perf E2E (llama/15715)
* Add fastdiv, use it in modulo and use modulo in rms_norm_f32

Fastdiv is much faster way to do integer division, which was identified
as bottleneck in rms_norm_f32

* Support more `block_size` values in `rms_norm_f32`

This makes us more flexible in selecting the optimal threads w.r.t
paralellizing across a col vs. launch-overheads of threads and mio
throttles

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

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

* Replace modulo with fastmodulo in `rms_norm_f32`

* Use `BinPackArguments=true` for formating function calls

Will file a separate PR to adjust .clang-format file

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

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

* Use uint3 for both `fastdiv` and `fastmodulo`

The compiler seems to reliably optimize away the unused .z component in
the fastdiv use-case, see https://godbolt.org/z/rx8KPrKr3

* More constrained type declarations

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

* Rename fastdiv and fastmodulo variables to shared variable name

As suggest by JohannesGaessler, this increases clarity of the intended
use

* Pack fastdiv/fastmodulo constants into uint2/uint3 objects

By packing constants to be used together into a struct, we are less
likely to make errors.

* Rename function parameter of fastmodulo

`modulo_consts` is more fitting/descriptive

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2025-09-20 13:42:48 +03:00
hipudding 51bc843f3a
CANN: Add RoPE contiguous check for 310I DUP device (llama/15735) 2025-09-20 13:42:48 +03:00
xctan 75f739c7c8
ggml-cpu : optimize RVV kernels (llama/15720)
* ggml-cpu : optimize rvv ggml_vec_dot_f32

* ggml-cpu : optimize 128-bit rvv ggml_vec_dot_q4_K_q8_K

* ggml-cpu : fix riscv arch flags

* ggml-cpu : add more rvv ops

* ggml-cpu : optimize rvv ggml_vec_dot_q4_K_q8_K

* ggml-cpu : optimize rvv ggml_vec_dot_q6_K_q8_K

* ggml-cpu : minor rvv adjustments

* ggml-cpu : fix riscv include
2025-09-20 13:42:48 +03:00
hipudding 91e9e72ecd
CANN: Mask unsupported TRANSPOSE_1D operator (llama/15733)
CANN currently does not support kernels larger than 255.
This change disables such cases.
2025-09-20 13:42:48 +03:00
Chenguang Li d84b96d9d0
CANN: Fix type float_t to float (llama/15736)
Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:48 +03:00
Ruben Ortlam e584edb5ba
vulkan: fix shaders gen when no integer dot is available (llama/15740) 2025-09-20 13:42:48 +03:00
hipudding 5aee53c40f
CANN: Resolve soft_max precision issue (llama/15730)
Previously, the slope tensor was set to fp16 to improve efficiency.
While this worked correctly in FA, it caused precision issues in soft_max.
This change applies different data types for different operators
to balance both accuracy and performance.
2025-09-20 13:42:47 +03:00
Jeff Bolz 1e03aa66f7
vulkan: Fix macro parameter order for f32 matmul shaders (llama/15716) 2025-09-20 13:42:47 +03:00
rmatif fb37f91163
opencl: add attn sinks support for FA kernels (llama/15706) 2025-09-20 13:42:47 +03:00
Chenguang Li 3db49c1c26
CANN: Support eager execution mode under ACL graph compilation (llama/15712)
* [CANN] Support eager execution mode under ACL graph compilation

Add support for running operators in eager mode while ACL graph
compilation is enabled. This allows bypassing graph execution
and directly submitting ops, which is useful for debugging and
reducing graph build overhead in certain scenarios.

Signed-off-by: noemotiovon <757486878@qq.com>

* fix typo

Signed-off-by: noemotiovon <757486878@qq.com>

* rename to acl_graph_mode

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:47 +03:00
hipudding 13d3963f71
CANN: Support ext_factor in rope (llama/15710) 2025-09-20 13:42:47 +03:00
Johannes Gäßler f20a7b0e99
ggml-backend: raise GGML_MAX_SPLIT_INPUTS (llama/15722) 2025-09-20 13:42:47 +03:00
Gilad S 9e3600e569
vulkan: use memory budget extension to read memory usage (llama/15545)
* vulkan: use memory budget extension to read memory usage

* fix: formatting and names

* formatting

* fix: detect and cache memory budget extension availability on init

* fix: read `budgetprops.heapBudget` instead of `heap.size` when memory budget extension is available

* style: lints
2025-09-20 13:42:47 +03:00
Jeff Bolz 7a5e7368a3
vulkan: add missing clamps in new mul_mat_id paths (llama/15702)
This is a missing interaction between #15546 and #15652
2025-09-20 13:42:46 +03:00
Ruben Ortlam d5f80a2982
vulkan: disable large mmv subgroups on older Nvidia GPUs (llama/15717) 2025-09-20 13:42:46 +03:00
s-goto-11 8218dc609c
ggml: SVE support for exponential functions (llama/15145)
* SVE support for exponential functions

Add const notation to variable pg

* Update ggml/src/ggml-cpu/vec.cpp

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

* Add const

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2025-09-20 13:42:46 +03:00
Prashant Vithule 31840a3a56
ggml: aarch64: Implement SVE F16 kernels for vector functions (llama/15115)
* Added sve implementation for vec_dot_fp16 Kernel

* removed white spaces

* Added comment

* removed white spaces

* changed GGML_F16x_VEC_FMA for code consistency

* Update vec.h

---------

Co-authored-by: vithulep <p.m.vithule1517@gmail.com>
2025-09-20 13:42:46 +03:00
Ruben Ortlam 5e70d901b0
Vulkan: Add Integer Dot Product mul_mat_vec shader for legacy quants (llama/14903)
* vulkan: Add Integer Dot Product mul_mat_vec shader for legacy quants

* vulkan: use subgroup operations for quantize_q8_1 shader

* vulkan: add q8_1_x4 type with 128-bit alignment, use in mul_mat_vecq shader

* vulkan: use q8_1_x4 blocks in mul_mmq shader

* vulkan: do 8 calculations per invocation instead of 32 in mul_mat_vecq, similar to mul_mat_vec

* vulkan: tune mul_mat_vecq performance for Intel

* vulkan: fix quantizing issue when tensor is not divisible by 128

* vulkan: adapt integer dot mmv to mmv small m optimization (llama/15355)

* vulkan: allow all subgroup modes for mmv and mmvq

* vulkan: use prealloc intermediate reuse for mmvq path

* vulkan: tune mmvq for Intel, AMD GCN and Nvidia RTX 3090

* vulkan: adapt mmv quantize_y path to conditional sync logic

* vulkan: disable q8_0 mmvq on Nvidia

* vulkan: enable q8_0 on Nvidia pre-turing

* fix prealloc sync condition

* fix llvmpipe subgroup 8 issue
2025-09-20 13:42:46 +03:00
Daniel Bevenius c5f511e697
ggml : WebGPU add TRANSPOSE and RESHAPE to supported ops (llama/15695)
* ggml : WebGPU add TRANSPOSE and RESHAPE to supported ops

This commit adds support for the TRANSPOSE and RESHAPE operations in the
ggml webgpu backend.

Co-authored-by: Diego Devesa <slarengh@gmail.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
2025-09-20 13:42:46 +03:00
Akarshan Biswas 2ba5e0cb47
CUDA: fix build error from ambiguous __half conversions in conv2d (llama/15690)
* CUDA: fix build error from ambiguous __half conversions in conv2d

Building conv2d with half precision failed because `__half` defines
multiple implicit conversion operators (to float, int, short, etc.),
causing ambiguous overload resolution when multiplying with float.

Introduce a templated `to_float` helper that explicitly converts
`__half` via `__half2float`, while passing through float unchanged.
Use this helper in conv2d accumulation to ensure unambiguous and
correct promotion to float.

Fixes some build errors with half-precision kernels on CUDA.

ggml-ci

* CUDA: Replace custom to_float helper with unified ggml_cuda_cast and add half‑>float conversion

* CUDA: Add missing convert.cuh header

* CUDA: remove unnecessary extension in ggml_cuda_cast

* CUDA: Address review comment, remove second type template argument
2025-09-20 13:42:46 +03:00
hipudding bb5f844ec7
CANN: Optimize MUL_MAT_ID (llama/15658) 2025-09-20 13:42:46 +03:00
hipudding ed7ebdc757
CANN: fix RoPE cache issue on multi-device (llama/15629)
* CANN: fix RoPE cache issue on multi-device

RoPE cache only needs to be computed once per token.
However, in multi-device scenarios, not every device starts
computation from layer 0, which may lead to unallocated memory
issues and precision errors.

This commit records the first layer of each device to avoid
the above issues.

* CANN: Optimize first-layer detection method

* CANN: Remove trailing whitespace

* CANN: Only cache the data that can be determined as unchanged through the parameters.

* CANN: Update function comment
2025-09-20 13:42:45 +03:00
Georgi Gerganov 3d470687de
metal : fix checks for available FA kernels (llama/15700)
* metal : fix checks for available FA kernels

ggml-ci

* cont : fix comment [no ci]
2025-09-20 13:42:45 +03:00
Diego Devesa b11c972b88
llama : separate compute buffer reserve from fattn check (llama/15696)
Exposes ggml_backend_sched_split_graph() to allow splitting the graph without allocating compute buffers and uses it to split the graph for the automatic Flash Attention check.
2025-09-20 13:42:45 +03:00
Jeff Bolz db7ecfb61d
vulkan: handle large sizes for get_rows (llama/15686) 2025-09-20 13:42:45 +03:00
Jeff Bolz 191def71ce
vulkan: mul_mat_id coopmat2 optimizations (llama/15546)
* vulkan: mul_mat_id coopmat2 optimizations

Add a path for when the tile fits in BN/2, similar to what we have for mul_mat.

Only call fetch_scales/store_scales once per QUANT_K block, and once at the
beginning in case start_k is not aligned.

* Also add a path for BN/4 - worth a couple more percent
2025-09-20 13:42:45 +03:00
Daniel Bevenius b092e95aaa
vulkan : remove unused portability_enumeration_ext variable (llama/15679)
This commit removes the portability_enumeration_ext variable from the
ggml_vk_instance_portability_enumeration_ext_available function as it
is initialized to false but never modified, making it redundant.
2025-09-20 13:42:45 +03:00
Jeff Bolz 20ce6fcf6a
vulkan: Allow fallback to sysmem memory when vidmem is full (llama/15649)
* vulkan: Allow fallback to sysmem memory when vidmem is full

* vulkan: Add env var GGML_VK_ALLOW_SYSMEM_FALLBACK
2025-09-20 13:42:45 +03:00
Jeff Bolz 71f0ee70bf
vulkan: clamp matmul and FA results to the max finite value (llama/15652)
* vulkan: clamp matmul and FA results to the max finite value

* only clamp for fp16
2025-09-20 13:42:45 +03:00
Charles Xu 74583845b6
ggml: update kleidiai to v1.13.0 (llama/15663) 2025-09-20 13:42:44 +03:00
Johannes Gäßler f6ba3949b6
llama: use FA + max. GPU layers by default (llama/15434)
* llama: use max. GPU layers by default, auto -fa

* ggml-backend: abort instead of segfault
2025-09-20 13:42:44 +03:00
Johannes Gäßler b7809c401b
CUDA: use FP32 arithmetic for conv2d (llama/15683) 2025-09-20 13:42:44 +03:00
Jeff Bolz a6dec4f49d
vulkan: Skip syncing for prealloc_y when it is reused (llama/15544) 2025-09-20 13:42:44 +03:00
Chenguang Li d629af157e
CANN: FIx compiler warnings (llama/15661)
Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:44 +03:00
Aman Gupta 82ce91e7d2
CUDA: fix bug in rms_norm fusion (llama/15660)
* CUDA: fix bug in rms_norm fusion

* Fix bug for OP_REPEAT

* Fix index for add
2025-09-20 13:42:44 +03:00
Aman Gupta 6d7ddaf793
CUDA: fuse adds, fuse add with rms norm (llama/15631)
* CUDA: fused add with rms_norm_mul

* Non-broadcast fuse works

* Add fused adds

* format

* Remove n_fuse from template params

* Address review comments

* Move template inside binbcast
2025-09-20 13:42:44 +03:00
mnehete32 dc9f55bbb0
CUDA: add conv2d (llama/15635)
* CUDA: add conv2d

* CUDA: conv2d - correct formatting and added const
2025-09-20 13:42:44 +03:00
Aaron Teo 6287027a2c
ggml-cpu: fix invalid hsum build in debug s390x (llama/15634)
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2025-09-20 13:42:43 +03:00
compilade 6dffbaa0cb
ggml : fix SSM_SCAN for n_groups > 1 (llama/15625) 2025-09-20 13:42:43 +03:00
Georgi Gerganov cac6253744
kv-cache : remove LLAMA_SET_ROWS checks (llama/15505)
ggml-ci
2025-09-20 13:42:43 +03:00
matiaslin 88c0582b61
cuda: Add cublasLt_static linking when GGML_STATIC is enabled (llama/15622)
Prior to this change, we faced undefined cublasLt references when
attempting to compile 'llama-cli' with GGML_STATIC=ON on Linux.

We add linking with CUDA::cublasLt_static when CUDA version is greater
than 10.1.
2025-09-20 13:42:43 +03:00
uvos 65fa2c0c1a
HIP: Enable support for ggml_backend_cuda_register_host_buffer (llama/15615) 2025-09-20 13:42:43 +03:00
Chenguang Li 02e8b23137
CANN: refactor mask handling and improve performance in FA (llama/15561)
* CANN(flash-attn): refactor mask handling and improve performance

1. Refactored the mask computation in Flash Attention, unified the logic without separating prefill and decode.
2. Optimized performance in non-alibi scenarios by reducing one repeat operation.
3. Updated operator management to explicitly mark unsupported cases on 310P devices and when dim is not divisible by 16.

Signed-off-by: noemotiovon <757486878@qq.com>

* [CANN]: fix review

Signed-off-by: noemotiovon <757486878@qq.com>

* [CANN]: Optimization FA BNSD to BSND

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:43 +03:00
xctan ece1bdfe7e
ggml-cpu : add basic RVV support for vector f32 ops (llama/15057)
* ggml-cpu : add basic RVV support for vector f32 ops

* ggml-cpu : add RVV support for f32 softmax
2025-09-20 13:42:43 +03:00
rmatif a6ec224efa
OpenCL: add fused group_norm/norm, mul, add (llama/15314)
* add fused group_norm/norm, mul, add

* fix spacing

* revert rms_norm logic

* fix trailing whitespace
2025-09-20 13:42:43 +03:00
Akarshan Biswas 94fa9f63b3
SYCL: fix rms_norm_mul_add for tensor dim not a multiple of sg_size (llama/15592)
The original implementation unconditionally returned true for this operation, leading to a failure when the tensor's first dimension (ne[0]) was not a multiple of WARP_SIZE. This caused an GGML_ASSERT(ncols % WARP_SIZE == 0) failure in ggml-sycl/norm.cpp.

This change updates the ggml_backend_sycl_device_supports_op check to correctly return true for GGML_OP_RMS_NORM only when the first dimension of the tensor is a multiple of WARP_SIZE, ensuring the operation can be performed without error.
2025-09-20 13:42:42 +03:00
shalinib-ibm 31c7784e09
llamafile: PowerPC Sgemm Optimization (llama/15558)
This patch improves GEMM for FP32 Data Type on PowerPC

Implements GEMM on large blocks with configurable block size mc, nc, kc
(default: 256, 256, 256).
Packing Function optimized to access blocks as per memory layout.
GEMM Optimized to work on larger blocks.
Isolated Packing from GEMM Operations for better MMA utilization.

Verified functionality and correctness uing llama-cli and stand alone
test case (performs matmul and compares final mattrix C result with base).

Minor code refactoring changes:
Replace macro with inline function
Code Indent made consistent with 4 spaces

Performance Testing:

Observed 50% ~ 70% improvement in Prompt Processing Speed mesured using
llama-bench with Meta-Llama3-8B FP32 Model.  Similar gains observed with
Mistral-7b-Instruct-v0.3 Model.

model                   Size                Params     Backend       Threads   Test    Patch   Base
llama 8B all F32        29.92 GiB           8.03 B      CPU           20       pp512   98.58   60.3
llama 8B all F32        29.92 GiB           8.03 B      CPU           20       pp1024  95.88   57.36
llama 8B all F32        29.92 GiB           8.03 B      CPU           20       pp2048  85.46   53.26
llama 8B all F32        29.92 GiB           8.03 B      CPU           20       pp4096  68.66   45.78
llama 8B all F32        29.92 GiB           8.03 B      CPU           20       pp6144  57.35   40.44

25 ~ 30% improvement in llama-batched-bench with Metla-Llama3-8B in
Prompt Processing Speed for large prompts (256, 512, 1024, 2048, 4096)tokens with various batch
sizes ( 1, 2, 4, 8, 16)

Signed-off-by: Shalini Salomi Bodapati <Shalini.Salomi.Bodapati@ibm.com>
2025-09-20 13:42:42 +03:00
Johannes Gäßler 53010199a1
CUDA: return -1 for nonexistent compiled arch (llama/15587) 2025-09-20 13:42:42 +03:00
Georgi Gerganov 1c21a850be
metal : optimize FA vec for large sequences and BS <= 8 (llama/15566)
* metal : optmize FA vec for large heads and sequences

* metal : adjust small-batch mul mv kernels

ggml-ci

* batched-bench : fix total speed computation

ggml-ci

* cont : add comments

ggml-ci
2025-09-20 13:42:42 +03:00
Georgi Gerganov dc693ca8c9
metal : improve `MUL_MAT_ID` (llama/15541)
* metal : mul_mm_id remove hdst

* metal : remove mul_mm_id hsrc1

* metal : mul_mm_id simplify + add test

* metal : opt mul_mm_id map0

* metal : optimize mul_mm_id id gathering

* metal : mul/div opt

* metal : optimize mul_mm_id_map0

ggml-ci
2025-09-20 13:42:42 +03:00
Sigbjørn Skjæret 3bb52acb46
metal : remove contiguous assertion for src0 in IM2COL (llama/15577)
* remove contiguous assertion for src0 in IM2COL

* add contiguous check in supports_op
2025-09-20 13:42:42 +03:00
Yoshi_likes_e4 9828caafb5
Add a warning for special devices (llama/15563)
* Add warning

* Print the devices names

* Add newlines

* Apply suggestions from code review

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

* Fix vector names

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2025-09-20 13:42:42 +03:00
Jeff Bolz 79e2bd5ea8
vulkan: Remove splitting for mul_mat_id (llama/15568)
row_ids only needs to hold the BN rows for the current tile.
2025-09-20 13:42:42 +03:00
Qeeweew 2468074e91
CUDA: Accelerate MXFP4 table lookup using `__byte_perm` (llama/15451)
* CUDA: optimize get_int_from_table_16

* CUDA: use v_perm_b32 to replace byte_perm on AMD GPUs

* revise documentation

---------

Co-authored-by: xix <xiapc@outlook.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2025-09-20 13:42:41 +03:00
lhez 582ef379ab
opencl: fix support ops condition for `rms_norm` (llama/15560) 2025-09-20 13:42:41 +03:00
Ruben Ortlam 335d2a5405
vulkan: fix min subgroup 16 condition for mmid subgroup optimization (llama/15565) 2025-09-20 13:42:41 +03:00
Ihar Hrachyshka 8851ef5463
metal: fix regression when no metal devices are present (llama/15531) 2025-09-20 13:42:41 +03:00
Johannes Gäßler 1e856b2919
CUDA: MoE helper in device code, better tile sizes (llama/15525)
* CUDA: MoE helper in device code, better tile sizes

* reduce superfluous CUDA blocks
2025-09-20 13:42:41 +03:00
Georgi Gerganov 54be54f4ce
metal : add FA kernels for HS=40 (llama/15559)
ggml-ci
2025-09-20 13:42:41 +03:00
Chenguang Li 86331f74e0
CANN: ROPE cache sin/cos repeat (llama/15501)
Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:41 +03:00
Ruben Ortlam ee11ed42a9
vulkan: apply MUL_MAT_ID subgroup optimization to non-coopmat devices (llama/15524)
* vulkan: use subgroup function for mul_mat_id shader even without coopmat

* vulkan: fix compile warnings

* vulkan: properly check for subgroup size control and require full subgroups for subgroup mul_mat_id

* vulkan: disable subgroup mul_mat_id on devices with subgroups < 16
2025-09-20 13:42:41 +03:00
Jeff Bolz 85d4d2c875
vulkan: Support FA with any multiple of 8 head sizes (llama/15537)
The scalar FA shader already handled multiples of 8. The coopmat1 FA
shader assumed 16x16x16 and the shared memory allocations need the HSK
dimensions padded to a multiple of 16. NVIDIA's coopmat2 implementation
requires multiples of 16 for N and K, and needs the matrix dimensions
padded and loads clamped.

Store the FA pipelines in a map, indexed by the pipeline state.
2025-09-20 13:42:40 +03:00
Ruben Ortlam 8c7872d6ed
vulkan: enable Conv2D for Apple after MoltenVK fixed the bug (llama/15526) 2025-09-20 13:42:40 +03:00
Jeff Bolz 27817867cc
vulkan: workaround MoltenVK compile failure in multi_add (llama/15506)
* vulkan: workaround MoltenVK compile failure in multi_add

* Update ggml/src/ggml-vulkan/vulkan-shaders/multi_add.comp

Co-authored-by: 0cc4m <picard12@live.de>
2025-09-20 13:42:40 +03:00
Johannes Gäßler b0d15e1eb6
CUDA: fix half2 -> half conversion for HIP (llama/15529) 2025-09-20 13:42:40 +03:00
Jeff Bolz 2f6288c33c
vulkan: optimize rms_norm, and allow the work to spread across multiple SMs (llama/15281)
* vulkan: optimize rms_norm, and allow the work to spread across multiple SMs

There are really two parts to this change:
(1) Some optimizations similar to what we have in soft_max, to unroll with
different numbers of iterations.
(2) A fusion optimization where we detect add followed by rms_norm, and make
the add shader atomically accumulate the values^2 into memory. Then the
rms_norm shader can just load that sum. This allows the rms_norm to be
parallelized across multiple workgroups, it just becomes a simple per-element
multiply.

The fusion optimization is currently only applied when the rms_norm is on a
single vector. This previously always ran on a single SM. It could apply more
broadly, but when there are other dimensions the work can already spread across
SMs, and there would be some complexity to tracking multiple atomic sums.

* Change add+rms_norm optimization to write out an array of partial sums
rather than using atomic add, to make it deterministic. The rms_norm
shader fetches a subgroup's worth in parallel and uses subgroupAdd to
add them up.

* complete rebase against fused adds - multi_add shader can also compute partial sums

* fix validation errors

* disable add_rms_fusion for Intel due to possible driver bug

* resolve against #15489, sync after clearing partial sums
2025-09-20 13:42:40 +03:00
Jeff Bolz d8eb9f7d67
vulkan: Rewrite synchronization to allow some overlap between nodes (llama/15489)
Track a list of nodes that need synchronization, and only sync if the new node
depends on them (or overwrites them). This allows some overlap which can
improve performance, and centralizes a big chunk of the synchronization logic.

The remaining synchronization logic involves writes to memory other than the
nodes, e.g. for dequantization or split_k. Each of these allocations has a bool
indicating whether they were in use and need to be synced. This should be
checked before they are written to, and set to true after they are done being
consumed.
2025-09-20 13:42:40 +03:00
Acly 5094171c37
vulkan : support ggml_mean (llama/15393)
* vulkan : support ggml_mean

* vulkan : support sum, sum_rows and mean with non-contiguous tensors

* vulkan : fix subbuffer size not accounting for misalign offset

* tests : add backend-op tests for non-contiguous sum_rows

* cuda : require contiguous src for SUM_ROWS, MEAN support
* sycl : require contiguous src for SUM, SUM_ROWS, ARGSORT support

* require ggml_contiguous_rows in supports_op and expect nb00=1 in the shader
2025-09-20 13:42:40 +03:00
Jeff Bolz 485c5c3b3b
vulkan: optimize mul_mat_id loading row ids into shared memory (llama/15427)
- Spread the work across the whole workgroup. Using more threads seems to
far outweigh the synchronization overhead.
- Specialize the code for when the division is by a power of two.
2025-09-20 13:42:40 +03:00
Reese Levine bb5d7e2c31
ggml WebGPU: add support for quantization types (llama/15440)
* Begin work on set_rows

* Work on set rows

* Add error buffers for reporting unsupported SET_ROWS indices

* Remove extra comments

* Work on templating for different types in shaders

* Work on shader type generation

* Working q4_0 mul_mat and some templating for different types

* Add q4_0_f16 matmul and fix device init

* Add matmul support for basic quantization types

* Add q2_k and q3_k quantization

* Add rest of k-quants

* Get firt i-quant working

* Closer to supporting all i-quants

* Support rest of i-quants

* Cleanup code

* Fix python formatting

* debug

* Bugfix for memset

* Add padding to end of buffers on creation

* Simplify bit-shifting

* Update usage of StringView
2025-09-20 13:42:39 +03:00
rmatif d7b7498e76
ggml: add `conv3d` op (llama/15182)
* add conv3d

* bump GGML_OP_COUNT
2025-09-20 13:42:39 +03:00
Yavor Ivanov 18ca4e8f63
cuda : add Pad Reflect 1D support (llama/14659)
* Add Pad Reflect 1D CUDA support

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

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
2025-09-20 13:42:39 +03:00
Aaron Teo 380d3db216
ggml-cpu: Support Q5_0 and Q5_1 on s390x (llama/15486)
* ggml-cpu: initial q5_0 impl for s390x

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: updated q5_0 code for better performance

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: use optimised hsum for better performance

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: introduce q5_1 simd + refactor q5_0

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: fix incorrect return type vec_hsum

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: q5_0 incomplete refactor + table_b2b_0 activation

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: refactor q5_1

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: q5_1 update loop unroll to 4

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: update q5_0 unroll to 4

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: update build-s390x docs

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: update unused variables q5_0

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* docs: update the last update date

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
2025-09-20 13:42:39 +03:00
Chenguang Li be841c3f6e
CANN: Optimize RMS_NORM using cache (llama/15419)
* [CANN] Optimize RMS_NORM using cache

Signed-off-by: noemotiovon <757486878@qq.com>

* fix typo

Signed-off-by: noemotiovon <757486878@qq.com>

* fix review comment

Signed-off-by: noemotiovon <757486878@qq.com>

* codestyle adjustment

Signed-off-by: noemotiovon <757486878@qq.com>

---------

Signed-off-by: noemotiovon <757486878@qq.com>
2025-09-20 13:42:39 +03:00
Diego Devesa 554f96f385
sched : fix possible use of wrong ids tensor when offloading moe prompt processing (llama/15488) 2025-09-20 13:42:39 +03:00
Acly 9dd5039968
vulkan : support conv_2d_dw with f16 weights (llama/15392) 2025-09-20 13:42:39 +03:00
Dong Won Kim 7eebd498ff
vulkan: add exp operation (llama/15456)
Co-authored-by: aeseulgi <kim2h7903@gmail.com>
2025-09-20 13:42:39 +03:00
Jeff Bolz 04d0f9a066
vulkan: Reuse conversion results in prealloc_y (llama/15410)
* vulkan: Reuse conversion results in prealloc_y

Cache the pipeline and tensor that were most recently used to fill prealloc_y,
and skip the conversion if the current pipeline/tensor match.

* don't use shared pointer for prealloc_y_last_pipeline_used
2025-09-20 13:42:38 +03:00
Xuan-Son Nguyen c5874bcf42
ggml : fix condition of im2col on Metal backend (llama/15460) 2025-09-20 13:42:38 +03:00
R0CKSTAR 7c077845fd
musa: add GGML_UNUSED_VARS (llama/15446)
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
2025-09-20 13:42:38 +03:00
Diego Devesa 622dec5bf6
sched : copy only the used experts when offloading prompt processing (llama/15346) 2025-09-20 13:42:38 +03:00
Johannes Gäßler 8f0579a33d
CUDA: refactor FA support/selection code (llama/15454) 2025-09-20 13:42:38 +03:00
Johannes Gäßler 316ed78d68
CUDA: replace GGML_CUDA_F16 with CUDA arch checks (llama/15433) 2025-09-20 13:42:38 +03:00