From fcb5b86659b6173631127cdddf4899b86968b7d0 Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Fri, 31 Jul 2026 14:20:28 +0800 Subject: [PATCH] sycl : support dev2dev memcpy by DEV2DEV_MEMCPY_FORWARD (llama/26234) Co-authored-by: Neo Zhang Jianyu --- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/ggml-sycl.cpp | 8 +++++++- 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index f27ec5dd6..160331191 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -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"); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 49bdc97d2..07ddfda88 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -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();