llama.cpp/ggml/src
deepsek 66906cd82a
HIP: Enable Matrix cores for MMQ Kernels, Enable stream-K for CDNA 3 (#14624)
This commit adds support for MFMA instructions to MMQ. CDNA1/GFX908 CDNA2/GFX90a and CDNA3/GFX942 are supported by the MFMA-enabled code path added by this commit. The code path and stream-k is only enabled on CDNA3 for now as it fails to outperform blas in all cases on the other devices.
Blas is currently only consistently outperformed on CDNA3 due to issues in the amd-provided blas libraries.
This commit also improves the awareness of MMQ towards different warp sizes and as a side effect improves the performance of all quant formats besides q4_0 and q4_1, which regress slightly, on GCN gpus.
2025-07-27 00:28:14 +02:00
..
ggml-blas cmake : Fix broken CMake error messages (ggml/1252) 2025-06-01 13:43:57 +03:00
ggml-cann CANN: Implement GLU ops (#14884) 2025-07-26 17:56:18 +08:00
ggml-cpu ggml-cpu : disable GGML_NNPA by default due to instability (#14880) 2025-07-25 19:09:03 +02:00
ggml-cuda HIP: Enable Matrix cores for MMQ Kernels, Enable stream-K for CDNA 3 (#14624) 2025-07-27 00:28:14 +02:00
ggml-hip HIP: disable rocwmma on gfx12 by default until rocm 7.0 (#14202) 2025-06-16 13:47:38 +02:00
ggml-metal metal: SSM_SCAN performance (#14743) 2025-07-25 10:47:39 -06:00
ggml-musa musa: upgrade musa sdk to rc4.2.0 (#14498) 2025-07-24 20:05:37 +01:00
ggml-opencl opencl: add fused `rms_norm_mul` (#14841) 2025-07-25 17:12:13 +02:00
ggml-rpc rpc : check for null buffers in get/set/copy tensor endpoints (#14868) 2025-07-25 12:17:02 +02:00
ggml-sycl sycl: fixed semantics of block offset calculation (#14814) 2025-07-24 11:09:57 +01:00
ggml-vulkan vulkan: fix rms_norm_mul to handle broadcasting dim0 (#14817) 2025-07-22 17:35:21 +02:00
ggml-webgpu ggml: Add initial WebGPU backend (#14521) 2025-07-16 18:18:51 +03:00
CMakeLists.txt ggml: Add initial WebGPU backend (#14521) 2025-07-16 18:18:51 +03:00
ggml-alloc.c metal : fuse add, mul + add tests (#14596) 2025-07-18 20:37:26 +03:00
ggml-backend-impl.h ggml : upgrade init_tensor API to return a ggml_status (#11854) 2025-02-28 14:41:47 +01:00
ggml-backend-reg.cpp ggml: Add initial WebGPU backend (#14521) 2025-07-16 18:18:51 +03:00
ggml-backend.cpp sched : fix multiple evaluations of the same graph with pipeline parallelism (#14855) 2025-07-25 11:07:26 +03:00
ggml-common.h ggml-cpu : split arch-specific implementations (#13892) 2025-06-09 16:47:13 +02:00
ggml-impl.h metal : fuse add, mul + add tests (#14596) 2025-07-18 20:37:26 +03:00
ggml-opt.cpp mnist: fix segmentation fault (ggml/1227) 2025-05-19 13:29:56 +03:00
ggml-quants.c ggml-quants : rename best_mad to best_error (ggml/1283) 2025-07-01 11:06:39 +03:00
ggml-quants.h ggml : build backends as libraries (#10256) 2024-11-14 18:04:35 +01:00
ggml-threading.cpp ggml : build backends as libraries (#10256) 2024-11-14 18:04:35 +01:00
ggml-threading.h remove CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS (#10797) 2024-12-12 19:02:49 +01:00
ggml.c ggml : remove invalid portPos specifiers from dot files (#14838) 2025-07-25 14:29:57 +03:00
ggml.cpp ggml : Print backtrace on uncaught C++ exceptions (ggml/1232) 2025-06-01 13:43:57 +03:00
gguf.cpp ggml : prevent integer overflow in gguf tensor size calculation (#14595) 2025-07-09 14:33:53 +02:00