From: Sigbjørn Skjæret <redacted>
Date: Tue, 9 Dec 2025 11:23:47 +0000 (+0100)
Subject: ggml : allow fill node alloc inplace (llama/17870)
X-Git-Tag: upstream/0.9.4.395~38
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=106c2686f8008496707ddc34c352fb37b7150307;p=pkg%2Fggml%2Fsources%2Fggml

ggml : allow fill node alloc inplace (llama/17870)
---

diff --git a/src/ggml-alloc.c b/src/ggml-alloc.c
index 218222ec..a5995fdc 100644
--- a/src/ggml-alloc.c
+++ b/src/ggml-alloc.c
@@ -25,6 +25,7 @@ static bool ggml_is_view(const struct ggml_tensor * t) {
 // ops that return true for this function must not use restrict pointers for their backend implementations
 bool ggml_op_can_inplace(enum ggml_op op) {
     switch (op) {
+        case GGML_OP_FILL:
         case GGML_OP_SCALE:
         case GGML_OP_DIAG_MASK_ZERO:
         case GGML_OP_DIAG_MASK_INF:
diff --git a/src/ggml-cuda/fill.cu b/src/ggml-cuda/fill.cu
index eb8ccb78..739062c4 100644
--- a/src/ggml-cuda/fill.cu
+++ b/src/ggml-cuda/fill.cu
@@ -4,7 +4,7 @@
 #define CUDA_FILL_BLOCK_SIZE 256
 
 template <typename T>
-static __global__ void fill_kernel(T * __restrict__ dst, const int64_t k, const T value) {
+static __global__ void fill_kernel(T * dst, const int64_t k, const T value) {
     const int64_t i = (int64_t)blockDim.x * blockIdx.x + threadIdx.x;
     if (i >= k) {
         return;