diff --git a/ggml/src/ggml-cuda/cpy.cu b/ggml/src/ggml-cuda/cpy.cu
index fd7ffc0bc..7b78b5083 100644
--- a/ggml/src/ggml-cuda/cpy.cu
+++ b/ggml/src/ggml-cuda/cpy.cu
@@ -386,6 +386,101 @@ static void ggml_cpy_f32_iq4_nl_cuda(
(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
}
+
+#if defined(GGML_USE_HIP)
+template <typename T>
+static __global__ void cpy_memcpy_2d(const char * GGML_CUDA_RESTRICT src, char * GGML_CUDA_RESTRICT dst,
+ const int64_t nchunks_row, const int64_t nchunks_total,
+ const int64_t spitch, const int64_t dpitch) {
+ ggml_cuda_pdl_sync();
+
+ const int64_t stride = (int64_t) gridDim.x*blockDim.x;
+
+ for (int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; i < nchunks_total; i += stride) {
+ const int64_t row = i / nchunks_row;
+ const int64_t col = i - row*nchunks_row;
+
+ const T * src_row = (const T *) (src + row*spitch);
+ T * dst_row = (T *) (dst + row*dpitch);
+
+ dst_row[col] = src_row[col];
+ }
+}
+
+template <typename T>
+static void ggml_cpy_memcpy_2d_cuda(const char * src, char * dst,
+ const size_t width, const size_t height, const size_t spitch, const size_t dpitch, cudaStream_t stream) {
+ const int64_t nchunks_row = width/sizeof(T);
+ const int64_t nchunks_total = nchunks_row*(int64_t) height;
+
+ GGML_ASSERT(nchunks_total > 0);
+
+ constexpr int block_size = 256;
+
+ // the kernel uses a grid-stride loop, so the number of blocks can be capped
+ const int64_t num_blocks = std::min<int64_t>((nchunks_total + block_size - 1)/block_size, 65536);
+
+ const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params((dim3) (unsigned int) num_blocks, (dim3) block_size, 0, stream);
+ ggml_cuda_kernel_launch(cpy_memcpy_2d<T>, launch_params,
+ src, dst, nchunks_row, nchunks_total, (int64_t) spitch, (int64_t) dpitch);
+}
+#endif // defined(GGML_USE_HIP)
+
+void ggml_cuda_memcpy_2d_d2d_async(void * dst, const size_t dpitch, const void * src, const size_t spitch,
+ const size_t width, const size_t height, cudaStream_t stream) {
+ if (width == 0 || height == 0) {
+ return;
+ }
+
+ GGML_ASSERT(spitch >= width && dpitch >= width);
+
+ if (height == 1 || (spitch == width && dpitch == width)) {
+ // the rows are adjacent in both src and dst, this is a plain 1D copy
+ CUDA_CHECK(cudaMemcpyAsync(dst, src, width*height, cudaMemcpyDeviceToDevice, stream));
+ return;
+ }
+
+#if !defined(GGML_USE_HIP)
+ CUDA_CHECK(cudaMemcpy2DAsync(dst, dpitch, src, spitch, width, height, cudaMemcpyDeviceToDevice, stream));
+#else
+ // On HIP, a device-to-device hipMemcpy2DAsync produces wrong results when the buffers were allocated with
+ // hipMallocManaged (GGML_CUDA_ENABLE_UNIFIED_MEMORY). Instead of relying on cudaMemcpy2DAsync use a strided
+ // cpy_memcpy_2d kernel to copy 2D data
+ static const bool unified_memory = getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr;
+ if (!unified_memory) {
+ CUDA_CHECK(cudaMemcpy2DAsync(dst, dpitch, src, spitch, width, height, cudaMemcpyDeviceToDevice, stream));
+ return;
+ }
+
+ const char * src_c = (const char *) src;
+ char * dst_c = (char *) dst;
+
+ const auto is_multiple_of = [&](const size_t n) {
+ return width % n == 0 && spitch % n == 0 && dpitch % n == 0 &&
+ ((uintptr_t) src_c) % n == 0 && ((uintptr_t) dst_c) % n == 0;
+ };
+
+ if (is_multiple_of(sizeof(int4))) {
+ ggml_cpy_memcpy_2d_cuda<int4>(src_c, dst_c, width, height, spitch, dpitch, stream);
+ } else if (is_multiple_of(sizeof(int32_t))) {
+ ggml_cpy_memcpy_2d_cuda<int32_t>(src_c, dst_c, width, height, spitch, dpitch, stream);
+ } else {
+ ggml_cpy_memcpy_2d_cuda<char>(src_c, dst_c, width, height, spitch, dpitch, stream);
+ }
+#endif // !defined(GGML_USE_HIP)
+}
+
+
// check if a same-type copy reduces to a 2D strided copy (height rows of width
// contiguous bytes), so it can use cudaMemcpy2DAsync instead of the scalar kernel
static bool ggml_cuda_cpy_as_memcpy_2d(const ggml_tensor * src0, const ggml_tensor * src1,
@@ -474,8 +576,7 @@ 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)) {
- CUDA_CHECK(cudaMemcpy2DAsync(src1_ddc, mc_dpitch, src0_ddc, mc_spitch,
- mc_width, mc_height, cudaMemcpyDeviceToDevice, main_stream));
+ ggml_cuda_memcpy_2d_d2d_async(src1_ddc, mc_dpitch, src0_ddc, mc_spitch, mc_width, mc_height, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
if (can_be_transposed) {
ggml_cpy_scalar_cuda<float, float, true>
diff --git a/ggml/src/ggml-cuda/cpy.cuh b/ggml/src/ggml-cuda/cpy.cuh
index a7a87d8fc..dc4b91090 100644
--- a/ggml/src/ggml-cuda/cpy.cuh
+++ b/ggml/src/ggml-cuda/cpy.cuh
@@ -5,3 +5,5 @@
void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, ggml_tensor * src1);
void ggml_cuda_dup(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
+
+void ggml_cuda_memcpy_2d_d2d_async(void * dst, size_t dpitch, const void * src, size_t spitch, size_t width, size_t height, cudaStream_t stream);
diff --git a/ggml/src/ggml-cuda/top-k.cu b/ggml/src/ggml-cuda/top-k.cu
index 4f41e071f..4f331271a 100644
--- a/ggml/src/ggml-cuda/top-k.cu
+++ b/ggml/src/ggml-cuda/top-k.cu
@@ -1,5 +1,6 @@
#include "argsort.cuh"
#include "top-k.cuh"
+#include "cpy.cuh"
#ifdef GGML_CUDA_USE_CUB
# ifdef GGML_CUDA_CUB_IS_HIPCUB
@@ -256,8 +257,7 @@ void ggml_cuda_op_top_k(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
} else {
argsort_f32_i32_cuda_cub(pool, src0_d, tmp_dst, ncols, iter_nrows, GGML_SORT_ORDER_DESC, stream);
}
- CUDA_CHECK(cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), iter_nrows,
- cudaMemcpyDeviceToDevice, stream));
+ ggml_cuda_memcpy_2d_d2d_async(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), iter_nrows, stream);
src0_d += ncols * iter_nrows;
dst_d += k * iter_nrows;
@@ -271,8 +271,7 @@ void ggml_cuda_op_top_k(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
ggml_cuda_pool_alloc<int> temp_dst_alloc(pool, ncols * nrows);
int * tmp_dst = temp_dst_alloc.get();
argsort_f32_i32_cuda_bitonic(src0_d, tmp_dst, ncols, nrows, GGML_SORT_ORDER_DESC, stream);
- CUDA_CHECK(cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows,
- cudaMemcpyDeviceToDevice, stream));
+ ggml_cuda_memcpy_2d_d2d_async(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, stream);
#if defined(GGML_USE_HIP)
}
#endif // defined(GGML_USE_HIP)
Name and Version
./llama-cli --version
version: 0.4.0-dev (build 11109, commit 6994643)
built with Clang 23.0.0 for Linux x86_64
Operating systems
Linux
GGML backends
HIP
Hardware
Ryzen AI Max 395+ (Strix Halo, Framework desktop)
Models
DeepSeek V4 Flash
Ling 3.0
Qwen 3.8 Flash
Problem description & steps to reproduce
Enabling GGML_CUDA_ENABLE_UNIFIED_MEMORY flag causes garbled output.
This flag is especially important for Strix Halo as it allows us to push our hardware further.
For example, with this flag I can run bullerwins DS4 Flash IQ3_XSS with Q8_0 experts with full context (262k). It crashes even with 16k context without the flag.
Same goes for Qwen 3.8 Flash. This flag allows jumping to Q5_K_XL quant (with n-gram offload of course) and still run it with full context, which is again impossible without it.
Unfortunately, by default, the output of these specific (and Ling 3) models is garbled from second token and onward.
First token is correct.
I've created the same bug in the upstream llama, but I feel like there are higher chances of this being fixed here, as this seems to be strix halo specific issue.
The problem seems to be with cudaMemcpy2DAsync when the memory was allocated with hipMallocManaged which is what is used when GGML_CUDA_ENABLE_UNIFIED_MEMORY flag is set. When the memory is copied that way (in top-k operation specifically), this results in corruption, resulting in garbled output.
I've made a crude fix for this, adding a helper function that switches to a strided copy when input is not a plane 1D data.
This works for all models I've tested that had issues with UM before.
I've created a crude (probably) fix, but I'm not really comfortable in creating a PR, so if someone can take a look at this and create a PR - feel free.
Diff
Relevant log output
Logs