From 6c6da89266ba7839d825c9997782af4f4d26b81b Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Mon, 31 Aug 2026 07:11:41 +0000 Subject: [PATCH] ggml-cuda: use cudaMemcpyDefault in the ggml_cuda_cpy 2D fast path The 2D fast path added in ggml-org#25057 copies with cudaMemcpyDeviceToDevice. That argument is not a hint, it is an assertion that both pointers are device resident. Under GGML_CUDA_ENABLE_UNIFIED_MEMORY the allocator returns hipMallocManaged pages, which migrate to the host once the working set exceeds VRAM, so at high footprint the assertion is false and the copy reads the wrong memory. cudaMemcpyDefault resolves residency per pointer instead, and is a no-op when both pointers really are on the device. Neither vendors/hip.h nor vendors/musa.h mapped cudaMemcpyDefault, so the one copy kind that resolves residency was unreachable from this file on non-CUDA backends. Both shims now carry the mapping. Measured on gfx1151, 45.1 GiB model at 73% of RAM, in one job on one host: this build produced identical tokens with the variable unset and set, while the stock build in the same job reproduced the reported corruption. A revert of 25057 is also clean but drops same-type strided copies outside F32/F16/BF16/I32 onto GGML_ABORT at cpy.cu:609, losing the coverage 25057 added for the GDN recurrent snapshot with -np 4. This keeps it. Not measured against a discrete GPU or the CUDA path. --- ggml/src/ggml-cuda/cpy.cu | 9 ++++++++- ggml/src/ggml-cuda/vendors/hip.h | 1 + ggml/src/ggml-cuda/vendors/musa.h | 1 + 3 files changed, 10 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/cpy.cu b/ggml/src/ggml-cuda/cpy.cu index fd7ffc0bc55..3b1af1ca48d 100644 --- a/ggml/src/ggml-cuda/cpy.cu +++ b/ggml/src/ggml-cuda/cpy.cu @@ -474,8 +474,15 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg CUDA_CHECK(cudaMemcpyAsync(src1_ddc, src0_ddc, ggml_nbytes(src0), cudaMemcpyDeviceToDevice, main_stream)); } } else if (ggml_cuda_cpy_as_memcpy_2d(src0, src1, mc_width, mc_height, mc_spitch, mc_dpitch)) { + // cudaMemcpyDefault, not DeviceToDevice: under + // GGML_CUDA_ENABLE_UNIFIED_MEMORY these buffers come from + // hipMallocManaged, and managed pages are not guaranteed device + // resident. DeviceToDevice asserts that they are, so the runtime skips + // residency resolution and a migrated page is read as garbage. Default + // lets it infer per pointer, and is a no-op when both really are on the + // device. CUDA_CHECK(cudaMemcpy2DAsync(src1_ddc, mc_dpitch, src0_ddc, mc_spitch, - mc_width, mc_height, cudaMemcpyDeviceToDevice, main_stream)); + mc_width, mc_height, cudaMemcpyDefault, main_stream)); } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) { if (can_be_transposed) { ggml_cpy_scalar_cuda diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index 9aa558f3f4c..466823101c9 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -87,6 +87,7 @@ #define cudaMemcpyAsync hipMemcpyAsync #define cudaMemcpyPeerAsync hipMemcpyPeerAsync #define cudaMemcpy2DAsync hipMemcpy2DAsync +#define cudaMemcpyDefault hipMemcpyDefault #define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice #define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost #define cudaMemcpyHostToDevice hipMemcpyHostToDevice diff --git a/ggml/src/ggml-cuda/vendors/musa.h b/ggml/src/ggml-cuda/vendors/musa.h index 6d725c7ec19..9160d88cb08 100644 --- a/ggml/src/ggml-cuda/vendors/musa.h +++ b/ggml/src/ggml-cuda/vendors/musa.h @@ -74,6 +74,7 @@ #define cudaMemcpyAsync musaMemcpyAsync #define cudaMemcpyPeerAsync musaMemcpyPeerAsync #define cudaMemcpy2DAsync musaMemcpy2DAsync +#define cudaMemcpyDefault musaMemcpyDefault #define cudaMemcpyDeviceToDevice musaMemcpyDeviceToDevice #define cudaMemcpyDeviceToHost musaMemcpyDeviceToHost #define cudaMemcpyHostToDevice musaMemcpyHostToDevice