upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-08-03 00:20:50 -05:00
Code Issues Packages Projects Releases Wiki Activity

sycl : support dev2dev memcpy by DEV2DEV_MEMCPY_FORWARD (#26234)

Browse Source 
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 GitHub
parent a2be61dc87
commit 1c5b89ff63
3 changed files with 9 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
docs/backend/SYCL.md
View File

@@ -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.<br>Value: <br>* 0: SYCL API (default)<br>* 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.<br>Value: <br>* 0: SYCL API (default), only support dGPUs.<br>* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case. <br>* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).<br>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. |

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();