Commit f7b41cca4 for llama.cpp
commit f7b41cca4fecc8164c7ea45bfb55dfff7633842f
Author: Hoang Vien Duy <hoangduy.cqb.2k@gmail.com>
Date: Thu Nov 5 15:58:41 2026 +0700
metal : use mul_mm for 32 <= K < 64 (#30982)
* metal : use mul_mm for 32 <= K < 64
* metal : add bf16 x bf16 mul mat kernel
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
diff --git a/ggml/src/ggml-metal/ggml-metal-common.cpp b/ggml/src/ggml-metal/ggml-metal-common.cpp
index 70fd551e3..f36312254 100644
--- a/ggml/src/ggml-metal/ggml-metal-common.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-common.cpp
@@ -38,7 +38,7 @@ bool ggml_metal_op_mul_mat_use_mm(const struct ggml_tensor * op, bool has_simdgr
return !ggml_is_transposed(op->src[0]) &&
!ggml_is_transposed(op->src[1]) &&
- has_simdgroup_mm && ne00 >= 64 && ne11 > 8;
+ has_simdgroup_mm && ne11 > 8 && ne00 >= 32;
}
bool ggml_metal_op_mul_mat_id_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm) {
diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal
index edcc061d5..fab720097 100644
--- a/ggml/src/ggml-metal/kernels/mul_mm.metal
+++ b/ggml/src/ggml-metal/kernels/mul_mm.metal
@@ -895,7 +895,8 @@ typedef decltype(kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4,
template [[host_name("kernel_mul_mm_f32_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_f16_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, half4x4, 1, dequantize_f16, half, half4x4, float, float2x4>;
#if defined(GGML_METAL_HAS_BF16)
-template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_mm<bfloat, bfloat4x4, simdgroup_bfloat8x8, bfloat, bfloat2x4, simdgroup_bfloat8x8, bfloat4x4, 1, dequantize_bf16, bfloat, bfloat4x4, float, float2x4>;
+template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_mm<bfloat, bfloat4x4, simdgroup_bfloat8x8, bfloat, bfloat2x4, simdgroup_bfloat8x8, bfloat4x4, 1, dequantize_bf16, bfloat, bfloat4x4, float, float2x4>;
+template [[host_name("kernel_mul_mm_bf16_bf16")]] kernel mul_mm_t kernel_mul_mm<bfloat, bfloat4x4, simdgroup_bfloat8x8, bfloat, bfloat2x4, simdgroup_bfloat8x8, bfloat4x4, 1, dequantize_bf16, bfloat, bfloat4x4, bfloat, bfloat2x4>;
#endif
template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q1_0, 8, dequantize_q1_0, float, float4x4, float, float2x4>;
template [[host_name("kernel_mul_mm_q2_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_0, 4, dequantize_q2_0, float, float4x4, float, float2x4>;
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 987d5488b..89050ed92 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -11429,6 +11429,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_mul_mat_id_w4a8(GGML_TYPE_MXFP4, GGML_TYPE_F32, 8, 2, false, 32, 32, 256));
test_cases.emplace_back(new test_mul_mat_id_w4a4(GGML_TYPE_MXFP4, GGML_TYPE_F32, 8, 2, false, 32, 32, 256));
+ // small K at batch > 8, e.g. LoRA B matrices with K = rank
+ for (ggml_type type_a : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
+ for (int64_t k : {16, 32, 48}) {
+ for (int64_t n : {9, 33, 128}) {
+ test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 100, n, k, {1, 1}, {1, 1}));
+ }
+ }
+ }
#if 0
// > 2^31-element dst (8.9 GB), about 45 GB peak memory. Too large to be enabled by default.
// kernel_mul_mm dst offset of the last batch is past INT32_MAX: rows 0-31 use full tiles, rows 32-39 use edge tiles.
@@ -12921,6 +12929,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
// delta-net conv input: conv state followed by the transposed qkv of the ubatch
test_cases.emplace_back(new test_concat(GGML_TYPE_F32, {3, 8192, 1, 1}, 512, 0, 16));
+ // LoRA at prefill: B is [rank, n_out], so K = rank
+ for (int64_t k : {16, 32, 48, 64}) {
+ test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 2048, 512, k, {1, 1}, {1, 1}));
+ }
+
// DSV4_HC_PRE: n_tokens 1 is decode, 128 and 512 are prompt batches
for (int64_t n_tokens : {1, 128, 512}) {
test_cases.emplace_back(new test_dsv4_hc_pre(4096, 4, n_tokens)); // DeepSeek-V4-Flash