From f288ae1a14da254a384d99b7c6f93e4cabd81749 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Thu, 4 Jul 2024 16:55:57 +0100 Subject: [PATCH 1/2] SYCL : Reenabled mmvq path for the SYCL Nvidia Backend --- 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 21006cd7bebf7..ff43accacf60d 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 Nvidia backend but slower on the Intel 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); From 0a00d6e4b8cc9b9c97194beed1255f4bbce6deb1 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Tue, 9 Jul 2024 10:47:43 +0100 Subject: [PATCH 2/2] Reduced verbosity of comment --- ggml/src/ggml-sycl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl.cpp b/ggml/src/ggml-sycl.cpp index ff43accacf60d..9c419ba896706 100644 --- a/ggml/src/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl.cpp @@ -3658,7 +3658,7 @@ 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 Nvidia backend but slower on the Intel backend + // 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;