386 Commits

Author	SHA1	Message	Date
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