From 1c5b89ff63151b5388cb5abb9357cebe446997c7 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 (#26234) Co-authored-by: Neo Zhang Jianyu --- docs/backend/SYCL.md | 2 +- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/ggml-sycl.cpp | 8 +++++++- 3 files changed, 9 insertions(+), 2 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index c72dc3b3ee..914d63af22 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -788,7 +788,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | Name | Value | Function | |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------| | GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG | -| GGML_SYCL_DEV2DEV_MEMCPY | 0 (default) or 1 | Choose the SYCL or L0 API in dev2dev memory copy.
Value:
* 0: SYCL API (default)
* 1: L0 API -- L0 API is found to lead to abnormal crash in some case. This debug flag is used to check the issue.| +| GGML_SYCL_DEV2DEV_MEMCPY | 0 (default), 1, 2 | Choose the method of dev2dev memory copy.
Value:
* 0: SYCL API (default), only support dGPUs.
* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case.
* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).
SYCL & L0 API are easy to be impacted by Intel GPU driver issue. When you meet the garbled output or crash issues in multiple GPUs case, try with this debug flag to work around or check the issue.| | GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.| | GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) | | GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. | diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index f27ec5dd62..160331191d 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 49bdc97d23..07ddfda887 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();