CUDA: fix kernel selection logic for tile FA (llama/19686)

* CUDA: fix kernel selection logic for tile FA

* add comment

llamafile: powerpc: add FP16 MMA path for Q4/Q8 matmul (llama/19709)

Avoid xvi8ger4pp signed→unsigned bias correction by dequantizing Q4/Q8
inputs to FP16 and using FP16×FP16→FP32 MMA. This removes
post-processing overhead and improves performance.

Performance Impact:
1.5 ~ 2x improvement in PP_Speed for Q4 and Q8 Models,
measured with llama-bench and llama-batched-bench.
Q8 Model: granite-4.0-h-micro-Q8_0.gguf (from huggingface)
Q4 Model: Meta-Llama3-8b Q4 model (generated with llama-quantize from
f32 model)

llama-bench Q8 Model Results:
 model                                  size       params   backend      threads              test  Base t/s Patch t/s
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10               pp8           64.48 ± 4.72           73.99 ± 0.27
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp16           80.11 ± 0.32          112.53 ± 0.40
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp32           89.10 ± 0.27          152.95 ± 0.68
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp64           93.65 ± 0.25          187.83 ± 0.83
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp128           99.93 ± 0.02          201.32 ± 0.11
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp256          102.32 ± 0.40          208.32 ± 0.41
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp512          103.42 ± 0.40          209.98 ± 0.14
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             tg128           20.35 ± 0.01           19.57 ± 0.01

llama-bench Q4 Model Results:
 model                                  size       params   backend      threads              test                Base    t/s                 Patch   t/s
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10               pp8           34.77 ± 0.10           41.23 ± 0.08
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp16           40.81 ± 0.04           64.55 ± 0.15
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp32           44.65 ± 0.05           90.84 ± 0.22
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp64           47.49 ± 0.03          114.39 ± 0.11
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp128           49.29 ± 0.24          120.13 ± 0.19
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp256           49.77 ± 0.23          121.51 ± 0.11
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp512           49.89 ± 0.23          117.52 ± 0.10
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             tg128           13.40 ± 0.01           13.37 ± 0.00

Llama perplexity Results:

Model                     Base Final PPL Estimate Patch Final PPL Estimate
granite-4.0-h-micro-Q8_0    1.3862 +/- 0.04424         1.3868 +/- 0.04432
Meta-Llama3-8b Q4     1.3801 +/- 0.04116         1.3803 +/- 0.04116

Signed-off-by: Shalini.Salomi.Bodapati <redacted>

ggml webgpu: Fix bug in dispatching large matrix-vector multiplication (llama/19535)

* Fix bug in dispatching large matrix-vector multiplication

sync : llama.cpp

ggml webgpu: shader library organization (llama/19530)

* Basic JIT compilation for mul_mat, get_rows, and scale (#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <redacted>
Co-authored-by: Reese Levine <redacted>
* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* flashattention and matrix multiplication moved to new format

* clean up preprocessing

* Formatting

* remove duplicate constants

* Split large shaders into multiple static strings

---------

Co-authored-by: neha-ha <redacted>

sync : llama.cpp

vulkan: split mul_mat into multiple dispatches to avoid overflow (llama/19509)

* vulkan: split mul_mat into multiple dispatches to avoid overflow

The batch dimensions can be greater than the max workgroup count limit,
in which case we need to split into multiple dispatches and pass the base
index through a push constant.

Fall back for the less common p021 and nc variants.

* address feedback

opencl: refactor expm1 and softplus (llama/19404)

* opencl: refactor expm1

* opencl: refactor softplus

* opencl: use h for half literals

---------

Co-authored-by: Li He <redacted>

opencl: optimize mean and sum_row kernels (llama/19614)

* opencl: optimize mean and sum_row kernels

* opencl: add comment for max subgroups

* opencl: format

---------

Co-authored-by: Li He <redacted>

ggml: ggml-cpu: force-no-lto-for-cpu-feats (llama/19609)

When LTO enabled in build environments it forces all builds to have LTO
in place. But feature detection logic is fragile, and causing Illegal
instruction errors with lto. This disables LTO for the feature
detection code to prevent cross-module optimization from inlining
architecture-specific instructions into the score function. Without this,
LTO can cause SIGILL when loading backends on older CPUs (e.g., loading
power10 backend on power9 crashes before feature check runs).

cuda : enable CUDA graphs for MMID 1 <= BS <= 4 (llama/19645)

* cuda : enable CUDA graphs for MMID BS <= 4

* cont : add stream capture check

Co-authored-by: Oliver Simons <redacted>
* cont : add MMVQ_MMID_MAX_BATCH_SIZE

---------

Co-authored-by: Oliver Simons <redacted>

ggml : make `ggml_is_view` as API (llama/19539)

* make `ggml_is_view` as API

* introduce `ggml_aux_is_view` as inline version for internal use.

* change `ggml_aux_is_view` to  `ggml_impl_is_view`

Adjust workaround for ROCWMMA_FATTN/GFX9 to only newer ROCm veresions (llama/19591)

Avoids issues with ROCm 6.4.4.

Closes: https://github.com/ggml-org/llama.cpp/issues/19580
Fixes: 6845f7f87 ("Add a workaround for compilation with ROCWMMA_FATTN and gfx9 (#19461)")
Signed-off-by: Mario Limonciello (AMD) <redacted>

ggml: aarch64: Implement SVE in Gemm q4_k 8x8 q8_k Kernel (llama/19132)

* Updated repack.cpp

* Updated repack.cpp

* Updated repack.cpp

* Added if condition to support only vector length 256.

* Changed the format removed comments and duplicate variable

* If SVE 256 not present then was using generic function to compute, hence slowing the performance.

So added code if SVE 256 is not present then use NEON code.

* Code format change suggestion

---------

Co-authored-by: Vithule, Prashant <redacted>

cuda: optimize iq2xxs/iq2xs/iq3xxs dequantization (llama/19624)

* cuda: optimize iq2xxs/iq2xs/iq3xxs dequantization

- load all 8 int8 for a grid position in one load
- calculate signs via popcnt instead of fetching from ksigns table
- broadcast signs to drop individual shift/mask

* cuda: iq2xxs: simplify sum scaling

express `(sum * scale + sum / 2) / 4` as `(sum * (scale * 2 + 1)) / 8`
express `((aux32 >> 28) * 2 + 1)` as `(aux32 >> 27 | 1)`

saves 3 registers for mul_mat_vec_q (152 -> 149) according to nsight
AFAICT no overflow can occur here as iq2xxs values are far too small

* uint -> uint32_t

error: identifier "uint" is undefined

cmake : check if KleidiAI API has been fetched (llama/19640)

This commit addresses a build issue with the KleidiAI backend when
building multiple cpu backends. Commmit
3a00c98584e42a20675b6569d81beadb282b0952 ("cmake : fix KleidiAI install
target failure with EXCLUDE_FROM_ALL") introduced a change where
FetchContent_Populate is called instead of FetchContent_MakeAvailable,
where the latter does handle this case (it is idempotent but
FetchContent_Populate is not).

I missed this during my review and I should not have commited without
verifying the CI failure, sorry about that.

ggml : avoid UB in gemm ukernel (llama/19642)

ggml-cpu: optimize ggml_vec_dot_bf16 for s390x (llama/19399)

ggml-cpu: FA add GEMM microkernel (llama/19422)

* ggml-cpu: FA add GEMM microkernel

* add guard for sizeless vector types

* fix case where DV % GGML_F32_EPR !=0

* move memset out of the loop

* move another memset out of the loop

* use RM=4 for arm

* simd_gemm: convert everything to int

* convert everything to size_t to avoid warnings

* fixup

* add pragma for ignoring aggressive loop optimizations

cmake : fix KleidiAI install target failure with EXCLUDE_FROM_ALL (llama/19581)

* cmake: fix KleidiAI install target failure with EXCLUDE_FROM_ALL

Fix for the bug #19501 by adding EXCLUDE_FROM_ALL to FetchContent_Declare. This properly excludes KleidiAI from both build and install targets, preventing install failures when GGML_CPU_KLEIDIAI=ON is used.

The KleidiAI source files are still compiled into libggml-cpu.so, preserving all functionality.

* addressed code review comments

ci : Upgrade GitHub Actions for Node 24 compatibility (#1426)

Signed-off-by: Salman Muin Kayser Chishti <redacted>

ggml : bump version to 0.9.7 (#1425)

sync : whisper.cpp

sync : llama.cpp

models : optimize qwen3next graph (llama/19375)

* models : optimizing qwen3next graph

* cont

* wip

* wip

* wip

* wip

* wip

* wip

* wip

* wip

* wip

* wip

* cont : remove redundant q, g chunking

* minor

* minor

* avoid passing masks around

* avoid concats during chunking

* naming + shapes

* update names and use prefix to disable CUDA graphs

ggml : fix GGML_DEBUG with OpenMP (llama/19599)

last_graph is only available without OpenMP, but
ggml_graph_compute_thread() is called in both cases.

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

metal : fix ACC op (llama/19427)

vulkan: support L2_NORM with contiguous rows (llama/19604)

vulkan: support GGML_OP_SET (llama/19584)

vulkan: Add vendor id for Qualcomm drivers (llama/19569)

This commit allows Qualcomm native vulkan driver to be used on Windows
instead of Mesa Dozen.

hexagon: further optimizations and refactoring for flash attention (llama/19583)

* ggml-hexagon: fa improvements

ggml-hexagon: optimize flash attention calculations with improved variable handling

ggml-hexagon: streamline flash attention operations by removing redundant checks for FP32

ggml-hexagon: optimize hvx_dot_f16_f16_aa_rx2 by simplifying variable handling for unused elements

ggml-hexagon: optimize flash attention by changing slope vector type to F16

* hexfa: fixed test-backend-ops failurs due to leftover element handling

* hexagon: refactor and optimize fa to use local context struct

* ggml-hexagon: optimize flash-attention using hvx_vec_expf

Use HVX for online softmax.

---------

Co-authored-by: chraac <redacted>

vulkan: restore -inf check in FA shaders (llama/19582)

Fix wrong memcpy length for block_interleave == 4 (llama/19575)

fix vulkan ggml_acc only works in 3d but not 4d (llama/19426)

* fix vulkan ggml_acc only works in 3d but not 4d

* removed clamp in test_acc_block

* use the correct stride and its test case

* cuda : fix "supports op" condition

* change src0 to src1 in ggml_vk_acc. Update acc.comp with jeffbolznv\'s suggestion except to keep the boundary check

* version without boundary check

* revert back to boundary check version

---------

Co-authored-by: Georgi Gerganov <redacted>

CUDA: loop over ne2*ne3 in case it overflows (llama/19538)

* CUDA: loop over ne2*ne3 in case it overflows

* use fastdiv

CUDA: Do not mutate cgraph for fused ADDs (llama/19566)

* Do not mutate cgraph for fused ADDs

1. We should try to minimize in-place changes to the incoming
   ggml_cgraph where possible (those should happen in graph_optimize)
2. Modifying in-place leads to an additional, unnecessary graph capture
   step as we store the properties before modifying the graph in-place
   in the cuda-backend

* Assert ggml_tensor is trivially copyable

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

Co-authored-by: Aman Gupta <redacted>
---------

Co-authored-by: Aman Gupta <redacted>

metal : improve concurrency (llama/19555)

metal : support GGML_OP_SET (llama/19548)

hexagon: fix typo in vtcm_needs_release (llama/19545)

opencl: add basic support for q4_1 (llama/19534)

* opencl: add q4_1 mv

* opencl: clean up

* opencl: add flattened q4_1 mv

* opencl: clean up

* opencl: add basic q4_1 mm

* opencl: fix whitespace

* opencl: add general q4_0 mm

metal : update sum_rows kernel to support float4 (llama/19524)

Add a workaround for compilation with ROCWMMA_FATTN and gfx9 (llama/19461)

There is an upstream problem [1] with AMD's LLVM 22 fork and
rocWMMA 2.2.0 causing compilation issues on devices without
native fp16 support (CDNA devices).

The specialized types aren't resolved properly:
```
/opt/rocm/include/rocwmma/internal/mfma_impl.hpp:2549:37: error: ambiguous partial specializations of 'amdgcn_mfma<__half, __half, __half, 16, 16, 16>'
 2549 |             using ARegsT = typename Impl::ARegsT;
```

Add a workaround to explicitly declare the types and cast when
compiling with HIP and ROCWMMA_FATTN [2].  When this is actually
fixed upstream some guards can be used to detect and wrap the
version that has the fix to only apply when necessary.

Link: https://github.com/ROCm/rocm-libraries/issues/4398
Link: https://github.com/ggml-org/llama.cpp/issues/19269
Signed-off-by: Mario Limonciello <redacted>

hexagon: further optimization and tuning of matmul and dot kernels (llama/19407)

* ggml-hexagon: implement 2x2 matmul kernel

* hexmm: implement vec_dot_rx2x2 for Q8_0 and MXFP4

* hexagon: fix editor config failures

* hexagon: refactor matmul ops to use context struct and remove wrappers

Also implement vec_dot_f16 2x2

* hexagon: refactor dyn quantizers to use mmctx

* hexagon: remove mm fastdiv from op_ctx

* hexagon: refactor matmul entry point to reduce code duplication

---------

Co-authored-by: Trivikram Reddy <redacted>

opencl: add general Q6_K mm and Q4_K mv (llama/19347)

* opencl: add general q6_k mm

* opencl: refine condition for q6_K mm

* opencl: add general q4_K mv

* opencl: fix whitespace

ggml : unary ops support non-cont src0 + metal F16 unary ops (llama/19511)

* ggml : unary ops support non-cont src0

* metal : support F16 unary ops + fix ELU

metal : extend l2_norm support for non-cont src0 (llama/19502)

hexagon: Add ARGSORT, DIV, SQR, SQRT, SUM_ROWS, GEGLU (llama/19406)

* hexagon: add ARGSORT op

Co-authored-by: Yarden Tal <redacted>
* hexagon: argsort reject tensors with huge rows for now

* Adding support for DIV,SQR,SQRT,SUM_ROWS ops in hexagon backend

* hexagon : Add GEGLU op

* hexagon: fix editor config check

* hexagon: rewrite and optimize binary ops ADD/SUB/MUL/DIV/ADD_ID to use DMA

---------

Co-authored-by: Yarden Tal <redacted>
Co-authored-by: Manohara Hosakoppa Krishnamurthy <redacted>

ggml : extend bin bcast for permuted src1 (llama/19484)

* tests : extend bin bcast for permuted src1

* cont : extend bin support

* cont : s0 is always 1

* tests : simplify

metal : consolidate unary ops (llama/19490)

CUDA : Update CCCL-tag for 3.2 to final release from RC (llama/19486)

CCCL 3.2 has been released since it was added to llama.cpp as part of
the backend-sampling PR, and it makes sense to update from RC to final
released version.

https://github.com/NVIDIA/cccl/releases/tag/v3.2.0

Plug memory leaks and free resources on shutdown (llama/19315)

* Fix memory leaks in shader lib, backend, backend_context, buffer_context, and webgpu_buf_pool

* Free pools

* Cleanup

* More cleanup

* Run clang-format

* Fix arg-parser and tokenizer test errors that free an unallocated buffer

* Fix device lost callback to not print on device teardown

* Fix include and run clang-format

* remove unused unused

* Update binary ops

---------

Co-authored-by: Reese Levine <redacted>

test: fix IMROPE perf test case (llama/19465)

ggml-cpu: arm64: q6_K repack gemm and gemv (and generic) implementations (dotprod) (llama/19360)

* First working version of GEMM and GEMV

* interleave loads and compute

* Clang-format

* Added missing fallback. Removed tested TODO.

* Swap M and N to be consistent with the repack template convention

ggml : use noexcept overload for is_regular_file in backend registration (llama/19452)

using noexcept std::filesystem::directory_entry::is_regular_file
overload prevents abnormal termination upon throwing an error
(as caused by symlinks to non-existent folders on linux)

Resolves: #18560

CANN: Remove unnecessary wrapper for `gml_backend_buft_is_cann` (llama/18968)

CANN: implement quantized MUL_MAT_ID for MoE models (llama/19228)

Implement ggml_cann_mul_mat_id_quant function to support quantized matrix
multiplication for Mixture of Experts (MoE) architectures on CANN backend.

Key features:
- Support Q4_0 and Q8_0 quantized weight formats
- Use IndexSelect to dynamically route expert-specific weights based on indices
- Leverage WeightQuantBatchMatmulV2 for efficient quantized computation
- Handle automatic F16 type conversion for hardware compatibility
- Support both per-expert and broadcast input modes

Implementation details:
- Extract expert weights and scales using CANN IndexSelect operation
- Process each batch and expert combination independently
- Create proper tensor views with correct stride for matmul operations
- Automatic input/output type casting to/from F16 as needed

Testing: All test cases passed for supported types (F32, F16, Q4_0, Q8_0).

cuda : extend GGML_OP_PAD to work with non-cont src0 (llama/19429)

* cuda : extend GGML_OP_PAD to work with non-cont src0

* tests : add permuted pad

CUDA: Fix non-contig rope (llama/19338)

* Rename variables + fix rope_neox

Seems memory layout is shared with Vulkan so we can port fix from
https://github.com/ggml-org/llama.cpp/pull/19299

* Fix rope_multi

* Fix rope_vision

* Fix rope_norm

* Rename ne* to ne0* for consistent variable naming

* cont : consistent stride names

---------

Co-authored-by: Georgi Gerganov <redacted>

sync : llama.cpp

metal : consolidate bin kernels (llama/19390)

* metal : refactor bin kernels

* cont

* cont : fix cv

metal : fix event synchronization in cpy_tensor_async (llama/19402)

ggml-webgpu: JIT compile binary operators and handle binding overlaps (llama/19310)

* ggml webgpu: port binary operators to use pre-wgsl

* Add binary.wgsl: unified shader with conditionals for all 4 ops

* Add gen_binary_shaders.cpp: build tool for using pre_wgsl preprocessor

* Remove bin_op.tmpl.wgsl and binary.wgsl (Python template)

* Update CMake to generate binary operator shaders at build time

* ggml-webgpu: migrate binary ops to JIT compilation with overlap handling

* port binary operators from AOT to pre-wgsl JIT compilation

* add src1=dst overlap handling for binary ops

* use compile-time workgroup size defines instead of runtime overrides

* ggml-webgpu: complete overlap handling for binary ops

* add support for inplace & overlap case in binding setup

* restructure conditional logic to handle all overlap cases

* ensure all buffer bindings are correctly assigned for edge cases

* ggml-webgpu: remove unused binary overlap cases

Remove src0==src1 binary overlap case that never occurs in practice.

* keep INPLACE (src0==dst), OVERLAP (src1==dst), DEFAULT

* remove unused src0==src1 and all-same variant

* refactor wgsl to eliminate duplication

sync : llama.cpp

sycl: add F16 support for GGML_OP_CEIL (llama/19306)

* Fix SYCL CEIL operator

* sycl: implement GGML_OP_CEIL

tests: reduce number of FA test permutations (llama/19381)

Only test non-F16 for head size 64 and 72 (one a multiple of QK, one not).

vulkan: For coopmat2 FA, use fp16 accumulators for the final result (llama/19376)

The cpu and cuda backends use fp16 for the VKQ accumulator type, this change
does the same for vulkan. This helps particularly with large head sizes which
are very register-limited.

I tried this for the coopmat1 path and it slowed down a bit. I didn't try for
scalar.

I applied the softmax bias that the cuda backend uses to avoid overflow,
although I was not able to reproduce the original bug without it.

vulkan: make FA mask/softcap enables spec constants (llama/19309)

* vulkan: make FA mask/softcap enables spec constants

* don't specialize for sinks

* bump timeout a little bit

metal : skip loading all-zero mask (llama/19337)

* metal : skip loading all-zero mask

* cont : minor

cuda : cuda graphs now compare all node params (llama/19383)

metal : adaptive CPU/GPU interleave based on number of nodes (llama/19369)

vulkan: Preprocess FA mask to detect all-neg-inf and all-zero. (llama/19281)

Write out a 2-bit code per block and avoid loading the mask when it
matches these two common cases.

Apply this optimization when the mask is relatively large (i.e. prompt
processing).

metal : add diag (llama/19330)

vulkan: fix GPU deduplication logic. (llama/19222)

* vulkan: fix GPU deduplication logic.

As reported in https://github.com/ggml-org/llama.cpp/issues/19221, the
(same uuid, same driver) logic is problematic for windows+intel igpu.

Let's just avoid filtering for MoltenVK which is apple-specific, and
keep the logic the  same as before 88d23ad5 - just dedup based on UUID.

Verified that MacOS + 4xVega still reports 4 GPUs with this version.

* vulkan: only skip dedup when both drivers are moltenVk

vulkan: Set k_load_shmem to false when K is too large (llama/19301)

vulkan: fix non-contig rope (llama/19299)

metal : add missing includes (llama/19348)

tests : add non-cont, inplace rope tests (llama/19296)

* tests : add non-cont, inplace rope tests

* cont : exercise dim 3

Co-authored-by: Jeff Bolz <redacted>
* cont : more dim3 exercises

---------

Co-authored-by: Jeff Bolz <redacted>

ggml-virtgpu: make the code thread safe (llama/19204)

* ggml-virtgpu: regenerate_remoting.py: add the ability to deprecate a function

* ggml-virtgpu: deprecate buffer_type is_host remoting

not necessary

* ggml-virtgpu: stop using static vars as cache

The static init isn't thread safe.

* ggml-virtgpu: protect the use of the shared memory to transfer data

* ggml-virtgpu: make the remote calls thread-safe

* ggml-virtgpu: backend: don't continue if couldn't allocate the tensor memory

* ggml-virtgpu: add a cleanup function for consistency

* ggml-virtgpu: backend: don't crash if buft->iface.get_max_size is missing

* fix style and ordering

* Remove the static variable in apir_device_get_count

* ggml-virtgpu: improve the logging

* fix review minor formatting changes

ggml-cpu: use LUT for converting e8->f32 scales on x86 (llama/19288)

* ggml-cpu: use LUT for converting e8->f32 scales on x86

* add dispatch based on macro

metal : add solve_tri (llama/19302)

vulkan: disable coopmat1 fa on Nvidia Turing (llama/19290)

CUDA: use mmvq for mul-mat-id for small batch sizes (llama/18958)

* CUDA: use mmvq for mul-mat-id for small batch sizes

* add mmvq too

* Fix perf issue on ampere. Use mmvf mm-id only for non-nvidia GPUs

* templatize multi_token_path

metal : minor cleanup (llama/19251)

CUDA: Fix loop unrolling for BW in mul_mat_q_stream_k_fixup (llama/19053)

By providing stride_* variables as size_t (i.e., 64-bit) the compiler can
correctly unroll the [two for-loops](https://github.com/ggml-org/llama.cpp/blob/557515be1e93ed8939dd8a7c7d08765fdbe8be31/ggml/src/ggml-cuda/mmq.cuh#L3789-L3816)
on BW. This gives some perf for prefill/pp phase on BW, while not affecting
other SMs:

| GPU                                                     | Model                 | Test   |   t/s master |   t/s osimons/fix_bw_mmq_fixup_kernel |   Speedup |
|:--------------------------------------------------------|:----------------------|:-------|-------------:|--------------------------------------:|----------:|
| NVIDIA RTX 6000 Ada Generation                          | gpt-oss 20B MXFP4 MoE | pp8096 |      8404.05 |                               8375.79 |      1.00 |
| NVIDIA RTX 6000 Ada Generation                          | llama 3B Q4_K_M       | pp8096 |     16148.93 |                              16019.60 |      0.99 |
| NVIDIA RTX 6000 Ada Generation                          | llama 8B Q4_0         | pp8096 |      8008.29 |                               7978.80 |      1.00 |
| NVIDIA RTX 6000 Ada Generation                          | nemotron_h 9B BF16    | pp8096 |      4263.16 |                               4248.53 |      1.00 |
| NVIDIA RTX 6000 Ada Generation                          | nemotron_h 9B Q4_K_M  | pp8096 |      5165.11 |                               5157.43 |      1.00 |
| NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition | gpt-oss 20B MXFP4 MoE | pp8096 |     12582.80 |                              12758.37 |      1.01 |
| NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition | llama 3B Q4_K_M       | pp8096 |     16879.10 |                              17619.47 |      1.04 |
| NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition | llama 8B Q4_0         | pp8096 |     10649.90 |                              10982.65 |      1.03 |
| NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition | nemotron_h 9B BF16    | pp8096 |      7717.73 |                               7716.22 |      1.00 |
| NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition | nemotron_h 9B Q4_K_M  | pp8096 |      7301.90 |                               7370.38 |      1.01 |

ggml: added cleanups in ggml_quantize_free (llama/19278)

Add missing cleanup calls for IQ2_S, IQ1_M quantization types and IQ3XS with 512 blocks during quantization cleanup.

cuda : revert CUDA_SCALE_LAUNCH_QUEUES override until investigated (llama/19227)

Hangs were reported on Jetson Orin AGX if we set CUDA_SCALE_LAUNCH_QUEUES=4x. Reverting the previous PR (#19042) and updating the document to consider setting CUDA_SCALE_LAUNCH_QUEUES=4x for faster throughput on multi-GPU systems.

opencl: refactor some ops, concat, repeat, tanh and scale (llama/19226)

* opencl: refactor concat

* opencl: refactor repeat

* opencl: refactor tanh

* opencl: enable fp16 for tanh

* opencl: refactor scale

* opencl: fix unused variables

ggml-cpu: FA split across kv for faster TG (llama/19209)

* ggml-cpu: split across kv for faster TG

* simplify sinks application

* add ref impl

Remove support for Nvidia & AMD GPU, because the oneAPI plugin for Nvidia & AMD GPU is unavailable: download/installation channels are out of work. (llama/19246)

User can't build up the software for Nvidia & AMD GPU.
rm the oneMath since it is only used in NV and AMD code path.

sycl: implement GGML_OP_TOP_K (llama/19242)

metal : support virtual devices (llama/18919)

* metal : support virtual devices

* cont : manage buffer type context memory

* metal : add events

* cont : implement cpy_tensor_async

ggml-backend: fix async set/get fallback sync (llama/19179)

docs : Minor cleanups (llama/19252)

* Update old URLs to github.com/ggml-org/

* Bump copyrights

Remove pipeline cache mutexes (llama/19195)

* Remove mutex for pipeline caches, since they are now per-thread.

* Add comment

* Run clang-format

* Cleanup

* Run CI again

* Run CI once more

* Run clang-format

Bump cmake max version (needed for Windows on Snapdragon builds) (llama/19188)

* Bump max cmake version (needed for Windows on Snapdragon builds)

* cmake: move max version setting into ggml/CMakeLists

ggml-hexagon: flash-attention and reduce-sum optimizations (llama/19141)

* wip

* ggml-hexagon: add vectorized dot product function for FP32 and FP16 accumulation

* ggml-hexagon: optimize dot product functions for FP16 and FP32 with new vectorized implementations

* wip

* ggml-hexagon: optimize hvx_vec_dump_f32_n and hvx_vec_reduce_sum_qf32x2 functions for improved performance

* ggml-hexagon: refactor dot product functions to use a common loading function for improved readability

* optimize vector dot product functions to use unified reduction for improved performance

* wip

* ggml-hexagon: add vectorized dot product function for FP32 and FP16 accumulation

* ggml-hexagon: optimize dot product functions for FP16 and FP32 with new vectorized implementations

* wip

* ggml-hexagon: optimize hvx_vec_dump_f32_n and hvx_vec_reduce_sum_qf32x2 functions for improved performance

* ggml-hexagon: refactor dot product functions to use a common loading function for improved readability

* optimize vector dot product functions to use unified reduction for improved performance

* hexagon: optimize reduce-sum for v75+

* hexagon: always keep row_sums in sf/fp32

* ggml-hexagon: enhance directory checks for HEXAGON_SDK_ROOT and HEXAGON_TOOLS_ROOT

* fix compiling error after rebase

---------

Co-authored-by: Max Krasnyansky <redacted>

opencl: add optimized q8_0 mm kernel for adreno (llama/18871)

* Add Q8_0 OpenCL kernel

Co-authored-by: yunjie <redacted>
* opencl: fix build for non-adreno

* opencl: refactor q8_0

* opencl: enforce subgroup size of 64 for adreno for q8_0

* For A750 and older generations, subgroup size can be 64 or 128.
  This kernel assumes subgroup size 64.

* opencl: suppress warning when adreno kernels are disabled

---------

Co-authored-by: yunjie <redacted>
Co-authored-by: Li He <redacted>

Correctly fetch q8_1 quantize pipeline in test as needed by 8a3519b (llama/19194)

tests : add GQA=20 FA test (llama/19095)

ci : remove "Release" word from the title of the release