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

diff --git a/src/ggml-alloc.c b/src/ggml-alloc.c
index 218222ece8726d6bbe3bcaa1ee5d6f03f07ccdc6..a5995fdc2c3979630f3334ab5499effaf5444077 100644 (file)
--- 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 eb8ccb7802bd6ce3c10e395009d2415c1efb9893..739062c4057af274b25a065e10800bad91cb5835 100644 (file)
--- 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;