vulkan: small mul_mat_vec optimizations (llama/10665)

* double the number of rows per workgroup

* Update ggml-vulkan.cpp

* Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats

* only increase the number of rows for amd and subgroup size 64

* fix missing NUM_ROWS for mul_mat_vec_iq4_nl_f16_f32, untested

* use subgroup min and max to check for gcn (requires https://github.com/ggerganov/llama.cpp/pull/10721)

* manual merge ggml-vulkan.cpp

* set min and max subgroup size in any case

* Also double the number of rows for Intel GPUs

SYCL: Reduce most of the compiler warnings (llama/10748)

* Try to reduce some unused and typecast warnings

* Reduce compiler warnings step 2

* add a newline at the end of the file

* Initialize nreduce as size_t

* [SYCL] Remove pragma directives from mmq.cpp

* SYCL: mmq add condition to prevent blocks_per_tile_x_row variable from becoming 0

* SYCL softmax: Initialize nreduce as size_t

* ggml-sycl.cpp: fix some trailing whitespaces

* SYCL: remove the unused variables instead of commenting it out

* SYCL poo2d kernel: set NAN for invalid pooling op

* SYCL gemm.hpp: remove pragma directives

* SYCL gemm.hpp: use const cast to properly support dnnl::memory

* SYCL: wkv6 remove a comment

* SYCL: clean comments step 2

* SYCL: clean comments and variables step 3

* SYCL: Use GGML_UNUSED for unused variables

* SYCL: remove extra empty lines and a comment

* Remove TODO

* cleanup spaces

* add a stdout for unsupported op

* use sycl printf over fprintf

* remove prints for CI

* SYCL ggml-sycl: pool2D use sycl::nan and remove if-else block

---------

Co-authored-by: Abhilash Majumder <redacted>

ggml : Fix compilation issues on ARM platform when building without fp16 (llama/10811)

CUDA: faster non-contiguous concat (llama/10760)

* faster uncontiguous concat

* Use a lambda to avoid code duplication

Co-authored-by: Diego Devesa <redacted>
* Update ggml/src/ggml-cuda/concat.cu

* add constexpr  and static assert

---------

Co-authored-by: Diego Devesa <redacted>

remove CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS (llama/10797)

other windows build fixes

Vulkan: Use improved q4_k and q5_k dequant code in dequant shaders (llama/10798)

Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats (llama/10721)

* Vulkan: Add VK_EXT_subgroup_size_control support to ensure full subgroups for coopmats

* Fix subgroup size control extension support check

Add accf32 and accf16 checks for coopmats

* Also disable coopmats on amdvlk

ggml: load all backends from a user-provided search path (llama/10699)

* feat: load all backends from a user-provided search path

* fix: Windows search path

* refactor: rename `ggml_backend_load_all_in_search_path` to `ggml_backend_load_all_from_path`

* refactor: rename `search_path` to `dir_path`

* fix: change `NULL` to `nullptr`

Co-authored-by: Diego Devesa <redacted>
* fix: change `NULL` to `nullptr`

---------

Co-authored-by: Diego Devesa <redacted>

vulkan: request round-to-even for fp16 in im2col/rope_head (llama/10767)

Vulkan doesn't mandate a specific rounding mode, but the shader_float_controls
feature allows rounding mode to be requested if the implementation supports it.

vulkan: dynamic subgroup size for the remaining k quants (llama/10745)

* q5_k

q4_k

q3_k

q2_k

q6_k multi row example

* revert as multi row isnt faster for k quants

CUDA: rename macros to avoid conflicts with WinAPI (llama/10736)

* Renames NVIDIA GPU-architecture flags to avoid name clashes with WinAPI. (e.g. CC_PASCAL, GPU architecture or WinAPI pascal compiler flag?)

* Reverts erroneous rename in SYCL-code.

* Renames GGML_CUDA_MIN_CC_DP4A to GGML_CUDA_CC_DP4A.

* Renames the rest of the compute capability macros for consistency.

vulkan: disable spirv-opt for coopmat shaders (llama/10763)

There are some bugs in the 1.3.296 SDK, so disable this. It isn't strictly
necessary anyway.

Add missing dependency on vulkan-shaders-gen, so shaders get recompiled when it
changes.

Fix coopmat support reporting when glslc doesn't support NV_coopmat2.

ggml : remove return from ggml_gallocr_allocate_node (ggml/1048)

This commit removes the return statement from ggml_gallocr_allocate_node
function.

The motivation behind this change is to make the code more readable and
consistent.

ggml : add check for grad_accs (ggml/1046)

* ggml : add check for grad_accs

This commit adds a check for grad_accs in ggml_graph_get_grad and
ggml_graph_get_grad_acc functions. This is necessary to avoid segfaults
when grad_accs is not initialized.

The motivation for this change is that I find it nice to be able to
print out a computation graph using ggml_graph_print but this function
segfaults when grad_accs is not initialized:
```console
(gdb) p g1
$2 = (ggml_cgraph *) 0x7ffff66004b0
(gdb) p *g1
$3 = {size = 2048, n_nodes = 1, n_leafs = 2, nodes = 0x7ffff6600500,
grads = 0x0, grad_accs = 0x0, leafs = 0x7ffff6604500,
visited_hash_set = {size = 4099, used = 0x7ffff6610518,
keys = 0x7ffff6608500}, order = GGML_CGRAPH_EVAL_ORDER_LEFT_TO_RIGHT}
(gdb) p ggml_graph_print(g1)
=== GRAPH ===
n_nodes = 1

Program received signal SIGSEGV, Segmentation fault.
0x0000555555579775 in ggml_graph_get_grad
(cgraph=0x7ffff66004b0,node=0x7ffff6600340)
    at /ggml/ggml/src/ggml.c:5990
5990  return igrad != GGML_HASHSET_FULL &&
          ggml_bitset_get(cgraph->visited_hash_set.used, igrad) ?
          cgraph->grads[igrad] : NULL;
```

* squash! ggml : add check for grad_accs

Fix the check in ggml_graph_get_grad. The check was incorrectly using
cgraph->grad_accs instead of cgraph->grads.

common : remove old types

ggml-ci

CUDA: fix shared memory access condition for mmv (llama/10740)

vulkan: fix compile warnings (llama/10731)

Vulkan: fix NaN in tanh.comp with AMD proprietary driver on Windows (llama/10723)

* Vulkan: fix NaN in tanh.comp

* Faster NaN-free tanh

vulkan: compile a test shader in cmake to check for coopmat2 support (llama/10713)

ggml : disable iq4_nl interleave size 8 (llama/10709)

ggml-ci

ggml : refactor online repacking (llama/10446)

* rename ggml-cpu-aarch64.c to .cpp

* reformat extra cpu backend.

- clean Q4_0_N_M and IQ4_0_N_M
  - remove from "file" tensor type
  - allow only with dynamic repack

- extract cpu extra bufts and convert to C++
  - hbm
  - "aarch64"

- more generic use of extra buffer
  - generalise extra_supports_op
  - new API for "cpu-accel":
     - amx
     - aarch64

* clang-format

* Clean Q4_0_N_M ref

Enable restrict on C++

* add op GGML_OP_MUL_MAT_ID for Q4_0_N_M with runtime repack

* added/corrected control on tensor size for Q4 repacking.

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

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp

Co-authored-by: Georgi Gerganov <redacted>
* add debug logs on repacks.

---------

Co-authored-by: Georgi Gerganov <redacted>

Vulkan: VK_KHR_cooperative_matrix support to speed up prompt processing (llama/10597)

* Vulkan: Implement VK_KHR_cooperative_matrix support in the matrix matrix multiplication shader

* Improve performance with better q4_k and q5_k dequant and store unrolling

* Add Vulkan MUL_MAT and MUL_MAT_ID accumulator precision selection

* Rework mulmat shader selection and compilation logic, avoid compiling shaders that won't get used by device

* Vulkan: Implement accumulator switch for specific mul mat mat shaders

* Vulkan: Unroll more loops for more mul mat mat performance

* Vulkan: Add VK_AMD_shader_core_properties2 support to read Compute Unit count for split_k logic

* Disable coopmat support on AMD proprietary driver

* Remove redundant checks

* Add environment variable GGML_VK_DISABLE_COOPMAT to disable VK_KHR_cooperative_matrix support

* Fix rebase typo

* Fix coopmat2 MUL_MAT_ID pipeline selection

metal : Extend how Llama.cpp locates metal resources (llama/10676)

* metal : Extend how Llama.cpp locates metal resources (llama/10675)

  * It searches the resource file in the directory where the current
    binary is located as well.
  * Resolves symbolic links.

Rationale:

When we plug this dependency into a Bazel build and run it in the
context of Bazel (e.g. testing):

  * the execution directory is often very different from where the files
    are located and no direct control over this (Bazel sandboxing),
  * the Bazel sandbox often use symbolic links to make files available.

With this patch, we can have the resource file added to the target,
can build and run tests in the context of Bazel.

* Update ggml/src/ggml-metal/ggml-metal.m

Co-authored-by: Georgi Gerganov <redacted>
* Update ggml/src/ggml-metal/ggml-metal.m

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

Co-authored-by: Georgi Gerganov <redacted>

vulkan: Add VK_NV_cooperative_matrix2 support for mul_mat and flash attention (llama/10206)

ruby : Add no_speech_thold (#2641)

* Remove Whisper::Model.[]

* Fix Whisper::Model::URI#request

* Make Whisper::Context#initialize accept pre-converted model name

* Use downloading pre-converted model feature for testing

* Update README

* Remove unnecessary task

* Move whisper/model.rb -> whisper/model/uri.rb

* Update document comment of Whisper::Context#initialize

* Don't show download progress when not tty

* Pass String to raise

* Use cache model file if download fails

* Add test for auto download

* Specify required Ruby version

* Fix a typo

* Remove unnecessary flags

* Initialize Whisper::Params#diarize explicitely

* Remove redundant code from README for simplicity

* Add Whisper::Params#no_speech_thold attribute

* Add test for Whisper::Params#no_speech_thold

stream : improve consistency in README (#2642)

whisper : support no_speech_thold (#2625)

* Implement no_speech_thold

no_speech_thold functionality is on par with OpenAI's whisper

* Addressed review comments

whisper : add single-timestamp logic (#2629)

* Fix hallucinations during silence

When the predicted tokens end with a single timestamp the the entire 30 segment should be considered as done, to avoid hallucinations for the remaining part of segment.
This behaviour is on par with openai's whisper. Refer to logic related to `single_timestamp_ending` in https://github.com/openai/whisper/blob/main/whisper/transcribe.py

* Accept review comments related to formatting.

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

Co-authored-by: Georgi Gerganov <redacted>

readme : fix typo (#2637)

cmake : fix "amd64" processor string (#2638)

vulkan : fix soft_max.comp division by zero (#2633)

This change prevents a division by zero error when p.KY is 0.

common : add cstdio header

stream : update build instructions

android : fix build and ci (#2624)

* Adding missing CMakeLists.txt include for ggm-cpu needed by whisper.android

* attempt to re-enable CI for JNI android

---------

Co-authored-by: Your Name <redacted>

models : fix typo in download-ggml-model.sh (#2623)

Introduced in #2589

ruby : Sync whisper.cpp and model download feature (#2617)

* Use C++17

* Add test for Pathname of model

* Make Whisper::Context#initialize accept Pathname

* Add shorthand for pre-converted models

* Update documents

* Add headings to API section in README [skip ci]

* Remove unused function

* Don't care about no longer included file

* Cosmetic fix

* Use conditional get when get model files

scripts : update to new build system

devops : add cmake

devops : update make commands

ggml : remove old files (skip) (#0)

ggml : sync remnants (skip) (#0)

scripts : remove amx from sync

ci : disable freeBSD builds [no ci]

readme : update build instructions

ci : disable CUDA and Android builds

ci : disable Obj-C build + fixes

make : shim cmake

talk-llama : sync llama.cpp

sync : ggml

ggml : add predefined list of CPU backend variants to build (llama/10626)

* ggml : add predefined list of CPU backend variants to build

* update CPU dockerfiles

ggml-cpu : fix HWCAP2_I8MM value (llama/10646)

vulkan: Implement "fast divide" (mul+shift) for unary ops like copy (llama/10642)

SYCL : Move to compile time oneMKL interface backend selection for NVIDIA backend (llama/10584)

* [SYCL] Move to Compile Time backend selection on oneMKL Interface for NVIDIA backend

Move to compile time selection to backend to avoid latency at run time.
Add it to all mkl gemm calls and only for NVIDIA backend.

Signed-off-by: nscipione <redacted>
* Formatting

* Address PR comments to increase readibility

---------

Signed-off-by: nscipione <redacted>

Avoid using __fp16 on ARM with old nvcc (llama/10616)

vulkan: optimize and reenable split_k (llama/10637)

Use vector loads when possible in mul_mat_split_k_reduce. Use split_k
when there aren't enough workgroups to fill the shaders.

ggml: add `GGML_SET` Metal kernel + i32 CPU kernel (ggml/1037)

* implemented cpu kernel

* add i32 test cases in test-backend-ops

* typedef `ggml_metal_kargs_set`

* implemented `kernel_set`

* memcpy

ggml : add `GGML_PAD_REFLECT_1D` operation (ggml/1034)

* ggml_pad_reflect_1d defined in header

* implemented on CPU

* called the forward pass

* impl Metal kernel

* added Metal kernel

* added OP_PAD_REFLECT_1D in test-backend-ops.cpp

* add test-pad-reflect-1d test case

* test case support multiple backend

files : remove make artifacts

common : fix compile warning

ggml-ci

ggml : move AMX to the CPU backend (llama/10570)

ggml : automatic selection of best CPU backend (llama/10606)

metal : small-batch mat-mul kernels (llama/10581)

* metal : small-batch mat-mul kernels

ggml-ci

* metal : add rest of types

ggml-ci

* metal : final adjustments

ggml-ci

* metal : add comments

ggml-ci

SYCL: Fix and switch to GGML_LOG system instead of fprintf (llama/10579)

* Switched to GGML_LOG

* Fix missing semicolon

ggml-cpu: replace AArch64 NEON assembly with intrinsics in ggml_gemv_q4_0_4x4_q8_0() (llama/10567)

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

vulkan: Dynamic subgroup size support for Q6_K mat_vec (llama/10536)

* subgroup 64 version with subgroup add. 15% faster

scalable version

tested for subgroup sizes 16-128

* check for subgroup multiple of 16 and greater than 16

* subgroup sizes are always a power of 2 (https://github.com/KhronosGroup/GLSL/issues/45)

* force 16 sequential threads per block

* make 16 subgroup size a constant

ggml : fix I8MM Q4_1 scaling factor conversion (llama/10562)

ggml-ci

ggml-cpu: fix typo in gemv/gemm iq4_nl_4_4 (llama/10580)

sycl : offload of get_rows set to 0 (llama/10432)

sycl : Reroute permuted mul_mats through oneMKL (llama/10408)

This PR fixes the failing MUL_MAT tests for the sycl backend.

CANN: RoPE operator optimization (llama/10563)

* [cann] RoPE operator optimization

* [CANN]Code Formatting

---------

Co-authored-by: noemotiovon <redacted>

vulkan: get the first command buffer submitted sooner (llama/10499)

This is an incremental improvement over #9118 to get work to the GPU a bit
sooner. The first part is to start with a smaller number of nodes before
the first submit, and ramp it up to the current 100 nodes/submit. The
second part is to reduce the dryrun overhead for all the nodes that just
need to request descriptor space.

With these changes I get around 1-2% speedup on RTX 4070 combined with my
old Haswell-era CPU.

ggml : remove redundant copyright notice + update authors

ggml : fix row condition for i8mm kernels (llama/10561)

ggml-ci

cmake : fix ARM feature detection (llama/10543)

ggml-ci

ggml-cpu: support IQ4_NL_4_4 by runtime repack (llama/10541)

* ggml-cpu: support IQ4_NL_4_4 by runtime repack

* ggml-cpu: add __ARM_FEATURE_DOTPROD guard

kompute : improve backend to pass test_backend_ops (llama/10542)

* kompute: op_unary: reject unsupported parameters

Signed-off-by: Sergio Lopez <redacted>
* kompute: softmax: implement ALiBi support

Signed-off-by: Sergio Lopez <redacted>
* kompute: rope: implement neox and phi3 support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_q4_k permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_[q4_0|q4_1|q8_0] permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_f16 permutted support

Signed-off-by: Sergio Lopez <redacted>
* kompute: op_mul_mat_q6_k permutted support

Signed-off-by: Sergio Lopez <redacted>
---------

Signed-off-by: Sergio Lopez <redacted>

CANN: Fix SOC_TYPE compile bug (llama/10519)

* CANN: Fix the bug build fail on Ascend310P under two cases:
1) Manual specify SOC_TYPE
2) Under some unusual compile environment

* Update the cann backend News content: Support F16 and F32 data type model for Ascend 310P NPU.

* fix CANN  compile fail bug: the assert in ascend kernel function doesn't supportted on some CANN version

CANN: ROPE operator optimization (llama/10540)

* [cann] ROPE operator optimization

Co-authored-by: noemotiovon <redacted>

Add some minimal optimizations for CDNA (llama/10498)

* Add some minimal optimizations for CDNA

* ggml_cuda: set launch bounds also for GCN as it helps there too

metal : fix group_norm support condition (llama/0)

vulkan: define all quant data structures in types.comp (llama/10440)

vulkan: Handle GPUs with less shared memory (llama/10468)

There have been reports of failure to compile on systems with <= 32KB
of shared memory (e.g. #10037). This change makes the large tile size
fall back to a smaller size if necessary, and makes mul_mat_id fall
back to CPU if there's only 16KB of shared memory.

vulkan: further optimize q5_k mul_mat_vec (llama/10479)

vulkan: skip integer div/mod in get_offsets for batch_idx==0 (llama/10506)

vulkan: optimize Q2_K and Q3_K mul_mat_vec (llama/10459)

mtgpu: Add MUSA_DOCKER_ARCH in Dockerfiles && update cmake and make (llama/10516)

Signed-off-by: Xiaodong Ye <redacted>

vulkan: fix group_norm (llama/10496)

Fix bad calculation of the end of the range. Add a backend test that
covers the bad case (taken from stable diffusion).

Fixes https://github.com/leejet/stable-diffusion.cpp/issues/439.

cmake : enable warnings in llama (llama/10474)

* cmake : enable warnings in llama

ggml-ci

* cmake : add llama_get_flags and respect LLAMA_FATAL_WARNINGS

* cmake : get_flags -> ggml_get_flags

* speculative-simple : fix warnings

* cmake : reuse ggml_get_flags

ggml-ci

* speculative-simple : fix compile warning

ggml-ci

ggml-cpu: cmake add arm64 cpu feature check for macos (llama/10487)

* ggml-cpu: cmake add arm64 cpu feature check for macos

* use vmmlaq_s32 for compile option i8mm check

CANN: Improve the Inferencing Performance for Ascend NPU Device (llama/10454)

* improve inferencing performance for ascend npu.

Co-authored-by: Frank Mai <redacted>
* some modification after review

* some modifications after review

* restore some modifications

* restore some modifications

---------

Co-authored-by: shanshan shen <redacted>
Co-authored-by: Frank Mai <redacted>

CANN: RoPE and CANCAT operator optimization (llama/10488)

Co-authored-by: noemotiovon <redacted>

vulkan: Fix a vulkan-shaders-gen arugment parsing error (llama/10484)

The vulkan-shaders-gen was not parsing the --no-clean argument correctly.
Because the previous code was parsing the arguments which have a value only
and the --no-clean argument does not have a value, it was not being parsed
correctly. This commit can now correctly parse arguments that don't have values.

metal : enable mat-vec kernels for bs <= 4 (llama/10491)

llama : accept a list of devices to use to offload a model (llama/10497)

* llama : accept a list of devices to use to offload a model

* accept `--dev none` to completely disable offloading

* fix dev list with dl backends

* rename env parameter to LLAMA_ARG_DEVICE for consistency

ggml : add support for dynamic loading of backends (llama/10469)

* ggml : add support for dynamic loading of backends

---------

Co-authored-by: Georgi Gerganov <redacted>

metal : minor code formatting

ggml : do not use ARM features not included in the build (llama/10457)

CANN: Support Ascend310P to accelerate F32 and F16 Model (llama/10216)

* CANN Support Ascend310P to accelerate F32 and F16 Model

* Add compile option soc type macro ASCEND_310P to ggml-cann lib

* Remove unused code

* Remove the ascend soc_type hard code compile option in CMakelist.txt

cuda : optimize argmax (llama/10441)

* cuda : optimize argmax

* remove unused parameter

ggml-ci

* fixup : use full warps

ggml-ci

* Apply suggestions from code review

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

* ggml : check ne00 <= INT32_MAX in argmax and argsort

---------

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

vulkan: predicate max operation in soft_max shaders/soft_max (llama/10437)

Fixes #10434

vulkan: copy iq4_nl LUT into shared memory (llama/10409)