e4abf643fa
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# ggml/src/ggml-hexagon/htp/act-ops.c
# ggml/src/ggml-rpc/ggml-rpc.cpp
# src/CMakeLists.txt
# src/llama-vocab.cpp
2026-01-03 15:37:30 +08:00
bcfc8c3cec
ggml-hexagon: optimize activation function ( #18393 )
...
Check Pre-Tokenizer Hashes / pre-tokenizer-hashes (push) Failing after 6s
Python check requirements.txt / check-requirements (push) Failing after 8s
Python Type-Check / pyright type-check (push) Failing after 7s
* refactor: refactor silu
* refactor: optimize swiglu
* refactor: remove unncessary if in swiglu
* refactor: refactor swiglu_oai
* chore: fix formatting issue
2026-01-02 21:24:24 -08:00
18ddaea2ae
vulkan: Optimize GGML_OP_CUMSUM ( #18417 )
...
* vulkan: Optimize GGML_OP_CUMSUM
There are two paths: The preexisting one that does a whole row per workgroup
in a single shader, and one that splits each row into multiple blocks and does
two passes. The first pass computes partials within a block, the second adds
the block partials to compute the final result. The multipass shader is used
when there are a small number of large rows.
In the whole-row shader, handle multiple elements per invocation.
* use 2 ELEM_PER_THREAD for AMD/Intel
* address feedback
2026-01-02 15:32:30 -06:00
706e3f93a6
vulkan: Implement mmvq for iq1_s/iq1_m ( #18450 )
2026-01-02 20:19:04 +01:00
f38de16341
metal : adjust extra size for FA buffer to avoid reallocations ( #18545 )
2026-01-02 19:02:18 +02:00
c6f0e832da
rpc : use unordered_map::reserve and emplace ( #18513 )
2026-01-02 12:09:36 +02:00
7e1ae49e7d
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# ggml/src/ggml-cuda/ggml-cuda.cu
# tests/test-backend-ops.cpp
# tools/mtmd/CMakeLists.txt
2026-01-02 11:05:20 +08:00
e86f3c2221
cuda : fix copy of large tensors (ggml_nbytes <= INT_MAX assertion) ( #18433 )
...
Check Pre-Tokenizer Hashes / pre-tokenizer-hashes (push) Failing after 7s
Python check requirements.txt / check-requirements (push) Failing after 6s
Python Type-Check / pyright type-check (push) Failing after 8s
* ggml-cuda: fixed assertion in ggml_cuda_cpy (#18140 )
* ggml-cuda: changes in data types to int64_t
* ggml-cuda: added asserts for CUDA block numbers
* ggml-cuda: changed the condition for y and z dimension
2026-01-02 00:24:20 +01:00
26831bded9
ggml-cuda: remove unneccesary prints on ggml_cuda_init ( #18502 )
2026-01-01 19:18:43 +08:00
be47fb9285
vulkan: extend topk_moe to handle sigmoid w/exp_probs_b for nemotron ( #18295 )
...
* vulkan: extend topk_moe to handle sigmoid w/exp_probs_b for nemotron
Also handle GGML_OP_SCALE at the end (nemotron, deepseek2).
Fewer pipeline variants and spec constants, just use push constants.
In test_topk_moe, change exp_probs_b to be 1D, matching real networks.
Update test-backend-ops and ggml-backend to allow verifying multiple outputs
in a fusion test (topk_moe has two outputs). Previously only the final node
was verified.
* change test_topk_moe to allow results in arbitrary order
* disable sigmoid fusion for moltenvk
2026-01-01 08:58:27 +01:00
54e419f587
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/workflows/docker.yml
# docs/ops.md
# docs/ops/Metal.csv
# ggml/CMakeLists.txt
# ggml/src/ggml-sycl/CMakeLists.txt
# grammars/README.md
# models/templates/llama-cpp-deepseek-r1.jinja
# scripts/sync-ggml.last
# tests/test-chat.cpp
2026-01-01 15:34:10 +08:00
66ccf8f6b8
Merge commit 'f14f4e421b2177fadcf9d15ebccb0492e5464d86' into concedo_experimental
...
# Conflicts:
# .github/workflows/docker.yml
# AGENTS.md
# CONTRIBUTING.md
# docs/build.md
# examples/llama.android/app/build.gradle.kts
# examples/llama.android/app/src/main/java/com/example/llama/MainActivity.kt
# examples/llama.android/app/src/main/res/layout/activity_main.xml
# examples/llama.android/gradle/libs.versions.toml
# examples/llama.android/lib/src/main/cpp/ai_chat.cpp
# examples/llama.android/lib/src/main/java/com/arm/aichat/InferenceEngine.kt
# examples/llama.android/lib/src/main/java/com/arm/aichat/internal/InferenceEngineImpl.kt
# examples/model-conversion/scripts/causal/compare-embeddings-logits.sh
# examples/model-conversion/scripts/embedding/run-original-model.py
# examples/retrieval/retrieval.cpp
# ggml/src/CMakeLists.txt
# ggml/src/ggml-cpu/CMakeLists.txt
# ggml/src/ggml-cpu/kleidiai/kernels.cpp
# ggml/src/ggml-cpu/kleidiai/kleidiai.cpp
# ggml/src/ggml-cuda/CMakeLists.txt
# ggml/src/ggml-cuda/mmq.cu
# ggml/src/ggml-cuda/mmq.cuh
# src/CMakeLists.txt
# tools/llama-bench/llama-bench.cpp
# tools/server/CMakeLists.txt
2026-01-01 15:20:56 +08:00
54f67b9b66
ggml : bump version to 0.9.5 (ggml/1410)
2025-12-31 18:54:43 +02:00
9a6369bb60
metal : add count_equal op ( #18314 )
...
* add count equal for metal
* remove trailing whitespace
* updated doc ops table
* changed shmem to i32
* added multi tg and templating
* removed BLAS support from Metal docs
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com >
* add memset to set dst to 0
* metal : cleanup
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com >
2025-12-31 10:39:48 +02:00
ecc343de63
CUDA: fix KQ max calculation ( #18487 )
2025-12-31 09:37:00 +01:00
01ade96e71
metal : remove BF16 x F16 kernels ( #18456 )
2025-12-31 09:53:48 +02:00
7bcaf815c2
sycl: add newline at the end of CMakeLists.txt ( #18503 )
2025-12-31 14:23:44 +08:00
c8a3798041
Work around broken IntelSYCLConfig.cmake in Intel oneAPI 2025.x ( #18345 )
...
* cmake: work around broken IntelSYCLConfig.cmake in oneAPI 2025.x
* [AI] sycl: auto-detect and skip incompatible IntelSYCL package
Automatically detect compiler versions with incompatible IntelSYCL
CMake configuration files and fall back to manual SYCL flags instead
of requiring users to set options manually.
Fixes build failures with oneAPI 2025.x where IntelSYCLConfig.cmake
has SYCL_FEATURE_TEST_EXTRACT invocation errors.
* refactor: improve SYCL provider handling and error messages in CMake configuration
* refactor: enhance SYCL provider validation and error handling in CMake configuration
* ggml-sycl: wrap find_package(IntelSYCL) to prevent build crashes
2025-12-31 09:08:44 +08:00
2d6c00a9b8
kleidiai: add and integrate SVE 256-bit vector-length kernel ( #18458 )
...
* kleidiai: add and integrate SVE 256-bit vector-length kernel
* updated for review comments
2025-12-30 14:04:53 +02:00
d77d7c5c06
CUDA: add log line when mxfp4 acceleration is used ( #18483 )
...
* CUDA: add log line when mxfp4 acceleration is used
* add in backend_get_features
2025-12-30 17:40:46 +08:00
0bd1212a43
CUDA: fix replacment of bad archs in CMake ( #18457 )
2025-12-29 17:58:20 +01:00
e70e640db3
CUDA: Blackwell features for non-native builds ( #18436 )
2025-12-29 09:35:42 +01:00
5fa66c6e67
cuda: fix race condition in cumsum ( #18448 )
...
* ggml-cuda: fix race condition in cumsum
* remove unneccesary sync_threads
2025-12-29 14:07:17 +08:00
4ffc47cb20
HIP: Use mmq on MFMA devices for MUL_MAT_ID in cases where a lot of splits would be generated ( #18202 )
2025-12-28 20:12:55 +01:00
0e26e4d354
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/ISSUE_TEMPLATE/010-bug-compilation.yml
# .github/ISSUE_TEMPLATE/011-bug-results.yml
# .github/ISSUE_TEMPLATE/019-bug-misc.yml
# ggml/CMakeLists.txt
# ggml/src/CMakeLists.txt
# ggml/src/ggml-cuda/CMakeLists.txt
# ggml/src/ggml-opencl/ggml-opencl.cpp
# ggml/src/ggml-rpc/ggml-rpc.cpp
2025-12-28 23:47:55 +08:00
82d562ad7b
unstable merge
2025-12-28 23:03:03 +08:00
9082403a43
disable vk events until directio pr or jeff's fix is added. (+1 squashed commits)
...
Squashed commits:
[4796db21a] disable vk events until directio pr or jeff's fix is added.
2025-12-28 21:54:25 +08:00
a94d5ffbec
Revert "Triage: revert https://github.com/ggml-org/llama.cpp/pull/18047 and https://github.com/ggml-org/llama.cpp/pull/18302 "
...
This reverts commit dfa1b72d2f
2025-12-28 21:48:55 +08:00
07a0c4ba92
Revert "ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON ( #18413 )" ( #18426 )
2025-12-28 20:53:36 +08:00
60f17f56da
rpc: fix segfault on invalid endpoint format ( #18387 )
...
* rpc: fix segfault on invalid endpoint format
* rpc: add error log for failed endpoint connection
2025-12-28 12:34:41 +02:00
94de74e7b1
cmake: Added more x86_64 CPU backends when building with GGML_CPU_ALL_VARIANTS=On ( #18186 )
...
* minor: Consolidated `#include <immintrin.h>` under `ggml-cpu-impl.h`
* cmake: Added more x86-64 CPU backends when building with `GGML_CPU_ALL_VARIANTS=On`
- `ivybridge`
- `piledriver`
- `cannonlake`
- `cascadelake`
- `cooperlake`
- `zen4`
Resolves : #17966
2025-12-28 09:33:29 +02:00
4fd59e8427
ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON ( #18413 )
2025-12-28 09:33:14 +08:00
08566977a7
opencl: allow resizing transpose buffers ( #18384 )
...
* opencl: allow resizing transpose buffers instead of using fixed sizes
* opencl: remove commented code
2025-12-27 15:51:14 -08:00
06705fdcb3
ggml-cuda: Use same regex for GGML_NATIVE=OFF ( #18407 )
2025-12-27 19:56:27 +08:00
c9ced4910b
vulkan: preprocess mul_mat_id experts and discard workgroups more quickly ( #18352 )
...
Run a preprocess to count how many times each expert is used, and use this to
quickly discard workgroups that aren't needed.
2025-12-26 16:12:58 -06:00
7ac8902133
vulkan: optimize decodeFuncB in coopmat2 mul_mat_id shader ( #18349 )
...
* vulkan: Use BK=32 for coopmat2 mul_mat_id
* vulkan: optimize decodeFuncB in coopmat2 mul_mat_id shader
Disable robustness, remove the OOB check in decodeFuncB, and initialize the
row_ids to zero to avoid OOB access.
Don't slice/offset the B matrix to ic * BN, only to adjust the coord back down
to the range [0, BN) in decodeFuncB. Instead just slice with a row offset of
zero and remove the '& (BN - 1)'. This allows the compiler to common some of
the shared memory loads.
2025-12-26 18:15:50 +01:00
9bf20d8ac3
vulkan: Use BK=32 for coopmat2 mul_mat_id ( #18332 )
2025-12-26 18:15:02 +01:00
cb999704fb
vulkan: small dequantization improvements ( #18380 )
...
* iq4_xs
* quants
2025-12-26 18:12:11 +01:00
b96b82fc85
vulkan: Support UPSCALE w/antialias ( #18327 )
2025-12-26 17:00:57 +01:00
10dc500bdb
vulkan: handle rope with large number of rows ( #18306 )
2025-12-26 16:53:46 +01:00
b07cda687c
CANN: implement the SSM_CONV operator ( #17737 )
...
* CANN: implement SSM_CONV operator
Co-authored-by: Aleksei Lobanov, <zeromarblectm@gmail.com >
Co-authored-by: Sujin Kang, <waterjin326@gmail.com >
* CANN: remove custom error limit for SSM_CONV
* CANN: merge SSM_CONV tensor shape/strides into one line
---------
Co-authored-by: Sujin Kang, <waterjin326@gmail.com >
2025-12-26 09:12:04 +08:00
85c40c9b02
ggml-cuda: fix regex for arch list ( #18371 )
...
* ggml-cuda: fix regex for arch list
* make regex exact
2025-12-26 01:35:14 +08:00
dfa1b72d2f
Triage: revert https://github.com/ggml-org/llama.cpp/pull/18047 and https://github.com/ggml-org/llama.cpp/pull/18302
...
Revert "vulkan: Implement set_tensor_async and the event interfaces (#18047 )"
This reverts commit e1f15b454f#18302 )"
This reverts commit 2a9ea2020c
2025-12-26 01:20:31 +08:00
83b3b1c271
cuda: optimize cumsum cub path ( #18362 )
...
Check Pre-Tokenizer Hashes / pre-tokenizer-hashes (push) Successful in 55m39s
Python check requirements.txt / check-requirements (push) Failing after 16s
Python Type-Check / pyright type-check (push) Failing after 1m23s
* cuda: optimize cumsum cub path
* remove heavy perf test
2025-12-25 23:55:38 +08:00
b0fb0f0aee
ggml-cuda: fix blackwell native builds ( #18361 )
...
* ggml-cuda: fix blackwell native builds
Replace 12x in native architectures by 12xa
* replace for GGML_NATIVE=OFF too
* only replace for native
* remove 120f-virtual for default compilation
---------
Co-authored-by: Aman Gupta <aman>
2025-12-25 22:12:11 +08:00
e68c19b0fd
CANN: Add support for CONV_TRANSPOSE_1D when kernel size > 255 ( #17934 )
...
* CONV_TRANSPOSE_1D kernel_size>255
* remove condition check
* fix the bug of type conversion
* removing trailing whitespaces
* fix: return true in the switch case
2025-12-25 16:46:09 +08:00
c54bba869d
ggml : optimize cuda cumsum fallback kernel ( #18343 )
2025-12-25 12:11:13 +08:00
6cc71db85a
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# docs/backend/SYCL.md
# examples/model-conversion/Makefile
# examples/model-conversion/scripts/causal/run-org-model.py
# ggml/src/ggml-cann/aclnn_ops.cpp
# ggml/src/ggml-cann/common.h
# ggml/src/ggml-cann/ggml-cann.cpp
# ggml/src/ggml-cuda/CMakeLists.txt
2025-12-25 00:06:27 +08:00
3589a5e136
Merge commit '12ee1763a6f6130ce820a366d220bbadff54b818' into concedo_experimental
...
# Conflicts:
# docs/backend/hexagon/README.md
# docs/backend/hexagon/developer.md
# examples/gen-docs/gen-docs.cpp
# examples/model-conversion/scripts/embedding/run-original-model.py
# examples/model-conversion/scripts/utils/semantic_check.py
# examples/sycl/run-llama2.sh
# examples/sycl/run-llama3.sh
# examples/sycl/win-run-llama2.bat
# examples/sycl/win-run-llama3.bat
# ggml/src/ggml-hexagon/ggml-hexagon.cpp
# ggml/src/ggml-hexagon/htp-utils.h
# ggml/src/ggml-hexagon/htp/act-ops.c
# ggml/src/ggml-hexagon/htp/htp-dma.c
# ggml/src/ggml-hexagon/htp/htp-dma.h
# ggml/src/ggml-hexagon/htp/hvx-utils.h
# ggml/src/ggml-hexagon/htp/main.c
# ggml/src/ggml-hexagon/htp/matmul-ops.c
# ggml/src/ggml-opencl/ggml-opencl.cpp
# ggml/src/ggml-opencl/kernels/cvt.cl
# ggml/src/ggml-opencl/kernels/transpose.cl
# ggml/src/ggml-rpc/ggml-rpc.cpp
# scripts/snapdragon/adb/run-cli.sh
# src/CMakeLists.txt
# tests/test-backend-ops.cpp
# tools/cli/README.md
# tools/completion/README.md
# tools/server/README.md
2025-12-24 23:57:41 +08:00
d1983959d2
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/workflows/release.yml
# AGENTS.md
# common/CMakeLists.txt
# docs/development/parsing.md
# ggml/src/ggml-rpc/ggml-rpc.cpp
# ggml/src/ggml-vulkan/ggml-vulkan.cpp
# tests/test-arg-parser.cpp
# tests/test-backend-ops.cpp
# tests/test-grammar-llguidance.cpp
# tests/test-tokenizer-0.cpp
# tests/test-tokenizer-1-bpe.cpp
# tests/test-tokenizer-1-spm.cpp
# tools/batched-bench/batched-bench.cpp
# tools/cli/cli.cpp
# tools/llama-bench/llama-bench.cpp
# tools/server/README.md
2025-12-24 23:42:28 +08:00
Concedo
Shouyu
Jeff Bolz
Georgi Gerganov
Chris Rohlf
MeeMin
Aman Gupta
gatbontonpc
Johannes Gäßler
Rahul Sathe
Charles Xu
uvos
o7si
Boian Berberov
QDelta
lhez
Eve
0Marble
Penglin Cai
Aadeshveer Singh