From: Alberto Cabrera Pérez <redacted>
Date: Tue, 9 Jul 2024 14:03:15 +0000 (+0100)
Subject: sycl : Reenabled mmvq path for the SYCL Nvidia Backend (llama/8372)
X-Git-Tag: upstream/1.7.4~564
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=2af4a52c392bbdd5a73140d56e0edb9c24b75539;p=pkg%2Fggml%2Fsources%2Fwhisper.cpp

sycl : Reenabled mmvq path for the SYCL Nvidia Backend (llama/8372)

* SYCL : Reenabled mmvq path for the SYCL Nvidia Backend

* Reduced verbosity of comment
---

diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp
index 21006cd7..9c419ba8 100644
--- a/ggml/src/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl.cpp
@@ -3658,6 +3658,10 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
     use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
 #endif // SYCL_USE_XMX
 
+    // mmvq path is faster in the CUDA backend.
+    if (ctx.stream()->get_backend() == sycl::backend::ext_oneapi_cuda)
+        use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && !use_mul_mat_vec_q;
+
     if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
         // KQ single-batch
         ggml_sycl_mul_mat_vec_p021(ctx, src0, src1, dst);