ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-02-24 15:14:10 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add check if device to device access is enabled

Browse Source
This commit is contained in:
Iwan Kawrakow
2025-05-22 09:30:29 +03:00
parent e87bfe39ba
commit d6c3d4b4a8
1 changed files with 15 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
ggml/src/ggml-cuda.cu
View File

@@ -3050,7 +3050,14 @@ GGML_CALL static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_
return false;
}
int cur_device;
CUDA_CHECK(cudaGetDevice(&cur_device));
if (backend_src != backend_dst) {
if (cuda_ctx_src->device != cur_device) {
GGML_CUDA_LOG_WARN("%s: attempt to copy on device %d while current device is %d\n", __func__, cuda_ctx_src->device, cur_device);
CUDA_CHECK(cudaSetDevice(cuda_ctx_src->device));
}
// copy on src stream
if (cuda_ctx_src->device == cuda_ctx_dst->device) {
CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream()));
@@ -3058,7 +3065,14 @@ GGML_CALL static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_
#ifdef GGML_CUDA_NO_PEER_COPY
return false;
#else
CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream()));
int can_access_peer;
CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, cuda_ctx_src->device, cuda_ctx_dst->device));
if (can_access_peer) {
CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream()));
} else {
GGML_CUDA_LOG_WARN("%s: attempt to copy from device %d to device %d without access enabled\n", __func__, cuda_ctx_src->device, cuda_ctx_dst->device);
return false;
}
#endif
}