439 Commits

Author	SHA1	Message	Date
Diego Devesa	2803540814	ggml-cpu : fix HWCAP2_I8MM value (#10646)	2024-12-04 14:40:44 +01:00
Jeff Bolz	2759916d86	vulkan: Implement "fast divide" (mul+shift) for unary ops like copy (#10642)	2024-12-04 08:28:59 +01:00
Nicolò Scipione	40c6d79fb5	SYCL : Move to compile time oneMKL interface backend selection for NVIDIA backend (#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 <nicolo.scipione@codeplay.com> * Formatting * Address PR comments to increase readibility --------- Signed-off-by: nscipione <nicolo.scipione@codeplay.com>	2024-12-04 09:29:20 +08:00
Frankie Robertson	cd2f37b304	Avoid using __fp16 on ARM with old nvcc (#10616)	2024-12-04 01:41:37 +01:00
Jeff Bolz	cc98896db8	vulkan: optimize and reenable split_k (#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.	2024-12-03 20:29:54 +01:00
mahorozte	e9e661bd59	CUDA: remove unnecessary warp reduce in FA (ggml/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 <zhao.xiaoyu@zte.com.cn>	2024-12-03 20:04:49 +02:00
PAB	efb6ae9630	feat: add GGML_UNARY_OP_ARGMAX Metal kernel (ggml/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 <slarengh@gmail.com> --------- Co-authored-by: Diego Devesa <slarengh@gmail.com>	2024-12-03 20:04:49 +02:00
PAB	667d70d170	metal : add GGML_OP_CONV_TRANSPOSE_1D kernels (ggml/1026) ... * wip * wip implementation f32 * kernel conv transpose 1d f32 working * initial commit	2024-12-03 20:04:49 +02:00
Georgi Gerganov	0115df2f65	metal : small-batch mat-mul kernels (#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	2024-12-03 11:52:33 +02:00
Akarshan Biswas	991f8aabee	SYCL: Fix and switch to GGML_LOG system instead of fprintf (#10579) ... * Switched to GGML_LOG * Fix missing semicolon	2024-12-02 15:04:11 +08:00
Diego Devesa	3420909dff	ggml : automatic selection of best CPU backend (#10606) ... * ggml : automatic selection of best CPU backend * amx : minor opt * add GGML_AVX_VNNI to enable avx-vnni, fix checks	2024-12-01 16:12:41 +01:00
Adrien Gallouët	0c39f44d70	ggml-cpu: replace AArch64 NEON assembly with intrinsics in ggml_gemv_q4_0_4x4_q8_0() (#10567) ... Signed-off-by: Adrien Gallouët <angt@huggingface.co>	2024-11-30 09:13:18 -08:00
Eve	0533e7fb38	vulkan: Dynamic subgroup size support for Q6_K mat_vec (#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	2024-11-30 08:00:02 +01:00
Diego Devesa	7cc2d2c889	ggml : move AMX to the CPU backend (#10570) ... * ggml : move AMX to the CPU backend --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2024-11-29 21:54:58 +01:00
Georgi Gerganov	f0678c5ff4	ggml : fix I8MM Q4_1 scaling factor conversion (#10562) ... ggml-ci	2024-11-29 16:25:39 +02:00
Shupei Fan	4b3242bbea	ggml-cpu: fix typo in gemv/gemm iq4_nl_4_4 (#10580)	2024-11-29 14:49:02 +01:00
Alberto Cabrera Pérez	0f77aae560	sycl : offload of get_rows set to 0 (#10432)	2024-11-29 20:38:45 +08:00
Alberto Cabrera Pérez	266b8519ee	sycl : Reroute permuted mul_mats through oneMKL (#10408) ... This PR fixes the failing MUL_MAT tests for the sycl backend.	2024-11-29 09:49:43 +00:00
Chenguang Li	938f608742	CANN: RoPE operator optimization (#10563) ... * [cann] RoPE operator optimization * [CANN]Code Formatting --------- Co-authored-by: noemotiovon <noemotiovon@gmail.com>	2024-11-29 14:46:55 +08:00
Jeff Bolz	f095a649ec	vulkan: get the first command buffer submitted sooner (#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.	2024-11-29 07:18:02 +01:00
Georgi Gerganov	dc22344088	ggml : remove redundant copyright notice + update authors	2024-11-28 20:46:40 +02:00
Georgi Gerganov	76b27d29c2	ggml : fix row condition for i8mm kernels (#10561) ... ggml-ci	2024-11-28 14:56:37 +02:00
Georgi Gerganov	eea986f215	cmake : fix ARM feature detection (#10543) ... ggml-ci	2024-11-28 14:56:23 +02:00
Shupei Fan	c202cef168	ggml-cpu: support IQ4_NL_4_4 by runtime repack (#10541) ... * ggml-cpu: support IQ4_NL_4_4 by runtime repack * ggml-cpu: add __ARM_FEATURE_DOTPROD guard	2024-11-28 13:52:03 +01:00
Sergio López	2025fa67e9	kompute : improve backend to pass test_backend_ops (#10542) ... * kompute: op_unary: reject unsupported parameters Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: softmax: implement ALiBi support Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: rope: implement neox and phi3 support Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: op_mul_mat_q4_k permutted support Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: op_mul_mat_[q4_0|q4_1|q8_0] permutted support Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: op_mul_mat_f16 permutted support Signed-off-by: Sergio Lopez <slp@redhat.com> * kompute: op_mul_mat_q6_k permutted support Signed-off-by: Sergio Lopez <slp@redhat.com> --------- Signed-off-by: Sergio Lopez <slp@redhat.com>	2024-11-28 12:51:38 +01:00
leo-pony	605fa66c50	CANN: Fix SOC_TYPE compile bug (#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	2024-11-28 15:25:24 +08:00
Chenguang Li	b7420131bf	CANN: ROPE operator optimization (#10540) ... * [cann] ROPE operator optimization Co-authored-by: noemotiovon <noemotiovon@gmail.com>	2024-11-28 14:24:46 +08:00
uvos	3ad5451f3b	Add some minimal optimizations for CDNA (#10498) ... * Add some minimal optimizations for CDNA * ggml_cuda: set launch bounds also for GCN as it helps there too	2024-11-27 17:10:08 +01:00
Georgi Gerganov	9e2301f4a4	metal : fix group_norm support condition (#0)	2024-11-27 11:22:14 +02:00
Frankie Robertson	9150f8fef9	Do not include arm_neon.h when compiling CUDA code (ggml/1028)	2024-11-27 11:10:27 +02:00
Jeff Bolz	c31ed2abfc	vulkan: define all quant data structures in types.comp (#10440)	2024-11-27 08:32:54 +01:00
Jeff Bolz	5b3466bedf	vulkan: Handle GPUs with less shared memory (#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.	2024-11-27 08:30:27 +01:00
Jeff Bolz	249a7902ec	vulkan: further optimize q5_k mul_mat_vec (#10479)	2024-11-27 08:21:59 +01:00
Jeff Bolz	71a64989a5	vulkan: skip integer div/mod in get_offsets for batch_idx==0 (#10506)	2024-11-27 08:08:54 +01:00
Jeff Bolz	4a57d362e1	vulkan: optimize Q2_K and Q3_K mul_mat_vec (#10459)	2024-11-27 08:00:50 +01:00
R0CKSTAR	249cd93da3	mtgpu: Add MUSA_DOCKER_ARCH in Dockerfiles && update cmake and make (#10516) ... Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>	2024-11-26 17:00:41 +01:00
Jeff Bolz	904109ed0d	vulkan: fix group_norm (#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.	2024-11-26 16:45:05 +01:00
Georgi Gerganov	ab96610b1e	cmake : enable warnings in 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	2024-11-26 14:18:08 +02:00
Charles Xu	25669aa92c	ggml-cpu: cmake add arm64 cpu feature check for macos (#10487) ... * ggml-cpu: cmake add arm64 cpu feature check for macos * use vmmlaq_s32 for compile option i8mm check	2024-11-26 13:37:05 +02:00
Shanshan Shen	9a4b79bcfa	CANN: Improve the Inferencing Performance for Ascend NPU Device (#10454) ... * improve inferencing performance for ascend npu. Co-authored-by: Frank Mai <thxCode@thxcode0824@gmail.com> * some modification after review * some modifications after review * restore some modifications * restore some modifications --------- Co-authored-by: shanshan shen <shanshanshen333@gmail.com> Co-authored-by: Frank Mai <thxCode@thxcode0824@gmail.com>	2024-11-26 18:08:37 +08:00
Chenguang Li	7066b4cce2	CANN: RoPE and CANCAT operator optimization (#10488) ... Co-authored-by: noemotiovon <noemotiovon@gmail.com>	2024-11-26 17:31:05 +08:00
Junil Kim	0eb4e12bee	vulkan: Fix a vulkan-shaders-gen arugment parsing error (#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.	2024-11-26 01:47:20 +00:00
Georgi Gerganov	106964e3d2	metal : enable mat-vec kernels for bs <= 4 (#10491)	2024-11-25 21:49:31 +02:00
Diego Devesa	10bce0450f	llama : accept a list of devices to use to offload a model (#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	2024-11-25 19:30:06 +01:00
Diego Devesa	5931c1f233	ggml : add support for dynamic loading of backends (#10469) ... * ggml : add support for dynamic loading of backends --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2024-11-25 15:13:39 +01:00
Georgi Gerganov	b756441104	metal : minor code formatting	2024-11-25 15:08:04 +02:00
Diego Devesa	55ed008b2d	ggml : do not use ARM features not included in the build (#10457)	2024-11-23 14:41:12 +01:00
leo-pony	c18610b4ee	CANN: Support Ascend310P to accelerate F32 and F16 Model (#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	2024-11-22 14:07:20 +08:00
Diego Devesa	a5e47592b6	cuda : optimize argmax (#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 <johannesg@5d6.de> * fix ub * ggml : check ne00 <= INT32_MAX in argmax and argsort --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de>	2024-11-21 18:18:50 +01:00
slaren	59b9172822	ggml/sched : do not skip views in pre-assignments	2024-11-21 09:22:05 +02:00
Johannes Gäßler	02e4eaf22f	ggml-opt: fix data corruption (ggml/1022)	2024-11-21 09:22:02 +02:00
Jeff Bolz	9abe9eeae9	vulkan: predicate max operation in soft_max shaders/soft_max (#10437) ... Fixes #10434	2024-11-20 20:47:36 +01:00
Jeff Bolz	8fd4b7fa29	vulkan: copy iq4_nl LUT into shared memory (#10409)	2024-11-20 08:40:18 +01:00
Jeff Bolz	1bacb9f625	vulkan: further optimize mul_mat_vec using larger loads (#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.	2024-11-20 08:11:00 +01:00
haopeng	42ae10bbcd	add cmake rvv support (#10411)	2024-11-19 21:10:31 +01:00
Plamen Minev	611fabd792	metal : fox offset integer overflows in im2col (ggml/1015) ... -- While running StableDiffusion.cpp locally with Metal some offsets overflow and results in incorrect calculations	2024-11-19 20:03:21 +02:00
PAB	12b0ad953a	metal : add GGML_UNARY_OP_ELU kernel (ggml/1018)	2024-11-19 20:03:21 +02:00
bandoti	2a11b6b094	Add required ggml-base and backend libs to cmake pkg (#10407)	2024-11-19 17:10:30 +01:00
Diego Devesa	3ee6382d48	cuda : fix CUDA_FLAGS not being applied (#10403)	2024-11-19 14:29:38 +01:00
Romain Biessy	2a1507c162	sycl : Add option to set the SYCL architecture for all targets (#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	2024-11-19 08:02:23 +00:00
Jeff Bolz	b3e585988f	vulkan: Optimize soft_max (#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.	2024-11-19 08:25:17 +01:00
Alberto Cabrera Pérez	557924f222	sycl: Revert MUL_MAT_OP support changes (#10385)	2024-11-19 08:50:04 +08:00
Diego Devesa	d3481e6316	cuda : only use native when supported by cmake (#10389)	2024-11-18 18:43:40 +01:00
Jeff Bolz	f139d2ea61	vulkan: remove use of null initializer (#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?	2024-11-18 08:28:42 -06:00
0cc4m	9b75f03cd2	Vulkan: Fix device info output format specifiers (#10366) ... * Vulkan: Fix device info output format specifiers * Vulkan: Use zu printf specifier for size_t instead of ld	2024-11-18 11:02:43 +01:00
Johannes Gäßler	76e9e58b78	CUDA: fix MMV kernel being used for FP16 src1 (#10357)	2024-11-17 23:20:42 +01:00
Johannes Gäßler	ce2e59ba10	CMake: fix typo in comment [no ci] (#10360)	2024-11-17 12:59:38 +01:00
Diego Devesa	be5caccef9	llama : only use default buffer types for the KV cache (#10358)	2024-11-17 12:25:45 +01:00
Georgi Gerganov	cf32a9b93a	metal : refactor kernel args into structs (#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	2024-11-17 11:23:01 +02:00
FirstTimeEZ	a43178299c	ggml : fix undefined reference to 'getcpu' (#10354) ... https://github.com/ggerganov/llama.cpp/issues/10352	2024-11-17 10:39:22 +02:00
Johannes Gäßler	c3ea58aca4	CUDA: remove DMMV, consolidate F16 mult mat vec (#10318)	2024-11-17 09:09:55 +01:00
Johannes Gäßler	467576b6cc	CMake: default to -arch=native for CUDA build (#10320)	2024-11-17 09:06:34 +01:00
Diego Devesa	eda7e1d4f5	ggml : fix possible buffer use after free in sched reserve (#9930)	2024-11-17 08:31:17 +02:00
Georgi Gerganov	24203e9dd7	ggml : inttypes.h -> cinttypes (#0) ... ggml-ci	2024-11-17 08:30:29 +02:00
Georgi Gerganov	5d9e59979c	ggml : adapt AMX to tensor->grad removal (#0) ... ggml-ci	2024-11-17 08:30:29 +02:00
Georgi Gerganov	68fcb4759c	ggml : fix compile warnings (#0) ... ggml-ci	2024-11-17 08:30:29 +02:00
Johannes Gäßler	8a43e940ab	ggml: new optimization interface (ggml/988)	2024-11-17 08:30:29 +02:00
Georgi Gerganov	db4cfd5dbc	llamafile : fix include path (#0) ... ggml-ci	2024-11-16 20:36:26 +02:00
Jeff Bolz	772703c8ff	vulkan: Optimize some mat-vec mul quant shaders (#10296) ... Compute two result elements per workgroup (for Q{4,5}_{0,1}). This reuses the B loads across the rows and also reuses some addressing calculations. This required manually partially unrolling the loop, since the compiler is less willing to unroll outer loops. Add bounds-checking on the last iteration of the loop. I think this was at least partly broken before. Optimize the Q4_K shader to vectorize most loads and reduce the number of bit twiddling instructions.	2024-11-16 07:26:57 +01:00
Dan Johansson	1e58ee1318	ggml : optimize Q4_0 into Q4_0_X_Y repack (#10324)	2024-11-16 01:53:37 +01:00
Srihari-mcw	74d73dc85c	Make updates to fix issues with clang-cl builds while using AVX512 flags (#10314)	2024-11-15 22:27:00 +01:00
slaren	883d206fbd	ggml : fix some build issues	2024-11-15 21:45:32 +02:00
Georgi Gerganov	09ecbcb596	cmake : fix ppc64 check (whisper/0) ... ggml-ci	2024-11-15 15:44:06 +02:00
thewh1teagle	3225008973	ggml : vulkan logs (whisper/2547)	2024-11-15 15:44:06 +02:00
Eve	18429220bd	AVX BF16 and single scale quant optimizations (#10212) ... * use 128 bit loads (i've tried 256->128 to death and its slower) * double accumulator * avx bf16 vec dot * +3% q4_0 inference * +7% tg +5% pp compared to master * slower f16c version, kep for reference * 256b version, also slow. i tried :) * revert f16 * faster with madd * split to functions * Q8_0 and IQ4_NL, 5-7% faster * fix potential overflow (performance reduced) * 16 bit add for q4_0 only * merge	2024-11-15 12:47:58 +01:00
Romain Biessy	5a54af4d4f	sycl: Use syclcompat::dp4a (#10267) ... * sycl: Use syclcompat::dp4a * Using the syclcompat version allow the compiler to optimize the operation with native function * Update news section * Update CI Windows oneAPI version to 2025.0 * Reword doc * Call syclcompat::dp4a inside dpct::dp4a This reverts commit 90cb61d692.	2024-11-15 11:09:12 +08:00
Charles Xu	1607a5e5b0	backend cpu: add online flow for aarch64 Q4_0 GEMV/GEMM kernels (#9921) ... * backend-cpu: add online flow for aarch64 Q4_0 GEMV/GEMM kernels --------- Co-authored-by: Diego Devesa <slarengh@gmail.com>	2024-11-15 01:28:50 +01:00
Diego Devesa	ae8de6d50a	ggml : build backends as libraries (#10256) ... * ggml : build backends as libraries --------- Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: R0CKSTAR <xiaodong.ye@mthreads.com>	2024-11-14 18:04:35 +01:00
Johannes Gäßler	4a8ccb37ad	CUDA: no -sm row for very small matrices (#10185)	2024-11-14 13:00:15 +01:00
Jeff Bolz	af148c9386	vulkan: Optimize binary ops (#10270) ... Reuse the index calculations across all of src0/src1/dst. Add a shader variant for when src0/src1 are the same dimensions and additional modulus for src1 aren't needed. Div/mod are slow, so add "fast" div/mod that have a fast path when the calculation isn't needed or can be done more cheaply.	2024-11-14 06:22:55 +01:00
Jeff Bolz	66798e42fb	vulkan: Use macros to make the mat mul pipeline creation more concise (#10259) ... Also add vk_matmul_pipeline2 to hold f16/f32 accumulator versions of a pipeline. This isn't really used yet.	2024-11-13 21:59:47 +01:00
Alberto Cabrera Pérez	2e82ffa4af	sycl : Fixes to broken builds and test-backend-ops (#10257) ... * Fixes broken build for the SYCL CUDA backend caused by non-explicit gemm call in outprod (merged in with RWKV6 in Optimize RWKV6 Operator Naming and Implement Multi-core CPU/ SYCL Acceleration #10133) * Marks permuted MUL_MAT as unsupported to be able to run test-backend-ops * Fixes asserts in norm to fix debug builds.	2024-11-13 09:40:57 +00:00
Jeff Bolz	80dd7ff22f	vulkan: Optimize contiguous copies (#10254) ... * tests: Fix memory bandwidth calculation for perf tests Add a flops calculation for flash attention. Add one GGML_OP_CPY perf test. * vulkan: Optimize contiguous copies Add a variant of the copy shader for when the tensors are contiguous. Avoid the complex addressing calculations, and do four elements per invocation to hide some other overhead. Apply similar changes to the scale shader, since scale is always contiguous. Add a "progress bar" for shader compiles.	2024-11-13 07:58:57 +01:00
Jeff Bolz	54ef9cfc72	vulkan: Throttle the number of shader compiles during the build step. (#10222) ... Fixes #9582 Spawning too many concurrent copies of glslc leads to "Failed to create pipes" errors on Linux. This change applies the same throttling we use for multithreaded pipeline creation.	2024-11-11 18:13:51 +01:00
Georgi Gerganov	b0cefea58a	metal : more precise Q*K in FA vec kernel (#10247)	2024-11-11 08:39:13 +02:00
Jeff Bolz	160687b3ed	vulkan: Fix newly added tests for permuted mul_mat and 1D im2col (#10226)	2024-11-10 12:37:56 +01:00
Georgi Gerganov	6423c65aa8	metal : reorder write loop in mul mat kernel + style (#10231) ... * metal : reorder write loop * metal : int -> short, style ggml-ci	2024-11-09 11:53:13 +02:00
Georgi Gerganov	39a334a9aa	metal : fix build and some more comments (#10229)	2024-11-09 11:53:02 +02:00
Georgi Gerganov	bb38cdd8ba	metal : fix F32 accumulation in FA vec kernel (#10232)	2024-11-09 11:52:45 +02:00
Georgi Gerganov	46323fa9ef	metal : hide debug messages from normal log	2024-11-09 11:21:49 +02:00