]> git.djapps.eu Git - pkg/ggml/sources/ggml/commit
CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)
authorAman Gupta <redacted>
Sat, 27 Sep 2025 16:49:32 +0000 (00:49 +0800)
committerGeorgi Gerganov <redacted>
Mon, 29 Sep 2025 09:41:09 +0000 (12:41 +0300)
commitf5301a67865a40a464698ce6c5f122c3a72ca3b7
treed674877e5a64da7bebcb188895c2caed3b9ab903
parent4977596ca0efb781339a111feb1d334efada08c2
CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)

* CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32

This commit adds mul_mat_id support for ncols_dst >= 16. It does this by
packing ncols_dst tiles into the blockDim.y.

My tests on a RTX 3090 show that this is faster than the cuBLAS fallback
for f16 till bs=64, and for f32 till bs=32

* Review: refactor if statement
src/ggml-cuda/ggml-cuda.cu
src/ggml-cuda/mmf.cu
src/ggml-cuda/mmf.cuh
tests/test-backend-ops.cpp