Add support for Jinja-based tool calling (MiniCPM5 and other recent models).

llama.cpp.patches/README.md

@@ -205,6 +205,7 @@ can change on a llama.cpp bump — when it does, the asset list in
 | `ggml_src_ggml-vulkan_ggml-vulkan.cpp.patch` | Fixes unsigned integer underflow in `ggml_backend_vk_get_device_memory` where Vulkan's `heapUsage` can exceed `heapBudget` (clamps to zero instead of wrapping) |
 | `src_models_t5.cpp.patch` | Forward-declares the `graph<false>`/`graph<true>` explicit specializations before `build_arch_graph` so clang's `-std=gnu++23` doesn't reject them as specializations after implicit instantiation |
 | `src_models_eagle3.cpp.patch` | Moves `build_arch_graph` to the end of the file, after the `graph<true>`/`graph<false>` constructor specializations, so clang's `-std=gnu++23` doesn't reject them as explicit specializations appearing after the `make_unique<graph<...>>` implicit instantiation point |
+| `src_models_dflash.cpp.patch` | Same cosmocc/clang `build_arch_graph` ordering fix as `eagle3` for the new DFlash draft model |
 
 ## Creating New Patches
 

llama.cpp.patches/llamafile-files/BUILD.mk

@@ -79,6 +79,7 @@ LLAMA_SRCS_CPP := \
 	llama.cpp/src/models/delta-net-base.cpp \
 	llama.cpp/src/models/dots1.cpp \
 	llama.cpp/src/models/dream.cpp \
+	llama.cpp/src/models/dflash.cpp \
 	llama.cpp/src/models/eagle3.cpp \
 	llama.cpp/src/models/ernie4-5-moe.cpp \
 	llama.cpp/src/models/ernie4-5.cpp \
@@ -240,7 +241,6 @@ COMMON_SRCS_CPP := \
 	llama.cpp/common/jinja/runtime.cpp \
 	llama.cpp/common/jinja/string.cpp \
 	llama.cpp/common/jinja/value.cpp \
-	llama.cpp/common/json-partial.cpp \
 	llama.cpp/common/json-schema-to-grammar.cpp \
 	llama.cpp/common/license.cpp \
 	llama.cpp/common/llguidance.cpp \
@@ -420,6 +420,7 @@ TOOL_SERVER_SRCS := \
 	llama.cpp/tools/server/server-models.cpp \
 	llama.cpp/tools/server/server-queue.cpp \
 	llama.cpp/tools/server/server-schema.cpp \
+	llama.cpp/tools/server/server-stream.cpp \
 	llama.cpp/tools/server/server-task.cpp \
 	llama.cpp/tools/server/server-tools.cpp
 

llama.cpp.patches/patches/common_common.cpp.patch

@@ -71,8 +71,8 @@ diff --git a/common/common.cpp b/common/common.cpp
  #else
 @@ -1201,10 +1246,31 @@ common_init_result::common_init_result(common_params & params, bool model_only)
      if (params.fit_params) {
-         LOG_INF("%s: fitting params to device memory ...\n", __func__);
-         LOG_INF("%s: (for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n", __func__);
+         COM_TRC("%s", "fitting params to device memory ...\n");
+         COM_TRC("%s", "(for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n");
 +
 +        // if a multimodal projection model is specified and will be loaded on GPU,
 +        // add its estimated size to the fit margins so the fit leaves enough room for it

llama.cpp.patches/patches/ggml_src_ggml-backend.cpp.patch

@@ -43,7 +43,7 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
  };
 
  ggml_backend_buffer_t ggml_backend_multi_buffer_alloc_buffer(ggml_backend_buffer_t * buffers, size_t n_buffers) {
-@@ -2210,7 +2218,7 @@ bool ggml_backend_compare_graph_backend(ggml_backend_t backend1, ggml_backend_t
+@@ -2214,7 +2222,7 @@ bool ggml_backend_compare_graph_backend(ggml_backend_t backend1, ggml_backend_t
 
  // CPU backend - buffer
 
@@ -52,7 +52,7 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
      GGML_ASSERT(buffer);
      uintptr_t data = (uintptr_t)buffer->context;
 
-@@ -2222,33 +2230,33 @@ static void * ggml_backend_cpu_buffer_get_base(ggml_backend_buffer_t buffer) {
+@@ -2226,33 +2234,33 @@ static void * ggml_backend_cpu_buffer_get_base(ggml_backend_buffer_t buffer) {
      return (void *)data;
  }
 
@@ -91,7 +91,7 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
      GGML_ASSERT(src);
      if (ggml_backend_buffer_is_host(src->buffer)) {
          memcpy(dst->data, src->data, ggml_nbytes(src));
-@@ -2259,7 +2267,7 @@ static bool ggml_backend_cpu_buffer_cpy_tensor(ggml_backend_buffer_t buffer, con
+@@ -2263,7 +2271,7 @@ static bool ggml_backend_cpu_buffer_cpy_tensor(ggml_backend_buffer_t buffer, con
      GGML_UNUSED(buffer);
  }
 
@@ -100,15 +100,15 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
      GGML_ASSERT(buffer);
      memset(buffer->context, value, buffer->size);
  }
-@@ -2276,6 +2284,7 @@ static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_i = {
+@@ -2280,6 +2288,7 @@ static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_i = {
      /* .cpy_tensor      = */ ggml_backend_cpu_buffer_cpy_tensor,
      /* .clear           = */ ggml_backend_cpu_buffer_clear,
      /* .reset           = */ NULL,
 +    /* .free_struct     = */ NULL,
  };
 
  static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_from_ptr_i = {
-@@ -2290,19 +2299,20 @@ static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_from_ptr_i = {
+@@ -2294,19 +2303,20 @@ static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_from_ptr_i = {
      /* .cpy_tensor      = */ ggml_backend_cpu_buffer_cpy_tensor,
      /* .clear           = */ ggml_backend_cpu_buffer_clear,
      /* .reset           = */ NULL,
@@ -131,7 +131,7 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
      void * data = ggml_aligned_malloc(size);
 
      if (data == NULL) {
-@@ -2313,13 +2323,13 @@ static ggml_backend_buffer_t ggml_backend_cpu_buffer_type_alloc_buffer(ggml_back
+@@ -2317,13 +2327,13 @@ static ggml_backend_buffer_t ggml_backend_cpu_buffer_type_alloc_buffer(ggml_back
      return ggml_backend_buffer_init(buft, ggml_backend_cpu_buffer_i, data, size);
  }
 
@@ -147,7 +147,7 @@ diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
      return true;
 
      GGML_UNUSED(buft);
-@@ -2342,7 +2352,7 @@ ggml_backend_buffer_type_t ggml_backend_cpu_buffer_type(void) {
+@@ -2346,7 +2356,7 @@ ggml_backend_buffer_type_t ggml_backend_cpu_buffer_type(void) {
      return &ggml_backend_cpu_buffer_type;
  }
 

llama.cpp.patches/patches/ggml_src_ggml-cpu_ops.cpp.patch

@@ -20,7 +20,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
  // ggml_compute_forward_dup
 
  static void ggml_compute_forward_dup_same_cont(
-@@ -8399,6 +8411,15 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
+@@ -8426,6 +8438,15 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
      GGML_ASSERT((                            q_to_vec_dot) && "fattn: unsupported K-type");
      GGML_ASSERT((v->type == GGML_TYPE_F32 || v_to_float  ) && "fattn: unsupported V-type");
 
@@ -36,7 +36,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
      int ith = params->ith;
 
      for (int ir = ir0; ir < ir1; ++ir) {
-@@ -8418,10 +8439,22 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
+@@ -8445,10 +8466,22 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
          ggml_fp16_t * VKQ16 = (ggml_fp16_t *) (VKQ32 + 1*DV); // (temporary) FP16 VKQ accumulator
          ggml_fp16_t * Q_q   = (ggml_fp16_t *) (VKQ32 + 2*DV); // (temporary) buffer for Q converted to quantized/FP16
 
@@ -62,7 +62,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
          }
 
          const ggml_fp16_t * mp = mask ? (ggml_fp16_t *)((char *) mask->data + iq1*mask->nb[1] + (iq2%mask->ne[2])*mask->nb[2] + (iq3%mask->ne[3])*mask->nb[3]) : NULL;
-@@ -8450,7 +8483,13 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
+@@ -8477,7 +8510,13 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
              float s; // KQ value
 
              const char * k_data = (const char *) k->data + ( ic*nbk1 + ik2*nbk2 + ik3*nbk3);
@@ -76,7 +76,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
 
              s = s*scale; // scale KQ value
 
-@@ -8467,21 +8506,36 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
+@@ -8494,21 +8533,36 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
 
              const char * v_data = ((const char *) v->data + (ic*nbv1 + iv2*nbv2 + iv3*nbv3));
 
@@ -116,7 +116,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
              } else {
                  if (s > M) {
                      // s is new maximum, ms < 1.0f, vs == expf(s - s) == 1.0f
-@@ -8489,26 +8543,20 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
+@@ -8516,26 +8570,20 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
                      ms = expf(Mold - M);
 
                      // V = V*expf(Mold - M)
@@ -146,7 +146,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
              for (int64_t d = 0; d < DV; ++d) {
                  VKQ32[d] = GGML_CPU_FP16_TO_FP32(VKQ16[d]);
              }
-@@ -8729,8 +8777,24 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
+@@ -8756,8 +8804,24 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
                  const char * k_data = (const char *)k->data + (ic + tk)*nbk1 + ik2*nbk2 + ik3*nbk3;
                  if (kv_type == GGML_TYPE_F16) {
                      const ggml_fp16_t * k_f16 = (const ggml_fp16_t *)k_data;
@@ -173,7 +173,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
                      }
                  } else {
                      const float * k_f32_src = (const float *)k_data;
-@@ -8740,7 +8804,18 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
+@@ -8767,7 +8831,18 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
                  }
              }
              memset(KQ, 0, Q_TILE_SZ * KV_TILE_SZ * sizeof(float));
@@ -192,7 +192,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
              ggml_vec_scale_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, scale);
 
              // Set padded KQ entries to -inf so softmax gives them zero weight
-@@ -8793,7 +8868,15 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
+@@ -8820,7 +8895,15 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
              for (int tk = 0; tk < kv_tile; tk++) {
                  const char * v_data = (const char *)v->data + (ic + tk)*nbv1 + iv2*nbv2 + iv3*nbv3;
                  if (kv_type == GGML_TYPE_F16) {
@@ -208,7 +208,7 @@ diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
                  } else {
                      memcpy(V32 + tk * DV, v_data, DV * sizeof(float));
                  }
-@@ -8803,7 +8886,16 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
+@@ -8830,7 +8913,16 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
                      memset(KQ + tq * KV_TILE_SZ, 0, KV_TILE_SZ * sizeof(float));
                  }
              }

llama.cpp.patches/patches/ggml_src_ggml-cuda_cpy.cu.patch

@@ -1,23 +1,23 @@
 diff --git a/ggml/src/ggml-cuda/cpy.cu b/ggml/src/ggml-cuda/cpy.cu
 --- a/llama.cpp/ggml/src/ggml-cuda/cpy.cu
 +++ b/llama.cpp/ggml/src/ggml-cuda/cpy.cu
-@@ -368,6 +368,7 @@ static void ggml_cpy_q5_1_f32_cuda(
+@@ -374,6 +374,7 @@ static void ggml_cpy_q5_1_f32_cuda(
          ne10, ne11, ne12, nb10, nb11, nb12, nb13);
  }
 
 +#ifndef GGML_CUDA_NO_IQ_QUANTS
  static void ggml_cpy_f32_iq4_nl_cuda(
      const char * cx, char * cdst, const int64_t ne,
      const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02,
-@@ -379,6 +380,7 @@ static void ggml_cpy_f32_iq4_nl_cuda(
+@@ -385,6 +386,7 @@ static void ggml_cpy_f32_iq4_nl_cuda(
      cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL><<<num_blocks, 1, 0, stream>>>
          (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
  }
 +#endif // GGML_CUDA_NO_IQ_QUANTS
 
- void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, ggml_tensor * src1) {
-     const int64_t ne = ggml_nelements(src0);
-@@ -473,9 +475,11 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
+ // 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
+@@ -524,9 +526,11 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
      } else if (src0->type == GGML_TYPE_Q5_0 && src1->type == GGML_TYPE_F32) {
          ggml_cpy_q5_0_f32_cuda
                  (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);

llama.cpp.patches/patches/ggml_src_ggml-cuda_ggml-cuda.cu.patch

@@ -372,7 +372,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_buffer_t buf_src = src->view_src ? src->view_src->buffer : src->buffer;
      ggml_backend_buffer_t buf_dst = dst->view_src ? dst->view_src->buffer : dst->buffer;
 
-@@ -3243,7 +3266,7 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_
+@@ -3259,7 +3282,7 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_
      return true;
  }
 
@@ -381,7 +381,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context;
 
      CUDA_CHECK(cudaStreamSynchronize(cuda_ctx->stream()));
-@@ -4465,7 +4488,7 @@ static bool ggml_cuda_graph_set_enabled(ggml_backend_cuda_context * cuda_ctx, co
+@@ -4481,7 +4504,7 @@ static bool ggml_cuda_graph_set_enabled(ggml_backend_cuda_context * cuda_ctx, co
  }
  #endif // USE_CUDA_GRAPH
 
@@ -390,7 +390,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context;
 
      ggml_cuda_set_device(cuda_ctx->device);
-@@ -4524,13 +4547,13 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend,
+@@ -4540,13 +4563,13 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend,
      return GGML_STATUS_SUCCESS;
  }
 
@@ -406,7 +406,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context;
 
      if (ggml_backend_is_cuda(backend)) {
-@@ -4549,7 +4572,7 @@ static void ggml_backend_cuda_event_wait(ggml_backend_t backend, ggml_backend_ev
+@@ -4565,7 +4588,7 @@ static void ggml_backend_cuda_event_wait(ggml_backend_t backend, ggml_backend_ev
      }
  }
 
@@ -415,7 +415,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context;
 
  #ifdef USE_CUDA_GRAPH
-@@ -4837,7 +4860,7 @@ void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * tot
+@@ -4853,7 +4876,7 @@ void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * tot
      CUDA_CHECK(cudaMemGetInfo(free, total));
  }
 
@@ -424,7 +424,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      if (getenv("GGML_CUDA_REGISTER_HOST") == nullptr) {
          return false;
      }
-@@ -4860,7 +4883,7 @@ bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) {
+@@ -4876,7 +4899,7 @@ bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) {
  #endif // CUDART_VERSION >= 11010 || defined(GGML_USE_MUSA)
  }
 
@@ -433,7 +433,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      if (getenv("GGML_CUDA_REGISTER_HOST") == nullptr) {
          return;
      }
-@@ -4883,12 +4906,12 @@ struct ggml_backend_cuda_device_context {
+@@ -4899,12 +4922,12 @@ struct ggml_backend_cuda_device_context {
      int op_offload_min_batch_size;
  };
 
@@ -448,7 +448,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context;
      return ctx->description.c_str();
  }
-@@ -4969,7 +4992,7 @@ static bool ggml_backend_cuda_get_available_uma_memory(long * available_memory_k
+@@ -4985,7 +5008,7 @@ static bool ggml_backend_cuda_get_available_uma_memory(long * available_memory_k
  }
  #endif // defined(__linux__)
 
@@ -457,7 +457,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context;
      ggml_cuda_set_device(ctx->device);
      CUDA_CHECK(cudaMemGetInfo(free, total));
-@@ -4999,7 +5022,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t *
+@@ -5015,7 +5038,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t *
 
  }
 
@@ -466,7 +466,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *) dev->context;
 
      cudaDeviceProp prop;
-@@ -5010,7 +5033,7 @@ static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend
+@@ -5026,7 +5049,7 @@ static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend
          : GGML_BACKEND_DEVICE_TYPE_GPU;
  }
 
@@ -475,7 +475,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *)dev->context;
 
      props->name        = ggml_backend_cuda_device_get_name(dev);
-@@ -5034,24 +5057,24 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back
+@@ -5050,24 +5073,24 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back
      };
  }
 
@@ -504,7 +504,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *) dev->context;
 
      // split buffers can only be used with GGML_OP_MUL_MAT
-@@ -5455,7 +5478,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
+@@ -5471,7 +5494,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
      }
  }
 
@@ -513,7 +513,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      ggml_backend_cuda_device_context * dev_ctx = (ggml_backend_cuda_device_context *) dev->context;
      const bool integrated = ggml_cuda_info().devices[dev_ctx->device].integrated;
      return (((ggml_backend_buft_is_cuda(buft) || ggml_backend_buft_is_cuda_split(buft)) && buft->device == dev) || (integrated && ggml_backend_buft_is_cuda_host(buft)));
-@@ -5476,13 +5499,13 @@ static int64_t get_op_batch_size(const ggml_tensor * op) {
+@@ -5492,13 +5515,13 @@ static int64_t get_op_batch_size(const ggml_tensor * op) {
      }
  }
 
@@ -529,7 +529,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
  #ifdef GGML_CUDA_NO_PEER_COPY
      return nullptr;
  #else
-@@ -5500,14 +5523,14 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_
+@@ -5516,14 +5539,14 @@ static ggml_backend_event_t ggml_backend_cuda_device_event_new(ggml_backend_dev_
  #endif
  }
 
@@ -546,7 +546,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      GGML_UNUSED(dev);
      CUDA_CHECK(cudaEventSynchronize((cudaEvent_t)event->context));
  }
-@@ -5536,23 +5559,23 @@ struct ggml_backend_cuda_reg_context {
+@@ -5552,23 +5575,23 @@ struct ggml_backend_cuda_reg_context {
      std::vector<ggml_backend_dev_t> devices;
  };
 
@@ -574,7 +574,7 @@ diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
      static std::vector<ggml_backend_feature> features = []() {
          std::vector<ggml_backend_feature> features;
      #define _STRINGIFY(...) #__VA_ARGS__
-@@ -5613,7 +5636,7 @@ static ggml_backend_feature * ggml_backend_cuda_get_features(ggml_backend_reg_t
+@@ -5629,7 +5652,7 @@ static ggml_backend_feature * ggml_backend_cuda_get_features(ggml_backend_reg_t
      GGML_UNUSED(reg);
  }