From 5b0b8d8cfb5ddf2118f686ba6c30fab3f71b384b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Alberto=20Cabrera=20P=C3=A9rez?= Date: Tue, 9 Jul 2024 15:03:15 +0100 Subject: [PATCH] sycl : Reenabled mmvq path for the SYCL Nvidia Backend (#8372) * SYCL : Reenabled mmvq path for the SYCL Nvidia Backend * Reduced verbosity of comment --- ggml/src/ggml-sycl.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp index 21006cd7b..9c419ba89 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);