Alberto Cabrera Pérez commited on
Commit
b969571
·
1 Parent(s): 011fbfd

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

Browse files

* SYCL : Reenabled mmvq path for the SYCL Nvidia Backend

* Reduced verbosity of comment

Files changed (1) hide show
  1. ggml/src/ggml-sycl.cpp +4 -0
ggml/src/ggml-sycl.cpp CHANGED
@@ -3658,6 +3658,10 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
3658
  use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
3659
  #endif // SYCL_USE_XMX
3660
 
 
 
 
 
3661
  if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
3662
  // KQ single-batch
3663
  ggml_sycl_mul_mat_vec_p021(ctx, src0, src1, dst);
 
3658
  use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
3659
  #endif // SYCL_USE_XMX
3660
 
3661
+ // mmvq path is faster in the CUDA backend.
3662
+ if (ctx.stream()->get_backend() == sycl::backend::ext_oneapi_cuda)
3663
+ use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && !use_mul_mat_vec_q;
3664
+
3665
  if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
3666
  // KQ single-batch
3667
  ggml_sycl_mul_mat_vec_p021(ctx, src0, src1, dst);