]> git.djapps.eu Git - pkg/ggml/sources/ggml/log
pkg/ggml/sources/ggml
5 weeks agoopencl: add q8_0 mm support (llama/16469)
lhez [Wed, 15 Oct 2025 17:51:04 +0000 (10:51 -0700)]
opencl: add q8_0 mm support (llama/16469)

* opencl: add mm_q8_0_f32

* opencl: fix data loading for incomplete tile

* opencl: use q8_0 mm for larger matrix

* opencl: add some tests to cover the path

5 weeks agoopencl: fix FA for f32 (llama/16584)
lhez [Wed, 15 Oct 2025 17:48:28 +0000 (10:48 -0700)]
opencl: fix FA for f32 (llama/16584)

5 weeks agometal: optimise `GGML_OP_SUM` (llama/16559)
Sam/Samuel [Wed, 15 Oct 2025 14:05:56 +0000 (23:05 +0900)]
metal: optimise `GGML_OP_SUM` (llama/16559)

* optimise GGML_OP_SUM

* add non-contiguous tests by permuting the input

* change tests to require full contiguity of OP_SUM

* cuda : add check GGML_OP_SUM

---------

Co-authored-by: Georgi Gerganov <redacted>
5 weeks agoCUDA: Changing the CUDA scheduling strategy to spin (llama/16585)
Julius Tischbein [Wed, 15 Oct 2025 11:54:15 +0000 (13:54 +0200)]
CUDA: Changing the CUDA scheduling strategy to spin (llama/16585)

* CUDA set scheduling strategy to spinning for cc121

* Using prop.major and prop.minor, include HIP and MUSA

* Exclude HIP and MUSA

* Remove trailing whitespace

Co-authored-by: Johannes Gäßler <redacted>
* Remove empty line

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
5 weeks agometal : avoid using Metal's gpuAddress property (llama/16576)
Georgi Gerganov [Tue, 14 Oct 2025 17:33:05 +0000 (20:33 +0300)]
metal : avoid using Metal's gpuAddress property (llama/16576)

* metal : avoid using Metal's gpuAddress property

* metal : fix rope kernels buffer check

6 weeks agosync : llama.cpp upstream/latest upstream/0.9.4.58
Georgi Gerganov [Tue, 14 Oct 2025 17:24:43 +0000 (20:24 +0300)]
sync : llama.cpp

6 weeks agovulkan: Add ACC_TYPE_VEC2 implementation (llama/16203)
SavicStefan [Tue, 14 Oct 2025 17:18:05 +0000 (19:18 +0200)]
vulkan: Add ACC_TYPE_VEC2 implementation (llama/16203)

Signed-off-by: Stefan Savic <redacted>
Co-authored-by: Stefan Savic <redacted>
6 weeks agoCUDA + openCL: fix bug in accessing rms_norm->src while doing fusion (llama/16577)
Aman Gupta [Tue, 14 Oct 2025 14:48:08 +0000 (22:48 +0800)]
CUDA + openCL: fix bug in accessing rms_norm->src while doing fusion (llama/16577)

6 weeks agovulkan: Support FA with K/V in F32 (llama/16543)
Jeff Bolz [Tue, 14 Oct 2025 13:53:37 +0000 (08:53 -0500)]
vulkan: Support FA with K/V in F32 (llama/16543)

6 weeks agovulkan: Improve build time for MSVC (llama/16545)
Jeff Bolz [Tue, 14 Oct 2025 12:51:36 +0000 (07:51 -0500)]
vulkan: Improve build time for MSVC (llama/16545)

Enable CMP0147 so custom build steps (invoking vulkan-shader-gen) are run in parallel.

Enable /MP so source files are compiled in parallel.

6 weeks agoCUDA: enable FA for FP32 KV cache (llama/16546)
Johannes Gäßler [Tue, 14 Oct 2025 12:22:47 +0000 (14:22 +0200)]
CUDA: enable FA for FP32 KV cache (llama/16546)

6 weeks agoCUDA: use fastdiv + ggml_cuda_mad for mmvf (llama/16557)
Aman Gupta [Tue, 14 Oct 2025 11:16:21 +0000 (19:16 +0800)]
CUDA: use fastdiv + ggml_cuda_mad for mmvf (llama/16557)

* CUDA: use fastdiv + ggml_cuda_mad for mmvf

* use bf16 directly + fix formatting

* Add exception for HIP code

6 weeks agoCUDA: add fp kernel for larger batch size MoE (llama/16512)
Aman Gupta [Tue, 14 Oct 2025 11:15:15 +0000 (19:15 +0800)]
CUDA: add fp kernel for larger batch size MoE (llama/16512)

* CUDA: kernel for larger batch sizes for MoE

* WIP

* WIP

* WIP

* WIP

* WIP

* WIP

* fixup

* tests

* Move mmq_ids_helper to mmid

* cleanup

* Remove redundant checks

6 weeks agocuda : remove legacy copy-op pointer indirection code (llama/16485)
Anav Prasad [Tue, 14 Oct 2025 09:53:49 +0000 (09:53 +0000)]
cuda : remove legacy copy-op pointer indirection code (llama/16485)

* remove legacy copy-op pointer indirection code

* further removal of copy-op indirection code

* renamed check_node_graph_compatibility_and_refresh_copy_ops function

6 weeks agometal : FA support F32 K and V and head size = 32 (llama/16531)
Georgi Gerganov [Mon, 13 Oct 2025 20:07:57 +0000 (23:07 +0300)]
metal : FA support F32 K and V and head size = 32 (llama/16531)

* metal : FA support F32 K and V and head size = 32

* graph : remove obsolete comment [no ci]

6 weeks agoopencl: fix build targeting CL 2 (llama/16554)
lhez [Mon, 13 Oct 2025 18:50:37 +0000 (11:50 -0700)]
opencl: fix build targeting CL 2 (llama/16554)

6 weeks agoCUDA: fix numerical issues in tile FA kernel (llama/16540)
Johannes Gäßler [Mon, 13 Oct 2025 14:29:45 +0000 (16:29 +0200)]
CUDA: fix numerical issues in tile FA kernel (llama/16540)

6 weeks agoggml : fix build broken with -march=armv9-a on MacOS (llama/16520)
Jie Fu (傅杰) [Mon, 13 Oct 2025 12:48:47 +0000 (20:48 +0800)]
ggml : fix build broken with -march=armv9-a on MacOS (llama/16520)

* ggml : fix build broken with -march=armv9-a on MacOS

Signed-off-by: Jie Fu <redacted>
* Add #pragma message

Signed-off-by: Jie Fu <redacted>
* Address review comment.

Signed-off-by: Jie Fu <redacted>
* Update ggml/src/ggml-cpu/ggml-cpu.c

---------

Signed-off-by: Jie Fu <redacted>
Co-authored-by: Diego Devesa <redacted>
6 weeks agoCANN: fix CPU memory leak in CANN backend (llama/16549)
Chenguang Li [Mon, 13 Oct 2025 09:01:24 +0000 (17:01 +0800)]
CANN: fix CPU memory leak in CANN backend (llama/16549)

This commit fixes a CPU-side memory leak issue in the CANN backend,
which occurred when intermediate aclTensorList objects were not properly
released after operator execution. The leak happened during repeated
invocations of CANN ops (e.g., FlashAttention), leading to increasing
host memory usage over time.

Proper resource cleanup (aclDestroyTensorList and related release logic)
has been added to ensure that all temporary tensors are correctly freed.

6 weeks agometal: add support for opt_step_sgd (llama/16539)
Sam/Samuel [Mon, 13 Oct 2025 08:25:02 +0000 (16:25 +0800)]
metal: add support for opt_step_sgd (llama/16539)

* metal: add support for opt_step_sgd

* add newline to pass EditorConfig check

6 weeks agoggml : fix scalar path for computing norm (llama/16558)
Georgi Gerganov [Mon, 13 Oct 2025 08:22:27 +0000 (11:22 +0300)]
ggml : fix scalar path for computing norm (llama/16558)

6 weeks agoCANN: Update several operators to support FP16 data format (llama/16251)
hipudding [Mon, 13 Oct 2025 00:52:22 +0000 (08:52 +0800)]
CANN: Update several operators to support FP16 data format (llama/16251)

Many Ascend operators internally use FP16 precision for computation.
If input data is in FP32, it must first be cast to FP16 before
computation, and then cast back to FP32 after computation, which
introduces unnecessary cast operations. Moreover, FP16 computation
requires significantly less workload compared to FP32, leading to
noticeable efficiency improvements.

In this change, `get_rows`, `rms_norm`, and `flash_attn_ext` are extended
to support multiple data types. Validation on the Qwen2 0.5b model shows
correct accuracy and about 10% performance gain in concurrent scenarios.

Co-authored-by: noemotiovon <redacted>
6 weeks agometal : add opt_step_adamw and op_sum (llama/16529)
Sam/Samuel [Sun, 12 Oct 2025 18:43:14 +0000 (02:43 +0800)]
metal : add opt_step_adamw and op_sum (llama/16529)

* scaffold to support opt step adamw on metal (not written so far)

* add opt-step-adamw kernel for metal

* pass op->src[4] as a separate buffer to the pipeline

* add bounds check to opt-step-adamw kernel

* complete scaffold for GGML_OP_SUM

* naive GGML_OP_SUM kernel

* remove unwanted comment

* change OP_SUM capability gate

* Add has_simdgroup_reduction to both ops to pass CI

6 weeks agofix UT fault cases: count-equal, argsort, pad OPs (llama/16521)
Neo Zhang Jianyu [Sun, 12 Oct 2025 13:53:35 +0000 (21:53 +0800)]
fix UT fault cases: count-equal, argsort, pad OPs (llama/16521)

* fix/refactor OP argsort, pad

* fix count-equal op

* update SYCL OP list

* fix format issue

---------

Co-authored-by: Zhang Jianyu <redacted>
6 weeks agoggml : Fix FP16 ELU positive branch (llama/16519)
sirus20x6 [Sun, 12 Oct 2025 05:25:37 +0000 (00:25 -0500)]
ggml : Fix FP16 ELU positive branch (llama/16519)

Co-authored-by: Aaron <redacted>
6 weeks agoggml: Correct SVE implementation in ggml_vec_dot_f16_unroll (llama/16518)
sirus20x6 [Sun, 12 Oct 2025 05:15:00 +0000 (00:15 -0500)]
ggml: Correct SVE implementation in ggml_vec_dot_f16_unroll (llama/16518)

The previous SVE implementation for `ggml_vec_dot_f16_unroll` contained a bug due to a copy-paste error. The wrong variable was used in an FMA instruction, leading to incorrect results. This commit corrects the variable usage and improves the clarity of the code by renaming variables to avoid confusion.

Co-authored-by: Aaron <redacted>
6 weeks agoCUDA: faster tile FA, add oob checks, more HSs (llama/16492)
Johannes Gäßler [Sat, 11 Oct 2025 18:54:32 +0000 (20:54 +0200)]
CUDA: faster tile FA, add oob checks, more HSs (llama/16492)

7 weeks agosync : llama.cpp
Georgi Gerganov [Sat, 11 Oct 2025 15:01:19 +0000 (18:01 +0300)]
sync : llama.cpp

7 weeks agometal : fix mul-mm condition + fix mul-mv permuted kernels (llama/16494)
Georgi Gerganov [Sat, 11 Oct 2025 13:54:10 +0000 (16:54 +0300)]
metal : fix mul-mm condition + fix mul-mv permuted kernels (llama/16494)

7 weeks agocuda : avoid initializing unused devices (llama/16510)
Diego Devesa [Sat, 11 Oct 2025 11:02:26 +0000 (04:02 -0700)]
cuda : avoid initializing unused devices (llama/16510)

7 weeks agocmake : Dont define XOPENSOURCE on AIX (llama/16481)
Prajwal B Mehendarkar [Fri, 10 Oct 2025 08:15:46 +0000 (13:45 +0530)]
cmake : Dont define XOPENSOURCE on AIX (llama/16481)

7 weeks agocpu : optimize the ggml NORM operation (llama/15953)
duduta [Thu, 9 Oct 2025 19:11:15 +0000 (22:11 +0300)]
cpu : optimize the ggml NORM operation (llama/15953)

* ggml-cpu: optimize norm operation to use intrinsics or Accelerate

          rename function

          add endif macro comment

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Aaron Teo <redacted>
* implement s390x SIMD suggested by @taronaeo

* add TODO comment

* tidy up spaces

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Aaron Teo <redacted>
7 weeks agoCANN: Improve ACL graph matching (llama/16166)
Chenguang Li [Thu, 9 Oct 2025 07:50:25 +0000 (15:50 +0800)]
CANN: Improve ACL graph matching (llama/16166)

* CANN: improve ACL graph matching

Record `ne` and `nb` information for src tensors and include them in the
graph matching check. This enhances the robustness of ACL graph matching
by preventing incorrect matches when src tensors share the same data
address but differ in shape or stride.

* CANN: add op_params match

7 weeks agokleidiai: kernel interface refactoring (llama/16460)
Charles Xu [Thu, 9 Oct 2025 07:29:17 +0000 (09:29 +0200)]
kleidiai: kernel interface refactoring (llama/16460)

7 weeks agorefactor soft_max, add soft_max_back (llama/16472)
Neo Zhang Jianyu [Thu, 9 Oct 2025 07:25:11 +0000 (15:25 +0800)]
refactor soft_max, add soft_max_back (llama/16472)

* refactor to support soft_max_ext

* fix error and support soft_max_back

* rm unused functions

* fix format issue

---------

Co-authored-by: Zhang Jianyu <redacted>
7 weeks agoDisable CUDA host buffers on integrated GPUs (llama/16308)
ai-fonsi [Wed, 8 Oct 2025 18:21:46 +0000 (20:21 +0200)]
Disable CUDA host buffers on integrated GPUs (llama/16308)

7 weeks agometal : mark FA blocks (llama/16372)
Georgi Gerganov [Wed, 8 Oct 2025 07:57:53 +0000 (10:57 +0300)]
metal : mark FA blocks (llama/16372)

* metal : better unroll in the FA kernels

* metal : index FA blocks

* tests : restore [no ci]

* metal : prevent division by zero in FA kernels

* metal : fix -INF detection logic

7 weeks agoggml webgpu: profiling, CI updates, reworking of command submission (llama/16452)
Reese Levine [Tue, 7 Oct 2025 20:48:56 +0000 (13:48 -0700)]
ggml webgpu: profiling, CI updates, reworking of command submission (llama/16452)

* Add profiling

* More detailed profiling

* Rework command submission to avoid global locks

* Update wait handling

* try new method of waiting on futures

* Add serializing of command submission in some cases

* Add new pool for timestamp queries and clean up logging

* Serialize command submission in CI and leave a TODO note

* Update webgpu CI

* Add myself as WebGPU codeowner

* Deadlock avoidance

* Leave WebGPU/Vulkan CI serialized

* Fix divide by 0

* Fix logic in division by inflight_threads

* Update CODEOWNERS and remove serialize submit option

7 weeks agometal : add support for non-padded FA KV (llama/16148)
Georgi Gerganov [Tue, 7 Oct 2025 05:23:30 +0000 (08:23 +0300)]
metal : add support for non-padded FA KV (llama/16148)

* metal : pad K, V and Mask when needed

* cont : simplify

* cuda : add TODO about KV padding requirement

* metal : add comments

* metal : remove mask padding requirement

7 weeks agotests : add -INF blocks to the KQ mask in the FA tests (llama/16380)
Georgi Gerganov [Tue, 7 Oct 2025 05:22:35 +0000 (08:22 +0300)]
tests : add -INF blocks to the KQ mask in the FA tests (llama/16380)

* tests : add -INF blocks to the KQ mask in the FA tests

* cont : bump -INF block size to 64

Co-authored-by: Jeff Bolz <redacted>
* ggml : prevent division by zero in FA CPU op

---------

Co-authored-by: Jeff Bolz <redacted>
7 weeks agometal : various optimizations + refactoring (llama/16446)
Georgi Gerganov [Tue, 7 Oct 2025 05:21:40 +0000 (08:21 +0300)]
metal : various optimizations + refactoring (llama/16446)

* metal : ssm_scan minor opts

* metal : get_rows optimize

* metal : cpy optimize

* metal : ssm_conv opt

* metal : ssm_scan simplify

* metal : ssm_Scan opt

7 weeks agoggml : fix unaligned access in AMX code (llama/16315)
Georgi Gerganov [Mon, 6 Oct 2025 13:05:27 +0000 (16:05 +0300)]
ggml : fix unaligned access in AMX code (llama/16315)

7 weeks agoggml-cpu : fix leftover handling in ggml_vec_scale_f32 for SVE (llama/16443)
Daniel Bevenius [Mon, 6 Oct 2025 12:17:12 +0000 (14:17 +0200)]
ggml-cpu : fix leftover handling in ggml_vec_scale_f32 for SVE (llama/16443)

This commit updates the leftover handling in ggml_vec_scale_f32.

The motivation for this is that the code currently incorrectly assumes
there would be fewer than ggml_f32_epr leftover elements. However,
since the main loop processes 2*ggml_f32_epr elements per iteration
, there can be up to (2*ggml_f32_epr - 1) leftover elements.

The original single-pass leftover code could only process ggml_f32_epr
elements, leaving some elements unscaled.

Example scenario with 256-bit SVE:
```
ggml_f32_epr  = 8 (elements per register)
ggml_f32_step = 16 (two registers per iteration)
n             = 25
np            = 16
leftovers     = 9 elements (16-24)

Original    : processes only elements 16-23, misses element 24
This commit : loop processes elements 16-23, then element 24
```

Refs: https://github.com/ggml-org/llama.cpp/actions/runs/18070620247/job/51419855630

7 weeks agoggml webgpu: actually add softmax, fix rms_norm offset (llama/16400)
Reese Levine [Sun, 5 Oct 2025 03:59:31 +0000 (20:59 -0700)]
ggml webgpu: actually add softmax, fix rms_norm offset (llama/16400)

* implement soft_max

* Fix soft_max data race

* Temporary fix, wait on each submit

7 weeks agovulkan: use a more appropriate amount of threads when generating shaders (llama/16418)
Eve [Sat, 4 Oct 2025 20:04:27 +0000 (20:04 +0000)]
vulkan: use a more appropriate amount of threads when generating shaders (llama/16418)

* use a more flexible amount of threads

* fix windows compile and 0 thread case

* nominmax

7 weeks agorpc : check src buffer when copying tensor (llama/16421)
Radoslav Gerganov [Sat, 4 Oct 2025 13:22:45 +0000 (16:22 +0300)]
rpc : check src buffer when copying tensor (llama/16421)

Only dst buffer is guaranteed to be an RPC buffer. Add check for the src
one.

7 weeks agorpc : add support for multiple devices (llama/16276)
Radoslav Gerganov [Sat, 4 Oct 2025 09:49:16 +0000 (12:49 +0300)]
rpc : add support for multiple devices (llama/16276)

* rpc : add support for multiple devices

Allow rpc-server to expose multiple devices from a single endpoint.
Change RPC protocol to include device identifier where needed.

closes: #15210

* fixes

* use ggml_backend_reg_t

* address review comments

* fix llama-bench backend report

* address review comments, change device naming

* fix cmd order

7 weeks agosync : llama.cpp
Georgi Gerganov [Sat, 11 Oct 2025 15:01:03 +0000 (18:01 +0300)]
sync : llama.cpp

7 weeks agovulkan : incremental shader builds (llama/16341)
Acly [Sat, 11 Oct 2025 14:59:36 +0000 (17:59 +0300)]
vulkan : incremental shader builds (llama/16341)

* vulkan (DRAFT): split shader generation by GLSL source file, to improve incremental build times

* support dep-files so shaders are recompiled if their included files change

* rename shader files which are used as "headers" to use .glsl extension
* move glslc extension detection shaders to separate folders
* the above is to prevent them from getting glob'd with the actual compute shaders that need to be compiled

* vulkan : only write embedded shader .hpp/.cpp when they change

* avoid recompiling ggml-vulkan.cpp when editing shaders
* pass single --source argument instead of --input-dir & --filter to shader gen
* check for source file match earlier

* fix hang in vulkan-shaders-gen when there are compilation errors

* early out did not decrement compile_count

* clean up

* fix glslc integer dot product test

* unconditionally write the embedded shader cpp output

* replace output filepath in generated dep-files to match output in CMakeLists

---------

Co-authored-by: Jeff Bolz <redacted>
7 weeks agosync : llama.cpp
Georgi Gerganov [Sat, 11 Oct 2025 14:58:06 +0000 (17:58 +0300)]
sync : llama.cpp

7 weeks agometal : fix loop bound in ggml_mem_ranges (llama/16412)
Georgi Gerganov [Fri, 3 Oct 2025 16:18:56 +0000 (19:18 +0300)]
metal : fix loop bound in ggml_mem_ranges (llama/16412)

7 weeks agoggml : fix graph reallocation with multiple chunks (llama/16396)
Acly [Fri, 3 Oct 2025 11:49:08 +0000 (13:49 +0200)]
ggml : fix graph reallocation with multiple chunks (llama/16396)

reallocation is needed if a single chunk grows in size,
even if total allocation size stays the same or is lower

7 weeks agovulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE (llama/16354)
Jeff Bolz [Fri, 3 Oct 2025 10:50:46 +0000 (05:50 -0500)]
vulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE (llama/16354)

* vulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE

Replace maxMemoryAllocationSize check with maxBufferSize when creating buffers.
The maxMemoryAllocationSize limit is a "soft" limit and allocations can succeed
beyond that limit. This allows > 4GB buffers to be allocated on some
implementations (e.g. NVIDIA) and tensors this large can be used for im2col
and mul_mat.

For temporary buffers (prealloc_x/y/etc) check against maxStorageBufferRange.
I'm not sure this check is ideal, but we always use these buffers as a single
full size binding and the limit may be smaller than maxMemoryAllocationSize
or maxBufferSize, so I think this is reasonable.

Replace descriptor range uses of VK_WHOLE_SIZE with a manually computed range.
The maxStorageBufferRange may be smaller than the maxBufferSize or
maxMemoryAllocationSize (and the Vulkan spec warns about this in a note) and
it's invalid usage if VK_WHOLE_SIZE computes a range larger than
maxStorageBufferRange.

With this change, it should be possible to generate videos using wan networks
in stable-diffusion.cpp.

* vulkan: Add env var GGML_VK_FORCE_MAX_BUFFER_SIZE and use stoull

7 weeks agovulkan: Fix FA coopmat1 invalid array indexing (llama/16365)
Jeff Bolz [Fri, 3 Oct 2025 09:52:46 +0000 (04:52 -0500)]
vulkan: Fix FA coopmat1 invalid array indexing (llama/16365)

When computing sinks, the cm1 shader was looping r from 0 to Br rather than
to rows_per_thread. I must have copied this from the scalar path (where it is
correct), and somehow it wasn't causing failures on current drivers.

7 weeks agovulkan: in flash attention, bounds check against nem1 (don't rely on GGML_KQ_MASK_PAD...
Jeff Bolz [Fri, 3 Oct 2025 08:33:08 +0000 (03:33 -0500)]
vulkan: in flash attention, bounds check against nem1 (don't rely on GGML_KQ_MASK_PAD) (llama/16316)

7 weeks agoggml webgpu: add support for soft_max, optimize rms_norm (llama/16357)
Reese Levine [Thu, 2 Oct 2025 18:00:31 +0000 (11:00 -0700)]
ggml webgpu: add support for soft_max, optimize rms_norm (llama/16357)

* Add inplace softmax

* Move rms_norm to split row approach

* Update debug for supports_op

* clean up debug statements

* Update tests/test-backend-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
---------

Co-authored-by: Georgi Gerganov <redacted>
7 weeks agomodel : Apertus model implementation (llama/15852)
Piotr Wilkin (ilintar) [Thu, 2 Oct 2025 17:43:22 +0000 (19:43 +0200)]
model : Apertus model implementation (llama/15852)

* First attempt

* No permute during convert (fixes qk tensors), proper norm application.

* RoPE = NeoX

* Coherence!

* Migrate xielu params from tensors to hyperparameters

* Simple CUDA kernel

* Revert stupid LLM refactorings

* Chat template support

* configchecker / flake8 errors

* Reorder unary.cu

* I do conclude that LLMs are, in fact, stupid.

* Fix after merge

* Final newline

* Make xIELU an UNARY_OP

* Final newline

* Correctly account for parameter shift

* Argh.

* Update ggml/src/ggml-cpu/unary-ops.cpp

Co-authored-by: Georgi Gerganov <redacted>
* Refactor: remove unused methods, inline and factorize softplus, add const modifiers

* Revert CUDA changes, implement xIELU as a separate OP

* Pesky newline

* Add float2half / half2float for F16 inputs/outputs

* CUDA variants, attempt 2

* Actually, attempt 3

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

Co-authored-by: Johannes Gäßler <redacted>
* Missing convert header

* Proper formula and reference for xIELU in the comments.

* Modify unary-ops.cpp to add the functor-based logic besides the template system to retain optimizations

* Apply suggestions from code review

Co-authored-by: Sigbjørn Skjæret <redacted>
* Add tensor mappings for Apertus to global list instead

* Fix lazy on scalars

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

Co-authored-by: Johannes Gäßler <redacted>
* Add comment about the constraints on positive/negative alpha

* Change `softplus` to `ggml_softplus`

---------

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Johannes Gäßler <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
7 weeks agomusa: update compile flags (llama/16265)
R0CKSTAR [Thu, 2 Oct 2025 13:29:56 +0000 (21:29 +0800)]
musa: update compile flags (llama/16265)

Signed-off-by: Xiaodong Ye <redacted>
7 weeks agoHIP: Disable ROCWMMA fattn on CDNA when compiled against ROCWMMA 2.0.0 (llama/16221)
uvos [Wed, 1 Oct 2025 21:09:25 +0000 (23:09 +0200)]
HIP: Disable ROCWMMA fattn on CDNA when compiled against ROCWMMA 2.0.0 (llama/16221)

* HIP: Disable ROCWMMA fatt on CDNA when compiled against ROCWMMA 2.0.0

rocwmma 2.0.0 includes a bug in the code fakeing fp16 accumulation on CDNA

* CUDA: Fix volta condition in ggml_cuda_should_use_wmma_fattn

7 weeks agovulkan: make ggml_vk_default_dispatcher support older vulkan headers (llama/16345)
Eve [Wed, 1 Oct 2025 07:56:36 +0000 (07:56 +0000)]
vulkan: make ggml_vk_default_dispatcher support older vulkan headers (llama/16345)

* make ggml_vk_default_dispatcher support older vulkan headers

* simpilfy with using

7 weeks agoopencl: support pad_ext (llama/15888)
lhez [Tue, 30 Sep 2025 17:45:45 +0000 (10:45 -0700)]
opencl: support pad_ext (llama/15888)

7 weeks agoggml webgpu: support for rope,div,sub,glu,scale,cont operators (llama/16187)
Reese Levine [Tue, 30 Sep 2025 16:57:51 +0000 (09:57 -0700)]
ggml webgpu: support for rope,div,sub,glu,scale,cont operators (llama/16187)

* Work on rope

* Simplify inplace operation generation and combine mul/add generation

* Work on rope variants

* implement neox rope

* rope complete

* Add sub,div,glu operators

* implement scale op

* Update cpy shader to handle cont/more types

* formatting

* Update test vars printing for rope,rms_norm

* Avoid ROPE hardcoded constants

* Add TODO to change ROPE constants to enum

Co-authored-by: Georgi Gerganov <redacted>
* fix TODO comment

---------

Co-authored-by: Georgi Gerganov <redacted>
7 weeks agoopencl: support ne3 in get_rows (llama/15866)
lhez [Tue, 30 Sep 2025 16:55:13 +0000 (09:55 -0700)]
opencl: support ne3 in get_rows (llama/15866)

2 months agoggml : bump version to 0.9.4 (#1363) upstream/0.9.4 v0.9.4
Georgi Gerganov [Tue, 30 Sep 2025 10:42:39 +0000 (13:42 +0300)]
ggml : bump version to 0.9.4 (#1363)

2 months agosync : whisper.cpp [no ci]
Georgi Gerganov [Tue, 30 Sep 2025 10:39:06 +0000 (13:39 +0300)]
sync : whisper.cpp [no ci]

2 months agosync : llama.cpp
Georgi Gerganov [Tue, 30 Sep 2025 08:18:45 +0000 (11:18 +0300)]
sync : llama.cpp

2 months agocuda : Enable CUDA Graph usage for Nemotron Nano v2 (NemotronH) (llama/16328)
anavp-nvidia [Tue, 30 Sep 2025 08:13:22 +0000 (08:13 +0000)]
cuda : Enable CUDA Graph usage for Nemotron Nano v2 (NemotronH) (llama/16328)

* Fix Nemotron Nano v2 9B not executing as CUDA Graph on NVIDIA GPUs

* fix to ensure test-backend-ops check passes

2 months agometal : dynamic simdgroups for MV kernels (llama/16340)
Georgi Gerganov [Tue, 30 Sep 2025 08:03:23 +0000 (11:03 +0300)]
metal : dynamic simdgroups for MV kernels (llama/16340)

* metal : dynamic simdgroups for MV kernels

* cont : minor

2 months agokleidiai : fix work size and threads sync for fp16 (llama/16246)
Charles Xu [Tue, 30 Sep 2025 07:07:20 +0000 (09:07 +0200)]
kleidiai : fix work size and threads sync for fp16 (llama/16246)

2 months agotests: override test_set_rows::max_nmse_err to allow for occasional rounding differen...
Jeff Bolz [Tue, 30 Sep 2025 00:26:34 +0000 (19:26 -0500)]
tests: override test_set_rows::max_nmse_err to allow for occasional rounding differences (llama/16295)

* tests: override test_set_rows::max_nmse_err to allow for occasional rounding differences

* apply similar error bounds to test_cpy

2 months agosync : llama.cpp
Georgi Gerganov [Mon, 29 Sep 2025 14:53:28 +0000 (17:53 +0300)]
sync : llama.cpp

2 months agoggml: riscv: add riscv spacemit backend (llama/15288)
alex-spacemit [Mon, 29 Sep 2025 14:50:44 +0000 (22:50 +0800)]
ggml: riscv: add riscv spacemit backend (llama/15288)

* ggml: add spacemit backend

Change-Id: I249bdc043485d815a9c351867137bc1e27cc2e23

* add new line at end of file

Change-Id: I889ed1c85fb45e62350ecde0c06f70450cadfbe2

* add riscv zba extension limit

Change-Id: I321eb200f859751727afe5cae13074dfce2bb0ce

* fixed for review comments, file renamed and format

Change-Id: Ia20b6ec24a36638e62e0fe07cf100916a7cce3ce

* fixed for code format, after clang-format

Change-Id: I5dc33a0412da3d3f2d77075d8939185d3009eca2

* use _Float16 instead of __fp16

Change-Id: I039fb02bb95270e641bc4442204e658735859d43

* add ci for riscv64-spacemit-ime-native

Change-Id: I711c1033061df1a289ea77891b2997599dfe8279

* update debian-13-riscv64-spacemit-ime-native ci label

Change-Id: Ifb2b891e2fca57b5da604fce2ac255f27731179a

* remove license comment for spacemit ime

Change-Id: If0dc3ca30a958631ccca0a28b62e0b825f9fb0c3

* upgrade binutils for gcc ime

Change-Id: Ibf2fa74c1064408974cb5b45f044d40987e5fb45

* add spacemit ime cross jobs

Change-Id: I80d74909941d41cb9cd09e51d8baf01c985cbfc6

* remove native compile for riscv64-spacemit-ime

Change-Id: I01920afafdc73fa7424014fd648d243f8ec9e25e

* ci : add caching for spacemit ime cross toolchain

Change-Id: Ic54a192019a2fd982bbd58225ce3bbc38f4053de

* ci: bug fixed for cache path and env

Change-Id: I28c42e10b6fff053bb6580926ca2353448cb042a

* Update .github/workflows/build-linux-cross.yml for cache path

Co-authored-by: Sigbjørn Skjæret <redacted>
* bugfixed for  build-linux-cross.yml,  syntax error

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Co-authored-by: cailinxi <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agoggml-backend : add root cause in error message if loading backend library fails ...
Rafal Lewczuk [Mon, 29 Sep 2025 11:17:09 +0000 (13:17 +0200)]
ggml-backend : add root cause in error message if loading backend library fails (llama/16172)

This PR adds additional information to an error message when loading backend library via ld_load_library() fails. This helps spotting why backend library did not load (missing library, missing dependency or unresolved symbol etc.).

2 months agosync : whisper.cpp (#1359)
Georgi Gerganov [Mon, 29 Sep 2025 13:49:11 +0000 (16:49 +0300)]
sync : whisper.cpp (#1359)

* ggml : Fix MKL detection by quoting BLAS_INCLUDE_DIRS (whisper/3426)

* sync : whisper.cpp

2 months agoci : print results [no ci] (#1358)
Georgi Gerganov [Mon, 29 Sep 2025 13:20:52 +0000 (16:20 +0300)]
ci : print results [no ci] (#1358)

2 months agoci : add self-hosted workflows (#1357)
Georgi Gerganov [Mon, 29 Sep 2025 12:15:13 +0000 (15:15 +0300)]
ci : add self-hosted workflows (#1357)

* ci : add self-hosted workflows

* ci : sync env vars with llama.cpp

* cont : always install python deps

* cont : cancel ongoing runs

* cont : exclude test-backend-ops from debug build

2 months agocmake : remove metal flag (llama/0)
Georgi Gerganov [Mon, 29 Sep 2025 09:33:38 +0000 (12:33 +0300)]
cmake : remove metal flag (llama/0)

2 months agosync : llama.cpp
Georgi Gerganov [Mon, 29 Sep 2025 09:31:53 +0000 (12:31 +0300)]
sync : llama.cpp

2 months agoggml : check cuda and metal argsort limits and add test (llama/16323)
Sigbjørn Skjæret [Mon, 29 Sep 2025 09:09:00 +0000 (11:09 +0200)]
ggml : check cuda and metal argsort limits and add test (llama/16323)

* check cuda argsort limits and add test

* add metal check

2 months agoggml : fix dependencies for ggml_set_rows (llama/16318)
Georgi Gerganov [Mon, 29 Sep 2025 05:41:28 +0000 (08:41 +0300)]
ggml : fix dependencies for ggml_set_rows (llama/16318)

2 months agovulkan: Fix validation failure in quantized flash attention (llama/16292)
Jeff Bolz [Mon, 29 Sep 2025 04:50:37 +0000 (23:50 -0500)]
vulkan: Fix validation failure in quantized flash attention (llama/16292)

2 months agoggml : fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32 (llama/16307)
Sigbjørn Skjæret [Sun, 28 Sep 2025 21:15:03 +0000 (23:15 +0200)]
ggml : fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32 (llama/16307)

* fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32

* add test that fails on simd

2 months agovulkan: 64-bit im2col (llama/16135)
Jeff Bolz [Sun, 28 Sep 2025 06:38:37 +0000 (01:38 -0500)]
vulkan: 64-bit im2col (llama/16135)

* vulkan: 64-bit im2col

Add variants of the im2col shaders that use buffer_device_address/buffer_reference,
and use 64-bit address calculations. This is needed for large convolutions used in
stable-diffusion.cpp.

* fix validation error for large im2col

2 months agometal : extend mat-mat multiplication support (llama/16225)
Georgi Gerganov [Sun, 28 Sep 2025 06:34:44 +0000 (09:34 +0300)]
metal : extend mat-mat multiplication support (llama/16225)

* metal : support mul_mm with src1->type == GGML_TYPE_F16

* metal : support mul_mm_id with src1->type == GGML_TYPE_F16

[no ci]

* metal : mul_mm support ne00 % 32 != 0

* metal : support mul_mm_id with ne00 % 32 != 0

* cont : remove unnecessary unrolls

* cont : simplify data loading

* metal : optimize mul_mm when output bounds checks are not needed

2 months agometal : fuse non-sequential nodes (llama/16102)
Georgi Gerganov [Sun, 28 Sep 2025 06:34:05 +0000 (09:34 +0300)]
metal : fuse non-sequential nodes (llama/16102)

* metal : fuse non-sequential nodes

* cont : add comment

* cont : simplify bounds checks

2 months agovulkan: handle mat_mul with A matrix > 4GB (llama/16176)
Jeff Bolz [Sun, 28 Sep 2025 01:36:34 +0000 (20:36 -0500)]
vulkan: handle mat_mul with A matrix > 4GB (llama/16176)

* vulkan: handle mat_mul with A matrix > 4GB

This change splits mat_mul operations with huge A matrix into chunks in the M
dimension. This works well for stable-diffusion use cases where the im2col
matrix has very large M.

Fix the order of setting the stride in mul_mm_cm2 - setting the dimension
clobbers the stride, so stride should be set after.

* build fixes

2 months agovulkan: support arbitrary KV dimension in flash attention (llama/16160)
Jeff Bolz [Sat, 27 Sep 2025 20:43:39 +0000 (16:43 -0400)]
vulkan: support arbitrary KV dimension in flash attention (llama/16160)

The "Clamp" spec constant is already based on whether KV is a multiple of Bc,
so use that to control whether bounds checking is performed. Add bounds checking
to the scalar and coopmat1 paths. Coopmat2 didn't need any changes (the K/V
tensors are already optionally clamped, nothing else needed to be changed).

2 months agovulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)
Acly [Sat, 27 Sep 2025 20:41:03 +0000 (22:41 +0200)]
vulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)

* don't use VULKAN_HPP_DEFAULT_DISPATCH_LOADER_DYNAMIC_STORAGE which can cause conflicts if application or other libraries do the same

2 months agoCUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)
Aman Gupta [Sat, 27 Sep 2025 16:49:32 +0000 (00:49 +0800)]
CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)

* CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32

This commit adds mul_mat_id support for ncols_dst >= 16. It does this by
packing ncols_dst tiles into the blockDim.y.

My tests on a RTX 3090 show that this is faster than the cuBLAS fallback
for f16 till bs=64, and for f32 till bs=32

* Review: refactor if statement

2 months agoCUDA: refactor and deduplicate vector FA kernels (llama/16208)
Johannes Gäßler [Sat, 27 Sep 2025 16:45:07 +0000 (18:45 +0200)]
CUDA: refactor and deduplicate vector FA kernels (llama/16208)

* CUDA: refactor and deduplicate vector FA kernels

2 months agovulkan: throw system error instead of SIGABRT during init on older devices (llama...
Dmytro Minochkin [Sat, 27 Sep 2025 16:26:46 +0000 (19:26 +0300)]
vulkan: throw system error instead of SIGABRT during init on older devices (llama/16156)

* Throw system error on old Vulkan driver rather than SIGABRT

* Optionally handle any potential error in vulkan init

2 months agovulkan: support GET_ROWS for k-quants (llama/16235)
Jeff Bolz [Sat, 27 Sep 2025 10:36:11 +0000 (06:36 -0400)]
vulkan: support GET_ROWS for k-quants (llama/16235)

The dequantize functions are copy/pasted from mul_mm_funcs.comp with very few
changes - add a_offset and divide iqs by 2. It's probably possible to call
these functions from mul_mm_funcs and avoid the duplication, but I didn't go
that far in this change.

2 months agodevops: add s390x & ppc64le CI (llama/15925)
Aaron Teo [Fri, 26 Sep 2025 18:03:33 +0000 (02:03 +0800)]
devops: add s390x & ppc64le CI (llama/15925)

* devops: move s390x and ppc64le ci build

we have access to ubuntu-24.04-s390x and ppc64le images now

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le for now since they have compiler errors

Signed-off-by: Aaron Teo <redacted>
* devops: stop warnings as errors

Signed-off-by: Aaron Teo <redacted>
* devops: switch to non-macro flag

Signed-off-by: Aaron Teo <redacted>
* devops: going the llama macro route

Signed-off-by: Aaron Teo <redacted>
* devops: add big-endian gguf test models

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le to test s390x, check test build

Signed-off-by: Aaron Teo <redacted>
* devops: dup .gguf.inp files for big-endian tests

Signed-off-by: Aaron Teo <redacted>
* devops: dup .gguf.out files for big-endian too

Signed-off-by: Aaron Teo <redacted>
* devops: add python setup and endian byteswap

Signed-off-by: Aaron Teo <redacted>
* devops: pooring thing does not have s390x python3

Signed-off-by: Aaron Teo <redacted>
* devops: add missing rust compiler for s390x

Signed-off-by: Aaron Teo <redacted>
* devops: try rust actions runner

Signed-off-by: Aaron Teo <redacted>
* Revert "devops: try rust actions runner"

This reverts commit 3f8db04356033d6c1d7eccc75ca396bc5298250c.

Signed-off-by: Aaron Teo <redacted>
* devops: try a different path for rust

Signed-off-by: Aaron Teo <redacted>
* devops: dump home directory and user info

Signed-off-by: Aaron Teo <redacted>
* devops: install gguf-py only

Signed-off-by: Aaron Teo <redacted>
* devops: missed relative path

Signed-off-by: Aaron Teo <redacted>
* devops: remove big-endian files since local swapping is working

Signed-off-by: Aaron Teo <redacted>
* devops: revert test-tokenizer-0 cmakelists

Signed-off-by: Aaron Teo <redacted>
* Fix unicode flags conversion from and to uint16_t

Bitfields are allocated in different order on s390x

Signed-off-by: Aaron Teo <redacted>
* Simplify byteswap command

Signed-off-by: Aaron Teo <redacted>
* Add byteswapping and git-lfs for test-tokenizers-ggml-vocabs

Signed-off-by: Aaron Teo <redacted>
* Fix endianness detection in vocab loader

Signed-off-by: Aaron Teo <redacted>
* Disable test-thread-safety on s390x

In this test a model is downloaded,
then immediately loaded to check if more downloads are needed,
and then used for test.

There is no clean way to separate all those steps
 to add byteswapping between them, so just skip this test.

Signed-off-by: Aaron Teo <redacted>
* Fix q8_0 test in test-quantize-fns

vec_signed uses unexpected rounding mode.
Explicitly use different rounding function.

Signed-off-by: Aaron Teo <redacted>
* devops: add big-endian stories260K

Signed-off-by: Aaron Teo <redacted>
* devops: add s390x test-eval-callback

Signed-off-by: Aaron Teo <redacted>
* devops: fix test does not exist

Signed-off-by: Aaron Teo <redacted>
* devops: fix model not found llama-eval-callback

Signed-off-by: Aaron Teo <redacted>
* Fix q3_K dot product error in test-quantize-fns on s390x

Array q8bytes had only 4 elements allocated, but 8 elements accessed.
This lead to write out of bounds and later read of overwritten values out of bounds
and incorrect result.

Signed-off-by: Aaron Teo <redacted>
* devops: re-enable ppc64le for testing

Signed-off-by: Aaron Teo <redacted>
* devops: activate test-thread-safety for s390x

Signed-off-by: Aaron Teo <redacted>
* devops: disable ppc64le tests

for some reason it keeps failing test-thread-safety tests and I do not
    have a machine that is able to replicate the tests.

Signed-off-by: Aaron Teo <redacted>
* devops: LLAMA_FATAL_WARNINGS=ON

Signed-off-by: Aaron Teo <redacted>
* Correct repository URL for s390x for test-thread-safety model

Signed-off-by: Aaron Teo <redacted>
* Fix fs_get_cache_directory

Ensure it works even if both XDG_CACHE_HOME and HOME are unset.
This might happen in containers.

Signed-off-by: Aaron Teo <redacted>
* Re-enable CI for ppc64le

Signed-off-by: Aaron Teo <redacted>
* Fortify ggml_rope_impl

Only memcpy data from sections argument if it's non-NULL.

Signed-off-by: Aaron Teo <redacted>
* Add TODO in struct unicode_cpt_flags to reimplement it in endian-independent way

* Update URL for big-endian model

* Update .github/workflows/build.yml

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update remaining mentions of BE models to ggml-org/models repo

---------

Signed-off-by: Aaron Teo <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Aleksei Nikiforov <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
2 months agometal : report OOM errors (llama/16274)
Georgi Gerganov [Fri, 26 Sep 2025 11:14:28 +0000 (14:14 +0300)]
metal : report OOM errors (llama/16274)

2 months agocommon : use cpp-httplib as a cURL alternative for downloads (llama/16185)
Adrien Gallouët [Fri, 26 Sep 2025 11:12:19 +0000 (13:12 +0200)]
common : use cpp-httplib as a cURL alternative for downloads (llama/16185)

* vendor : update httplib

Signed-off-by: Adrien Gallouët <redacted>
* common : use cpp-httplib as a cURL alternative for downloads

The existing cURL implementation is intentionally left untouched to
prevent any regressions and to allow for safe, side-by-side testing by
toggling the `LLAMA_CURL` CMake option.

Signed-off-by: Adrien Gallouët <redacted>
* ggml : Bump to Windows 10

Signed-off-by: Adrien Gallouët <redacted>
---------

Signed-off-by: Adrien Gallouët <redacted>
2 months agoggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)
Aaron Teo [Fri, 26 Sep 2025 10:27:25 +0000 (18:27 +0800)]
ggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)

* ggml-cpu: impl mxfp4 s390x

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: missing s = sumf

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix incorrect kval_mxfp4 type

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: rework mxfp4

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: missing delta calc

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix typo

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: fix typo for vec_splats

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: expand to 2 blocks per loop

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: add unroll to boost perf

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: back to 1 block per loop to test perf

Signed-off-by: Aaron Teo <redacted>
* Revert "ggml-cpu: back to 1 block per loop to test perf"

This reverts commit 1fe55724e2dc295701101bf838bdd4a512237492.

Signed-off-by: Aaron Teo <redacted>
* ggml-cpu: rm unroll from single block

Signed-off-by: Aaron Teo <redacted>
---------

Signed-off-by: Aaron Teo <redacted>
2 months agomusa: fix build warnings (llama/15611)
R0CKSTAR [Fri, 26 Sep 2025 00:56:10 +0000 (08:56 +0800)]
musa: fix build warnings (llama/15611)

Signed-off-by: Xiaodong Ye <redacted>
2 months agoCUDA: add a fused top-K MoE kernel (llama/16130)
Aman Gupta [Thu, 25 Sep 2025 14:35:05 +0000 (22:35 +0800)]
CUDA: add a fused top-K MoE kernel (llama/16130)

* CUDA: add a fused top-K MoE kernel

This kernel does the following:
1. softmax over the logits per token [n_experts, n_tokens]
2. argmax reduce over the top-k (n_experts_used) logits
3. write weights + ids to global memory

It is intended as fusion of softmax->top-k->get_rows pipeline for MoE models

* Refactor into ggml_cuda_should_use_topk_moe

* Review: Use better coalescing pattern, use WARP_SIZE, store logits into registers before

* Review: format + micro-optimizations

* Fix bug: fix tie breakers

* Add optional norm + clean-up code

* Use smem for final write

* Add bounds check

* Use better memory pattern for writeback

2 months agoggml : fix loongarch lsx compilation error (llama/15864)
junchao-zhao [Thu, 25 Sep 2025 09:22:55 +0000 (17:22 +0800)]
ggml : fix loongarch lsx compilation error (llama/15864)

2 months agoggml : remove -dev suffix from release version (#1355)
Daniel Bevenius [Fri, 26 Sep 2025 15:34:42 +0000 (17:34 +0200)]
ggml : remove -dev suffix from release version (#1355)

This commit removes the `-dev` suffix from the version string in
CMakeLists.txt and the release script. The version will now be
just be formatted as `MAJOR.MINOR.PATCH`.