chore: update mllama to use ollama engine (#10637)

llama/llama.cpp/src/llama-quant.cpp

@@ -639,9 +639,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
         if (llama_model_has_encoder(&model)) {
             n_attn_layer *= 3;
         }
-        if (qs.n_attention_wv != n_attn_layer) {
-            LLAMA_LOG_WARN("%s: n_attention_wv is unexpected, expected: %d, found: %d\n", __func__, n_attn_layer, qs.n_attention_wv);
-        }
+        GGML_ASSERT((qs.n_attention_wv == n_attn_layer) && "n_attention_wv is unexpected");
     }
 
     size_t total_size_org = 0;

llama/llama.cpp/tools/mtmd/llava.cpp

@@ -462,7 +462,7 @@ struct llava_embd_batch {
     std::vector<llama_seq_id *> seq_ids;
     std::vector<int8_t>         logits;
     llama_batch batch;
-    llava_embd_batch(float * embd, int32_t n_embd, int32_t n_tokens, llama_pos pos_0, llama_seq_id seq_id) {
+    llava_embd_batch(float * embd, int32_t n_tokens, llama_pos pos_0, llama_seq_id seq_id) {
         pos     .resize(n_tokens);
         n_seq_id.resize(n_tokens);
         seq_ids .resize(n_tokens + 1);
@@ -474,7 +474,6 @@ struct llava_embd_batch {
             /*n_tokens       =*/ n_tokens,
             /*tokens         =*/ nullptr,
             /*embd           =*/ embd,
-            /*n_embd         =*/ n_embd,
             /*pos            =*/ pos.data(),
             /*n_seq_id       =*/ n_seq_id.data(),
             /*seq_id         =*/ seq_ids.data(),
@@ -498,7 +497,7 @@ bool llava_eval_image_embed(llama_context * ctx_llama, const struct llava_image_
             n_eval = n_batch;
         }
         float * embd = image_embed->embed+i*n_embd;
-        llava_embd_batch llava_batch = llava_embd_batch(embd, n_embd, n_eval, *n_past, 0);
+        llava_embd_batch llava_batch = llava_embd_batch(embd, n_eval, *n_past, 0);
         if (llama_decode(ctx_llama, llava_batch.batch)) {
             LOG_ERR("%s : failed to eval\n", __func__);
             return false;

llama/llama.go

@@ -17,7 +17,6 @@ package llama
 #include "llava.h"
 #include "gguf.h"
 
-#include "mllama.h"
 #include "sampling_ext.h"
 
 extern bool llamaProgressCallback(float progress, void *user_data);
@@ -510,63 +509,6 @@ func (c *ClipContext) NewEmbed(llamaContext *Context, data []byte) ([][]float32,
 	return embed, nil
 }
 
-type MllamaContext struct {
-	c *C.struct_mllama_ctx
-}
-
-func NewMllamaContext(llamaContext *Context, modelPath string) (*MllamaContext, error) {
-	mp := C.CString(modelPath)
-	defer C.free(unsafe.Pointer(mp))
-	c := C.mllama_model_load(mp, 1)
-	if c == nil {
-		return nil, fmt.Errorf("unable to load mllama model: %v", modelPath)
-	}
-
-	projEmbedSize := int(C.mllama_n_embd(c))
-	modelEmbedSize := llamaContext.Model().NEmbd()
-	if projEmbedSize != modelEmbedSize {
-		return nil, fmt.Errorf("projector embedding size (%d) does not match model (%d)", projEmbedSize, modelEmbedSize)
-	}
-
-	return &MllamaContext{c: c}, nil
-}
-
-func (m *MllamaContext) Free() {
-	C.mllama_free(m.c)
-}
-
-func (m *MllamaContext) NewEmbed(llamaContext *Context, data []byte, aspectRatioId int) ([][]float32, error) {
-	img := C.mllama_image_init()
-	defer C.mllama_image_free(img)
-
-	ok := bool(C.mllama_image_load_from_data(unsafe.Pointer(&data[0]), C.int(len(data)), 560, 560, 3, 4, C.int(aspectRatioId), img))
-	if !ok {
-		return nil, errors.New("unable to load mllama image data")
-	}
-
-	rows := make([]float32, m.EmbedSize(llamaContext))
-	ok = bool(C.mllama_image_encode(m.c, C.int(llamaContext.numThreads), img, (*C.float)(unsafe.Pointer(&rows[0]))))
-	if !ok {
-		return nil, errors.New("unable to make mllama embedding from image")
-	}
-
-	embed := make([][]float32, 1)
-	embed[0] = rows
-
-	return embed, nil
-}
-
-func (m *MllamaContext) EmbedSize(llamaContext *Context) int {
-	numTokens := int(C.mllama_n_positions(m.c) * C.mllama_n_tiles(m.c))
-	numEmbed := llamaContext.Model().NEmbd()
-
-	return numTokens * numEmbed
-}
-
-func (c *Context) SetCrossAttention(state bool) {
-	C.llama_set_cross_attention(c.c, C.bool(state))
-}
-
 func (c *Context) Synchronize() {
 	C.llama_synchronize(c.c)
 }

llama/mllama.h

-#ifndef MLLAMA_H
-#define MLLAMA_H
-
-#include <stddef.h>
-#include <stdint.h>
-
-#ifdef LLAMA_SHARED
-#if defined(_WIN32) && !defined(__MINGW32__)
-#ifdef LLAMA_BUILD
-#define MLLAMA_API __declspec(dllexport)
-#else
-#define MLLAMA_API __declspec(dllimport)
-#endif
-#else
-#define MLLAMA_API __attribute__((visibility("default")))
-#endif
-#else
-#define MLLAMA_API
-#endif
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-struct mllama_ctx;
-
-struct mllama_image_batch {
-    struct mllama_image *data;
-    size_t size;
-};
-
-MLLAMA_API struct mllama_ctx *mllama_model_load(const char *fname, int verbosity);
-MLLAMA_API struct mllama_ctx *mllama_model_load_cpu(const char *fname, int verbosity);
-
-MLLAMA_API void mllama_free(struct mllama_ctx *ctx);
-
-MLLAMA_API int32_t mllama_image_size(const struct mllama_ctx *ctx);
-MLLAMA_API int32_t mllama_patch_size(const struct mllama_ctx *ctx);
-MLLAMA_API int32_t mllama_hidden_size(const struct mllama_ctx *ctx);
-
-MLLAMA_API int mllama_n_patches(const struct mllama_ctx *ctx);
-MLLAMA_API int mllama_n_positions(const struct mllama_ctx *ctx);
-MLLAMA_API int mllama_n_tiles(const struct mllama_ctx *ctx);
-MLLAMA_API int mllama_n_embd(const struct mllama_ctx *ctx);
-MLLAMA_API size_t mllama_n_embd_bytes(const struct mllama_ctx *ctx);
-
-MLLAMA_API struct mllama_image *mllama_image_init();
-
-MLLAMA_API void mllama_image_free(struct mllama_image *img);
-MLLAMA_API void mllama_image_batch_free(struct mllama_image_batch *batch);
-
-MLLAMA_API bool mllama_image_load_from_data(const void *data, const int n, const int nx, const int ny, const int nc, const int nt, const int aspect_ratio_id, struct mllama_image *img);
-
-MLLAMA_API bool mllama_image_encode(struct mllama_ctx *ctx, int n_threads, struct mllama_image *img, float *vec);
-MLLAMA_API bool mllama_image_batch_encode(struct mllama_ctx *ctx, int n_threads, const struct mllama_image_batch *imgs, float *vec);
-
-#ifdef __cplusplus
-}
-#endif
-
-#endif // MLLAMA_H

llama/patches/0005-solar-pro.patch

@@ -270,7 +270,7 @@ index 3a4e72a3..831b68c0 100644
 +            // self-attention
 +            {
 +                // rope freq factors for llama3; may return nullptr for llama2 and other models
-+                ggml_tensor * rope_factors = static_cast<const llama_kv_cache_unified *>(memory)->cbs.get_rope_factors(n_ctx_per_seq, il);
++                ggml_tensor * rope_factors = model.get_rope_factors(n_ctx_per_seq, il);
 +
 +                // compute Q and K and RoPE them
 +                ggml_tensor * Qcur = build_lora_mm(model.layers[il].wq, cur);

llama/patches/0007-add-unpad-operator.patch

-From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001
-From: jmorganca <jmorganca@gmail.com>
-Date: Sun, 13 Apr 2025 22:10:06 -0400
-Subject: [PATCH] add unpad operator
-
-adds the unpad operator to GGML
----
- ggml/include/ggml.h                  | 10 +++++
- ggml/src/ggml-cpu/ggml-cpu.c         |  5 +++
- ggml/src/ggml-cpu/ops.cpp            | 55 ++++++++++++++++++++++++++++
- ggml/src/ggml-cpu/ops.h              |  1 +
- ggml/src/ggml-cuda/ggml-cuda.cu      |  4 ++
- ggml/src/ggml-cuda/pad.cu            | 46 +++++++++++++++++++++++
- ggml/src/ggml-cuda/pad.cuh           |  1 +
- ggml/src/ggml-metal/ggml-metal.m     | 33 +++++++++++++++++
- ggml/src/ggml-metal/ggml-metal.metal | 45 +++++++++++++++++++++++
- ggml/src/ggml.c                      | 25 ++++++++++++-
- 10 files changed, 223 insertions(+), 2 deletions(-)
-
-diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
-index e91dedf1..8dc107ba 100644
---- a/ggml/include/ggml.h
-+++ b/ggml/include/ggml.h
-@@ -489,6 +489,7 @@ extern "C" {
-         GGML_OP_UPSCALE, // nearest interpolate
-         GGML_OP_PAD,
-         GGML_OP_PAD_REFLECT_1D,
-+        GGML_OP_UNPAD,
-         GGML_OP_ARANGE,
-         GGML_OP_TIMESTEP_EMBEDDING,
-         GGML_OP_ARGSORT,
-@@ -1781,6 +1782,15 @@ extern "C" {
-             int                   p0,
-             int                   p1);
- 
-+    // unpad each dimension: [x, ..., x, y, ..., y] -> [x, ..., x]
-+    GGML_API struct ggml_tensor * ggml_unpad(
-+            struct ggml_context * ctx,
-+            struct ggml_tensor  * a,
-+            int                  p0,
-+            int                  p1,
-+            int                  p2,
-+            int                  p3);
-+
-     // Ref: https://github.com/CompVis/stable-diffusion/blob/main/ldm/modules/diffusionmodules/util.py#L151
-     // timesteps: [N,]
-     // return: [N, dim]
-diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
-index a30e67f2..835e6495 100644
---- a/ggml/src/ggml-cpu/ggml-cpu.c
-+++ b/ggml/src/ggml-cpu/ggml-cpu.c
-@@ -1951,6 +1951,10 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
-             {
-                 ggml_compute_forward_pad_reflect_1d(params, tensor);
-             } break;
-+        case GGML_OP_UNPAD:
-+            {
-+                ggml_compute_forward_unpad(params, tensor);
-+            } break;
-         case GGML_OP_ARANGE:
-             {
-                 ggml_compute_forward_arange(params, tensor);
-@@ -2274,6 +2278,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
-         case GGML_OP_UPSCALE:
-         case GGML_OP_PAD:
-         case GGML_OP_PAD_REFLECT_1D:
-+        case GGML_OP_UNPAD:
-         case GGML_OP_ARANGE:
-         case GGML_OP_TIMESTEP_EMBEDDING:
-         case GGML_OP_ARGSORT:
-diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
-index 955fec59..1868a10c 100644
---- a/ggml/src/ggml-cpu/ops.cpp
-+++ b/ggml/src/ggml-cpu/ops.cpp
-@@ -6690,6 +6690,61 @@ void ggml_compute_forward_pad_reflect_1d(
-     }
- }
- 
-+// ggml_compute_forward_unpad
-+
-+static void ggml_compute_forward_unpad_f32(
-+    const struct ggml_compute_params *params,
-+    struct ggml_tensor *dst) {
-+
-+    const struct ggml_tensor * src0 = dst->src[0];
-+
-+    GGML_ASSERT(src0->nb[0] == sizeof(float));
-+    GGML_ASSERT( dst->nb[0] == sizeof(float));
-+
-+    const int ith = params->ith;
-+    const int nth = params->nth;
-+
-+    GGML_TENSOR_UNARY_OP_LOCALS
-+
-+    float * dst_ptr = (float *) dst->data;
-+
-+    // TODO: optimize
-+
-+    for (int64_t i2 = 0; i2 < ne2; ++i2) {
-+        for (int64_t i1 = ith; i1 < ne1; i1 += nth) {
-+            for (int64_t i0 = 0; i0 < ne0; ++i0) {
-+                for (int64_t i3 = 0; i3 < ne3; ++i3) {
-+                    const int64_t dst_idx = i3*(ne0*ne1*ne2) + i2*(ne0*ne1) + i1*ne0 + i0;
-+
-+                    const float * src_ptr = (const float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
-+
-+                    if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
-+                        dst_ptr[dst_idx] = *src_ptr;
-+                    }
-+                }
-+            }
-+        }
-+    }
-+}
-+
-+void ggml_compute_forward_unpad(
-+    const struct ggml_compute_params * params,
-+    struct ggml_tensor * dst) {
-+
-+    const struct ggml_tensor * src0 = dst->src[0];
-+
-+    switch (src0->type) {
-+        case GGML_TYPE_F32:
-+            {
-+                ggml_compute_forward_unpad_f32(params, dst);
-+            } break;
-+        default:
-+            {
-+                GGML_ABORT("fatal error");
-+            }
-+    }
-+}
-+
- // ggml_compute_forward_arange
- 
- static void ggml_compute_forward_arange_f32(
-diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h
-index dc081b9e..a7125555 100644
---- a/ggml/src/ggml-cpu/ops.h
-+++ b/ggml/src/ggml-cpu/ops.h
-@@ -72,6 +72,7 @@ void ggml_compute_forward_pool_2d_back(const struct ggml_compute_params * params
- void ggml_compute_forward_upscale(const struct ggml_compute_params * params, struct ggml_tensor * dst);
- void ggml_compute_forward_pad(const struct ggml_compute_params * params, struct ggml_tensor * dst);
- void ggml_compute_forward_pad_reflect_1d(const struct ggml_compute_params * params, struct ggml_tensor * dst);
-+void ggml_compute_forward_unpad(const struct ggml_compute_params * params, struct ggml_tensor * dst);
- void ggml_compute_forward_arange(const struct ggml_compute_params * params, struct ggml_tensor * dst);
- void ggml_compute_forward_timestep_embedding(const struct ggml_compute_params * params, struct ggml_tensor * dst);
- void ggml_compute_forward_argsort(const struct ggml_compute_params * params, struct ggml_tensor * dst);
-diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
-index cb0d8528..6fe86674 100644
---- a/ggml/src/ggml-cuda/ggml-cuda.cu
-+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
-@@ -2238,6 +2238,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
-         case GGML_OP_PAD:
-             ggml_cuda_op_pad(ctx, dst);
-             break;
-+        case GGML_OP_UNPAD:
-+            ggml_cuda_op_unpad(ctx, dst);
-+            break;
-         case GGML_OP_ARANGE:
-             ggml_cuda_op_arange(ctx, dst);
-             break;
-@@ -3212,6 +3215,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
-         case GGML_OP_UPSCALE:
-             return op->src[0]->type == GGML_TYPE_F32 && op->op_params[0] == GGML_SCALE_MODE_NEAREST;
-         case GGML_OP_PAD:
-+        case GGML_OP_UNPAD:
-         case GGML_OP_ARANGE:
-         case GGML_OP_TIMESTEP_EMBEDDING:
-         case GGML_OP_LEAKY_RELU:
-diff --git a/ggml/src/ggml-cuda/pad.cu b/ggml/src/ggml-cuda/pad.cu
-index 77432b04..7d45a7e1 100644
---- a/ggml/src/ggml-cuda/pad.cu
-+++ b/ggml/src/ggml-cuda/pad.cu
-@@ -47,3 +47,49 @@ void ggml_cuda_op_pad(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
-         src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
-         dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], stream);
- }
-+
-+static __global__ void unpad_f32(const float * x, float * dst, const int ne0, const int ne00, const int ne01, const int ne02, const int ne03) {
-+    // blockIdx.z: idx of ne2*ne3, aka ne02*ne03
-+    // blockIdx.y: idx of ne1
-+    // blockIDx.x: idx of ne0 / BLOCK_SIZE
-+    int nidx = threadIdx.x + blockIdx.x * blockDim.x;
-+    if (nidx >= ne0) {
-+        return;
-+    }
-+
-+    // operation
-+    int offset_dst =
-+        nidx +
-+        blockIdx.y * ne0 +
-+        blockIdx.z * ne0 * gridDim.y;
-+    if (nidx < ne00 && blockIdx.y < ne01 && blockIdx.z < ne02*ne03) {
-+        int offset_src =
-+            nidx +
-+            blockIdx.y * ne00 +
-+            blockIdx.z * ne00 * ne01;
-+        dst[offset_dst] = x[offset_src];
-+    }
-+}
-+
-+static void unpad_f32_cuda(const float * x, float * dst,
-+    const int ne00, const int ne01, const int ne02, const int ne03,
-+    const int ne0, const int ne1, const int ne2, const int ne3, cudaStream_t stream) {
-+    int num_blocks = (ne0 + CUDA_PAD_BLOCK_SIZE - 1) / CUDA_PAD_BLOCK_SIZE;
-+    dim3 gridDim(num_blocks, ne1, ne2*ne3);
-+    unpad_f32<<<gridDim, CUDA_PAD_BLOCK_SIZE, 0, stream>>>(x, dst, ne0, ne00, ne01, ne02, ne03);
-+}
-+
-+void ggml_cuda_op_unpad(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
-+    const ggml_tensor * src0 = dst->src[0];
-+    const float * src0_d = (const float *)src0->data;
-+    float * dst_d = (float *)dst->data;
-+    cudaStream_t stream = ctx.stream();
-+
-+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
-+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
-+    GGML_ASSERT(src0->ne[3] == 1 && dst->ne[3] == 1); // just 3D tensors
-+
-+    unpad_f32_cuda(src0_d, dst_d,
-+        src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
-+        dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], stream);
-+}
-\ No newline at end of file
-diff --git a/ggml/src/ggml-cuda/pad.cuh b/ggml/src/ggml-cuda/pad.cuh
-index 8fd386b0..e2ededc3 100644
---- a/ggml/src/ggml-cuda/pad.cuh
-+++ b/ggml/src/ggml-cuda/pad.cuh
-@@ -3,3 +3,4 @@
- #define CUDA_PAD_BLOCK_SIZE 256
- 
- void ggml_cuda_op_pad(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
-+void ggml_cuda_op_unpad(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
-diff --git a/ggml/src/ggml-metal/ggml-metal.m b/ggml/src/ggml-metal/ggml-metal.m
-index 1b56f858..7641247e 100644
---- a/ggml/src/ggml-metal/ggml-metal.m
-+++ b/ggml/src/ggml-metal/ggml-metal.m
-@@ -347,6 +347,7 @@ static void ggml_backend_metal_device_rel(struct ggml_backend_metal_device_conte
-     GGML_METAL_KERNEL_TYPE_UPSCALE_F32,
-     GGML_METAL_KERNEL_TYPE_PAD_F32,
-     GGML_METAL_KERNEL_TYPE_PAD_REFLECT_1D_F32,
-+    GGML_METAL_KERNEL_TYPE_UNPAD_F32,
-     GGML_METAL_KERNEL_TYPE_ARANGE_F32,
-     GGML_METAL_KERNEL_TYPE_TIMESTEP_EMBEDDING_F32,
-     GGML_METAL_KERNEL_TYPE_ARGSORT_F32_I32_ASC,
-@@ -1294,6 +1295,7 @@ @implementation GGMLMetalClass
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_UPSCALE_F32,                     upscale_f32,                     true);
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_PAD_F32,                         pad_f32,                         true);
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_PAD_REFLECT_1D_F32,              pad_reflect_1d_f32,              true);
-+        GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_UNPAD_F32,                       unpad_f32,                       true);
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_TIMESTEP_EMBEDDING_F32,          timestep_embedding_f32,          true);
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ARANGE_F32,                      arange_f32,                      true);
-         GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ARGSORT_F32_I32_ASC,             argsort_f32_i32_asc,             true);
-@@ -1655,6 +1657,7 @@ static bool ggml_metal_supports_op(const struct ggml_backend_metal_device_contex
-         case GGML_OP_POOL_2D:
-         case GGML_OP_PAD:
-         case GGML_OP_PAD_REFLECT_1D:
-+        case GGML_OP_UNPAD:
-         case GGML_OP_TIMESTEP_EMBEDDING:
-         case GGML_OP_ARGSORT:
-         case GGML_OP_LEAKY_RELU:
-@@ -4184,6 +4187,36 @@ static bool ggml_metal_encode_node(
- 
-                 const int nth = MIN(1024, ne0);
- 
-+                [encoder dispatchThreadgroups:MTLSizeMake(ne1, ne2, ne3) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
-+            } break;
-+        case GGML_OP_UNPAD:
-+            {
-+                GGML_ASSERT(src0->type == GGML_TYPE_F32);
-+
-+                id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_UNPAD_F32].pipeline;
-+
-+                [encoder setComputePipelineState:pipeline];
-+                [encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
-+                [encoder setBuffer:id_dst  offset:offs_dst  atIndex:1];
-+                [encoder setBytes:&ne00 length:sizeof(ne00) atIndex:2];
-+                [encoder setBytes:&ne01 length:sizeof(ne01) atIndex:3];
-+                [encoder setBytes:&ne02 length:sizeof(ne02) atIndex:4];
-+                [encoder setBytes:&ne03 length:sizeof(ne03) atIndex:5];
-+                [encoder setBytes:&nb00 length:sizeof(nb00) atIndex:6];
-+                [encoder setBytes:&nb01 length:sizeof(nb01) atIndex:7];
-+                [encoder setBytes:&nb02 length:sizeof(nb02) atIndex:8];
-+                [encoder setBytes:&nb03 length:sizeof(nb03) atIndex:9];
-+                [encoder setBytes:&ne0  length:sizeof(ne0)  atIndex:10];
-+                [encoder setBytes:&ne1  length:sizeof(ne1)  atIndex:11];
-+                [encoder setBytes:&ne2  length:sizeof(ne2)  atIndex:12];
-+                [encoder setBytes:&ne3  length:sizeof(ne3)  atIndex:13];
-+                [encoder setBytes:&nb0  length:sizeof(nb0)  atIndex:14];
-+                [encoder setBytes:&nb1  length:sizeof(nb1)  atIndex:15];
-+                [encoder setBytes:&nb2  length:sizeof(nb2)  atIndex:16];
-+                [encoder setBytes:&nb3  length:sizeof(nb3)  atIndex:17];
-+
-+                const int nth = MIN(1024, ne0);
-+
-                 [encoder dispatchThreadgroups:MTLSizeMake(ne1, ne2, ne3) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
-             } break;
-         case GGML_OP_ARANGE:
-diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal
-index 9cfddf45..080a943b 100644
---- a/ggml/src/ggml-metal/ggml-metal.metal
-+++ b/ggml/src/ggml-metal/ggml-metal.metal
-@@ -3121,6 +3121,51 @@ kernel void kernel_pad_reflect_1d_f32(
-     }
- }
- 
-+kernel void kernel_unpad_f32(
-+    device  const char * src0,
-+    device        char * dst,
-+    constant   int64_t & ne00,
-+    constant   int64_t & ne01,
-+    constant   int64_t & ne02,
-+    constant   int64_t & ne03,
-+    constant  uint64_t & nb00,
-+    constant  uint64_t & nb01,
-+    constant  uint64_t & nb02,
-+    constant  uint64_t & nb03,
-+    constant   int64_t & ne0,
-+    constant   int64_t & ne1,
-+    constant   int64_t & ne2,
-+    constant   int64_t & ne3,
-+    constant  uint64_t & nb0,
-+    constant  uint64_t & nb1,
-+    constant  uint64_t & nb2,
-+    constant  uint64_t & nb3,
-+    uint3 tgpig[[threadgroup_position_in_grid]],
-+    uint3 tpitg[[thread_position_in_threadgroup]],
-+    uint3   ntg[[threads_per_threadgroup]]) {
-+
-+    const int64_t i3 = tgpig.z;
-+    const int64_t i2 = tgpig.y;
-+    const int64_t i1 = tgpig.x;
-+
-+    const int64_t i03 = i3;
-+    const int64_t i02 = i2;
-+    const int64_t i01 = i1;
-+
-+    device const float * src0_ptr = (device const float *) (src0 + i03*nb03 + i02*nb02 + i01*nb01);
-+    device       float * dst_ptr  = (device       float *) (dst  +  i3*nb3  +  i2*nb2  +  i1*nb1);
-+
-+    if (i1 < ne01 && i2 < ne02 && i3 < ne03) {
-+        for (int i0 = tpitg.x; i0 < ne0; i0 += ntg.x) {
-+            if (i0 < ne00) {
-+                dst_ptr[i0] = src0_ptr[i0];
-+            }
-+        }
-+
-+        return;
-+    }
-+}
-+
- kernel void kernel_arange_f32(
-     device        char * dst,
-     constant   ggml_metal_kargs_arange & args,
-diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
-index 8a654624..6b034d35 100644
---- a/ggml/src/ggml.c
-+++ b/ggml/src/ggml.c
-@@ -923,6 +923,7 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
-     "UPSCALE",
-     "PAD",
-     "PAD_REFLECT_1D",
-+    "UNPAD",
-     "ARANGE",
-     "TIMESTEP_EMBEDDING",
-     "ARGSORT",
-@@ -953,7 +954,7 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
-     "OPT_STEP_ADAMW",
- };
- 
--static_assert(GGML_OP_COUNT == 82, "GGML_OP_COUNT != 82");
-+static_assert(GGML_OP_COUNT == 83, "GGML_OP_COUNT != 83");
- 
- static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
-     "none",
-@@ -1018,6 +1019,7 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
-     "upscale(x)",
-     "pad(x)",
-     "pad_reflect_1d(x)",
-+    "unpad(x)",
-     "arange(start, stop, step)",
-     "timestep_embedding(timesteps, dim, max_period)",
-     "argsort(x)",
-@@ -1048,7 +1050,7 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
-     "adamw(x)",
- };
- 
--static_assert(GGML_OP_COUNT == 82, "GGML_OP_COUNT != 82");
-+static_assert(GGML_OP_COUNT == 83, "GGML_OP_COUNT != 83");
- 
- static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2");
- 
-@@ -4274,6 +4276,25 @@ struct ggml_tensor * ggml_pad_reflect_1d(
-     return result;
- }
- 
-+// ggml_unpad
-+
-+struct ggml_tensor * ggml_unpad(
-+    struct ggml_context * ctx,
-+    struct ggml_tensor  * a,
-+    int p0, int p1, int p2, int p3) {
-+
-+    struct ggml_tensor * result = ggml_new_tensor_4d(ctx, a->type,
-+            a->ne[0] - p0,
-+            a->ne[1] - p1,
-+            a->ne[2] - p2,
-+            a->ne[3] - p3);
-+
-+    result->op = GGML_OP_UNPAD;
-+    result->src[0] = a;
-+
-+    return result;
-+}
-+
- // ggml_arange
- 
- struct ggml_tensor * ggml_arange(

llama/patches/0010-ensure-KV-cache-is-fully-defragmented.patch → llama/patches/0008-ensure-KV-cache-is-fully-defragmented.patch

@@ -58,7 +58,7 @@ index c22687e4..c5948e8f 100644
  
              auto * gf = graph_init();
 diff --git a/src/llama-context.h b/src/llama-context.h
-index c4ab242a..9970dfc6 100644
+index c0ceacb1..0264e937 100644
 --- a/src/llama-context.h
 +++ b/src/llama-context.h
 @@ -5,6 +5,7 @@
@@ -70,10 +70,10 @@ index c4ab242a..9970dfc6 100644
  #include "ggml-cpp.h"
  #include "ggml-opt.h"
 diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp
-index a7b0a7eb..1a50c034 100644
+index 3dcad65b..60e67b03 100644
 --- a/src/llama-kv-cache.cpp
 +++ b/src/llama-kv-cache.cpp
-@@ -372,8 +372,6 @@ void llama_kv_cache_unified::commit() {
+@@ -364,8 +364,6 @@ void llama_kv_cache_unified::commit() {
  }
  
  bool llama_kv_cache_unified::update(llama_context & lctx) {
@@ -82,7 +82,7 @@ index a7b0a7eb..1a50c034 100644
      auto * sched = lctx.get_sched();
  
      if (has_shift) {
-@@ -396,8 +394,6 @@ bool llama_kv_cache_unified::update(llama_context & lctx) {
+@@ -388,8 +386,6 @@ bool llama_kv_cache_unified::update(llama_context & lctx) {
              res->set_inputs(nullptr);
  
              lctx.graph_compute(gf, false);
@@ -91,7 +91,7 @@ index a7b0a7eb..1a50c034 100644
          }
  
          {
-@@ -411,27 +407,36 @@ bool llama_kv_cache_unified::update(llama_context & lctx) {
+@@ -403,27 +399,36 @@ bool llama_kv_cache_unified::update(llama_context & lctx) {
  
      if (do_defrag) {
          LLAMA_LOG_DEBUG("%s: defragmenting KV cache\n", __func__);
@@ -133,7 +133,7 @@ index a7b0a7eb..1a50c034 100644
  }
  
  void llama_kv_cache_unified::defrag_sched(float thold) {
-@@ -715,11 +720,10 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_shift(
+@@ -707,11 +712,10 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_shift(
  llm_graph_result_ptr llama_kv_cache_unified::build_graph_defrag(
          const llama_cparams & cparams,
                 ggml_context * ctx,
@@ -147,7 +147,7 @@ index a7b0a7eb..1a50c034 100644
  #if 0
      // CPU defrag
      //
-@@ -791,32 +795,20 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_defrag(
+@@ -783,32 +787,20 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_defrag(
          ggml_backend_tensor_set(v_l[il], buf_v.data(), 0, buf_v.size());
      }
  #else
@@ -185,7 +185,7 @@ index a7b0a7eb..1a50c034 100644
  
              ggml_tensor * view_v_src;
              ggml_tensor * view_v_dst;
-@@ -824,31 +816,29 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_defrag(
+@@ -816,31 +808,29 @@ llm_graph_result_ptr llama_kv_cache_unified::build_graph_defrag(
              if (cparams.flash_attn) {
                  // NOTE: the V cache is not transposed when using flash attention
                  view_v_src = ggml_view_2d(ctx, v_l[il],
@@ -225,7 +225,7 @@ index a7b0a7eb..1a50c034 100644
      }
  
      //LLAMA_LOG_INFO("gf->n_nodes = %d\n", gf->n_nodes);
-@@ -865,17 +855,7 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
+@@ -857,17 +847,7 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
  
      assert(n_used <= n_kv);
  
@@ -244,7 +244,7 @@ index a7b0a7eb..1a50c034 100644
  
      // determine which KV cells to move where
      //
-@@ -883,10 +863,7 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
+@@ -875,10 +855,7 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
      //
      //  if ids[i] == i || ids[i] == n_kv, then cell i is not moved
      //
@@ -256,7 +256,7 @@ index a7b0a7eb..1a50c034 100644
  
      for (uint32_t i0 = 0; i0 < n_used; ++i0) {
          const auto & cell0 = cells[i0];
-@@ -935,19 +912,11 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
+@@ -927,19 +904,11 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
          // are we moving a continuous block of memory?
          bool cont = false;
  
@@ -276,7 +276,7 @@ index a7b0a7eb..1a50c034 100644
                  cont = false;
                  continue;
              }
-@@ -963,8 +932,10 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
+@@ -955,8 +924,10 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
              head = n_used;
  
              if (!cont) {
@@ -288,7 +288,7 @@ index a7b0a7eb..1a50c034 100644
              }
  
              nf++;
-@@ -974,22 +945,16 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
+@@ -966,22 +937,16 @@ bool llama_kv_cache_unified::defrag_prepare(int32_t n_max_nodes) {
              }
          }
  

llama/patches/0011-sort-devices-by-score.patch → llama/patches/0009-sort-devices-by-score.patch

@@ -11,7 +11,7 @@ with the fastest acceleration is loaded
  1 file changed, 13 insertions(+), 8 deletions(-)
 
 diff --git a/ggml/src/ggml-backend-reg.cpp b/ggml/src/ggml-backend-reg.cpp
-index 82ae1b5b..1487f322 100644
+index 405d8e31..4e67d243 100644
 --- a/ggml/src/ggml-backend-reg.cpp
 +++ b/ggml/src/ggml-backend-reg.cpp
 @@ -157,7 +157,7 @@ struct ggml_backend_reg_entry {

llama/patches/0015-ollama-debug-tensor.patch → llama/patches/0013-ollama-debug-tensor.patch

@@ -8,7 +8,7 @@ Subject: [PATCH] ollama debug tensor
  1 file changed, 6 insertions(+)
 
 diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
-index 835e6495..3902894b 100644
+index a30e67f2..2462d2b8 100644
 --- a/ggml/src/ggml-cpu/ggml-cpu.c
 +++ b/ggml/src/ggml-cpu/ggml-cpu.c
 @@ -15,6 +15,8 @@
@@ -20,7 +20,7 @@ index 835e6495..3902894b 100644
  #if defined(_MSC_VER) || defined(__MINGW32__)
  #include <malloc.h> // using malloc.h with MSC/MINGW
  #elif !defined(__FreeBSD__) && !defined(__NetBSD__) && !defined(__OpenBSD__)
-@@ -2846,6 +2848,10 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
+@@ -2841,6 +2843,10 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
  
          ggml_compute_forward(&params, node);
  

llm/memory.go

@@ -111,9 +111,8 @@ func EstimateGPULayers(gpus []discover.GpuInfo, f *ggml.GGML, projectors []strin
 	slog.Debug("evaluating", "library", gpus[0].Library, "gpu_count", len(gpus), "available", availableList)
 
 	for _, projector := range projectors {
-		weight, graph := projectorMemoryRequirements(projector)
+		weight := projectorMemoryRequirements(projector)
 		projectorWeights += weight
-		projectorGraph += graph
 
 		// multimodal models require at least 2048 context
 		opts.NumCtx = max(opts.NumCtx, 2048)
@@ -409,51 +408,21 @@ func (m MemoryEstimate) LogValue() slog.Value {
 	return slog.GroupValue(attrs...)
 }
 
-func projectorMemoryRequirements(filename string) (weights, graphSize uint64) {
+func projectorMemoryRequirements(filename string) (weights uint64) {
 	file, err := os.Open(filename)
 	if err != nil {
-		return 0, 0
+		return 0
 	}
 	defer file.Close()
 
 	ggml, _, err := ggml.Decode(file, 1024)
 	if err != nil {
-		return 0, 0
+		return 0
 	}
 
 	for _, layer := range ggml.Tensors().GroupLayers() {
 		weights += layer.Size()
 	}
 
-	switch arch := ggml.KV().Architecture(); arch {
-	case "mllama":
-		kv := func(n string) uint64 {
-			if v, ok := ggml.KV()[arch+".vision."+n].(uint32); ok {
-				return uint64(v)
-			}
-
-			return 0
-		}
-
-		imageSize := kv("image_size")
-
-		maxNumTiles := kv("max_num_tiles")
-		embeddingLength := kv("embedding_length")
-		headCount := kv("attention.head_count")
-
-		numPatches := (imageSize / kv("patch_size")) * (imageSize / kv("patch_size"))
-		if _, ok := ggml.Tensors().GroupLayers()["v"]["class_embd"]; ok {
-			numPatches++
-		}
-
-		numPaddedPatches := numPatches + 8 - (numPatches%8)%8
-
-		graphSize = 4 * (8 +
-			imageSize*imageSize*kv("num_channels")*maxNumTiles +
-			embeddingLength*numPatches*maxNumTiles +
-			9*embeddingLength*numPaddedPatches*maxNumTiles +
-			numPaddedPatches*maxNumTiles*numPaddedPatches*maxNumTiles*headCount)
-	}
-
-	return weights, graphSize
+	return weights
 }

llm/server.go

@@ -679,9 +679,8 @@ ws ::= ([ \t\n] ws)?
 const maxBufferSize = 512 * format.KiloByte
 
 type ImageData struct {
-	Data          []byte `json:"data"`
-	ID            int    `json:"id"`
-	AspectRatioID int    `json:"aspect_ratio_id"`
+	Data []byte `json:"data"`
+	ID   int    `json:"id"`
 }
 
 type CompletionRequest struct {

ml/backend.go

@@ -161,7 +161,6 @@ type Tensor interface {
 	Set(ctx Context, t2 Tensor, offset int, strides ...int) Tensor
 
 	Pad(ctx Context, shape ...int) Tensor
-	Unpad(ctx Context, shape ...int) Tensor
 
 	Stack(ctx Context, dim int, s ...Tensor) Tensor