87616f0680
ggml : fix trailing whitespaces ( #0 )
2025-04-24 17:32:47 +03:00
c6e8cc28c1
ggml : Depthwise 2D convolution (ggml/1152)
...
* ggml-cpu : kernels for faster depthwise 2D convolution
* fix compile: remove static after moving to ops.cpp
* add dilation for depthwise_conv_2d
* review: rename to ggml_conv_2d_dw_direct, remove redundant struct keywords, pass by ref, whitespace
* review: rename depthwise_conv_2d -> conv_2d_dw everywhere
2025-04-24 17:32:47 +03:00
b10d8bfdb1
CUDA: use switch statements in constexpr functions ( #13095 )
2025-04-24 15:57:10 +02:00
c21c8cd00a
Merge branch 'upstream' into concedo_experimental
2025-04-24 18:00:29 +08:00
7604a7d6b8
metal : fix floating-point range of attention scores in FA kernels ( #13090 )
...
ggml-ci
2025-04-24 10:38:30 +03:00
28a2723100
merged pixtral support, not fully working
2025-04-24 15:27:02 +08:00
b3b6d862cf
vulkan: matmul gcn tuning ( #13016 )
...
* tune matmul for gcn
* this one is more power efficient
* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp
Co-authored-by: 0cc4m <picard12@live.de >
* disable this tune for the proprietary driver
---------
Co-authored-by: 0cc4m <picard12@live.de >
2025-04-24 09:18:33 +02:00
8f1edcbdac
Merge commit 'dc39a5e7a84815a90fa0c515ed8927870cf858c9' into concedo_experimental
...
# Conflicts:
# README.md
# SECURITY.md
# docs/multimodal/MobileVLM.md
# examples/llava/CMakeLists.txt
# examples/llava/README.md
# examples/llava/android/adb_run.sh
# ggml/CMakeLists.txt
# ggml/src/CMakeLists.txt
# ggml/src/ggml-cpu/CMakeLists.txt
# ggml/src/ggml-sycl/ggml-sycl.cpp
# ggml/src/ggml-sycl/rope.cpp
# ggml/src/ggml-sycl/rope.hpp
2025-04-24 11:49:08 +08:00
658987cfc9
CUDA: noncont MMVQ + batched bs1 MUL_MAT_ID ( #13014 )
...
* CUDA: noncont MMVQ + batched bs1 MUL_MAT_ID
* fix logic for RoPE support, CUDA graphs
2025-04-22 21:27:40 +02:00
7b53389c24
metal : add memory pool for temp allocs ( #12850 )
...
* metal : add memory pool for temp allocs (wip) [no ci]
* cont : free buffers from the heap
* cont : resize heap [no ci]
* cont : refactor heap [no ci]
* cont : heap for each cmd buffer [no ci]
* cont : fix free
* wip
* cont : fix alignment [no ci]
* cont : not working .. [no ci]
* cont : heap allocation now works [no ci]
* cont : use MTLHeapTypePlacement
ggml-ci
* metal : use dynamic MTLHeap allocations
ggml-ci
* metal : add comments
* metal : disable softmax use of mem_pool
ggml-ci
* metal : final touches
2025-04-22 16:15:51 +03:00
1d735c0b4f
ggml : add SSE 4.2 and x64 base variant for CPUs without AVX ( #12871 )
...
* ggml : add SSE 4.2 variant for CPUs without AVX
* ggml : add x64 base ABI variant
2025-04-21 18:13:51 +02:00
5368ddda7a
SYCL: Add non-contiguous support in ROPE ( #12993 )
...
ggml-ci
2025-04-21 19:13:30 +05:30
17360a3b32
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/workflows/build.yml
# examples/llava/clip.cpp
2025-04-20 17:59:58 +08:00
66168204be
vulkan: support noncontiguous rms_norm ( #13031 )
2025-04-20 10:50:02 +02:00
4ba9d711ba
metal: add neg operator ( #13029 )
2025-04-20 08:28:40 +03:00
95d1aaf4d4
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# examples/rpc/rpc-server.cpp
# ggml/src/ggml-rpc/ggml-rpc.cpp
# ggml/src/ggml-sycl/backend.hpp
# ggml/src/ggml-sycl/common.hpp
# ggml/src/ggml-sycl/element_wise.cpp
# ggml/src/ggml-sycl/element_wise.hpp
# ggml/src/ggml-sycl/ggml-sycl.cpp
# requirements/requirements-all.txt
2025-04-19 13:17:13 +08:00
4b0f63ed62
cleanup
2025-04-18 22:57:10 +08:00
8d66005763
SYCL: Refactor and enable FP16 in binary broadcast OPs ( #12975 )
...
* SYCL: refactor move to a separate file
* Fix binbcast
* Remove duplicates
* fix include formatting
* fix typo
2025-04-18 15:57:56 +02:00
2db9ba1464
rpc : add RPC_CMD_HELLO ( #12955 )
...
Add RPC_CMD_HELLO for getting the version of the protocol implemend by
the server. Follow the semantic versioning rules at https://semver.org
Hopefully this bring better user experience when we make breaking
changes at the protocol level and avoid issues like #12465
2025-04-18 10:13:42 +03:00
bce519cee7
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# ggml/src/ggml-cann/aclnn_ops.cpp
# ggml/src/ggml-cann/aclnn_ops.h
# ggml/src/ggml-cann/common.h
# ggml/src/ggml-cann/ggml-cann.cpp
# tests/test-backend-ops.cpp
2025-04-18 12:44:20 +08:00
2f74c354c0
graph : make FA compatible with MLA + add initial Metal kernels ( #12953 )
...
* graph : make mla compatible with FA
* metal : add exp FA kernels for DeepSeek models
ggml-ci
* llama : minor naming updates
ggml-ci
* ggml : disable FA for DS head sizes
* tests : add FA tests for MLA shapes
ggml-ci
2025-04-17 18:16:36 +03:00
207c22ec2d
ggml: Re-enable CUDA graphs in presence of CONT and DUP nodes ( #12970 )
2025-04-17 15:19:42 +02:00
7a395f67a7
CANN: Add support for async operator submission ( #12864 )
...
Submit operators using asynchronous threads to improve performance.
Use the environment variable GGML_CANN_ASYNC_MODE to control whether
asynchronous submission is enabled. It is disabled by default.
Testing shows a 10%–20% performance improvement in scenarios with
small parameter sizes, especially in quantized models.
2025-04-17 20:34:16 +08:00
12b17501e6
opencl: fix incorrect local_size index in profiling log ( #12868 )
2025-04-16 14:25:57 -07:00
015022bb53
vulkan: enable coopmat2 FA gqa and split_k optimizations more often ( #12931 )
...
The grouped query attention optmization doesn't require a power of two ratio,
the only thing relying on it was the modulo operation written as bitwise &.
split_k need not depend on gqa_ratio - enable it any time there's only one
workgroup in the X dimension. The shader gets the split index from the x coord,
and multiple workgroups in the X dimension (pre-split) indicates a larger
FA operation that wouldn't need splitting.
2025-04-16 20:37:25 +02:00
06159939d9
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/workflows/build.yml
# Makefile
# docs/build.md
# examples/rpc/rpc-server.cpp
# examples/sycl/build.sh
# ggml/CMakeLists.txt
# ggml/src/ggml-cann/aclnn_ops.cpp
# ggml/src/ggml-cann/ggml-cann.cpp
# ggml/src/ggml-hip/CMakeLists.txt
# scripts/sync-ggml.last
2025-04-17 00:52:37 +08:00
a5c39143a1
enabled coopmat2 for the price of 7mb
2025-04-17 00:21:03 +08:00
b43d89e311
CANN: Add 310P operator support check ( #12962 )
2025-04-16 16:21:05 +08:00
80f19b4186
opencl: split ggml-opencl.cl into multiple files and cleanup ( #12886 )
...
* opencl: refactor - split the kernel files
---------
Co-authored-by: Shangqing Gu <quic_shawngu@quicinc.com >
* opencl: split more kernels into separate files
* opencl: specify subgroup size instead of querying it
* opencl: refine Adreno cl compiler version parsing
* opencl: skip some kernels not used by Adreno on old compilers
* opencl: refine logic for selecting Adreno kernels
* opencl: refine Adreno cl compiler version
* opencl: cleanup preprocessor for kernels
* opencl: consider Adreno CL compiler on Windows
* opencl: add final newline for `mul_mv_f16_f16.cl`
---------
Co-authored-by: Shangqing Gu <quic_shawngu@quicinc.com >
2025-04-15 12:26:00 -07:00
f8f820cc4d
metal : add FA-vec kernels for head size 96 ( #12952 )
...
ggml-ci
2025-04-15 14:45:05 +03:00
54a7272043
CANN: Add x86 build ci ( #12950 )
...
* CANN: Add x86 build ci
* CANN: fix code format
2025-04-15 12:08:55 +01:00
84778e9770
CUDA/HIP: Share the same unified memory allocation logic. ( #12934 )
...
Replace compile-time `GGML_HIP_UMA` with environment variable `GGML_CUDA_ENABLE_UNIFIED_MEMORY`. This unifies the usage on NVIDIA and AMD GPUs, and allows a single binary to be shared between integrated and dedicated GPUs.
2025-04-15 11:20:38 +02:00
510676475f
SYCL: Add ROPE vision kernel ( #12887 )
...
* SYCL: Add ROPE vision kernel
* Add comment about rope mode
2025-04-15 10:37:42 +02:00
eccc7a1602
ggml : Add AVX512 implementation of GEMM - Q4_Kx8 ( #12829 )
...
* Add AVX512 implementation of GEMM - q4kx8
* Update changes to remove unnecessary whitespaces
2025-04-15 09:22:36 +03:00
0019279bb5
CANN: Opt ROPE optimization ( #12865 )
...
* [CANN]Opt ROPE optimization
* [CANN]Codestyle adjustment
* [CANN]Fix the ROPE precision issue
* [CANN]codestyle fix
* [CANN]add rope unsupport case
Signed-off-by: noemotiovon <noemotiovon@gmail.com >
2025-04-15 10:09:35 +08:00
b0c75ac9f9
CANN: Optimize CANN buffer pool memory management ( #12875 )
...
Multiple optional memory pools are provided for CANN, including VMM,
priority queue-based, and traditional memory pools.
1.When the memory pool is available and GGML_CANN_DISABLE_VMM_POOL
is not defined, the VMM pool is selected by default.
2.Otherwise, if GGML_CANN_ENABLE_BUF_PRIO_POOL is defined,
the priority queue-based memory pool is used.
3.If neither condition is met, the default memory pool is used.
2025-04-15 10:04:24 +08:00
75afa0ae31
SYCL: Fix im2col ( #12910 )
...
* SYCL: Fix im2col
* restore local workgroup size adjustments for large inputs
* restore format
2025-04-14 14:23:53 +02:00
c772d54926
rpc : use ggml_context_ptr ( #12938 )
2025-04-14 13:59:34 +03:00
a25355e264
cpu: fix cpu backend's supports-op for GET_ROWS_BACK. fixes a fatal when running test-backend-ops with only the CPU backend (ggml/1190)
2025-04-14 09:26:15 +03:00
e959d32b1c
ggml: use _mm[512/256]_dpbusd[_avx]_epi32 to directly accumulate into the result register ( #12773 )
...
* ggml: use _mm[512/256]_dpbusd[_avx]_epi32 to directly accumulate into the result register
* simplifies the codebase by removing redundant functions
2025-04-14 08:47:55 +03:00
e1ee857b1e
allow vulkan to be packaged without coopmat for noavx2
2025-04-14 12:40:00 +08:00
307bfa253d
ggml: disable CUDA graphs for unsupported DUP and CONT node types ( #12891 )
...
Fixes #12798
2025-04-13 23:12:21 +02:00
7b4254bef9
not working on cpu
2025-04-12 18:55:29 +08:00
a4837577aa
vulkan: use aligned loads for flash attention mask ( #12853 )
...
Rewrite the stride logic for the mask tensor in the FA shader to force the
stride to be aligned, to allow using more efficient loads.
2025-04-12 10:44:48 +02:00
a0ae187563
Merge branch 'upstream' into concedo_experimental
...
# Conflicts:
# .github/workflows/docker.yml
# README.md
# build-xcframework.sh
# examples/llava/CMakeLists.txt
# examples/llava/clip.cpp
# examples/rpc/rpc-server.cpp
# examples/run/run.cpp
# ggml/src/ggml-cann/ggml-cann.cpp
# scripts/sync-ggml-am.sh
# scripts/sync-ggml.last
# tests/test-backend-ops.cpp
# tests/test-chat.cpp
2025-04-12 10:06:47 +08:00
ea9bd61e47
Merge commit '64eda5deb9859e87a020e56bab5d2f9ca956f1de' into concedo_experimental
...
# Conflicts:
# .devops/cuda.Dockerfile
# .devops/intel.Dockerfile
# .devops/llama-cli-cann.Dockerfile
# .devops/musa.Dockerfile
# .devops/rocm.Dockerfile
# .devops/vulkan.Dockerfile
# .github/workflows/build.yml
# .github/workflows/docker.yml
# README.md
# docs/backend/SYCL.md
# examples/llava/clip.cpp
# examples/server_embd.py
# ggml/src/ggml-cann/acl_tensor.cpp
# ggml/src/ggml-cann/aclnn_ops.cpp
# ggml/src/ggml-cann/aclnn_ops.h
# ggml/src/ggml-cann/ggml-cann.cpp
# src/CMakeLists.txt
# tests/test-chat-template.cpp
2025-04-12 08:31:22 +08:00
578754b315
sycl: Support sycl_ext_oneapi_limited_graph ( #12873 )
...
The current usage of the SYCL-Graph extension checks for
the `sycl_ext_oneapi_graph` device aspect. However, it is also
possible to support `sycl_ext_oneapi_limied_graph` devices that
don't support update
2025-04-11 15:32:14 +02:00
fccf9cae83
SYCL: Add fp16 type support to unary op kernels ( #12788 )
...
* SYCL: Add fp16 support to some elementwise OP kernels
* remove comment
ggml-ci
* Use static_cast directly
* remove not needed cast from tanh
* Use static cast and remove unneeded castings
* Adjust device_support_op for unary OPs
* Use cast_data and typed_data struct to deduplicate casting code
2025-04-11 16:03:50 +08:00
0fed24c347
ggml: fix compilation error s390x ( #12848 )
...
* ggml: fixes #12846 compilation error
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com >
Co-authored-by: Aleksei Nikiforov <aleksei.nikiforov@ibm.com >
* ggml: add documentation for code change
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com >
Co-authored-by: Aleksei Nikiforov <aleksei.nikiforov@ibm.com >
* ggml: refactor to type-cast and update documentation
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com >
Co-authored-by: Aleksei Nikiforov <aleksei.nikiforov@ibm.com >
* ggml: update documentation to provide full issue link
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com >
Co-authored-by: Aleksei Nikiforov <aleksei.nikiforov@ibm.com >
---------
Co-authored-by: Aleksei Nikiforov <aleksei.nikiforov@ibm.com >
2025-04-11 08:20:07 +03:00
cb79c2e7fa
ggml: don't include arm_neon.h when using CUDA 12 with ARM Neon (ggml/1187)
...
fix #1186
2025-04-11 00:17:47 +03:00
Georgi Gerganov
Acly
Johannes Gäßler
Concedo
Eve
Diego Devesa
Akarshan Biswas
Jeff Bolz
Jeffrey Morgan
Radoslav Gerganov
Alan Gray
hipudding
kimminsu
Chenguang Li
lhez
David Huang
Srihari-mcw
Xinpeng Dou
cmdr2
SXX
Ewan Crawford
Aaron Teo