From: Johannes Gäßler <redacted>
Date: Wed, 12 Nov 2025 22:13:55 +0000 (+0100)
Subject: CUDA: static assert to prevent misuse of memcpy_1 (#17198)
X-Git-Tag: upstream/0.0.7446~405
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=5d6838b74f151443ef126a6efda46e93f3b721e3;p=pkg%2Fggml%2Fsources%2Fllama.cpp

CUDA: static assert to prevent misuse of memcpy_1 (#17198)
---

diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index ca876459..25e9308d 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -586,6 +586,12 @@ static __device__ __forceinline__ void ggml_cuda_mad(half2 & acc, const half2 v,
 //     If dst and src point at different address spaces then they are guaranteed to not be aliased.
 template <int nbytes, int alignment = 0>
 static __device__ __forceinline__ void ggml_cuda_memcpy_1(void * __restrict__ dst, const void * __restrict__ src) {
+    static_assert(
+        nbytes <= ggml_cuda_get_max_cpy_bytes() || alignment == 0,
+        "You are misusing the alignment parameter for ggml_cuda_memcpy_1. "
+        "The intent is for the parameter is only as a workaround if either one of the pointers is not properly aligned. "
+        "If you use it to do more bytes per copy than ggml_cuda_max_cpy_bytes() the reads and writes may not be coalesced. "
+        "Call ggml_cuda_memcpy_1 in a loop instead.");
     if constexpr (alignment != 0) {
         static_assert(nbytes % alignment == 0, "bad alignment");
     }