Merge branch 'upstream' into concedo_experimental

# Conflicts:
#	.github/workflows/build-openvino.yml
#	.github/workflows/server-sanitize.yml
#	AUTHORS
#	README.md
#	app/llama.cpp
#	ci/run.sh
#	common/build-info.h
#	docs/backend/SYCL.md
#	docs/build.md
#	docs/ops.md
#	ggml/CMakeLists.txt
#	ggml/src/ggml-cuda/CMakeLists.txt
#	ggml/src/ggml-musa/CMakeLists.txt
#	ggml/src/ggml-opencl/CMakeLists.txt
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-opencl/kernels/cvt.cl
#	ggml/src/ggml-opencl/kernels/gemm_noshuffle_q4_k_f32.cl
#	ggml/src/ggml-opencl/kernels/gemm_noshuffle_q6_k_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_1_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
#	ggml/src/ggml-opencl/kernels/mul_mm_f32_f32_l4_lm.cl
#	ggml/src/ggml-opencl/kernels/rms_norm.cl
#	ggml/src/ggml-sycl/binbcast.cpp
#	ggml/src/ggml-sycl/binbcast.hpp
#	ggml/src/ggml-sycl/common.hpp
#	ggml/src/ggml-sycl/fattn.cpp
#	ggml/src/ggml-sycl/fusion.cpp
#	ggml/src/ggml-sycl/ggml-sycl.cpp
#	ggml/src/ggml-sycl/norm.cpp
#	ggml/src/ggml-sycl/norm.hpp
#	scripts/snapdragon/build.py
#	scripts/snapdragon/qdc/run_qdc_jobs.py
#	scripts/snapdragon/qdc/tests/linux/run_linux.sh
#	scripts/snapdragon/qdc/tests/run_backend_ops_posix.py
#	scripts/snapdragon/qdc/tests/run_bench_tests_posix.py
#	scripts/snapdragon/qdc/tests/utils.py
#	scripts/snapdragon/run.py
#	src/CMakeLists.txt
#	src/llama.cpp
#	tests/test-backend-ops.cpp
#	tests/test-json-schema-to-grammar.cpp

common/build-info.cpp.in

@@ -29,7 +29,7 @@ const char * llama_build_info(void) {
     return s.c_str();
 }
 
-void llama_print_build_info(const char * llama_version) {
-    fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version, llama_build_number(), llama_commit());
-    fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target());
+void llama_print_build_info(const char * llama_version, FILE * stream) {
+    fprintf(stream, "version: %s (build %d, commit %s)\n", llama_version, llama_build_number(), llama_commit());
+    fprintf(stream, "built with %s for %s\n", llama_compiler(), llama_build_target());
 }

common/build-info.h

@@ -33,7 +33,7 @@ static inline const char * llama_build_info(void) {
     return s.c_str();
 }
 
-static inline void llama_print_build_info(const char *) {
+static inline void llama_print_build_info(const char *, FILE * = stderr) {
     fprintf(stderr, "%s: build = %d (%s)\n",      __func__, llama_build_number(), llama_commit());
     fprintf(stderr, "%s: built with %s for %s\n", __func__, llama_compiler(), llama_build_target());
 }

common/json-schema-to-grammar.cpp

@@ -748,6 +748,10 @@ private:
             optional_props.push_back("*");
         }
 
+        if (required_props.empty() && optional_props.empty()) {
+            return "\"{\" space \"}\"";
+        }
+
         std::string rule = "\"{\" space ";
         for (size_t i = 0; i < required_props.size(); i++) {
             if (i > 0) {

ggml/src/ggml-backend-reg.cpp

@@ -491,7 +491,13 @@ static ggml_backend_reg_t ggml_backend_load_best(const char * name, bool silent,
 #endif
         // default search paths: executable directory, current directory
         search_paths.push_back(get_executable_path());
-        search_paths.push_back(fs::current_path());
+        std::error_code cwd_ec;
+        const fs::path cwd = fs::current_path(cwd_ec);
+        if (cwd_ec) {
+            GGML_LOG_DEBUG("%s: current_path() failure, error-message: %s\n", __func__, cwd_ec.message().c_str());
+        } else {
+            search_paths.push_back(cwd);
+        }
     } else {
         search_paths.push_back(fs::u8path(user_search_path));
     }
@@ -509,8 +515,14 @@ static ggml_backend_reg_t ggml_backend_load_best(const char * name, bool silent,
             }
             continue;
         }
-        fs::directory_iterator dir_it(search_path, fs::directory_options::skip_permission_denied);
-        for (const auto & entry : dir_it) {
+        std::error_code dir_ec;
+        fs::directory_iterator dir_it(search_path, fs::directory_options::skip_permission_denied, dir_ec);
+        if (dir_ec) {
+            GGML_LOG_DEBUG("%s: failed to enumerate %s: %s\n", __func__, path_str(search_path).c_str(), dir_ec.message().c_str());
+            continue;
+        }
+        for (const fs::directory_iterator end; dir_it != end; dir_it.increment(dir_ec)) {
+            const auto & entry = *dir_it;
             if (entry.is_regular_file(ec)) {
                 auto filename = entry.path().filename();
                 auto ext = entry.path().extension();

ggml/src/ggml-cpu/arch/s390/quants.c

@@ -636,7 +636,7 @@ void ggml_vec_dot_q5_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const voi
         const float32x4_t v_xyf = vec_float(v_xy);
 
         const float32x4_t v_d = vec_splats(GGML_CPU_FP16_TO_FP32(x0->d) * GGML_CPU_FP16_TO_FP32(y0->d));
-        const float32x4_t v_acc = vec_madd(v_xyf, v_d, v_acc);
+        const float32x4_t v_acc = vec_madd(v_xyf, v_d, vec_splats(0.0f));
 
         sumf += vec_hsum_f32x4(v_acc) + summs;
     }

ggml/src/ggml-cuda/common.cuh

@@ -52,6 +52,7 @@
 #define GGML_CUDA_CC_VOLTA           700
 #define GGML_CUDA_CC_TURING          750
 #define GGML_CUDA_CC_AMPERE          800
+#define GGML_CUDA_CC_ORIN            870
 #define GGML_CUDA_CC_ADA_LOVELACE    890
 #define GGML_CUDA_CC_HOPPER          900
 // While BW spans CC 1000, 1100 & 1200, we are integrating Tensor Core instructions available to 1200 family, see

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -4555,10 +4555,12 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
     ggml_cuda_stream_context & stream_context = cuda_ctx->stream_context();
     stream_context.reset();
 
-    if (!use_cuda_graph || ggml_backend_cuda_get_device_count() != 1) {
+    if (!use_cuda_graph) {
         return;
     }
 
+    ggml_cuda_set_device(cuda_ctx->device);
+
     // number of out-degrees for a particular node
     std::unordered_map<const ggml_tensor *, int> fan_out;
     // reverse mapping of node to index in the cgraph

ggml/src/ggml-cuda/mmvq.cu

@@ -326,6 +326,18 @@ bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) {
                 return ne11 <= MMVQ_MAX_BATCH_SIZE;
         }
     }
+    if (GGML_CUDA_CC_IS_NVIDIA(cc) && cc == GGML_CUDA_CC_ORIN) {
+        switch (type) { // tuned for Jetson Orin
+            case GGML_TYPE_Q2_K:
+            case GGML_TYPE_Q3_K:
+            case GGML_TYPE_Q4_K:
+            case GGML_TYPE_Q5_K:
+            case GGML_TYPE_Q6_K:
+                return ne11 <= 1;
+            default:
+                return ne11 <= MMVQ_MAX_BATCH_SIZE;
+        }
+    }
     if (GGML_CUDA_CC_IS_CDNA(cc)) {
         if (GGML_CUDA_CC_IS_CDNA1(cc)) {
             switch (type) {

ggml/src/ggml-metal/ggml-metal-device.cpp

@@ -1577,6 +1577,26 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext(
     return res;
 }
 
+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_vec_idx(
+        ggml_metal_library_t lib,
+        const ggml_tensor * op) {
+    assert(op->op == GGML_OP_FLASH_ATTN_EXT);
+    assert(op->src[3]);
+
+    char name[256];
+
+    snprintf(name, 256, "kernel_flash_attn_ext_vec_idx");
+
+    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
+    if (!res.pipeline) {
+        res = ggml_metal_library_compile_pipeline(lib, name, name, nullptr);
+    }
+
+    GGML_UNUSED(op);
+
+    return res;
+}
+
 ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_vec(
         ggml_metal_library_t lib,
         const ggml_tensor * op,
@@ -1585,6 +1605,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_v
         bool    has_bias,
         bool    has_scap,
         bool    has_kvpad,
+        bool    has_sparse,
         int32_t nqpsg,
         int32_t ne,
         int32_t nsg,
@@ -1614,13 +1635,14 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_v
             dv,
             qne_suffix);
 
-    snprintf(name, 256, "%s_mask=%d_sink=%d_bias=%d_scap=%d_kvpad=%d_ns10=%d_ns20=%d_nsg=%d_nwg=%d",
+    snprintf(name, 256, "%s_mask=%d_sink=%d_bias=%d_scap=%d_kvpad=%d_sparse=%d_ns10=%d_ns20=%d_nsg=%d_nwg=%d",
             base,
             has_mask,
             has_sinks,
             has_bias,
             has_scap,
             has_kvpad,
+            has_sparse,
             ns10,
             ns20,
             nsg, nwg);
@@ -1633,7 +1655,8 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_v
         ggml_metal_cv_set_bool(cv, has_sinks, FC_FLASH_ATTN_EXT_VEC + 1);
         ggml_metal_cv_set_bool(cv, has_bias,  FC_FLASH_ATTN_EXT_VEC + 2);
         ggml_metal_cv_set_bool(cv, has_scap,  FC_FLASH_ATTN_EXT_VEC + 3);
-        ggml_metal_cv_set_bool(cv, has_kvpad, FC_FLASH_ATTN_EXT_VEC + 4);
+        ggml_metal_cv_set_bool(cv, has_kvpad,  FC_FLASH_ATTN_EXT_VEC + 4);
+        ggml_metal_cv_set_bool(cv, has_sparse, FC_FLASH_ATTN_EXT_VEC + 5);
 
         ggml_metal_cv_set_int32(cv, ns10, FC_FLASH_ATTN_EXT_VEC + 20);
         ggml_metal_cv_set_int32(cv, ns20, FC_FLASH_ATTN_EXT_VEC + 21);

ggml/src/ggml-metal/ggml-metal-device.h

@@ -201,6 +201,10 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_att
         int32_t ns10,
         int32_t ns20);
 
+struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_vec_idx(
+        ggml_metal_library_t lib,
+        const struct ggml_tensor * op);
+
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_attn_ext_vec(
         ggml_metal_library_t lib,
         const struct ggml_tensor * op,
@@ -209,6 +213,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_flash_att
         bool    has_bias,
         bool    has_scap,
         bool    has_kvpad,
+        bool    has_sparse,
         int32_t nqpsg,
         int32_t ne,
         int32_t nsg,

ggml/src/ggml-metal/ggml-metal-impl.h

@@ -458,8 +458,21 @@ typedef struct {
     float    m1;
     int32_t  n_head_log2;
     float    logit_softcap;
+    int32_t  n_kv_max_padded;
 } ggml_metal_kargs_flash_attn_ext_vec;
 
+typedef struct {
+    int32_t  ne30;
+    int32_t  ne31;
+    int32_t  ne32;
+    int32_t  ne33;
+    uint64_t nb31;
+    uint64_t nb32;
+    uint64_t nb33;
+    int32_t  n_kv_max;
+    int32_t  n_kv_max_padded;
+} ggml_metal_kargs_flash_attn_ext_vec_idx;
+
 typedef struct {
     int32_t  nrows;
 } ggml_metal_kargs_flash_attn_ext_vec_reduce;

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -917,7 +917,7 @@ int ggml_metal_op_glu(ggml_metal_op_t ctx, int idx) {
 
     const int64_t nrows = ggml_nrows(op->src[0]);
 
-    const int32_t nth = std::min(ggml_metal_pipeline_max_theads_per_threadgroup(pipeline), ne00/2);
+    const int32_t nth = std::max(1, std::min(ggml_metal_pipeline_max_theads_per_threadgroup(pipeline), ne00/2));
 
     ggml_metal_encoder_set_pipeline(enc, pipeline);
     ggml_metal_encoder_set_bytes   (enc, &args, sizeof(args), 0);
@@ -2857,6 +2857,65 @@ static bool ggml_metal_op_flash_attn_ext_use_kv_f16(const ggml_tensor * op) {
     }
 }
 
+// returns the n_kv_max hint if the sparse path is available for this op, or 0 otherwise
+// the mask (src[3]) remains the single source of truth: finite entries are the valid KV positions,
+// n_kv_max is only an upper bound on their number per mask row, used to size the index lists
+static int ggml_metal_op_flash_attn_ext_n_kv_max_sparse(const ggml_tensor * op) {
+    assert(op->op == GGML_OP_FLASH_ATTN_EXT);
+
+    int32_t n_kv_max = 0;
+    memcpy(&n_kv_max, ((const int32_t *) op->op_params) + 4, sizeof(n_kv_max));
+
+    if (n_kv_max <= 0) {
+        return 0;
+    }
+
+    // the sparse indices are gathered from the mask
+    if (!op->src[3]) {
+        return 0;
+    }
+
+    // bound the size of the index lists
+    if (n_kv_max > 4096) {
+        return 0;
+    }
+
+    // vec kernel instantiations exist for these (type, dk, dv) combinations only
+    const int64_t dk = op->src[1]->ne[0];
+    const int64_t dv = op->src[2]->ne[0];
+
+    const bool dk_dv_ok = (dk == 32  && dv == 32)  ||
+                          (dk == 64  && dv == 64)  ||
+                          (dk == 96  && dv == 96)  ||
+                          (dk == 128 && dv == 128) ||
+                          (dk == 192 && dv == 128) ||
+                          (dk == 192 && dv == 192) ||
+                          (dk == 256 && dv == 256) ||
+                          (dk == 320 && dv == 256) ||
+                          (dk == 512 && dv == 512) ||
+                          (dk == 576 && dv == 512);
+
+    if (!dk_dv_ok) {
+        return 0;
+    }
+
+    switch (op->src[1]->type) {
+        case GGML_TYPE_F16:
+        case GGML_TYPE_BF16:
+        case GGML_TYPE_F32:
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_Q4_1:
+        case GGML_TYPE_Q5_0:
+        case GGML_TYPE_Q5_1:
+        case GGML_TYPE_Q8_0:
+            break;
+        default:
+            return 0;
+    }
+
+    return n_kv_max;
+}
+
 // in some models (e.g. MLA-based), V is a view of K (the first ne20 elements of each K row);
 // the dequantized V is then a view of the dequantized K and does not need its own dequant or scratch
 // - ref: https://github.com/ggml-org/llama.cpp/pull/13435
@@ -3027,6 +3086,24 @@ size_t ggml_metal_op_flash_attn_ext_extra_kv_f16(const ggml_tensor * op) {
     return k_size + v_size;
 }
 
+// size of the sparse index lists: one list of KV indices per mask row,
+// padded with -1 up to a multiple of OP_FLASH_ATTN_EXT_VEC_NCPSG
+size_t ggml_metal_op_flash_attn_ext_extra_idx(const ggml_tensor * op) {
+    assert(op->op == GGML_OP_FLASH_ATTN_EXT);
+
+    GGML_TENSOR_LOCALS( int32_t, ne3, op->src[3], ne);
+
+    const int n_kv_max = ggml_metal_op_flash_attn_ext_n_kv_max_sparse(op);
+
+    if (n_kv_max <= 0) {
+        return 0;
+    }
+
+    const int n_kv_max_padded = GGML_PAD(n_kv_max, OP_FLASH_ATTN_EXT_VEC_NCPSG);
+
+    return GGML_PAD(sizeof(int32_t)*(size_t) n_kv_max_padded*ne31*ne32*ne33, 16);
+}
+
 int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
     ggml_tensor * op = ctx->node(idx);
 
@@ -3104,7 +3181,16 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
     ggml_metal_buffer_id bid_kv_f16 = bid_tmp;
     bid_kv_f16.offs += ggml_metal_op_flash_attn_ext_extra_tmp(op);
 
-    const bool use_kv_f16 = ggml_metal_op_flash_attn_ext_use_kv_f16(op);
+    // sparse path: gather the finite mask entries into index lists and run the vec kernels over them
+    const int n_kv_max_sparse = ggml_metal_op_flash_attn_ext_n_kv_max_sparse(op);
+    const bool use_sparse = n_kv_max_sparse > 0;
+    const int n_kv_max_padded = use_sparse ? GGML_PAD(n_kv_max_sparse, OP_FLASH_ATTN_EXT_VEC_NCPSG) : 0;
+
+    // the vec kernels dequantize the KV inline; no need for the F16 dequant pass in the sparse path
+    const bool use_kv_f16 = !use_sparse && ggml_metal_op_flash_attn_ext_use_kv_f16(op);
+
+    ggml_metal_buffer_id bid_idx = bid_kv_f16;
+    bid_idx.offs += ggml_metal_op_flash_attn_ext_extra_kv_f16(op);
 
     ggml_metal_buffer_id bid_k = bid_src1;
     ggml_metal_buffer_id bid_v = bid_src2;
@@ -3206,7 +3292,7 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
         }
     }
 
-    if (!ggml_metal_op_flash_attn_ext_use_vec(op)) {
+    if (!use_sparse && !ggml_metal_op_flash_attn_ext_use_vec(op)) {
         // half8x8 kernel
         const int nqptg = OP_FLASH_ATTN_EXT_NQPSG; // queries per threadgroup
         const int ncpsg = OP_FLASH_ATTN_EXT_NCPSG; // cache values per simdgroup
@@ -3378,13 +3464,18 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
 #undef FATTN_SMEM
     } else {
         // half4x4 kernel
-        auto cfg = ggml_metal_tuning::fa_vec_pick(
-                props_dev->device_id,
-                props_dev->gpu_family,
-                (int) op->src[1]->type,
-                (int) ne00, (int) ne20,   // dk, dv (ne00 == dk for FA)
-                ne11, ne01);
-        int nqptg = cfg.Q;                             // queries per threadgroup
+        // sparse: the index lists are per query row, so a threadgroup can share KV with Q == 1 only
+        auto cfg = use_sparse
+                ? ggml_metal_tuning::fa_vec_baseline_cfg((int) ne00, (int) ne20)
+                : ggml_metal_tuning::fa_vec_pick(
+                          props_dev->device_id,
+                          props_dev->gpu_family,
+                          (int) op->src[1]->type,
+                          (int) ne00, (int) ne20,   // dk, dv (ne00 == dk for FA)
+                          ne11, ne01);
+
+        int nqptg = cfg.Q; // queries per threadgroup
+
         const int ncpsg = OP_FLASH_ATTN_EXT_VEC_NCPSG; // cache values per simdgroup !! sync with kernel template arguments !!
         const int nhptg = 1;                           // heads per threadgroup
 
@@ -3394,7 +3485,39 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
 
         bool need_sync = false;
 
-        const bool has_kvpad = ne11 % ncpsg != 0;
+        const bool has_kvpad = !use_sparse && ne11 % ncpsg != 0;
+
+        if (use_sparse) {
+            assert(ggml_metal_op_flash_attn_ext_extra_idx(op) != 0);
+
+            GGML_ASSERT(ne30 == ne11);
+
+            ggml_metal_kargs_flash_attn_ext_vec_idx args0 = {
+                /*.ne30              =*/ ne30,
+                /*.ne31              =*/ ne31,
+                /*.ne32              =*/ ne32,
+                /*.ne33              =*/ ne33,
+                /*.nb31              =*/ nb31,
+                /*.nb32              =*/ nb32,
+                /*.nb33              =*/ nb33,
+                /*.n_kv_max          =*/ n_kv_max_sparse,
+                /*.n_kv_max_padded   =*/ n_kv_max_padded,
+            };
+
+            auto pipeline0 = ggml_metal_library_get_pipeline_flash_attn_ext_vec_idx(lib, op);
+
+            ggml_metal_encoder_set_pipeline(enc, pipeline0);
+            ggml_metal_encoder_set_bytes   (enc, &args0, sizeof(args0), 0);
+            ggml_metal_encoder_set_buffer  (enc, bid_src3, 1);
+            ggml_metal_encoder_set_buffer  (enc, bid_idx,  2);
+
+            int nth = std::min(ggml_metal_pipeline_max_theads_per_threadgroup(pipeline0), 256);
+            nth = std::max(32, (nth/32)*32);
+
+            ggml_metal_encoder_dispatch_threadgroups(enc, ne31, ne32, ne33, nth, 1, 1);
+
+            need_sync = true;
+        }
 
         if (has_kvpad) {
             assert(ggml_metal_op_flash_attn_ext_extra_pad(op) != 0);
@@ -3455,11 +3578,26 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
         // workgroups
         // each workgroup handles nsg*nkpsg cache values
         int32_t nwg = 1;
-        if (false) {
-            // for small KV caches, we could launch a single workgroup and write the results directly to dst/
-            // however, this does not lead to significant improvement, so disabled
-            nwg = 1;
-            nsg = 4;
+        if (use_sparse) {
+            if (ne01 > 32) {
+                // large sparse batch
+                nwg = 1;
+                nsg = 1;
+                if (n_kv_max_padded == 640) {
+                    nsg = 4; // 640 % (4*32) == 0
+                } else {
+                    while (2*nwg*nsg*ncpsg < n_kv_max_padded && nsg < 4) {
+                        nsg *= 2;
+                    }
+                }
+            } else {
+                // small sparse batch
+                nwg = 32;
+                nsg = 1;
+                while (2*nwg*nsg*ncpsg < n_kv_max_padded && nsg < 4) {
+                    nsg *= 2;
+                }
+            }
         } else {
             nwg = 32;
             nsg = 1;
@@ -3484,7 +3622,7 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
             /*.nb01          =*/ nb01,
             /*.nb02          =*/ nb02,
             /*.nb03          =*/ nb03,
-            /*.ne11          =*/ ne11,
+            /*.ne11          =*/ use_sparse ? n_kv_max_padded : ne11,
             /*.ne_12_2       =*/ ne12,
             /*.ne_12_3       =*/ ne13,
             /*.ns10          =*/ ns10,
@@ -3510,9 +3648,10 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
             /*.m1            =*/ m1,
             /*.n_head_log2   =*/ n_head_log2,
             /*.logit_softcap =*/ logit_softcap,
+            /*.n_kv_max_padded =*/ n_kv_max_padded,
         };
 
-        auto pipeline = ggml_metal_library_get_pipeline_flash_attn_ext_vec(lib, op, has_mask, has_sinks, has_bias, has_scap, has_kvpad, nqptg, cfg.NE, nsg, nwg, use_kv_f16, ns10, ns20);
+        auto pipeline = ggml_metal_library_get_pipeline_flash_attn_ext_vec(lib, op, has_mask, has_sinks, has_bias, has_scap, has_kvpad, use_sparse, nqptg, cfg.NE, nsg, nwg, use_kv_f16, ns10, ns20);
 
         GGML_ASSERT(nsg*32 <= ggml_metal_pipeline_max_theads_per_threadgroup(pipeline));
 
@@ -3523,6 +3662,7 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
         ggml_metal_encoder_set_buffer  (enc, bid_v,    3);
         ggml_metal_encoder_set_buffer  (enc, bid_src3, 4);
         ggml_metal_encoder_set_buffer  (enc, bid_src4, 5);
+        ggml_metal_encoder_set_buffer  (enc, use_sparse ? bid_idx : bid_src0, 8);
 
         const size_t smem = FATTN_SMEM(nsg);
 
@@ -3530,8 +3670,6 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
         GGML_ASSERT(smem <= props_dev->max_theadgroup_memory_size);
 
         if (nwg == 1) {
-            assert(ggml_metal_op_flash_attn_ext_extra_tmp(op) == 0);
-
             // using 1 workgroup -> write the result directly into dst
             ggml_metal_encoder_set_buffer(enc, bid_pad, 6);
             ggml_metal_encoder_set_buffer(enc, bid_dst, 7);

ggml/src/ggml-metal/ggml-metal-ops.h

@@ -43,6 +43,7 @@ size_t ggml_metal_op_flash_attn_ext_extra_pad(const struct ggml_tensor * op);
 size_t ggml_metal_op_flash_attn_ext_extra_blk(const struct ggml_tensor * op);
 size_t ggml_metal_op_flash_attn_ext_extra_tmp(const struct ggml_tensor * op);
 size_t ggml_metal_op_flash_attn_ext_extra_kv_f16(const struct ggml_tensor * op);
+size_t ggml_metal_op_flash_attn_ext_extra_idx(const struct ggml_tensor * op);
 
 int ggml_metal_op_concat            (ggml_metal_op_t ctx, int idx);
 int ggml_metal_op_repeat            (ggml_metal_op_t ctx, int idx);

ggml/src/ggml-metal/ggml-metal.cpp

@@ -232,6 +232,7 @@ static size_t ggml_backend_metal_buffer_type_get_alloc_size(ggml_backend_buffer_
                 res += ggml_metal_op_flash_attn_ext_extra_blk(tensor);
                 res += ggml_metal_op_flash_attn_ext_extra_tmp(tensor);
                 res += ggml_metal_op_flash_attn_ext_extra_kv_f16(tensor);
+                res += ggml_metal_op_flash_attn_ext_extra_idx(tensor);
             } break;
         case GGML_OP_CUMSUM:
         case GGML_OP_ARGSORT:

ggml/src/ggml-metal/kernels/fa.metal

@@ -1071,6 +1071,112 @@ constant int32_t FC_flash_attn_ext_vec_ns10 [[function_constant(FC_FLASH_ATTN_EX
 constant int32_t FC_flash_attn_ext_vec_ns20 [[function_constant(FC_FLASH_ATTN_EXT_VEC + 21)]];
 constant int32_t FC_flash_attn_ext_vec_nsg  [[function_constant(FC_FLASH_ATTN_EXT_VEC + 22)]];
 constant int32_t FC_flash_attn_ext_vec_nwg  [[function_constant(FC_FLASH_ATTN_EXT_VEC + 23)]];
+constant bool    FC_flash_attn_ext_vec_has_sparse [[function_constant(FC_FLASH_ATTN_EXT_VEC + 5)]];
+
+// compress the finite entries of each KQ mask row into a list of KV indices (ascending order),
+// padded with -1 up to n_kv_max_padded (a multiple of OP_FLASH_ATTN_EXT_VEC_NCPSG)
+// one threadgroup per mask row; the mask remains the single source of truth for the values
+kernel void kernel_flash_attn_ext_vec_idx(
+        constant ggml_metal_kargs_flash_attn_ext_vec_idx & args,
+        device const half * mask,
+        device       int  * idx,
+        uint3   tgpig[[threadgroup_position_in_grid]],
+        ushort  tiitg[[thread_index_in_threadgroup]],
+        ushort3 ntg[[threads_per_threadgroup]]) {
+    constexpr short NW = N_SIMDWIDTH;
+    constexpr short NLOCAL = 32; // max finite positions kept in registers per thread
+
+    const int i1 = tgpig[0];
+    const int i2 = tgpig[1];
+    const int i3 = tgpig[2];
+
+    device const half * pm  = (device const half *) ((device const char *) mask + i1*args.nb31 + i2*args.nb32 + i3*args.nb33);
+    device int * pidx = idx + (((int64_t)i3*args.ne32 + i2)*args.ne31 + i1)*args.n_kv_max_padded;
+
+    const int n  = args.ne30;
+    const int q  = n/ntg.x;
+    const int r  = n%ntg.x;
+
+    // each thread handles a contiguous slice of the mask row
+    const int r0 = q*tiitg + min((int) tiitg, r);
+    const int r1 = r0 + q + (tiitg < r ? 1 : 0);
+
+    // count the finite entries in the slice and keep their positions in registers (single mask read)
+    int cnt = 0;  // total finite entries in the slice
+    int nloc = 0; // finite entries kept in registers
+    int local[NLOCAL];
+    for (int i = r0; i < r1; ++i) {
+        if (isfinite((float) pm[i])) {
+            if (nloc < NLOCAL) {
+                local[nloc] = i;
+                nloc++;
+            }
+            cnt++;
+        }
+    }
+
+    const short sgitg = tiitg/NW;
+    const short tiisg = tiitg%NW;
+
+    threadgroup int tcount[8];
+
+    // simd_sum is a collective: all lanes must evaluate it
+    const int sg_sum = simd_sum(cnt);
+    if (tiisg == 0) {
+        tcount[sgitg] = sg_sum;
+    }
+
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+
+    int total = 0;
+    for (short s = 0; s < ntg.x/NW; ++s) {
+        total += tcount[s];
+    }
+
+    // base offset of this thread's slice in the output list (exclusive scan within the simdgroup)
+    int sg_base = 0;
+    for (short s = 0; s < sgitg; ++s) {
+        sg_base += tcount[s];
+    }
+
+    // exclusive prefix scan of the per-thread counts within the simdgroup
+    int incl = cnt;
+    for (int d = 1; d < NW; d <<= 1) {
+        const int v = simd_shuffle_up(incl, d);
+        if (tiisg >= d) {
+            incl += v;
+        }
+    }
+    const int base = sg_base + (incl - cnt);
+
+    // write the finite positions in order; if the hint is violated, keep only the first n_kv_max entries
+    int j = 0;
+    for (; j < nloc && base + j < args.n_kv_max; ++j) {
+        pidx[base + j] = local[j];
+    }
+
+    // a dense mask may have more than NLOCAL finite entries in a slice; re-read the mask to write the rest
+    if (cnt > nloc && base + nloc < args.n_kv_max) {
+        int j2 = 0;
+        for (int i = r0; i < r1; ++i) {
+            if (isfinite((float) pm[i])) {
+                if (j2 >= nloc) {
+                    pidx[base + j2] = i;
+                }
+                j2++;
+                if (base + j2 >= args.n_kv_max) {
+                    break;
+                }
+            }
+        }
+    }
+
+    // pad the tail of the list with -1
+    const int count = min(total, args.n_kv_max);
+    for (int i = count + tiitg; i < args.n_kv_max_padded; i += ntg.x) {
+        pidx[i] = -1;
+    }
+}
 
 template<
     typename q4_t,  // query types in shared memory
@@ -1091,6 +1197,7 @@ template<
     short NE = 4,   // head elements per thread
     short Q  = OP_FLASH_ATTN_EXT_VEC_NQPSG,  // queries per threadgroup
     short C  = OP_FLASH_ATTN_EXT_VEC_NCPSG>  // cache items per threadgroup
+
 kernel void kernel_flash_attn_ext_vec(
         constant ggml_metal_kargs_flash_attn_ext_vec & args,
         device const char * q,
@@ -1100,6 +1207,7 @@ kernel void kernel_flash_attn_ext_vec(
         device const char * sinks,
         device const char * pad,
         device       char * dst,
+        device const char * idx,
         threadgroup  half * shmem_f16 [[threadgroup(0)]],
         uint3   tgpig[[threadgroup_position_in_grid]],
         ushort  tiisg[[thread_index_in_simdgroup]],
@@ -1137,8 +1245,8 @@ kernel void kernel_flash_attn_ext_vec(
 
   //const short T = PK + NSG*SH; // shared memory size per query in (half)
 
-  //threadgroup q_t   * sq  = (threadgroup q_t   *) (shmem_f16 +                            0*PK); // holds the query data
-    threadgroup q4_t  * sq4 = (threadgroup q4_t  *) (shmem_f16 +                            0*PK); // same as above but in q4_t
+  //threadgroup q_t   * sq  = (threadgroup q_t   *) (shmem_f16 +                          0*PK); // holds the query data
+    threadgroup q4_t  * sq4 = (threadgroup q4_t  *) (shmem_f16 +                          0*PK); // same as above but in q4_t
     threadgroup s_t   * ss  = (threadgroup s_t   *) (shmem_f16 +   sgitg*SH         + Q*NSG*PK); // scratch buffer for attention
     threadgroup s4_t  * ss4 = (threadgroup s4_t  *) (shmem_f16 +   sgitg*SH         + Q*NSG*PK); // same as above but in s4_t
     threadgroup half  * sm  = (threadgroup half  *) (shmem_f16 +   sgitg*SH + 2*Q*C + Q*NSG*PK); // scratch buffer for mask
@@ -1207,6 +1315,14 @@ kernel void kernel_flash_attn_ext_vec(
         // pointer to the mask
         device const half * pm_base = (device const half *) (mask + iq1*Q*args.nb31 + (iq2%args.ne32)*args.nb32 + (iq3%args.ne33)*args.nb33);
 
+        // sparse indices: the list of finite mask entries per query row
+        // the sparse path requires Q == 1 (enforced by the host)
+        device const int * pidx = nullptr;
+        if (FC_flash_attn_ext_vec_has_sparse) {
+            pidx = (device const int *) idx +
+                ((int64_t)(iq3%args.ne33)*args.ne32 + (iq2%args.ne32))*args.ne31*args.n_kv_max_padded + (iq1%args.ne31)*args.n_kv_max_padded;
+        }
+
         float slope = 1.0f;
 
         // ALiBi
@@ -1265,11 +1381,22 @@ kernel void kernel_flash_attn_ext_vec(
             }
 
             if (FC_flash_attn_ext_vec_has_mask) {
-                FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
-                    if ((iq1*Q + qq) < args.ne01) {
-                        sm[qq*C + tiisg] = pm[qq][ic + tiisg];
-                    } else {
-                        sm[qq*C + tiisg] = -MAXHALF;
+                if (FC_flash_attn_ext_vec_has_sparse) {
+                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                        const int i11 = pidx[ic + tiisg];
+                        if ((iq1*Q + qq) < args.ne01 && i11 >= 0) {
+                            sm[qq*C + tiisg] = pm[qq][i11];
+                        } else {
+                            sm[qq*C + tiisg] = -MAXHALF;
+                        }
+                    }
+                } else {
+                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                        if ((iq1*Q + qq) < args.ne01) {
+                            sm[qq*C + tiisg] = pm[qq][ic + tiisg];
+                        } else {
+                            sm[qq*C + tiisg] = -MAXHALF;
+                        }
                     }
                 }
             } else {
@@ -1280,6 +1407,7 @@ kernel void kernel_flash_attn_ext_vec(
                 }
             }
 
+            // skip -INF mask
             {
                 bool any_finite = false;
                 FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
@@ -1294,9 +1422,13 @@ kernel void kernel_flash_attn_ext_vec(
 
             // Q*K^T
             {
-                device      const k4_t * pk4 = (device const k4_t *) (k + ic*args.nb11);
+                device      const k4_t * pk4 = nullptr;
 
-                pk4 += ty*NS10/4 + tx;
+                if (!FC_flash_attn_ext_vec_has_sparse) {
+                    pk4 = (device const k4_t *) (k + ic*args.nb11);
+
+                    pk4 += ty*NS10/4 + tx;
+                }
 
                 qk_t mqk[Q][C/NE];
                 FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
@@ -1307,7 +1439,35 @@ kernel void kernel_flash_attn_ext_vec(
 
                 // each simdgroup processes Q queries and NE (NW/NL) cache elements
                 FOR_UNROLL (short cc = 0; cc < C/NE; ++cc) {
-                    if (is_same<kd4_t, k4_t>::value) {
+                    if (FC_flash_attn_ext_vec_has_sparse) {
+                        // the KV rows are gathered from the index list; -1 entries are padding
+                        const int i11 = pidx[ic + NE*cc + ty];
+                        if (i11 >= 0) {
+                            if (is_same<kd4_t, k4_t>::value) {
+                                device const k4_t * pk4s = (device const k4_t *) (k + i11*args.nb11) + tx;
+                                FOR_UNROLL (short ii = 0; ii < DK4/NL; ++ii) {
+                                    const k4_t k_elem = pk4s[ii*NL];
+                                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                                        mqk[qq][cc] += dot((float4) k_elem, (float4) sq4[qq*PK4 + ii*NL + tx]);
+                                    }
+                                }
+                            } else {
+                                device const kd4_t * pk = (device const kd4_t *) (k + i11*args.nb11);
+
+                                k4_t mk;
+
+                                FOR_UNROLL (short ii = 0; ii < DK4/NL; ++ii) {
+                                    const short i = ii*NL + tx;
+
+                                    deq_k_t4(pk + i/nl_k, i%nl_k, mk);
+
+                                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                                        mqk[qq][cc] += dot((float4) mk, (float4) sq4[qq*PK4 + i]);
+                                    }
+                                }
+                            }
+                        }
+                    } else if (is_same<kd4_t, k4_t>::value) {
                         FOR_UNROLL (short ii = 0; ii < DK4/NL; ++ii) {
                             const k4_t k_elem = pk4[cc*NE*NS10/4 + ii*NL];
                             FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
@@ -1422,7 +1582,40 @@ kernel void kernel_flash_attn_ext_vec(
                     }
                 }
 
-                if (is_same<vd4_t, v4_t>::value) {
+                if (FC_flash_attn_ext_vec_has_sparse) {
+                    FOR_UNROLL (short cc = 0; cc < C/NE; ++cc) {
+                        // the KV rows are gathered from the index list; -1 entries are padding
+                        const int i11 = pidx[ic + NE*cc + ty];
+                        if (i11 >= 0) {
+                            if (is_same<vd4_t, v4_t>::value) {
+                                device const v4_t * pv4 = (device const v4_t *) (v + i11*args.nb21);
+
+                                pv4 += tx;
+
+                                FOR_UNROLL (short ii = 0; ii < DV4/NL; ++ii) {
+                                    const v4_t v_elem = pv4[ii*NL];
+                                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                                        lo[qq][ii] += o4_t(float4(v_elem)*float4(ss[qq*C + cc*NE + ty]));
+                                    }
+                                }
+                            } else {
+                                device const vd4_t * pv4 = (device const vd4_t *) (v + i11*args.nb21);
+
+                                FOR_UNROLL (short ii = 0; ii < DV4/NL; ++ii) {
+                                    const short i = ii*NL + tx;
+
+                                    v4_t mv;
+
+                                    deq_v_t4(pv4 + i/nl_v, i%nl_v, mv);
+
+                                    FOR_UNROLL (short qq = 0; qq < Q; ++qq) {
+                                        lo[qq][ii] += o4_t(float4(mv)*float4(ss[qq*C + cc*NE + ty]));
+                                    }
+                                }
+                            }
+                        }
+                    }
+                } else if (is_same<vd4_t, v4_t>::value) {
                     device const v4_t * pv4 = (device const v4_t *) (v + ic*args.nb21);
 
                     pv4 += ty*NS20/4 + tx;

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -10999,7 +10999,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
         return t->nb[0] == ggml_type_size(t->type) &&
                t->nb[2] == ggml_row_size(t->type, t->ne[0]) &&
                t->nb[1] == t->nb[2] * t->ne[2] &&
-               t->nb[3] == t->nb[1] * t->ne[1];
+               (t->ne[3] == 1 || t->nb[3] == t->nb[1] * t->ne[1]);
     };
     const bool k_quant = k->type != GGML_TYPE_F16 && k->type != GGML_TYPE_BF16 && k->type != GGML_TYPE_F32;
     const bool v_quant = v->type != GGML_TYPE_F16 && v->type != GGML_TYPE_BF16 && v->type != GGML_TYPE_F32;

json_to_gbnf.py

@@ -734,6 +734,9 @@ class SchemaConverter:
             )
             optional_props.append("*")
 
+        if not required_props and not optional_props:
+            return '"{" space "}"'
+
         rule = '"{" space '
         rule += ' "," space '.join(prop_kv_rule_names[k] for k in required_props)
 

src/llama-hparams.cpp

@@ -87,6 +87,15 @@ uint32_t llama_hparams::n_expert_used(uint32_t il) const {
     GGML_ABORT("fatal error");
 }
 
+uint32_t llama_hparams::n_expert_used_max() const {
+    uint32_t val = 0;
+    for (uint32_t il = 0; il < n_layer_all; ++il) {
+        val = std::max(val, n_expert_used(il));
+    }
+
+    return val;
+}
+
 uint32_t llama_hparams::n_gqa(uint32_t il) const {
     const uint32_t n_head    = this->n_head(il);
     const uint32_t n_head_kv = this->n_head_kv(il);

src/llama-hparams.h

@@ -392,6 +392,9 @@ struct llama_hparams {
 
     uint32_t n_expert_used(uint32_t il = 0) const;
 
+    // return the maximum n_expert_used across all layers
+    uint32_t n_expert_used_max() const;
+
     uint32_t n_gqa(uint32_t il = 0) const;
 
     uint32_t n_rot(uint32_t il = 0) const;

src/llama-model-loader.cpp

@@ -952,7 +952,7 @@ static bool weight_buft_supported(const llama_hparams & hparams, ggml_tensor * w
         case GGML_OP_MUL_MAT_ID:
             {
                 // Used for either MoE expert routing or embedded adapter routing
-                const int n_ids_used = hparams.router_layer >= 0 ? 1 : hparams.n_expert_used();
+                const int n_ids_used = hparams.router_layer >= 0 ? 1 : hparams.n_expert_used_max();
                 GGML_ASSERT(n_ids_used > 0);
                 ggml_tensor * b = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, w->ne[0], n_ids_used, 512);
                 ggml_tensor * ids = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_ids_used, 512);
@@ -965,7 +965,7 @@ static bool weight_buft_supported(const llama_hparams & hparams, ggml_tensor * w
             } break;
         case GGML_OP_ADD_ID:
             {
-                const int n_expert_used = hparams.n_expert_used();
+                const int n_expert_used = hparams.n_expert_used_max();
                 GGML_ASSERT(n_expert_used > 0);
                 ggml_tensor * a = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, w->ne[0], n_expert_used, 512);
                 ggml_tensor * c = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_expert_used, 512);

src/llama-model.cpp

@@ -1406,10 +1406,7 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
     }
 
     // models may route a different number of experts per layer, so validate the maximum
-    uint32_t n_expert_used_max = 0;
-    for (uint32_t il = 0; il < hparams.n_layer_all; ++il) {
-        n_expert_used_max = std::max(n_expert_used_max, hparams.n_expert_used(il));
-    }
+    uint32_t n_expert_used_max = hparams.n_expert_used_max();
 
     GGML_ASSERT(hparams.n_expert <= LLAMA_MAX_EXPERTS);
     GGML_ASSERT(n_expert_used_max <= hparams.n_expert);
@@ -1661,10 +1658,9 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         // TODO: move to a separate function
         const auto tn = LLM_TN(arch);
 
-        const int64_t n_expert      = hparams.n_expert;
-        const int64_t n_expert_used = hparams.n_expert_used();
+        const int64_t n_expert = hparams.n_expert;
 
-        if (n_expert > 0 && n_expert_used == 0) {
+        if (n_expert > 0 && hparams.n_expert_used_max() == 0) {
             throw std::runtime_error("model has expert layers but no expert layers are used");
         }
 

src/models/qwen4exp.cpp

@@ -744,6 +744,9 @@ ggml_tensor * llama_model_qwen4exp::graph::build_attn_qsa(
     ggml_tensor * k = mctx_cur->get_k(ctx0, il);
     ggml_tensor * v = mctx_cur->get_v(ctx0, il);
 
+    // TODO: enable sparse attention when we are ready
+    // ref: https://github.com/ggml-org/llama.cpp/pull/27970
+    //ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, top_k->ne[0], kq_scale, il);
     ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, 0, kq_scale, il);
     cb(cur, "kqv_out", il);
 

tools/mtmd/mtmd-audio.cpp

@@ -549,7 +549,7 @@ void mtmd_audio_preprocessor_whisper::initialize() {
 
 bool mtmd_audio_preprocessor_whisper::preprocess(const float *                 samples,
                                                  size_t                        n_samples,
-                                                 std::vector<mtmd_audio_mel> & output) {
+                                                 std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         // empty audio
         return false;
@@ -637,7 +637,7 @@ void mtmd_audio_preprocessor_qwen3a::initialize() {
 
 bool mtmd_audio_preprocessor_qwen3a::preprocess(const float *                 samples,
                                                  size_t                        n_samples,
-                                                 std::vector<mtmd_audio_mel> & output) {
+                                                 std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -739,7 +739,7 @@ void mtmd_audio_preprocessor_dots3note::initialize() {
 
 bool mtmd_audio_preprocessor_dots3note::preprocess(const float *                 samples,
                                                    size_t                        n_samples,
-                                                   std::vector<mtmd_audio_mel> & output) {
+                                                   std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -839,7 +839,7 @@ void mtmd_audio_preprocessor_mimo_audio::initialize() {
 
 bool mtmd_audio_preprocessor_mimo_audio::preprocess(const float *                 samples,
                                                     size_t                        n_samples,
-                                                    std::vector<mtmd_audio_mel> & output) {
+                                                    std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -898,7 +898,7 @@ void mtmd_audio_preprocessor_qwen3tts_spk::initialize() {
 
 bool mtmd_audio_preprocessor_qwen3tts_spk::preprocess(const float *                 samples,
                                                       size_t                        n_samples,
-                                                      std::vector<mtmd_audio_mel> & output) {
+                                                      std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -955,7 +955,7 @@ void mtmd_audio_preprocessor_conformer::initialize() {
 
 bool mtmd_audio_preprocessor_conformer::preprocess(const float *                 samples,
                                                    size_t                        n_samples,
-                                                   std::vector<mtmd_audio_mel> & output) {
+                                                   std::vector<mtmd_audio_mel> & output) const {
     // empty audio
     if (n_samples == 0) {
         return false;
@@ -1003,7 +1003,7 @@ void mtmd_audio_preprocessor_granite_speech::initialize() {
 
 bool mtmd_audio_preprocessor_granite_speech::preprocess(const float *                 samples,
                                                         size_t                        n_samples,
-                                                        std::vector<mtmd_audio_mel> & output) {
+                                                        std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -1117,7 +1117,7 @@ void mtmd_audio_preprocessor_gemma4a::initialize() {
 
 bool mtmd_audio_preprocessor_gemma4a::preprocess(const float *                 samples,
                                                   size_t                        n_samples,
-                                                  std::vector<mtmd_audio_mel> & output) {
+                                                  std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -1266,7 +1266,7 @@ void mtmd_audio_preprocessor_parakeet::initialize() {
 
 bool mtmd_audio_preprocessor_parakeet::preprocess(const float * samples,
                                                        size_t   n_samples_in,
-                                  std::vector<mtmd_audio_mel> & output) {
+                                  std::vector<mtmd_audio_mel> & output) const {
     if (n_samples_in == 0) {
         return false;
     }
@@ -1386,7 +1386,7 @@ void mtmd_audio_preprocessor_gemma4ua::initialize() {
 
 bool mtmd_audio_preprocessor_gemma4ua::preprocess(const float *                 samples,
                                                    size_t                        n_samples,
-                                                   std::vector<mtmd_audio_mel> & output) {
+                                                   std::vector<mtmd_audio_mel> & output) const {
     if (n_samples == 0) {
         return false;
     }
@@ -1527,7 +1527,7 @@ std::vector<float> mtmd_audio_streaming_istft::flush() {
 
 bool mtmd_audio_preprocessor_pockettts::preprocess(const float *                 samples,
                                                    size_t                        n_samples,
-                                                   std::vector<mtmd_audio_mel> & output) {
+                                                   std::vector<mtmd_audio_mel> & output) const {
     // the encoder needs whole frames, see pad_for_conv1d() in the reference
     const int64_t frame_size = (int64_t) hparams.mimi_downsample * 120;
     if (n_samples == 0 || frame_size <= 0) {

tools/mtmd/mtmd-audio.h

@@ -57,13 +57,13 @@ struct mtmd_audio_preprocessor {
 
     virtual ~mtmd_audio_preprocessor() = default;
     virtual void initialize() = 0; // NOT thread-safe
-    virtual bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) = 0;
+    virtual bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const = 0;
 };
 
 struct mtmd_audio_preprocessor_whisper : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_whisper(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -72,7 +72,7 @@ struct mtmd_audio_preprocessor_whisper : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_conformer : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_conformer(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -81,7 +81,7 @@ struct mtmd_audio_preprocessor_conformer : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_granite_speech : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_granite_speech(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -90,7 +90,7 @@ struct mtmd_audio_preprocessor_granite_speech : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_gemma4a : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_gemma4a(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -99,13 +99,13 @@ struct mtmd_audio_preprocessor_gemma4a : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_gemma4ua : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_gemma4ua(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 };
 
 struct mtmd_audio_preprocessor_qwen3a : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_qwen3a(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -114,7 +114,7 @@ struct mtmd_audio_preprocessor_qwen3a : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_dots3note : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_dots3note(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -123,7 +123,7 @@ struct mtmd_audio_preprocessor_dots3note : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_mimo_audio : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_mimo_audio(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -132,7 +132,7 @@ struct mtmd_audio_preprocessor_mimo_audio : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_qwen3tts_spk : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_qwen3tts_spk(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;
@@ -142,13 +142,13 @@ struct mtmd_audio_preprocessor_qwen3tts_spk : mtmd_audio_preprocessor {
 struct mtmd_audio_preprocessor_pockettts : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_pockettts(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
     void initialize() override {}
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 };
 
 struct mtmd_audio_preprocessor_parakeet : mtmd_audio_preprocessor {
     mtmd_audio_preprocessor_parakeet(clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) { }
     void initialize() override;
-    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
+    bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) const override;
 
   private:
     mtmd_audio_cache cache;

tools/mtmd/mtmd-helper.cpp

@@ -370,11 +370,11 @@ static bool is_webp_file(const unsigned char * buf, size_t len) {
 }
 
 #ifdef MTMD_VIDEO
-static mtmd_bitmap * decode_webp_with_ffmpeg(mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
+static mtmd_bitmap * decode_webp_with_ffmpeg(const mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
                                              const mtmd_helper_video_init_params & params);
 #endif
 
-mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(mtmd_context * ctx, const unsigned char * buf, size_t len, bool placeholder,
+mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(const mtmd_context * ctx, const unsigned char * buf, size_t len, bool placeholder,
                                                             mtmd_helper_init_opt opt) {
     // calculate the hash if needed
     std::string id;
@@ -460,7 +460,7 @@ mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(mtmd_context * ctx,
     return {nullptr, nullptr};
 }
 
-mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(mtmd_context * ctx, const char * fname, bool placeholder,
+mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(const mtmd_context * ctx, const char * fname, bool placeholder,
                                                              mtmd_helper_init_opt opt) {
 #ifdef _WIN32
     int wlen = MultiByteToWideChar(CP_UTF8, 0, fname, -1, NULL, 0);
@@ -505,7 +505,7 @@ mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(mtmd_context * ctx,
     return mtmd_helper_bitmap_init_from_buf(ctx, buf.data(), buf.size(), placeholder, opt);
 }
 
-bool mtmd_helper_support_video(mtmd_context * ctx) {
+bool mtmd_helper_support_video(const mtmd_context * ctx) {
 #ifdef MTMD_VIDEO
     return mtmd_support_vision(ctx);
 #else
@@ -521,7 +521,7 @@ bool mtmd_helper_support_video(mtmd_context * ctx) {
 #ifdef MTMD_VIDEO
 
 struct mtmd_helper_video {
-    mtmd_context * mctx;
+    const mtmd_context * mctx;
     std::string path;
     std::vector<uint8_t> input_buf; // non-empty when initialized from buffer
     std::string ffmpeg_bin;
@@ -887,7 +887,7 @@ static std::string video_resolve_bin(const char * bin_dir, const char * name) {
 }
 
 #ifdef MTMD_VIDEO
-static mtmd_bitmap * decode_webp_with_ffmpeg(mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
+static mtmd_bitmap * decode_webp_with_ffmpeg(const mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
                                              const mtmd_helper_video_init_params & params) {
     mtmd_helper_video vctx;
     vctx.mctx        = mctx;
@@ -914,7 +914,7 @@ static mtmd_bitmap * decode_webp_with_ffmpeg(mtmd_context * mctx, const unsigned
 #endif
 
 mtmd_helper_video * mtmd_helper_video_init(
-        mtmd_context * mctx,
+        const mtmd_context * mctx,
         const char * path,
         mtmd_helper_video_init_params params) {
 #ifdef MTMD_VIDEO
@@ -949,7 +949,7 @@ mtmd_helper_video * mtmd_helper_video_init(
 }
 
 mtmd_helper_video * mtmd_helper_video_init_from_buf(
-        mtmd_context * mctx,
+        const mtmd_context * mctx,
         const unsigned char * buf, size_t len,
         mtmd_helper_video_init_params params) {
 #ifdef MTMD_VIDEO
@@ -1017,7 +1017,7 @@ int32_t mtmd_helper_video_read_next(mtmd_helper_video * ctx,
 #endif
 }
 
-bool mtmd_helper_model_can_chat(llama_context * lctx, mtmd_context * mctx) {
+bool mtmd_helper_model_can_chat(const llama_context * lctx, const mtmd_context * mctx) {
     if (!mctx) {
         return true;
     }

tools/mtmd/mtmd-helper.h

@@ -46,7 +46,7 @@ MTMD_API struct mtmd_helper_init_opt mtmd_helper_init_opt_default(void);
 MTMD_API void mtmd_helper_log_set(ggml_log_callback log_callback, void * user_data);
 
 // Returns true if this build includes video support (MTMD_VIDEO was ON at compile time).
-MTMD_API bool mtmd_helper_support_video(mtmd_context * ctx);
+MTMD_API bool mtmd_helper_support_video(const mtmd_context * ctx);
 
 struct mtmd_helper_bitmap_wrapper {
     mtmd_bitmap * bitmap;
@@ -58,7 +58,7 @@ struct mtmd_helper_bitmap_wrapper {
 // returns nullptr on failure
 // this function is thread-safe
 MTMD_API struct mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(
-                    mtmd_context * ctx,
+                    const mtmd_context * ctx,
                     const char * fname,
                     bool placeholder,
                     struct mtmd_helper_init_opt opt);
@@ -75,7 +75,7 @@ MTMD_API struct mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(
 // returns nullptr on failure
 // this function is thread-safe
 MTMD_API struct mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(
-                    mtmd_context * ctx,
+                    const mtmd_context * ctx,
                     const unsigned char * buf, size_t len,
                     bool placeholder,
                     struct mtmd_helper_init_opt opt);
@@ -153,7 +153,7 @@ struct mtmd_helper_video_info {
 
 // returns NULL on failure (ffprobe not found, file unreadable, etc.)
 MTMD_API mtmd_helper_video * mtmd_helper_video_init(
-                    struct mtmd_context * mctx,
+                    const struct mtmd_context * mctx,
                     const char * path,
                     struct mtmd_helper_video_init_params params);
 
@@ -162,7 +162,7 @@ MTMD_API mtmd_helper_video * mtmd_helper_video_init(
 // Note: pipe input is not seekable, so seeking will use output-side seeking
 // (ffmpeg decodes and discards frames up to the target position).
 MTMD_API mtmd_helper_video * mtmd_helper_video_init_from_buf(
-                    struct mtmd_context * mctx,
+                    const struct mtmd_context * mctx,
                     const unsigned char * buf, size_t len,
                     struct mtmd_helper_video_init_params params);
 MTMD_API void mtmd_helper_video_free(mtmd_helper_video * ctx);
@@ -177,7 +177,7 @@ MTMD_API int32_t mtmd_helper_video_read_next(mtmd_helper_video * ctx,
             char ** out_text);
 
 // return true if model can be used for chat
-MTMD_API bool mtmd_helper_model_can_chat(struct llama_context * lctx, struct mtmd_context * mctx);
+MTMD_API bool mtmd_helper_model_can_chat(const struct llama_context * lctx, const struct mtmd_context * mctx);
 
 //
 // Audio generation helpers

tools/mtmd/mtmd-image.cpp

@@ -485,7 +485,7 @@ private:
 // mtmd_image_preprocessor_llava_uhd
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_llava_uhd::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_llava_uhd::preprocess(const clip_image_u8 & img) const {
     const clip_image_size original_size = img.get_size();
     auto const inst = get_slice_instructions(original_size);
     auto sliced = slice_image(img, inst);
@@ -499,7 +499,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_llava_uhd::preprocess(const clip_
     return output;
 }
 
-mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_llava_uhd::get_slice_instructions(const clip_image_size & original_size) {
+mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_llava_uhd::get_slice_instructions(const clip_image_size & original_size) const {
     mtmd_image_preprocessor_llava_uhd::slice_instructions res;
     // align slices by patch_size * n_merge so an integer number of merger output tokens fits per slice
     const int n_merge         = hparams.n_merge;
@@ -604,7 +604,7 @@ mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_ll
     return res;
 }
 
-mtmd_image_preprocessor_llava_uhd::slice_output mtmd_image_preprocessor_llava_uhd::slice_image(const clip_image_u8 & img, const mtmd_image_preprocessor_llava_uhd::slice_instructions & inst) {
+mtmd_image_preprocessor_llava_uhd::slice_output mtmd_image_preprocessor_llava_uhd::slice_image(const clip_image_u8 & img, const mtmd_image_preprocessor_llava_uhd::slice_instructions & inst) const {
     slice_output output;
 
     // resize to overview size
@@ -636,7 +636,7 @@ mtmd_image_preprocessor_llava_uhd::slice_output mtmd_image_preprocessor_llava_uh
     return output;
 }
 
-clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_resize(const clip_image_size & original_size, int scale_resolution, int patch_size, bool allow_upscale) {
+clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_resize(const clip_image_size & original_size, int scale_resolution, int patch_size, bool allow_upscale) const {
     int width  = original_size.width;
     int height = original_size.height;
     if ((width * height > scale_resolution * scale_resolution) || allow_upscale) {
@@ -650,7 +650,7 @@ clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_resize(const clip_im
     return res;
 }
 
-clip_image_size mtmd_image_preprocessor_llava_uhd::resize_maintain_aspect_ratio(const clip_image_size & orig, const clip_image_size & target_max) {
+clip_image_size mtmd_image_preprocessor_llava_uhd::resize_maintain_aspect_ratio(const clip_image_size & orig, const clip_image_size & target_max) const {
     float scale_width  = static_cast<float>(target_max.width)  / orig.width;
     float scale_height = static_cast<float>(target_max.height) / orig.height;
     float scale = std::min(scale_width, scale_height);
@@ -660,7 +660,7 @@ clip_image_size mtmd_image_preprocessor_llava_uhd::resize_maintain_aspect_ratio(
     };
 }
 
-clip_image_size mtmd_image_preprocessor_llava_uhd::select_best_resolution(const clip_image_size & original_size, const std::vector<clip_image_size> & possible_resolutions) {
+clip_image_size mtmd_image_preprocessor_llava_uhd::select_best_resolution(const clip_image_size & original_size, const std::vector<clip_image_size> & possible_resolutions) const {
     clip_image_size best_fit;
     int min_wasted_area = std::numeric_limits<int>::max();
     int max_effective_resolution = 0;
@@ -684,11 +684,11 @@ clip_image_size mtmd_image_preprocessor_llava_uhd::select_best_resolution(const
     return best_fit;
 }
 
-int mtmd_image_preprocessor_llava_uhd::ensure_divide(int length, int patch_size) {
+int mtmd_image_preprocessor_llava_uhd::ensure_divide(int length, int patch_size) const {
     return std::max(static_cast<int>(std::round(static_cast<float>(length) / patch_size) * patch_size), patch_size);
 }
 
-clip_image_size mtmd_image_preprocessor_llava_uhd::get_refine_size(const clip_image_size & original_size, const clip_image_size & grid, int scale_resolution, int patch_size, bool allow_upscale) {
+clip_image_size mtmd_image_preprocessor_llava_uhd::get_refine_size(const clip_image_size & original_size, const clip_image_size & grid, int scale_resolution, int patch_size, bool allow_upscale) const {
     int width  = original_size.width;
     int height = original_size.height;
     int grid_x = grid.width;
@@ -711,7 +711,7 @@ clip_image_size mtmd_image_preprocessor_llava_uhd::get_refine_size(const clip_im
     return refine_size;
 }
 
-clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_grid(const int max_slice_nums, const int multiple, const float log_ratio) {
+clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_grid(const int max_slice_nums, const int multiple, const float log_ratio) const {
     std::vector<int> candidate_split_grids_nums;
     for (int i : {multiple - 1, multiple, multiple + 1}) {
         if (i == 1 || i > max_slice_nums) {
@@ -747,7 +747,7 @@ clip_image_size mtmd_image_preprocessor_llava_uhd::get_best_grid(const int max_s
 // mtmd_image_preprocessor_fixed_size
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_fixed_size::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_fixed_size::preprocess(const clip_image_u8 & img) const {
     clip_image_u8 resized_image;
     int sz = hparams.image_size;
     img_tool::resize(img, resized_image, {sz, sz},
@@ -763,7 +763,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_fixed_size::preprocess(const clip
 // mtmd_image_preprocessor_dyn_size
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_dyn_size::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_dyn_size::preprocess(const clip_image_u8 & img) const {
     GGML_ASSERT(hparams.image_min_pixels > 0 && hparams.image_max_pixels > 0);
     clip_image_u8 resized_image;
     const clip_image_size original_size = img.get_size();
@@ -790,7 +790,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_dyn_size::preprocess(const clip_i
 // mtmd_image_preprocessor_longest_edge
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_longest_edge::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_longest_edge::preprocess(const clip_image_u8 & img) const {
     GGML_ASSERT(hparams.image_longest_edge > 0);
     clip_image_u8 resized_image;
     const clip_image_size original_size = img.get_size();
@@ -817,7 +817,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_longest_edge::preprocess(const cl
 // mtmd_image_preprocessor_minicpmv
 //
 
-mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_minicpmv::get_slice_instructions(const clip_image_size & original_size) {
+mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_minicpmv::get_slice_instructions(const clip_image_size & original_size) const {
     if (hparams.n_merge == 2) {
         const int   slice_size = hparams.image_size;
         const float ratio      = (float)original_size.width * original_size.height / (slice_size * slice_size);
@@ -837,7 +837,7 @@ mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_mi
 // mtmd_image_preprocessor_lfm2
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_lfm2::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_lfm2::preprocess(const clip_image_u8 & img) const {
     auto const inst = get_slice_instructions(img.get_size());
     if (!inst.slices.empty()) {
         return mtmd_image_preprocessor_llava_uhd::preprocess(img);
@@ -868,7 +868,7 @@ bool mtmd_image_preprocessor_lfm2::should_tile(
            static_cast<double>(hparams.image_max_pixels) * max_pixels_tolerance;
 }
 
-mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_lfm2::get_slice_instructions(const clip_image_size & original_size) {
+mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_lfm2::get_slice_instructions(const clip_image_size & original_size) const {
     mtmd_image_preprocessor_llava_uhd::slice_instructions inst;
     const int align_size = hparams.patch_size * hparams.n_merge;
     inst.overview_size = img_tool::calc_size_preserved_ratio(
@@ -914,7 +914,7 @@ mtmd_image_preprocessor_llava_uhd::slice_instructions mtmd_image_preprocessor_lf
 clip_image_size mtmd_image_preprocessor_lfm2::find_closest_aspect_ratio(
         float aspect_ratio,
         const std::vector<clip_image_size> & target_ratios,
-        int width, int height) {
+        int width, int height) const {
     float best_ratio_diff = std::numeric_limits<float>::max();
     clip_image_size best_ratio = {1, 1};
     const float area = static_cast<float>(width * height);
@@ -935,7 +935,7 @@ clip_image_size mtmd_image_preprocessor_lfm2::find_closest_aspect_ratio(
     return best_ratio;
 }
 
-std::vector<clip_image_size> mtmd_image_preprocessor_lfm2::get_target_ratios() {
+std::vector<clip_image_size> mtmd_image_preprocessor_lfm2::get_target_ratios() const {
     std::vector<clip_image_size> ratios;
     for (int n = min_tiles; n <= max_tiles; n++) {
         for (int w = 1; w <= n; w++) {
@@ -961,7 +961,7 @@ std::vector<clip_image_size> mtmd_image_preprocessor_lfm2::get_target_ratios() {
     return ratios;
 }
 
-clip_image_size mtmd_image_preprocessor_lfm2::get_grid_layout(int height, int width) {
+clip_image_size mtmd_image_preprocessor_lfm2::get_grid_layout(int height, int width) const {
     const float aspect_ratio = static_cast<float>(width) / height;
     const auto ratios = get_target_ratios();
     return find_closest_aspect_ratio(aspect_ratio, ratios, width, height);
@@ -971,7 +971,7 @@ clip_image_size mtmd_image_preprocessor_lfm2::get_grid_layout(int height, int wi
 // mtmd_image_preprocessor_idefics3
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_idefics3::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_idefics3::preprocess(const clip_image_u8 & img) const {
     // The refined size has two steps:
     // 1. Resize w/ aspect-ratio preserving such that the longer side is
     //      the preprocessor longest size
@@ -1071,7 +1071,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_idefics3::preprocess(const clip_i
 // mtmd_image_preprocessor_internvl
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_internvl::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_internvl::preprocess(const clip_image_u8 & img) const {
     GGML_ASSERT(!hparams.image_res_candidates.empty());
     const clip_image_size original_size = img.get_size();
     auto const inst = get_slice_instructions(original_size);
@@ -1206,7 +1206,7 @@ void mtmd_image_preprocessor_deepseek4v::safe_resize(int height, int width, int
 }
 
 // ref: load_image()
-mtmd_image_preproc_out mtmd_image_preprocessor_deepseek4v::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_deepseek4v::preprocess(const clip_image_u8 & img) const {
     mtmd_image_preproc_out out;
 
     const int p           = hparams.patch_size;
@@ -1244,7 +1244,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_deepseek4v::preprocess(const clip
     return out;
 }
 
-mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const clip_image_u8 & img) const {
     mtmd_image_preproc_out output;
     int grid_w = 0;
     int grid_h = 0;
@@ -1320,7 +1320,7 @@ void mtmd_image_preprocessor_step3vl::img_u8_resize_bilinear_to_f32(
         int target_width,
         int target_height,
         const float mean[3],
-        const float std[3]) {
+        const float std[3]) const {
     const auto src_size = src.get_size();
     if (src_size.width == target_width && src_size.height == target_height) {
         dst.from_u8(src);
@@ -1519,7 +1519,7 @@ mtmd_image_preprocessor_step3vl::slice_instructions mtmd_image_preprocessor_step
     return instructions;
 }
 
-mtmd_image_preproc_out mtmd_image_preprocessor_step3vl::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_step3vl::preprocess(const clip_image_u8 & img) const {
     clip_image_u8 prepared = prepare_image(img, hparams);
     const auto instructions = build_slice_instructions(hparams, prepared.get_size());
 
@@ -1573,7 +1573,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_step3vl::preprocess(const clip_im
 // mtmd_image_preprocessor_youtuvl
 //
 
-mtmd_image_preproc_out mtmd_image_preprocessor_youtuvl::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_youtuvl::preprocess(const clip_image_u8 & img) const {
     const int patch_size = hparams.patch_size;   // typically 16
     const int merge_size = hparams.n_merge;      // typically 2
     const int align_size = patch_size * merge_size;  // 32
@@ -1622,7 +1622,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_youtuvl::preprocess(const clip_im
     return output;
 }
 
-mtmd_image_preproc_out mtmd_image_preprocessor_granite::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_granite::preprocess(const clip_image_u8 & img) const {
     GGML_ASSERT(!hparams.image_res_candidates.empty());
 
     const clip_image_size orig_size = img.get_size();
@@ -1717,7 +1717,7 @@ static clip_image_size muse_glimmer_grid_size(int img_w, int img_h, int patch_hw
     return clip_image_size{ best_npw * patch_hw, best_nph * patch_hw };
 }
 
-mtmd_image_preproc_out mtmd_image_preprocessor_muse_glimmer::preprocess(const clip_image_u8 & img) {
+mtmd_image_preproc_out mtmd_image_preprocessor_muse_glimmer::preprocess(const clip_image_u8 & img) const {
     const int patch_hw   = hparams.patch_size * hparams.n_merge;
     const int patch_area = hparams.patch_size * hparams.patch_size * hparams.n_merge * hparams.n_merge;
     GGML_ASSERT(patch_area > 0 && hparams.image_max_pixels > 0);

tools/mtmd/mtmd-image.h

@@ -33,7 +33,7 @@ struct mtmd_image_preprocessor {
     mtmd_image_preprocessor(const clip_ctx * ctx): hparams(*clip_get_hparams(ctx)) {}
 
     virtual ~mtmd_image_preprocessor() = default;
-    virtual mtmd_image_preproc_out preprocess(const clip_image_u8 & img) = 0;
+    virtual mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const = 0;
 };
 
 /**
@@ -59,7 +59,7 @@ struct mtmd_image_preprocessor {
  */
 struct mtmd_image_preprocessor_llava_uhd : mtmd_image_preprocessor {
     mtmd_image_preprocessor_llava_uhd(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 
     struct slice_coordinates {
         int x;
@@ -74,16 +74,16 @@ struct mtmd_image_preprocessor_llava_uhd : mtmd_image_preprocessor {
         std::vector<slice_coordinates> slices;
     };
 
-    virtual slice_instructions get_slice_instructions(const clip_image_size & original_size);
+    virtual slice_instructions get_slice_instructions(const clip_image_size & original_size) const;
 
     struct slice_output {
         clip_image_u8 overview;
         std::vector<clip_image_u8> slices;
     };
-    slice_output slice_image(const clip_image_u8 & img, const slice_instructions & inst);
+    slice_output slice_image(const clip_image_u8 & img, const slice_instructions & inst) const;
 
 protected:
-    clip_image_size get_best_resize(const clip_image_size & original_size, int scale_resolution, int patch_size, bool allow_upscale = false);
+    clip_image_size get_best_resize(const clip_image_size & original_size, int scale_resolution, int patch_size, bool allow_upscale = false) const;
 
     /**
      * Selects the best resolution from a list of possible resolutions based on the original size.
@@ -100,19 +100,19 @@ protected:
      * @param possible_resolutions A list of possible resolutions
      * @return The best fit resolution
      */
-    clip_image_size select_best_resolution(const clip_image_size & original_size, const std::vector<clip_image_size> & possible_resolutions);
+    clip_image_size select_best_resolution(const clip_image_size & original_size, const std::vector<clip_image_size> & possible_resolutions) const;
 
 private:
-    clip_image_size resize_maintain_aspect_ratio(const clip_image_size & orig, const clip_image_size & target_max);
-    int ensure_divide(int length, int patch_size);
-    clip_image_size get_refine_size(const clip_image_size & original_size, const clip_image_size & grid, int scale_resolution, int patch_size, bool allow_upscale = false);
-    clip_image_size get_best_grid(const int max_slice_nums, const int multiple, const float log_ratio);
+    clip_image_size resize_maintain_aspect_ratio(const clip_image_size & orig, const clip_image_size & target_max) const;
+    int ensure_divide(int length, int patch_size) const;
+    clip_image_size get_refine_size(const clip_image_size & original_size, const clip_image_size & grid, int scale_resolution, int patch_size, bool allow_upscale = false) const;
+    clip_image_size get_best_grid(const int max_slice_nums, const int multiple, const float log_ratio) const;
 };
 
 // downscale or upscale the input image to fixed size
 struct mtmd_image_preprocessor_fixed_size : mtmd_image_preprocessor {
     mtmd_image_preprocessor_fixed_size(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // resize image to multiple of patch_size*n_merge, while preserving aspect ratio
@@ -120,19 +120,19 @@ struct mtmd_image_preprocessor_fixed_size : mtmd_image_preprocessor {
 // this is used by models with native support for dynamic image size, for example: Qwen-VL, Pixtral, Kimi-VL, etc
 struct mtmd_image_preprocessor_dyn_size : mtmd_image_preprocessor {
     mtmd_image_preprocessor_dyn_size(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // similar to mtmd_image_preprocessor_dyn_size, but resize the image to have longest edge equal to hparams.image_longest_edge, while preserving aspect ratio
 struct mtmd_image_preprocessor_longest_edge : mtmd_image_preprocessor {
     mtmd_image_preprocessor_longest_edge(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // ref: inference/image_processor.py in the HF repo (DeepSeek-V4-Flash-Vision)
 struct mtmd_image_preprocessor_deepseek4v : mtmd_image_preprocessor {
     mtmd_image_preprocessor_deepseek4v(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 
 private:
     struct grid_info {
@@ -148,7 +148,7 @@ private:
 // custom llava-uhd slicing logic for MiniCPM-V
 struct mtmd_image_preprocessor_minicpmv : mtmd_image_preprocessor_llava_uhd {
     using mtmd_image_preprocessor_llava_uhd::mtmd_image_preprocessor_llava_uhd;
-    slice_instructions get_slice_instructions(const clip_image_size & original_size) override;
+    slice_instructions get_slice_instructions(const clip_image_size & original_size) const override;
 };
 
 // custom llava-uhd slicing logic for LFM2
@@ -161,8 +161,8 @@ struct mtmd_image_preprocessor_lfm2 : mtmd_image_preprocessor_llava_uhd {
     static constexpr int   tile_size            = 512;
 
     using mtmd_image_preprocessor_llava_uhd::mtmd_image_preprocessor_llava_uhd;
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
-    slice_instructions get_slice_instructions(const clip_image_size & original_size) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
+    slice_instructions get_slice_instructions(const clip_image_size & original_size) const override;
 
     static bool should_tile(const clip_hparams & hparams, const clip_image_size & original_size);
 
@@ -170,19 +170,19 @@ private:
     clip_image_size find_closest_aspect_ratio(
             float aspect_ratio,
             const std::vector<clip_image_size> & target_ratios,
-            int width, int height);
-    std::vector<clip_image_size> get_target_ratios();
-    clip_image_size get_grid_layout(int height, int width);
+            int width, int height) const;
+    std::vector<clip_image_size> get_target_ratios() const;
+    clip_image_size get_grid_layout(int height, int width) const;
 };
 
 struct mtmd_image_preprocessor_idefics3 : mtmd_image_preprocessor_llava_uhd {
     mtmd_image_preprocessor_idefics3(const clip_ctx * ctx) : mtmd_image_preprocessor_llava_uhd(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 struct mtmd_image_preprocessor_internvl : mtmd_image_preprocessor_llava_uhd {
     mtmd_image_preprocessor_internvl(const clip_ctx * ctx) : mtmd_image_preprocessor_llava_uhd(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // DeepSeek-OCR (v1/v2) global view + optional local tile grid
@@ -194,7 +194,7 @@ struct mtmd_image_preprocessor_deepseekocr : mtmd_image_preprocessor {
           tile_size(hparams.preproc_tile_size),
           min_tiles(hparams.preproc_min_tiles),
           max_tiles(hparams.preproc_max_tiles) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 
 private:
     bool fuse_row; // v1 fuses a tile-row into one image; v2 keeps tiles separate
@@ -214,7 +214,7 @@ private:
 // ref: https://huggingface.co/stepfun-ai/Step3-VL-10B/blob/main/processing_step3.py
 struct mtmd_image_preprocessor_step3vl : mtmd_image_preprocessor_llava_uhd {
     mtmd_image_preprocessor_step3vl(const clip_ctx * ctx) : mtmd_image_preprocessor_llava_uhd(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
     static slice_instructions build_slice_instructions(const clip_hparams & params, const clip_image_size & prepared_size);
 
 private:
@@ -230,7 +230,7 @@ private:
             int target_width,
             int target_height,
             const float mean[3],
-            const float std[3]);
+            const float std[3]) const;
     static int get_image_longest_edge(const clip_hparams & params);
     static int determine_window_size(const clip_hparams & params, int longer, int shorter);
     static int calc_crop_extent(int length, int window_size);
@@ -241,17 +241,17 @@ private:
 
 struct mtmd_image_preprocessor_youtuvl : mtmd_image_preprocessor {
     mtmd_image_preprocessor_youtuvl(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // llava-next "anyres": stacks the overview and all tiles into one image, assembled by clip in a single graph
 struct mtmd_image_preprocessor_granite : mtmd_image_preprocessor_llava_uhd {
     mtmd_image_preprocessor_granite(const clip_ctx * ctx) : mtmd_image_preprocessor_llava_uhd(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };
 
 // pick the patch grid closest to the input aspect ratio under the per-image token cap, stretch-resize.
 struct mtmd_image_preprocessor_muse_glimmer : mtmd_image_preprocessor {
     mtmd_image_preprocessor_muse_glimmer(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
-    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) override;
+    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
 };

tools/mtmd/mtmd.cpp

@@ -1117,7 +1117,7 @@ std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::
 }
 
 struct mtmd_tokenizer {
-    mtmd_context * ctx;
+    const mtmd_context * ctx;
 
     std::string input_text; // note: can contain null bytes; do not use c_str()
     bool add_special;
@@ -1140,9 +1140,9 @@ struct mtmd_tokenizer {
         }
     }
 
-    mtmd_tokenizer(mtmd_context * ctx,
+    mtmd_tokenizer(const mtmd_context * ctx,
             const mtmd_input_text * text,
-            const mtmd_bitmap ** bmps,
+            const mtmd_bitmap * const * bmps,
             size_t n_bitmaps) : ctx(ctx) {
         add_special   = text->add_special;
         parse_special = text->parse_special;
@@ -1177,8 +1177,8 @@ struct mtmd_tokenizer {
         expand_lazy_bitmaps();
     }
 
-    mtmd_tokenizer(mtmd_context * ctx,
-            const mtmd_input_part ** input_parts,
+    mtmd_tokenizer(const mtmd_context * ctx,
+            const mtmd_input_part * const * input_parts,
             size_t n_parts,
             bool add_special) : ctx(ctx) {
         this->add_special = add_special;
@@ -1733,10 +1733,10 @@ struct mtmd_tokenizer {
     }
 };
 
-int32_t mtmd_tokenize(mtmd_context * ctx,
+int32_t mtmd_tokenize(const mtmd_context * ctx,
             mtmd_input_chunks * output,
             const mtmd_input_text * text,
-            const mtmd_bitmap ** bitmaps,
+            const mtmd_bitmap * const * bitmaps,
             size_t n_bitmaps) {
     try {
         mtmd_tokenizer tokenizer(ctx, text, bitmaps, n_bitmaps);
@@ -1747,9 +1747,9 @@ int32_t mtmd_tokenize(mtmd_context * ctx,
     }
 }
 
-int32_t mtmd_tokenize_from_parts(mtmd_context * ctx,
+int32_t mtmd_tokenize_from_parts(const mtmd_context * ctx,
             mtmd_input_chunks * output,
-            const mtmd_input_part ** parts,
+            const mtmd_input_part * const * parts,
             size_t n_parts,
             bool add_special) {
     for (size_t i = 0; i < n_parts; i++) {
@@ -2271,7 +2271,7 @@ void mtmd_bitmap_set_mergeable(mtmd_bitmap * bitmap, bool mergeable) {
     bitmap->mergeable = mergeable;
 }
 
-mtmd_bitmap * mtmd_bitmap_init_lazy(mtmd_context * ctx,
+mtmd_bitmap * mtmd_bitmap_init_lazy(const mtmd_context * ctx,
                                     const char * id,
                                     void * user_data,
                                     mtmd_bitmap_lazy_callback callback) {

tools/mtmd/mtmd.h

@@ -211,7 +211,7 @@ typedef int(* mtmd_bitmap_lazy_callback)(
     mtmd_bitmap ** out_bitmap,
     char ** out_text);
 
-MTMD_API mtmd_bitmap * mtmd_bitmap_init_lazy(mtmd_context * ctx,
+MTMD_API mtmd_bitmap * mtmd_bitmap_init_lazy(const mtmd_context * ctx,
                                              const char * id, // usually set to file hash
                                              void * user_data,
                                              mtmd_bitmap_lazy_callback callback);
@@ -299,10 +299,10 @@ MTMD_API struct mtmd_decoder_pos mtmd_image_tokens_get_decoder_pos(const mtmd_im
 //   0 on success
 //   1 on number of bitmaps not matching the number of markers
 //   2 on media preprocessing error
-MTMD_API int32_t mtmd_tokenize(mtmd_context * ctx,
+MTMD_API int32_t mtmd_tokenize(const mtmd_context * ctx,
                                mtmd_input_chunks * output,
                                const mtmd_input_text * text,
-                               const mtmd_bitmap ** bitmaps,
+                               const mtmd_bitmap * const * bitmaps,
                                size_t n_bitmaps);
 
 // same as mtmd_tokenize(), but takes an array of mtmd_input_part
@@ -311,9 +311,9 @@ MTMD_API int32_t mtmd_tokenize(mtmd_context * ctx,
 // - when you want to control parse_special for each text part
 // note: per-part add_special will be ignored
 // return 1 if a part has both text and bitmap set (or neither)
-MTMD_API int32_t mtmd_tokenize_from_parts(mtmd_context * ctx,
+MTMD_API int32_t mtmd_tokenize_from_parts(const mtmd_context * ctx,
                                           mtmd_input_chunks * output,
-                                          const mtmd_input_part ** parts,
+                                          const mtmd_input_part * const * parts,
                                           size_t n_parts,
                                           bool add_special);
 

tools/server/tests/conftest.py

@@ -1,7 +1,17 @@
+import os
 import pytest
+from filelock import FileLock
 from utils import *
 
 
+@pytest.fixture(scope="session", autouse=True)
+def configure_worker_port(request):
+    worker_id = getattr(request.config, "workerinput", {}).get("workerid", "master")
+    if worker_id != "master":
+        worker_num = int(worker_id[2:])
+        os.environ["PORT"] = str(8080 + worker_num * 10)
+
+
 # ref: https://stackoverflow.com/questions/22627659/run-code-before-and-after-each-test-in-py-test
 @pytest.fixture(autouse=True)
 def stop_server_after_each_test():
@@ -16,6 +26,10 @@ def stop_server_after_each_test():
 
 
 @pytest.fixture(scope="session", autouse=True)
-def load_server_presets():
+def load_server_presets(configure_worker_port, tmp_path_factory):
     # this will be run once per test session, before any tests
-    ServerPreset.load_all()
+
+    # serialize model downloads across parallel workers.
+    root_tmp_dir = tmp_path_factory.getbasetemp().parent
+    with FileLock(str(root_tmp_dir / "load_all.lock")):
+        ServerPreset.load_all()

tools/server/tests/requirements.txt

@@ -1,5 +1,7 @@
 aiohttp~=3.9.3
 pytest~=8.3.3
+pytest-xdist~=3.6
+filelock~=3.16
 numpy~=1.26.4
 openai~=2.14.0
 prometheus-client~=0.20.0

tools/server/tests/tests.sh

@@ -6,13 +6,15 @@ cd $SCRIPT_DIR
 
 set -eu
 
+WORKERS="${PYTEST_WORKERS:-auto}"
+
 if [ $# -lt 1 ]
 then
     if [[ "${SLOW_TESTS:-0}" == 1 ]]; then
-        pytest --durations=30 -v -x
+        pytest --durations=30 -v -x -n "${WORKERS}" --dist=worksteal
     else
-        pytest --durations=30 -v -x -m "not slow"
+        pytest --durations=30 -v -x -n "${WORKERS}" --dist=worksteal -m "not slow"
     fi
 else
-    pytest --durations=30 "$@"
+    pytest --durations=30 -n "${WORKERS}" --dist=worksteal "$@"
 fi

tools/server/tests/unit/test_compat_anthropic.py

@@ -21,7 +21,6 @@ def create_server():
     global server
     server = ServerPreset.tinyllama2()
     server.model_alias = "tinyllama-2-anthropic"
-    server.server_port = 8082
     server.n_slots = 1
     server.n_ctx = 8192
     server.n_batch = 2048
@@ -34,7 +33,6 @@ def vision_server():
     server = ServerPreset.tinygemma3()
     server.offline = False  # Allow downloading the model
     server.model_alias = "tinygemma3-anthropic"
-    server.server_port = 8083  # Different port to avoid conflicts
     server.n_slots = 1
     return server
 
@@ -1015,7 +1013,6 @@ def test_anthropic_thinking_with_reasoning_model(stream):
     server.jinja = True
     server.n_ctx = 8192
     server.n_predict = 1024
-    server.server_port = 8084
     server.start(timeout_seconds=600)  # large model needs time to download
 
     if stream:

tools/server/tests/unit/test_mcp_servers.py

@@ -37,7 +37,6 @@ def _start_server_with_mcp(mcp_json: str, **kwargs) -> ServerProcess:
     srv = ServerPreset.router()
     srv.server_tools = "all"
     srv.no_ui = True
-    srv.server_port = 8085  # avoid conflict with load_all() which uses 8080
     srv.mcp_servers_json = mcp_json
     for k, v in kwargs.items():
         setattr(srv, k, v)
@@ -183,7 +182,6 @@ def test_mcp_tools_not_listed_when_not_configured():
     server = ServerPreset.router()
     server.server_tools = "all"
     server.no_ui = True
-    server.server_port = 8085
     server.start()
 
     try:
@@ -250,7 +248,6 @@ def test_mcp_tools_via_json_config_file():
         server = ServerPreset.router()
         server.server_tools = "all"
         server.no_ui = True
-        server.server_port = 8085
         server.mcp_servers_config = config_path
         server.start()
 
@@ -468,7 +465,6 @@ def test_mcp_config_file_errors():
     server = ServerPreset.router()
     server.server_tools = "all"
     server.no_ui = True
-    server.server_port = 8085
     server.mcp_servers_json = "not valid json"
     try:
         server.start()
@@ -480,7 +476,6 @@ def test_mcp_config_file_errors():
     server = ServerPreset.router()
     server.server_tools = "all"
     server.no_ui = True
-    server.server_port = 8085
     server.mcp_servers_config = "/nonexistent/path.json"
     try:
         server.start()

tools/server/tests/unit/test_slot_save.py

@@ -10,10 +10,10 @@ STATE_FILE_HEADER_SIZE = 12
 server = ServerPreset.tinyllama2()
 
 @pytest.fixture(autouse=True)
-def create_server():
+def create_server(tmp_path):
     global server
     server = ServerPreset.tinyllama2()
-    server.slot_save_path = "./tmp"
+    server.slot_save_path = str(tmp_path)
     server.temperature = 0.0
 
 
@@ -94,7 +94,7 @@ def test_slot_restore_legacy_token_list():
     assert res.body["n_saved"] == 84
 
     # rewrite the token payload into a plain token list, as written by servers that predate the packed server_tokens format
-    path = os.path.join("tmp", "slot_legacy.bin")
+    path = os.path.join(server.slot_save_path, "slot_legacy.bin")
     with open(path, "rb") as f:
         data = bytearray(f.read())
 
@@ -462,7 +462,7 @@ def test_slot_save_restore_image_payload_larger_than_context(mmproj_server):
     })
     assert res.status_code == 200
 
-    path = os.path.join("tmp", "mm_slot_large_payload.bin")
+    path = os.path.join(server.slot_save_path, "mm_slot_large_payload.bin")
     with open(path, "rb") as f:
         data = bytearray(f.read())
     payload_size = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE - 4)[0]

tools/server/tests/unit/test_tool_call.py

@@ -21,7 +21,6 @@ def create_server():
     global server
     server = ServerPreset.tinyllama2()
     server.model_alias = "tinyllama-2-tool-call"
-    server.server_port = 8081
     server.n_slots = 1
     server.n_ctx = 8192
     server.n_batch = 2048

tools/server/tests/unit/test_tools_builtin.py

@@ -64,11 +64,11 @@ def test_tools_builtin_read_file():
     assert "def test_tools_builtin_read_file" in text
 
 
-def test_tools_builtin_write_then_edit_file():
+def test_tools_builtin_write_then_edit_file(tmp_path):
     global server
     server.start()
 
-    log_path = os.path.join(PROJECT_ROOT, "test.log")
+    log_path = str(tmp_path / "test.log")
     try:
         write_res = call_tool("write_file", {"path": log_path, "content": "line1\nline2\nline3\n"})
         assert write_res["result"] == "file written successfully"
@@ -93,11 +93,11 @@ def test_tools_builtin_write_then_edit_file():
             os.remove(log_path)
 
 
-def test_tools_builtin_edit_file_rejects_non_unique_old_text():
+def test_tools_builtin_edit_file_rejects_non_unique_old_text(tmp_path):
     global server
     server.start()
 
-    log_path = os.path.join(PROJECT_ROOT, "test.log")
+    log_path = str(tmp_path / "test.log")
     try:
         call_tool("write_file", {"path": log_path, "content": "dup\ndup\n"})
         err = call_tool_expect_error("edit_file", {
@@ -275,11 +275,11 @@ def test_tools_builtin_docker_runtime_cleans_up_spawned_container():
     assert leftover.returncode != 0, f"container {container_id} was not cleaned up after server exit"
 
 
-def test_tools_builtin_edit_file_rejects_overlapping_edits():
+def test_tools_builtin_edit_file_rejects_overlapping_edits(tmp_path):
     global server
     server.start()
 
-    log_path = os.path.join(PROJECT_ROOT, "test.log")
+    log_path = str(tmp_path / "test.log")
     try:
         call_tool("write_file", {"path": log_path, "content": "line1\nline2\n"})
         err = call_tool_expect_error("edit_file", {

tools/server/tests/utils.py

@@ -294,6 +294,7 @@ class ServerProcess:
             server_args.append("--backend_sampling")
         if self.gcp_compat:
             env["AIP_MODE"] = "PREDICTION"
+            env["AIP_HTTP_PORT"] = str(self.server_port)
 
         args = [str(arg) for arg in [server_path, *server_args]]
         print(f"tests: starting server with: {' '.join(args)}")

vendor/cpp-httplib/CMakeLists.txt

@@ -43,7 +43,7 @@ if (LLAMA_BUILD_BORINGSSL)
     set(FIPS OFF CACHE BOOL "Enable FIPS (BoringSSL)")
 
     set(BORINGSSL_GIT "https://boringssl.googlesource.com/boringssl" CACHE STRING "BoringSSL git repository")
-    set(BORINGSSL_VERSION "0.20260813.0" CACHE STRING "BoringSSL version")
+    set(BORINGSSL_VERSION "0.20260903.0" CACHE STRING "BoringSSL version")
 
     message(STATUS "Fetching BoringSSL version ${BORINGSSL_VERSION}")