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

sync : whisper.cpp

ggml-ci

scripts : update whisper sync

sync : llama.cpp

ggml-ci

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 (#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 (#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

ci : update requirements.txt

ggml-ci

ci : remove opencl workflow

ci : fix pip env

ggml-ci

files : remove make artifacts

common : fix compile warning

ggml-ci

authors : update

sync : llama.cpp

ggml-ci

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

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

scripts : remove amx from sync

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>

tests : fix compile warning

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)

vulkan: further optimize mul_mat_vec using larger loads (llama/10387)

* vulkan: Use pipeline_robustness to disable robustness in mul_mat_vec.

Add some early returns for nonexistent rows in mul_mat_vec shaders. These
can only be hit when dispatching a 2D grid of workgroups. Fix the logic
for the 2D grid of workgroups to round up.

Enable the pipeline robustness extension if it's available, and use it to
disable robustness for these pipelines. The instructions to do the bounds
checking contend for the same ALU resources as the bit twiddling dequant
instructions.

* vulkan: Add GLSL structure aliases for quant types to allow larger loads

In Vulkan it's not possible to cast pointer types, so instead you have to
declare an aliased binding for the memory with a different type. This
commit adds aliases for the quant formats using 16b ints, and in a few
places where the struct size is a multiple of 4 also using 32b ints.
Currently only q4_k's aliases are used, but others will be used in
subsequent commits.

* vulkan: use larger loads in q5_k and q6_k shaders.

Similar to the optimization I did in q4_k recently, this vectorizes some loads
and reduces the number of bit twiddling instructions.

* vulkan: use larger K step per iteration in mul_mat_vec.

Add vec4 dequantization functions, and use them to do K=8 per iteration in
mul_mat_vec. This uses 16b loads for the quant values and 128b loads for B
which helps reduce the load on the memory system.

The K_PER_ITER==2 logic is still there, just for F16/F32, and really only
because they support unaligned sizes.

Tweak the num_iters/unrolling logic to be simpler and catch a couple missed
unrolling opportunities.

add cmake rvv support (llama/10411)

CUDA: remove unnecessary warp reduce in FA (#1032)

* kqmax_new_j in every thread within warp is same after operate at line 199,this reduce can be omit

* same problem in vec32

---------

Co-authored-by: ZhaoXiaoYu <redacted>

feat: add `GGML_UNARY_OP_ARGMAX` Metal kernel (#1019)

* implemented argmax kernel

* tpig -> tgpig

* change to strides

* contiguous assertions

* kernel working and tested

* argmax simd parallel implementation

* added 2 new tests for argmax in test-backend-ops

* cosmit

* added 3 tests cases for perf eval

* add test_argmax in make_test_cases_perf

* Update test-backend-ops.cpp

Co-authored-by: Diego Devesa <redacted>
---------

Co-authored-by: Diego Devesa <redacted>

metal : add `GGML_OP_CONV_TRANSPOSE_1D` kernels (#1026)

* wip

* wip implementation f32

* kernel conv transpose 1d f32 working

* initial commit

examples: link to HuggingFace mirror of MNIST data (#1030)

Fix build docs for hip (#1029)

Do not include arm_neon.h when compiling CUDA code (#1028)

Create .gitmodules for the kompute backend (#1024)

sync : whisper.cpp

ggml/sched : do not skip views in pre-assignments

ggml-opt: fix data corruption (#1022)

sync : llama.cpp

Add required ggml-base and backend libs to cmake pkg (llama/10407)

sync : llama.cpp

cuda : fix CUDA_FLAGS not being applied (llama/10403)

sync : llama.cpp

sycl : Add option to set the SYCL architecture for all targets (llama/10266)

* Add option to set the SYCL architecture for all targets
* Convert GGML_SYCL_HIP_TARGET to the more generic GGML_SYCL_ARCH option
* Document that setting GGML_SYCL_ARCH can improve the performance

vulkan: Optimize soft_max (llama/10301)

* vulkan: Optimize soft_max

Large soft_max could already saturate memory, but small/medium sizes were
pretty slow. The bulk of the gains for them comes from using a smaller
workgroup size, and making the workgroup size match the subgroup size also
makes the barriers much cheaper.

Cache some values in locals to avoid refetching/recomputing. And stamp
out a few "template instantiations" so smaller cases will fully unroll.

Add a missing early return for OOB rows. This happens when there are more
than 512 rows and the dispatch is 512 x H.

* vulkan: Further soft_max optimizations

Restore the workgroup size of 512 case, use it for >1024.

Use unrollable loops for more iteration counts.

sycl: Revert MUL_MAT_OP support changes (llama/10385)

cuda : only use native when supported by cmake (llama/10389)

vulkan: remove use of null initializer (llama/10372)

Seems like this isn't working for vulkan-over-metal when the array is sized
by a spec constant. Maybe a spirv-cross limitation?

metal : fox offset integer overflows in im2col (#1015)

-- While running StableDiffusion.cpp locally with Metal some offsets overflow and results in incorrect calculations

sync : llama.cpp

ggml-ci

Vulkan: Fix device info output format specifiers (llama/10366)

* Vulkan: Fix device info output format specifiers

* Vulkan: Use zu printf specifier for size_t instead of ld

metal : add `GGML_UNARY_OP_ELU` kernel (#1018)

sync : llama.cpp

ggml-ci

CUDA: fix MMV kernel being used for FP16 src1 (llama/10357)

sync : llama.cpp

ggml-ci

CMake: fix typo in comment [no ci] (llama/10360)

llama : only use default buffer types for the KV cache (llama/10358)

metal : refactor kernel args into structs (llama/10238)

* metal : add kernel arg structs (wip)

* metal : fattn args

ggml-ci

* metal : cont + avoid potential int overflow [no ci]

* metal : mul mat struct (wip)

* cont : mul mat vec

* cont : pass by reference

* cont : args is first argument

* cont : use char ptr

* cont : shmem style

* cont : thread counters style

* cont : mul mm id

ggml-ci

* cont : int safety + register optimizations

ggml-ci

* metal : GGML_OP_CONCAT

ggml-ci

* metal : GGML_OP_ADD, GGML_OP_SUB, GGML_OP_MUL, GGML_OP_DIV

* metal : GGML_OP_REPEAT

* metal : GGML_OP_CPY

* metal : GGML_OP_RMS_NORM

* metal : GGML_OP_NORM

* metal : add TODOs for rest of ops

* ggml : add ggml-metal-impl.h

ggml-ci

ggml : fix undefined reference to 'getcpu' (llama/10354)

https://github.com/ggerganov/llama.cpp/issues/10352

CUDA: remove DMMV, consolidate F16 mult mat vec (llama/10318)

CMake: default to -arch=native for CUDA build (llama/10320)

ggml : fix possible buffer use after free in sched reserve (llama/9930)

ggml : inttypes.h -> cinttypes (llama/0)

ggml-ci

ggml : adapt AMX to tensor->grad removal (llama/0)

ggml-ci