6adcd0b5db
Merge commit '34df42f7bef5a711b2b40f5d2b6b78254def99c3' into concedo_experimental
...
# Conflicts:
# README.md
# ggml/src/ggml-hexagon/ggml-hexagon.cpp
# ggml/src/ggml-hexagon/htp/CMakeLists.txt
# ggml/src/ggml-hexagon/htp/act-ops.c
# ggml/src/ggml-hexagon/htp/binary-ops.c
# ggml/src/ggml-hexagon/htp/cpy-ops.c
# ggml/src/ggml-hexagon/htp/get-rows-ops.c
# ggml/src/ggml-hexagon/htp/htp-msg.h
# ggml/src/ggml-hexagon/htp/htp-ops.h
# ggml/src/ggml-hexagon/htp/hvx-arith.h
# ggml/src/ggml-hexagon/htp/hvx-base.h
# ggml/src/ggml-hexagon/htp/hvx-inverse.h
# ggml/src/ggml-hexagon/htp/hvx-utils.h
# ggml/src/ggml-hexagon/htp/main.c
# ggml/src/ggml-hexagon/htp/rope-ops.c
# ggml/src/ggml-hexagon/htp/set-rows-ops.c
# ggml/src/ggml-hexagon/htp/softmax-ops.c
# ggml/src/ggml-hexagon/htp/unary-ops.c
# ggml/src/ggml-opencl/CMakeLists.txt
# ggml/src/ggml-opencl/ggml-opencl.cpp
# tests/test-backend-ops.cpp
# tools/cli/cli.cpp
# tools/server/webui/src/lib/components/app/chat/ChatScreen/ChatScreen.svelte
2026-03-10 22:20:04 +08:00
746664fde6
Merge commit '2cd20b72ed3565ac6935911ca0d9b5d73ae70d0d' into concedo_experimental
...
# Conflicts:
# CONTRIBUTING.md
# docs/backend/CANN.md
# docs/backend/SYCL.md
# docs/backend/snapdragon/README.md
# docs/backend/snapdragon/windows.md
# docs/build.md
# docs/multimodal/MobileVLM.md
# docs/ops.md
# docs/ops/WebGPU.csv
# examples/debug/README.md
# examples/llama.vim
# examples/model-conversion/README.md
# examples/sycl/README.md
# ggml/src/ggml-cpu/amx/mmq.cpp
# ggml/src/ggml-cpu/arch/x86/repack.cpp
# ggml/src/ggml-hexagon/ggml-hexagon.cpp
# ggml/src/ggml-hexagon/htp-drv.cpp
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c
# ggml/src/ggml-hexagon/htp/hvx-base.h
# ggml/src/ggml-hexagon/htp/hvx-copy.h
# ggml/src/ggml-hexagon/htp/hvx-inverse.h
# ggml/src/ggml-hexagon/htp/hvx-reduce.h
# ggml/src/ggml-hexagon/htp/matmul-ops.c
# ggml/src/ggml-hexagon/htp/rope-ops.c
# ggml/src/ggml-hexagon/htp/worker-pool.c
# ggml/src/ggml-opencl/ggml-opencl.cpp
# ggml/src/ggml-opencl/kernels/cpy.cl
# ggml/src/ggml-sycl/common.hpp
# ggml/src/ggml-sycl/quants.hpp
# ggml/src/ggml-sycl/softmax.cpp
# ggml/src/ggml-vulkan/CMakeLists.txt
# ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp
# ggml/src/ggml-webgpu/ggml-webgpu.cpp
# scripts/pr2wt.sh
# scripts/server-bench.py
# scripts/snapdragon/windows/run-cli.ps1
# tests/test-alloc.cpp
# tests/test-backend-ops.cpp
# tests/test-chat.cpp
# tools/cli/cli.cpp
# tools/completion/README.md
# tools/cvector-generator/cvector-generator.cpp
# tools/imatrix/README.md
# tools/perplexity/README.md
# tools/server/public_simplechat/readme.md
# tools/server/tests/README.md
2026-03-10 22:11:08 +08:00
a40038d8e6
further reverse the mxfp4 changes
2026-03-07 22:42:22 +08:00
34df42f7be
hexagon: add f32 ssm_conv op ( #20122 )
...
* hexagon: add ssm_conv op
* hexagon: hvx kernel is functional
* hexagon: improvements to ssm-conv hvx kernel
* hexagon: added dma to ssm-conv hvx kernel
* hexagon: ssm-conv dynamically compute gather scratchpad
* hex-ssm-conv: add local context and fix various issues (spad indexing, etc)
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com >
2026-03-06 09:59:26 -08:00
ba2fd11cdf
cpu: skip redudant ROPE cache updates ( #20149 )
2026-03-06 08:32:40 -08:00
d48e876467
ggml-cuda: add mem check for fusion ( #19916 )
...
* ggml-cuda: add mem check for fusion
* Replace NaNs with -FLT_MAX
* fix typo
Co-authored-by: Johannes Gäßler <johannesg@5d6.de >
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de >
2026-03-07 00:05:43 +08:00
ba2ff79e43
ggml: update comments for backends which have no memory to report ( #20157 )
...
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com >
2026-03-06 23:24:38 +08:00
c6980ff29d
ggml-cpu: Fix gcc 15 ICE on ppc64le ( #20083 ) ( #20130 )
...
This patch addresses an Internal Compiler Error (Segmentation fault)
observed with gcc 15 by replacing the intrinsic + cast by doing
a cat on the data first and then calling the intrinsic. This bypasses the
buggy compiler path while maintaining identical instruction selection.
Performance Verification:
Assembly analysis on RHEL 9 (GCC 15.1.1) confirms that both the original
code and this fix generate the identical Power10 prefixed load instruction:
`plxv 40, 2(14)`
This ensures zero performance regression while unblocking builds on
newer toolchains.
Reproduced on:
- Alpine Linux + GCC 15.2.0-r2
- RHEL 9 + GCC 15.1.1 (gcc-toolset-15)
Signed-off-by: Shalini Salomi Bodapati <Shalini.Salomi.Bodapati@ibm.com >
2026-03-06 23:22:39 +08:00
1e38a7a6fa
CUDA: use shared mem for ssm_conv ( #20128 )
...
* CUDA: use shared mem for ssm_conv
* fuse silu + ssm_conv
* fuse unary + mul
* enable for fp16
* formatting
Co-authored-by: Johannes Gäßler <johannesg@5d6.de >
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de >
2026-03-06 23:09:59 +08:00
d20e60ddd5
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# docs/build.md
# examples/batched/batched.cpp
# examples/convert-llama2c-to-ggml/convert-llama2c-to-ggml.cpp
# examples/deprecation-warning/deprecation-warning.cpp
# examples/eval-callback/eval-callback.cpp
# examples/gen-docs/gen-docs.cpp
# examples/gguf-hash/gguf-hash.cpp
# examples/gguf/gguf.cpp
# examples/lookahead/lookahead.cpp
# examples/lookup/lookup-create.cpp
# examples/lookup/lookup-merge.cpp
# examples/lookup/lookup-stats.cpp
# examples/lookup/lookup.cpp
# examples/parallel/parallel.cpp
# examples/passkey/passkey.cpp
# examples/retrieval/retrieval.cpp
# examples/save-load-state/save-load-state.cpp
# examples/simple-chat/simple-chat.cpp
# examples/simple/simple.cpp
# examples/speculative-simple/speculative-simple.cpp
# examples/speculative/speculative.cpp
# examples/sycl/ls-sycl-device.cpp
# examples/training/finetune.cpp
# ggml/src/ggml-cpu/CMakeLists.txt
# ggml/src/ggml-cpu/amx/common.h
# ggml/src/ggml-cpu/kleidiai/kernels.cpp
# ggml/src/ggml-opencl/CMakeLists.txt
# ggml/src/ggml-opencl/ggml-opencl.cpp
# ggml/src/ggml-opencl/kernels/cvt.cl
# ggml/src/ggml-opencl/kernels/gemv_noshuffle_general_q8_0_f32.cl
# ggml/src/ggml-opencl/kernels/transpose.cl
# ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp
# ggml/src/ggml-webgpu/ggml-webgpu.cpp
# ggml/src/ggml-webgpu/wgsl-shaders/mul_mat.wgsl
# ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_reg_tile.wgsl
# ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_subgroup_matrix.wgsl
# scripts/get-wikitext-2.sh
# tests/test-backend-ops.cpp
# tools/batched-bench/batched-bench.cpp
# tools/cvector-generator/cvector-generator.cpp
# tools/export-lora/export-lora.cpp
# tools/imatrix/imatrix.cpp
# tools/llama-bench/llama-bench.cpp
# tools/perplexity/perplexity.cpp
# tools/rpc/rpc-server.cpp
# tools/tokenize/tokenize.cpp
2026-03-06 21:19:49 +08:00
2c38638b3d
Merge commit '2afcdb9777b1bac79fa4bfe284b9bf23085b0469' into concedo_experimental
...
# Conflicts:
# scripts/sync_vendor.py
# tests/CMakeLists.txt
2026-03-06 21:13:15 +08:00
abcca8c0f9
do not use the mxfp4 repack - repack must be synced again from before this commit if it's ever to be used in future. this will break compilation with older w64devkit
2026-03-06 21:07:41 +08:00
2fbc3b2ae5
Adjust int types in format strings ( #2009 )
...
* tweak format sting types
This may not be all of them, but it's the ones which warn on OpenBSD
* complete the changes needed to fix the format string specifers
* avoid using inttypes, directly cast to size_t (u64 usually) instead
---------
Co-authored-by: Concedo <39025047+LostRuins@users.noreply.github.com >
2026-03-06 19:06:18 +08:00
2850bc6a13
ggml-cpu: fix data race for debug asserts ( #20148 )
2026-03-06 09:12:49 +01:00
6c97bffd65
opencl: add neg, exp and diag ( #20127 )
...
* opencl: add `neg`
* opencl: add `exp`
* opencl: add `diag`
2026-03-05 21:16:39 -08:00
2b10b62677
hexagon: add fp16 support for binary ops: add,sub,mul,div ( #20139 )
...
* hexagon: add fp16 support for binary ops: add,sub,mul,div
* hexagon: fix test-backend-ops failures for fp16 binary ops on older arches (<v79)
* hexagon: decide on n_threads (aka n_jobs) early to avoid overallocating scratchpad
* snapdragon: fix readme link
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com >
2026-03-05 18:29:13 -08:00
2cd20b72ed
CUDA: Improve performance via less synchronizations between token ( #17795 )
...
* Adds CPU-to-CUDA copy capability to
ggml_backend_cuda_cpy_tensor_async()
* Adds function to relax sync requirements between input copies on
supported backends (CUDA for now)
* Exchanges synchronous copy with async copy function.
* Adds macro guards to allow compilation in non-CUDA builds
* Reworked backend detection in ggml-backend.cpp to avoid linking
conflicts
* Relax requirement of checks in async CUDA copies from backend and buffer type to just buffer type, to avoid linking issues
* Minor cleanup
* Makes opt-in to relax use of explicit syncs more general. Backends like
vulkan which require a synchronization between HtoD copies and graph
execution could also adopt this change now.
* Reintroduces stricter check for CPU->CUDA backend async copy via
GGML_DEVICE_TYPE_CPU.
* Corrects initialization of ggml_backend_sync_mode in
ggml_backend_sched_split initialization
* Simplifies synchronizations to adhere to `saaasg` pattern.
* Apply suggestion from @ggerganov (src->buffer to buf_src)
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com >
* Apply suggestion from @ggerganov (src->buffer to buf_src) v2
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com >
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com >
2026-03-05 13:53:21 +02:00
92f7da00b4
chore : correct typos [no ci] ( #20041 )
...
* fix(docs): correct typos found during code review
Non-functional changes only:
- Fixed minor spelling mistakes in comments
- Corrected typos in user-facing strings
- No variables, logic, or functional code was modified.
Signed-off-by: Marcel Petrick <mail@marcelpetrick.it >
* Update docs/backend/CANN.md
Co-authored-by: Aaron Teo <taronaeo@gmail.com >
* Revert "Auxiliary commit to revert individual files from 846d1c301281178efbc6ce6060ad34c1ebe45af8"
This reverts commit 02fcf0c7db661d5ff3eff96b2b2db9fdb7213256.
* Update tests/test-backend-ops.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com >
* Update tests/test-backend-ops.cpp
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com >
---------
Signed-off-by: Marcel Petrick <mail@marcelpetrick.it >
Co-authored-by: Aaron Teo <taronaeo@gmail.com >
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com >
2026-03-05 08:50:21 +01:00
7a99dc85e2
hexagon: Flash Attention optimizations (dma, mpyacc, multi-row) and MatMul updates ( #20118 )
...
* ggml-hexagon: enhance hvx_dot_f16_f16_aa_rx4 for improved performance by expanding vector handling and optimizing accumulation
# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c
* ggml-hexagon: optimize hvx_dot_f16_f16_aa_rx4 and enhance hvx_vec_reduce_sum_f32x4 for improved performance and reduced complexity
* ggml-hexagon: add hvx_dot_f16_f16_aa_rx32 for enhanced vector processing in flash attention
# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c
* optimize hvx_dot_f16_f16_aa_rx4 and hvx_dot_f16_f16_aa_rx32 by removing unused scale parameter and improving vector accumulation
# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c
* ggml-hexagon: refactor hvx_dot_f16_f16_aa_rx4 for improved readability and return HVX_Vector for better integration
# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c
* ggml-hexagon: initialize sums variable in hvx_dot_f16_f16_aa_rx32 for clarity
* ggml-hexagon: fix compiling error
* fix hvx_dot_f16_f16_aa_rx4 to handle leftover elements correctly using masking
* refactor hvx_dot_f16_f16_aa_rx4 to accept vector and leftover element counts as parameters for improved clarity and flexibility
* wip
* fa: instrumentation and dma reordering
* hex-fa: use block-size 64 to improve DMA pipelining
* hex-fa: optimize vec-dot for v79 and above
* hex-fa: use block size 64
* hex-fa: avoid scalar fp32->fp16 conversions
* hex-fa: simplify dot_f16 functions using optimized vec_mpyacc
* hex-fa: rewrite mad_f32_f16 using hvx_vec_mpyacc
* hex-mm: use mpyacc in matmul dot functions
---------
Co-authored-by: chraac <chraac@gmail.com >
2026-03-04 21:55:29 -08:00
69fd345335
opencl: add SET, support i32 for CPY, minor refactor for cpy ( #20101 )
2026-03-04 21:32:26 -08:00
24d2ee0527
[WebGPU] Fix wait logic for inflight jobs ( #20096 )
...
* Enable tmate debugging for investigating thread safety issue
* Refactor wait and submit to operate on vector<wgpu::FutureWaitInfo>, and fix wait to delete only the future that is completed.
* Cleanup
* Remove clear change and run clang-format
* Cleanup
2026-03-04 11:54:55 -08:00
541bf37622
Add concat op to webgpu. ( #20068 )
2026-03-04 11:19:00 -08:00
7f5ee54968
ggml: fix ggml_is_contiguous_n for ne == 1 ( #20092 )
2026-03-04 12:04:31 +01:00
66199c9f03
ggml : use a simple std::thread in AMX without OpenMP ( #20074 )
...
Disabling OpenMP generally provides better inference performance (at
least in my testing) but the loading becomes slightly slower.
Benchmark results for `convert_B_packed_format()`:
Before this commit:
N K | No OpenMP OpenMP | Diff | Speedup
------------------------------------------------------------
512 2880 | 640.9us 263.5us | -58.9% | 0.41x
2880 4096 | 2.55ms 261.7us | -89.8% | 0.10x
201088 2880 | 256.44ms 21.61ms | -91.6% | 0.08x
------------------------------------------------------------
Total: 325.43ms vs 31.05ms
After:
N K | No OpenMP OpenMP | Diff | Speedup
------------------------------------------------------------
512 2880 | 1.49ms 263.5us | -82.3% | 0.18x
2880 4096 | 1.55ms 261.7us | -83.1% | 0.17x
201088 2880 | 24.03ms 21.61ms | -10.1% | 0.90x
------------------------------------------------------------
Total: 78.97ms vs 31.05ms
Tested with unsloth/gpt-oss-20b-GGUF:Q4_K_M.
Signed-off-by: Adrien Gallouët <angt@huggingface.co >
2026-03-04 11:57:09 +01:00
137435ff15
kleidiai : add sme fp16 compute path for q4_0 gemm on aarch64 ( #20043 )
2026-03-03 11:40:26 +02:00
24350fdf9b
opencl: add optimized q4_1 mm kernel for adreno ( #19840 )
...
* Add Q4_1 OpenCL Kernels
* opencl: refactor transpose
* opencl: format
* opencl: refactor q4_1 unpack
* opencl: move `ggml_cl_mul_mat_q4_1_f32_adreno`
* opencl: refactor `ggml_cl_mul_mat_q4_1_f32_adreno` and kernels
* opencl: rename kernel files and kernes
* opencl: fix build for non adreno
* opencl: move code around and format
---------
Co-authored-by: Li He <lih@qti.qualcomm.com >
2026-03-02 19:49:41 -08:00
49a7564ac1
ggml webgpu: fix workgroup dispatch limit for large batch sizes ( #19965 )
...
* ggml-webgpu: fix workgroup dispatch limit for large batch sizes
WebGPU limits workgroup sizes to 65535 per dimension. Large MUL_MAT
operations with batch sizes exceedeing this limi would fail.
* add compute_2d_workgroups() helper to split total workgroup ID across
X/Y dimensions
* update mul_mat_reg_tile.wgsl to reconstruct linear workgroup ID from 2D
dispatch
* update mul_mat_subgroup_matrix.wgsl to reconstruct linear workgroup ID
from 2D dispatch
* update mul_mat.wgsl to compute global index from 2D workgroup
coordinates
* refactor all three mul_mat dispatch paths to use the shared helper
* ggml-webgpu: add bounds checking for over-dispatched workgroups
2D workgroup dispatch can over-dispatch when total workgroups don't
divide evenly into the 65535 per-dimension limit. Extra workgroups
would compute invalid batch indices, causing memory corruption.
* add batch_idx bound check to mul_mat_reg_tile.wgsl and
mul_mat_subgroup_matrix.wgsl to prevent over-dispatched workgroups
from accessing invalid memory
* fixes test failures with large batch sizes (eg., bs=[128, 1024])
* ggml-webgpu: add back TODO for spliting large sizes into batches
* Optimize 2d workgroup provisioning
* Set some parameters that increase speed
---------
Co-authored-by: Reese Levine <reeselevine1@gmail.com >
2026-03-02 19:35:11 -08:00
4d828bd1ab
ggml webgpu: Clean up per-thread parameter buffer pool and job submission logic ( #19772 )
...
* Allow webgpu_buf_pool to resize if needed, remove inflight_threads, and replace inflight_threads with num_kernels for submission
* Run clang-format
* Keep track of num batched kernels that have not been submitted yet
* Run clang-format
* Increase buf pool max size
* Increase param buf pool init size
* Remove webgpu buf pool resizing
* Merge with master
* Add buffer pool growth
* Move buffer pool growth outside of lock
* Reduce max pool size to 32
* Run clang-format
* Only resize param buf pool
2026-03-02 10:23:34 -08:00
36a7a6589c
ggml-webgpu: Support non-contiguous src0 and overlapping src0/src1 in binary ops ( #19850 )
...
* ggml-webgpu: Add binary op support for overlapping and non-contiguous.
* Add newline to binary.wgsl
* Append the test of binary op for src overlapping to test_bin_bcast.
* Remove unnecessary newline.
2026-03-02 07:59:53 -08:00
feefb92836
vulkan: tune MMVQ for Intel Windows ( #19988 )
2026-03-02 15:58:25 +01:00
2afcdb9777
ggml-cpu: optimise s390x multiply extend instructions ( #20032 )
2026-03-02 16:23:56 +08:00
319146247e
vulkan: improve partial offloading performance on AMD ( #19976 )
...
* vulkan: fix and enable cpy_tensor_async function
* use transfer_queue for async transfers on AMD, synchronize with timeline semaphore
* update offload_op logic
* fix missing transfer submission
* disable async transfer queue on AMD GCN
* revert op batch size change
* fix cpy_tensor_async checks
2026-03-01 17:32:14 +01:00
66d65ec29b
cuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels ( #19999 )
2026-03-01 13:40:22 +08:00
c9e651f7e5
updated lite, fix some cuda spams, fix qwen3tts voice loading
2026-03-01 00:41:56 +08:00
4e358265a3
Merge commit '8387ffb28d3467b81bc73727588c3fde772f8ebe' into concedo_experimental
...
# Conflicts:
# docs/backend/VirtGPU.md
# docs/backend/ZenDNN.md
# ggml/src/ggml-cpu/amx/amx.cpp
# ggml/src/ggml-cpu/amx/mmq.cpp
# ggml/src/ggml-sycl/add-id.cpp
# ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp
# ggml/src/ggml-virtgpu/backend/backend-dispatched-buffer-type.cpp
# ggml/src/ggml-virtgpu/backend/backend-dispatched-buffer.cpp
# ggml/src/ggml-virtgpu/backend/backend-dispatched.cpp
# ggml/src/ggml-virtgpu/backend/backend-dispatched.gen.h
# ggml/src/ggml-virtgpu/backend/backend-dispatched.h
# ggml/src/ggml-virtgpu/backend/backend-virgl-apir.h
# ggml/src/ggml-virtgpu/backend/backend.cpp
# ggml/src/ggml-virtgpu/backend/shared/api_remoting.h
# ggml/src/ggml-virtgpu/backend/shared/apir_backend.gen.h
# ggml/src/ggml-virtgpu/backend/shared/apir_backend.h
# ggml/src/ggml-virtgpu/backend/shared/apir_cs.h
# ggml/src/ggml-virtgpu/backend/shared/apir_cs_ggml.h
# ggml/src/ggml-virtgpu/backend/shared/apir_cs_rpc.h
# ggml/src/ggml-virtgpu/ggml-backend-buffer-type.cpp
# ggml/src/ggml-virtgpu/ggml-backend-device.cpp
# ggml/src/ggml-virtgpu/ggml-backend-reg.cpp
# ggml/src/ggml-virtgpu/ggml-backend.cpp
# ggml/src/ggml-virtgpu/ggml-remoting.h
# ggml/src/ggml-virtgpu/include/apir_hw.h
# ggml/src/ggml-virtgpu/regenerate_remoting.py
# ggml/src/ggml-virtgpu/virtgpu-forward-backend.cpp
# ggml/src/ggml-virtgpu/virtgpu-forward-buffer-type.cpp
# ggml/src/ggml-virtgpu/virtgpu-forward-buffer.cpp
# ggml/src/ggml-virtgpu/virtgpu-forward-device.cpp
# ggml/src/ggml-virtgpu/virtgpu-forward-impl.h
# ggml/src/ggml-virtgpu/virtgpu-forward.gen.h
# ggml/src/ggml-virtgpu/virtgpu.cpp
# ggml/src/ggml-virtgpu/virtgpu.h
# ggml/src/ggml-zendnn/CMakeLists.txt
# ggml/src/ggml-zendnn/ggml-zendnn.cpp
# src/CMakeLists.txt
# tests/CMakeLists.txt
# tests/test-tokenizer-0.sh
# tools/cli/README.md
# tools/completion/README.md
# tools/imatrix/imatrix.cpp
# tools/server/README.md
2026-02-28 12:45:16 +08:00
ecbcb7ea9d
CUDA: add CDNA3 MFMA support for flash attention MMA kernel ( #19806 )
...
* CUDA: add CDNA3 MFMA support for flash attention MMA kernel
Add MI300X (gfx942) MFMA tensor core flash attention using
v_mfma_f32_16x16x16_f16 (FP16 in, FP32 accumulate).
- Add FATTN_WARP_SIZE=64 for CDNA wavefront64
- Add CDNA config for head sizes 64, 80, 96, 112, 128
- Add FP16 MFMA intrinsic path in mma.cuh
- Add manual V transpose load for MFMA register layout
- Route CDNA to MMA for prompt processing, VEC for token generation
- Fix Q loading and combine stride granularity for non-power-of-2 heads
Benchmarks (Qwen2.5-1.5B Q4_K_M, MI300X):
pp512 +7%, pp1024 +13%, pp2048 +23%, pp4096 +39%
tg128 -10% (FA overhead, VEC used for both)
All 2480 flash attention tests pass.
Ref: https://github.com/ggml-org/llama.cpp/issues/17917
* address review: replace FATTN_WARP_SIZE with constexpr, improve dispatch
- Replace #define FATTN_WARP_SIZE with constexpr int warp_size =
ggml_cuda_get_physical_warp_size() in each device function
- Use ne[1]*gqa_ratio threshold for MMA vs tile dispatch. Benchmarked
crossover on MI300X @ d32768 with power-of-2 GQA models:
hsk=64 (Llama 1B, gqa=4): MMA wins at eff >= 128 (+11%)
hsk=128 (Llama 3B, gqa=4): MMA wins at eff >= 128 (+4%)
Unified threshold: eff_nq >= 128 for all head sizes.
- Remove VEC fallback; small batches fall through to tile kernel
* Update ggml/src/ggml-cuda/fattn.cu
* use ggml_cuda_info().devices warp_size instead of hardcoded check
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de >
2026-02-27 19:37:26 +01:00
d903f30e25
ggml-cpu: add repack for mxfp4 ( #19738 )
2026-02-27 18:15:09 +08:00
c17dce4f5c
replace the magic nunber 768 by max work group size to support iGPU ( #19920 )
...
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com >
2026-02-27 09:26:07 +08:00
88cf781f51
ggml-zendnn: update code for latest ZenDNN API ( #19923 )
...
- adapt ggml-zendnn.cpp to the new lowoha::matmul interface
- update the ZenDNN git tag in CMake to the latest release (ZenDNN‑2026‑WW08)
- add static lib support in CMake
2026-02-27 08:43:41 +08:00
4e76d24f28
ggml : fix AMX and add batched support ( #19925 )
...
llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF:Q4_0 -f wikitext-2-raw/wiki.test.raw -c 2048 -b 2048 --chunks 2
before this commit:
```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.31 seconds per pass - ETA 0.07 minutes
[1]17.3868,[2]22.2199,
Final estimate: PPL = 22.2199 +/- 1.59692
llama_perf_context_print: load time = 878.56 ms
llama_perf_context_print: prompt eval time = 2037.82 ms / 4096 tokens ( 0.50 ms per token, 2009.99 tokens per second)
llama_perf_context_print: eval time = 0.00 ms / 1 runs ( 0.00 ms per token, inf tokens per second)
llama_perf_context_print: total time = 6403.17 ms / 4097 tokens
llama_perf_context_print: graphs reused = 0
llama_memory_breakdown_print: | memory breakdown [MiB] | total free self model context compute unaccounted |
llama_memory_breakdown_print: | - Host | 845 = 318 + 224 + 302 |
llama_memory_breakdown_print: | - CPU_REPACK | 288 = 288 + 0 + 0 |
llama_memory_breakdown_print: | - AMX | 31 = 31 + 0 + 0 |
```
after this commit:
```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 1.98 seconds per pass - ETA 0.05 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485
llama_perf_context_print: load time = 719.23 ms
llama_perf_context_print: prompt eval time = 1676.23 ms / 4096 tokens ( 0.41 ms per token, 2443.58 tokens per second)
llama_perf_context_print: eval time = 0.00 ms / 1 runs ( 0.00 ms per token, inf tokens per second)
llama_perf_context_print: total time = 4258.74 ms / 4097 tokens
llama_perf_context_print: graphs reused = 0
llama_memory_breakdown_print: | memory breakdown [MiB] | total free self model context compute unaccounted |
llama_memory_breakdown_print: | - Host | 845 = 318 + 224 + 302 |
llama_memory_breakdown_print: | - AMX | 319 = 319 + 0 + 0 |
```
(no more CPU_REPACK)
after this commit, disabling amx:
```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.34 seconds per pass - ETA 0.07 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485
llama_perf_context_print: load time = 841.91 ms
llama_perf_context_print: prompt eval time = 2057.28 ms / 4096 tokens ( 0.50 ms per token, 1990.98 tokens per second)
llama_perf_context_print: eval time = 0.00 ms / 1 runs ( 0.00 ms per token, inf tokens per second)
llama_perf_context_print: total time = 6454.51 ms / 4097 tokens
llama_perf_context_print: graphs reused = 0
llama_memory_breakdown_print: | memory breakdown [MiB] | total free self model context compute unaccounted |
llama_memory_breakdown_print: | - Host | 845 = 318 + 224 + 302 |
llama_memory_breakdown_print: | - CPU_REPACK | 319 = 319 + 0 + 0 |
```
=> same perplexity.
Signed-off-by: Adrien Gallouët <angt@huggingface.co >
2026-02-26 21:39:11 +01:00
723c71064d
vulkan: fix fp16 Flash Attention on Windows AMD RDNA2 and below ( #19921 )
2026-02-26 19:11:04 +01:00
ffaafde16f
ggml-virtgpu: improve the reliability of the code ( #19846 )
...
* ggml-virtgpu-backend: validate the consistency of the received objects
This patch adds consistency checks in the
ggml-virtgpu-backend (running on the host side) to ensure that the
data received from the guest is consistent (valid pointers, valid
sizes and offsets).
* ggml-virtgpu-backend: add fallback/skips for optional ggml backend methods
```
1. bck->iface.synchronize(bck)
2. buft->iface.get_alloc_size(buft, op)
3. buft->iface.get_max_size(buft)
```
these three methods are optional in the GGML interface. `get_max_size`
was already properly defaulted, but `backend sychronize` and `butf
get_max_size` would have segfaulted the backend if not implemented.
* ggml-virtgpu-backend: fix log format missing argument
* ggml-virtgpu-backend: improve the abort message
* ggml-virtgpu-backend: more safety checks
* ggml-virtgpu-backend: new error code
* ggml-virtgpu-backend: initialize all the error codes
* ggml-virtgpu: add a missing comment generated by the code generator
* ggml-virtgpu: add the '[virtgpu]' prefix to the device/buffer names
* ggml-virtgpu: apir_device_buffer_from_ptr: improve the error message
* ggml-virtgpu: shared: make it match the latest api_remoting.h of Virglrenderer APIR
(still unmerged)
* ggml-virtgpu: update the code generator to have dispatch_command_name in a host/guest shared file
* ggml-virtgpu: REMOTE_CALL: fail if the backend returns an error
* docs/backend/VirtGPU.md: indicate that the RAM+VRAM size is limed to 64 GB with libkrun
* ggml-virtgpu: turn off clang-format header ordering for some of the files
Compilation breaks when ordered alphabetically.
* ggml-virtgpu: clang-format
* ggml-virtgpu/backend/shared/api_remoting: better comments for the APIR return codes
2026-02-26 20:00:57 +08:00
05834eecb3
Merge commit '1ca3d1de153152645cf890b637c2c6450f1615e3' into concedo_experimental
...
# Conflicts:
# tools/server/README.md
2026-02-26 19:55:06 +08:00
1ca3d1de15
gguf : avoid too many file size calls ( #19919 )
2026-02-26 12:46:32 +02:00
749a606374
whisper broke
2026-02-26 16:45:04 +08:00
44182ebefe
Merge commit '8c2c0108dd77f90c0a5233420601ae05ac735d59' into concedo_experimental
...
# Conflicts:
# examples/model-conversion/Makefile
# examples/model-conversion/scripts/utils/inspect-org-model.py
# ggml/src/ggml-hexagon/ggml-hexagon.cpp
# ggml/src/ggml-hexagon/htp/act-ops.c
# ggml/src/ggml-hexagon/htp/get-rows-ops.c
# ggml/src/ggml-hexagon/htp/hex-dma.h
# ggml/src/ggml-hexagon/htp/htp-ops.h
# ggml/src/ggml-hexagon/htp/matmul-ops.c
# ggml/src/ggml-hexagon/htp/rope-ops.c
# ggml/src/ggml-hexagon/htp/set-rows-ops.c
# ggml/src/ggml-hexagon/htp/softmax-ops.c
# ggml/src/ggml-hexagon/htp/unary-ops.c
# scripts/snapdragon/adb/run-cli.sh
# scripts/snapdragon/adb/run-completion.sh
# scripts/snapdragon/adb/run-mtmd.sh
# scripts/snapdragon/windows/run-cli.ps1
# scripts/sync_vendor.py
# tests/test-backend-sampler.cpp
2026-02-26 16:30:37 +08:00
2943210c1e
support permuted, remove check s0/s10 ( #19889 )
...
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com >
2026-02-26 10:27:20 +08:00
3769fe6eb7
vulkan: check for memory overlap before doing fusion ( #19768 )
...
* vulkan: check for memory overlap before doing fusion
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
* address feedback
2026-02-25 18:25:38 +01:00
a96a1120b4
gguf : fix ftell/fseek for Windows ( #19870 )
2026-02-25 06:58:11 +02:00
418dea39ce
ggml/gguf : prevent integer overflows ( #19856 )
...
* gguf : prevent integer overflow for ggml_context mem size
* ggml : fix int overflows in ggml_new_object()
* gguf : prevent string exhaustion
* gguf : prevent array elements exhaustion
* ggml : fix negative tensor type oob
* py : assert that alignment is non-zero power of 2
* ggml : check int overflow in ggml_new_tensor_impl and ggml_new_object
* gguf-py : error on duplicate keys when reading
* py : restore tensor_fields
* enforce proper alignment in add_custom_alignment
* gguf : better name
* gguf : fix ctx size for no_alloc == true
* gguf : minor print fix
* ggml : print values when overflow
* ggml : remove deprecated ggml_type_sizef()
* ggml : relax ggml_type asserts to debug-only
* gguf : add mem_size overflow test
* gguf : add file size check for arrays
* ggml : relax asseerts for ggml_get_type_traits()
* flake8 fix
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com >
2026-02-24 20:17:11 +02:00
Concedo
Todor Boinovski
Max Krasnyansky
Aman Gupta
Aaron Teo
shalinib-ibm
JustCommitRandomness
Johannes Gäßler
lhez
YardenTal44
Andreas Kieslinger
Marcel Petrick
Nikhil Jain
Masashi Yoshimura
Adrien Gallouët
Charles Xu
shaofeiqi
Abhijit Ramesh
Ruben Ortlam
oobabooga
Jayant Lohia
Neo Zhang
Vishal Singh
Kevin Pouget
Georgi Gerganov
Jeff Bolz
Aldehir Rojas