From: Johannes Gäßler <redacted>
Date: Sat, 13 Apr 2024 22:21:55 +0000 (+0200)
Subject: CUDA: fix matrix multiplication logic for tests (llama/6667)
X-Git-Tag: upstream/1.7.4~818
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=66aaf03a7a06e595f43c95ef6e8cbfe846bfde09;p=pkg%2Fggml%2Fsources%2Fwhisper.cpp

CUDA: fix matrix multiplication logic for tests (llama/6667)
---

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 4a2bbdab..a3bbb920 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -1946,7 +1946,7 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
     } else if (!split && !fp16_performance_good && src0->type == GGML_TYPE_F16 && !ggml_is_contiguous(src0) && !ggml_is_transposed(src1) && src1->ne[1] == 1) {
         // KQV single-batch
         ggml_cuda_mul_mat_vec_nc(ctx, src0, src1, dst);
-    } else if (!split && fp16_performance_good && src0->type == GGML_TYPE_F16 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1) && src1->ne[2]*src1->ne[3] > 1) {
+    } else if (!split && src0->type == GGML_TYPE_F16 && (src1->type == GGML_TYPE_F16 || fp16_performance_good) && !ggml_is_transposed(src0) && !ggml_is_transposed(src1) && src1->ne[2]*src1->ne[3] > 1) {
         // KQ + KQV multi-batch
         ggml_cuda_mul_mat_batched_cublas(ctx, src0, src1, dst);
     } else if (use_dequantize_mul_mat_vec) {