From d6c3d4b4a817610d8a17babbd7903a9d4c0fb7f9 Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Thu, 22 May 2025 09:30:29 +0300 Subject: [PATCH] Add check if device to device access is enabled --- ggml/src/ggml-cuda.cu | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda.cu b/ggml/src/ggml-cuda.cu index 17b9731c..00a5466d 100644 --- a/ggml/src/ggml-cuda.cu +++ b/ggml/src/ggml-cuda.cu @@ -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 }