vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

SYCL : Reenabled mmvq path for the SYCL Nvidia Backend

Browse source
This commit is contained in:
Alberto Cabrera 2024-07-04 16:55:57 +01:00
parent 3f2d538b81
commit f288ae1a14
1 changed files with 4 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

4
ggml/src/ggml-sycl.cpp
View file

@ -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); use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
#endif // SYCL_USE_XMX #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) { if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
// KQ single-batch // KQ single-batch
ggml_sycl_mul_mat_vec_p021(ctx, src0, src1, dst); ggml_sycl_mul_mat_vec_p021(ctx, src0, src1, dst);