public_git
/
whisper.cpp
mirror of https://github.com/ggml-org/whisper.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

sycl : support dev2dev memcpy by DEV2DEV_MEMCPY_FORWARD (llama/26234) 

Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
This commit is contained in:
Neo Zhang 2026-07-31 14:20:28 +08:00 committed by Georgi Gerganov
parent 2dba496574
commit fcb5b86659
2 changed files with 8 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
ggml/src/ggml-sycl/common.hpp
View File

@ -133,6 +133,7 @@ enum ggml_sycl_backend_gpu_mode {
enum ggml_sycl_dev2dev_memcpy_mode {
DEV2DEV_MEMCPY_SYCL = 0,
DEV2DEV_MEMCPY_L0 = 1,
DEV2DEV_MEMCPY_FORWARD = 2
};
static_assert(sizeof(sycl::half) == sizeof(ggml_fp16_t), "wrong fp16 size");

8
ggml/src/ggml-sycl/ggml-sycl.cpp
View File

@ -274,6 +274,8 @@ static const char* dev2dev_int2str(int dev2dev) {
return "SYCL API";
} else if (dev2dev == DEV2DEV_MEMCPY_L0) {
return "Level Zero API";
} else if (dev2dev == DEV2DEV_MEMCPY_FORWARD) {
return "Host Forward";
} else {
return "Unknown";
}
@ -684,7 +686,11 @@ static void dev2dev_memcpy(int device_dst, sycl::queue &q_dst, int device_src, s
}
// Host-staged copy
GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward\n");
if(g_ggml_sycl_dev2dev_memcpy == DEV2DEV_MEMCPY_FORWARD) {
GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward for setting GGML_SYCL_DEV2DEV_MEMCPY=2\n");
} else {
GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward for SYCL/L0 fallback\n");
}
char *host_buf = (char *)malloc(size);
q_src.memcpy(host_buf, (const char *)ptr_src, size).wait();
q_dst.memcpy((char *)ptr_dst, host_buf, size).wait();