mtmd: fix LFM2 image tiling threshold (#27057)

* mtmd: fix LFM2 image tiling threshold

* refactor testing

* fix

* fix on windows

---------

Co-authored-by: Xuan Son Nguyen <son@huggingface.co>

.github/workflows/make-release.yml

@@ -49,6 +49,33 @@ jobs:
           git push origin "${VERSION}"
           echo "Created and pushed tag ${VERSION}"
 
+      - name: Generate release description
+        id: desc
+        run: bash scripts/make-release-desc.sh "${{ steps.checks.outputs.version }}"
+        env:
+          GITHUB_REPOSITORY: ${{ github.repository }}
+
+      - name: Create release
+        if: ${{ github.event.inputs.dry_run == 'false' }}
+        uses: ggml-org/action-create-release@v1
+        env:
+          GITHUB_TOKEN: ${{ github.token }}
+        with:
+          tag_name: ${{ steps.checks.outputs.version }}
+          # TODO: remove the prerelease flag once the semantic versioning workflow is ready
+          # ref: https://github.com/ggml-org/ggml/discussions/1579
+          prerelease: true
+          body: |
+            > [!NOTE]
+            > Semantic versioning is still work in progress.
+            > More info can be found in https://github.com/ggml-org/ggml/discussions/1579
+
+            ${{ steps.desc.outputs.nightly }}
+
+            ## ${{ steps.desc.outputs.changelog_title }}
+
+            ${{ steps.desc.outputs.changelog }}
+
       - name: Dry run summary
         if: ${{ github.event.inputs.dry_run == 'true' }}
         run: |

CMakeLists.txt

@@ -5,7 +5,7 @@ include(CheckIncludeFileCXX)
 ### llama.cpp version
 set(LLAMA_VERSION_MAJOR 0)
 set(LLAMA_VERSION_MINOR 1)
-set(LLAMA_VERSION_PATCH 1)
+set(LLAMA_VERSION_PATCH 2)
 set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
 
 # whether this is a development/nightly build
@@ -224,12 +224,10 @@ add_subdirectory(src)
 # utils, programs, examples and tests
 #
 
-# mtmd needs this even when common is not built
-add_subdirectory(vendor/hash)
+add_subdirectory(vendor)
 
 if (LLAMA_BUILD_COMMON)
     add_subdirectory(common)
-    add_subdirectory(vendor/cpp-httplib)
 endif()
 
 if (LLAMA_BUILD_COMMON AND LLAMA_BUILD_TESTS AND NOT CMAKE_JS_VERSION)

README.md

@@ -119,7 +119,7 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
 ## Acknowledgements
 
 - [yhirose/cpp-httplib](https://github.com/yhirose/cpp-httplib) - Single-header HTTP server, used by `llama-server` - MIT license
-- [stb-image](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
+- [nothings/stb](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
 - [nlohmann/json](https://github.com/nlohmann/json) - Single-header JSON library, used by various tools/examples - MIT License
-- [miniaudio.h](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
-- [subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
+- [mackron/miniaudio](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
+- [sheredom/subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain

build-xcframework.sh

@@ -290,6 +290,7 @@ combine_static_libraries() {
         "${base_dir}/${build_dir}/ggml/src/ggml-metal/${release_dir}/libggml-metal.a"
         "${base_dir}/${build_dir}/ggml/src/ggml-blas/${release_dir}/libggml-blas.a"
         "${base_dir}/${build_dir}/tools/mtmd/${release_dir}/libmtmd.a"
+        "${base_dir}/${build_dir}/vendor/hash/${release_dir}/libvendor-hash.a"
     )
 
     # Create temporary directory for processing

common/CMakeLists.txt

@@ -126,7 +126,8 @@ set_target_properties(${TARGET} PROPERTIES
     MACHO_CURRENT_VERSION 0 # keep macOS linker from seeing oversized version number
 )
 
-target_include_directories(${TARGET} PUBLIC . ../vendor)
+target_include_directories(${TARGET} PUBLIC .)
+target_link_libraries     (${TARGET} PUBLIC vendor::nlohmann vendor::sheredom)
 target_compile_features   (${TARGET} PUBLIC cxx_std_17)
 
 if (LLAMA_SUBPROCESS)

examples/gguf-hash/CMakeLists.txt

@@ -2,5 +2,5 @@ set(TARGET llama-gguf-hash)
 add_executable(${TARGET} gguf-hash.cpp)
 install(TARGETS ${TARGET} RUNTIME)
 
-target_link_libraries(${TARGET} PRIVATE vendor-hash ggml ${CMAKE_THREAD_LIBS_INIT})
+target_link_libraries(${TARGET} PRIVATE vendor::hash ggml ${CMAKE_THREAD_LIBS_INIT})
 target_compile_features(${TARGET} PRIVATE cxx_std_17)

examples/gguf-hash/gguf-hash.cpp

@@ -17,15 +17,15 @@
 extern "C" {
 #endif
 
-#include "xxhash/xxhash.h"
-#include "sha256/sha256.h"
+#include "hash/xxhash/xxhash.h"
+#include "hash/sha256/sha256.h"
 
 #ifdef __cplusplus
 }
 #endif
 
 // sha1 is compiled as C++ and lives in a namespace, see scripts/sync_vendor.py
-#include "sha1/sha1.h"
+#include "hash/sha1/sha1.h"
 using namespace vendor_hash;
 
 

ggml/CMakeLists.txt

@@ -5,7 +5,7 @@ project("ggml" C CXX ASM)
 ### GGML Version
 set(GGML_VERSION_MAJOR 0)
 set(GGML_VERSION_MINOR 20)
-set(GGML_VERSION_PATCH 1)
+set(GGML_VERSION_PATCH 2)
 set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
 
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")

ggml/src/ggml-cuda/mmvq.cu

@@ -4,6 +4,7 @@
 #include "vecdotq.cuh"
 
 #include <cstdint>
+#include <type_traits>
 
 typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
 
@@ -69,7 +70,8 @@ enum mmvq_parameter_table_id {
     MMVQ_PARAMETERS_GCN,
     MMVQ_PARAMETERS_RDNA2,
     MMVQ_PARAMETERS_RDNA3_0,
-    MMVQ_PARAMETERS_RDNA4
+    MMVQ_PARAMETERS_RDNA4,
+    MMVQ_PARAMETERS_GB10
 };
 
 static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
@@ -83,6 +85,8 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
     return MMVQ_PARAMETERS_GCN;
 #elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE
     return MMVQ_PARAMETERS_TURING;
+#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK
+    return MMVQ_PARAMETERS_GB10;
 #else
     return MMVQ_PARAMETERS_GENERIC;
 #endif
@@ -104,6 +108,9 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
     if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) {
         return MMVQ_PARAMETERS_TURING;
     }
+    if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARK) {
+        return MMVQ_PARAMETERS_GB10;
+    }
     return MMVQ_PARAMETERS_GENERIC;
 }
 
@@ -351,7 +358,7 @@ static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() {
 #endif
 }
 
-static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id) {
+static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id, bool small_k = false, bool halve_iters = false) {
     if (table_id == MMVQ_PARAMETERS_GENERIC) {
         switch (ncols_dst) {
             case 1:
@@ -454,11 +461,32 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
                 return 1;
         }
     }
+    if (table_id == MMVQ_PARAMETERS_GB10) {
+        const int generic = calc_nwarps(type, ncols_dst, MMVQ_PARAMETERS_GENERIC);
+        // Only worth the wider block when it actually retires the K loop in half the trips (Observation)
+        if (ncols_dst == 1 && !small_k && halve_iters) {
+            switch (type) {
+                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:
+                case GGML_TYPE_Q4_K:
+                case GGML_TYPE_Q5_K:
+                case GGML_TYPE_Q6_K:
+                case GGML_TYPE_IQ4_NL:
+                    return 2 * generic;
+                default:
+                    break;
+            }
+        }
+        return generic;
+    }
     return 1;
 }
 
 static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int table_id, bool small_k = false, int nwarps = 1) {
-    if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING) {
+    if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING || table_id == MMVQ_PARAMETERS_GB10) {
         switch (ncols_dst) {
             case 1:
                 return small_k ? nwarps : 1;
@@ -477,8 +505,8 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
     return 1;
 }
 
-template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
-__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
+template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false, bool halve_iters = false>
+__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id(), small_k, halve_iters)*ggml_cuda_get_physical_warp_size(), 1)
 static __global__ void mul_mat_vec_q(
         const void * vx_ptr, const void * vy_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
         const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -495,7 +523,7 @@ static __global__ void mul_mat_vec_q(
     constexpr int qi  = ggml_cuda_type_traits<type>::qi;
     constexpr int vdr = get_vdr_mmvq(type);
     constexpr mmvq_parameter_table_id table_id = get_device_table_id();
-    constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id);
+    constexpr int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
     constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();
 
@@ -773,8 +801,8 @@ static __global__ void mul_mat_vec_q_moe(
 template<ggml_type type>
 static std::pair<dim3, dim3> calc_launch_params(
         const int ncols_dst, const int nrows_x, const int nchannels_dst, const int nsamples_or_ntokens,
-        const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false) {
-    const int nwarps = calc_nwarps(type, ncols_dst, table_id);
+        const int warp_size, const mmvq_parameter_table_id table_id, const bool small_k = false, const bool halve_iters = false) {
+    const int nwarps = calc_nwarps(type, ncols_dst, table_id, small_k, halve_iters);
     const int rpb = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps);
     const int64_t nblocks = (nrows_x + rpb - 1) / rpb;
     const dim3 block_nums(nblocks, nchannels_dst, nsamples_or_ntokens);
@@ -782,7 +810,7 @@ static std::pair<dim3, dim3> calc_launch_params(
     return {block_nums, block_dims};
 }
 
-template<ggml_type type, int c_ncols_dst, bool small_k = false>
+template<ggml_type type, int c_ncols_dst, bool small_k = false, bool halve_iters = false>
 static void mul_mat_vec_q_switch_fusion(
         const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
         const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -797,7 +825,7 @@ static void mul_mat_vec_q_switch_fusion(
     if constexpr (c_ncols_dst == 1) {
         if (has_fusion) {
             const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
-            ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k>, launch_params,
+            ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, true, small_k, halve_iters>, launch_params,
                  vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
                  channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
                  sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -808,7 +836,7 @@ static void mul_mat_vec_q_switch_fusion(
     GGML_ASSERT(!has_fusion && "fusion only supported for ncols_dst=1");
 
     const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
-    ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k>, launch_params,
+    ggml_cuda_kernel_launch(mul_mat_vec_q<type, c_ncols_dst, false, small_k, halve_iters>, launch_params,
         vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
         channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
         sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride);
@@ -860,16 +888,18 @@ static void mul_mat_vec_q_switch_ncols_dst(
 
     const bool has_ids = ids != nullptr;
 
+    // How the K loop divides up at the baseline block width, both decisions below use these.
+    constexpr int qk                    = ggml_cuda_type_traits<type>::qk;
+    constexpr int qi                    = ggml_cuda_type_traits<type>::qi;
+    constexpr int vdr                   = get_vdr_mmvq(type);
+    const int     blocks_per_row_x      = ncols_x / qk;
+    const int     blocks_per_iter_1warp = vdr * warp_size / qi;
+
     const auto should_use_small_k = [&](int c_ncols_dst) {
         // When K is small, increase rows_per_block to match nwarps so each warp has more work to do
         // Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle.
-        constexpr int qk                    = ggml_cuda_type_traits<type>::qk;
-        constexpr int qi                    = ggml_cuda_type_traits<type>::qi;
-        constexpr int vdr                   = get_vdr_mmvq(type);
-        const int     blocks_per_row_x      = ncols_x / qk;
-        const int     blocks_per_iter_1warp = vdr * warp_size / qi;
-        const int     nwarps                = calc_nwarps(type, c_ncols_dst, table_id);
-        bool          use                   = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;
+        const int  nwarps = calc_nwarps(type, c_ncols_dst, table_id);
+        bool       use    = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp;
 
         constexpr std::array<ggml_type, 2> iq_slow_turing = {
             GGML_TYPE_IQ3_XXS,
@@ -902,6 +932,28 @@ static void mul_mat_vec_q_switch_ncols_dst(
         return use;
     };
 
+    // Whether doubling nwarps pays off on the ncols_dst == 1 path, where K sets the K loop trip count.
+    const auto should_halve_iters = [&] {
+        if (table_id != MMVQ_PARAMETERS_GB10) {
+            return false;
+        }
+
+        // Expert rows are gathered per token, so a wider block adds reduction work without reuse.
+        if (has_ids) {
+            return false;
+        }
+
+        const int blocks_per_iter = calc_nwarps(type, 1, table_id) * blocks_per_iter_1warp;
+        const int iters           = (blocks_per_row_x + blocks_per_iter - 1) /  blocks_per_iter;
+        const int iters_wide      = (blocks_per_row_x + blocks_per_iter * 2 - 1) / (blocks_per_iter * 2);
+
+        // An odd trip count leaves half the wider block idle for its last iteration, that tail is
+        // only affordable once the loop is long enough to dilute it to an eighth of the work (observation).
+        const int idle = iters_wide * 2 - iters;
+
+        return idle * 8 <= iters_wide * 2;
+    };
+
     if (has_ids && ncols_dst > 1) {
         // Multi-token MUL_MAT_ID path - dedicated MoE kernel
         mul_mat_vec_q_moe_launch<type>(
@@ -914,26 +966,34 @@ static void mul_mat_vec_q_switch_ncols_dst(
 
     switch (ncols_dst) {
         case 1: {
-            constexpr int c_ncols_dst = 1;
+            // static, else MSVC lambda capture breaks the constexpr uses below
+            static constexpr int c_ncols_dst = 1;
 
-            bool use_small_k = should_use_small_k(c_ncols_dst);
+            // Tag types keep the flags compile-time, so __launch_bounds__ matches what is launched.
+            const auto launch = [&](auto small_k_tag, auto halve_iters_tag) {
+                constexpr bool c_small_k = decltype(small_k_tag)::value;
+                // Types the table does not promote would compile a second, identical kernel.
+                constexpr bool c_promoted =
+                    calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, true) !=
+                    calc_nwarps(type, c_ncols_dst, MMVQ_PARAMETERS_GB10, false, false);
 
-            if (use_small_k) {
-                std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
-                                                                        nsamples_dst, warp_size, table_id, true);
-                mul_mat_vec_q_switch_fusion<type, c_ncols_dst, true>(
+                constexpr bool c_halve_iters = decltype(halve_iters_tag)::value && c_promoted;
+
+                const std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
+                                                                              nsamples_dst, warp_size, table_id, c_small_k, c_halve_iters);
+                mul_mat_vec_q_switch_fusion<type, c_ncols_dst, c_small_k, c_halve_iters>(
                     vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
                     channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
                     stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
                     stream);
+            };
+
+            if (should_use_small_k(c_ncols_dst)) {
+                launch(std::true_type{},  std::false_type{});
+            } else if (should_halve_iters()) {
+                launch(std::false_type{}, std::true_type{});
             } else {
-                std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
-                                                                        nsamples_dst, warp_size, table_id);
-                mul_mat_vec_q_switch_fusion<type, c_ncols_dst>(
-                    vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst,
-                    channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd,
-                    stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride,
-                    stream);
+                launch(std::false_type{}, std::false_type{});
             }
         } break;
         case 2: {

scripts/make-release-desc.sh

@@ -0,0 +1,87 @@
+#!/bin/bash
+# Generate the description of a release: the previous release version, the
+# change log and the link to the nightly release corresponding to the commit being released.
+#
+# Usage: make-release-desc.sh <version>
+#   <version>: current release version (v<maj>.<min>.<pat>, the leading v is optional)
+#
+# The previous version is the highest plain semver tag (v<maj>.<min>.<pat>)
+# strictly below <version>. The change log lists all commits between the
+# previous version tag and the release commit, one line per commit.
+#
+# The release commit is the commit <version> points at when the tag exists,
+# HEAD otherwise. The nightly release is the b* tag pointing at that commit
+# (release.yml tags the same commit); the link is only generated when that
+# tag exists.
+#
+# Env (when running in GitHub Actions):
+#   GITHUB_OUTPUT: previous_tag, changelog_title, changelog and nightly are written here
+#   GITHUB_REPOSITORY: owner/repo, used to build the nightly release URL (skipped when unset)
+set -euo pipefail
+
+if [[ $# -ne 1 ]]; then
+    echo "Usage: $(basename "$0") <version>"
+    exit 1
+fi
+VERSION="$1"
+
+# Accept the version with or without the leading v, reject anything else
+if [[ "${VERSION}" =~ ^[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
+    VERSION="v${VERSION}"
+elif [[ ! "${VERSION}" =~ ^v[0-9]+\.[0-9]+\.[0-9]+$ ]]; then
+    echo "Error: invalid version '${VERSION}' (expected v<maj>.<min>.<pat>)"
+    exit 1
+fi
+
+# Make sure all remote tags are available locally (skipped on local runs without origin)
+if ! git fetch --tags origin 2>/dev/null; then
+    echo "Warning: could not fetch tags from origin (local run?)"
+fi
+
+# Release commit: the commit <version> points at when the tag exists, HEAD otherwise.
+if ! RELEASE_COMMIT="$(git rev-parse -q --verify "refs/tags/${VERSION}^{commit}" 2>/dev/null)"; then
+    RELEASE_COMMIT="$(git rev-parse HEAD)"
+fi
+
+echo "Release commit: $(git rev-parse --short "${RELEASE_COMMIT}")"
+
+PREV="$( { git tag --list; echo "${VERSION}"; } \
+    | grep -E '^v[0-9]+\.[0-9]+\.[0-9]+$' \
+    | sort -V \
+    | awk -v cur="${VERSION}" '$0 == cur { exit } { prev = $0 } END { print prev }')"
+
+if [[ -n "${PREV}" ]]; then
+    CHANGELOG="$(git log --oneline "${PREV}..${RELEASE_COMMIT}")"
+    CHANGELOG_TITLE="Change log since ${PREV}"
+else
+    CHANGELOG="(no previous release tag found)"
+    CHANGELOG_TITLE="Change log"
+fi
+
+# Nightly release: the b* tag pointing at the release commit (|| true: no match is not an error)
+NIGHTLY_TAG="$(git tag --points-at "${RELEASE_COMMIT}" | grep -E '(^|-)b[0-9]+(-[0-9a-f]{7})?$' | head -n 1 || true)"
+
+NIGHTLY=""
+if [[ -n "${NIGHTLY_TAG}" ]]; then
+    if [[ -n "${GITHUB_REPOSITORY:-}" ]]; then
+        NIGHTLY_URL="https://github.com/${GITHUB_REPOSITORY}/releases/tag/${NIGHTLY_TAG}"
+        NIGHTLY="**Nightly build:** [${NIGHTLY_TAG}](${NIGHTLY_URL})"
+        echo "Nightly release: ${NIGHTLY_URL}"
+    fi
+else
+    echo "No nightly release found for commit $(git rev-parse --short "${RELEASE_COMMIT}")"
+fi
+
+echo "Previous version: ${PREV:-none}"
+echo "${CHANGELOG}"
+
+if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
+    {
+        echo "previous_tag=${PREV}"
+        echo "changelog_title=${CHANGELOG_TITLE}"
+        echo "nightly=${NIGHTLY}"
+        echo "changelog<<CHANGELOG_EOF"
+        echo "${CHANGELOG}"
+        echo "CHANGELOG_EOF"
+    } >> "${GITHUB_OUTPUT}"
+fi

scripts/sync-ggml.last

@@ -1 +1 @@
-3834fd814e74e8af277939dabd69ecc780affd21
+8c63e70982c95ceb862e3a1073a2c1beef75d60a

tests/CMakeLists.txt

@@ -310,6 +310,9 @@ llama_build_and_test(test-mtmd-c-api.c)
 target_link_libraries(${LLAMA_TEST_NAME} PRIVATE mtmd)
 unset(LLAMA_TEST_NAME)
 
+llama_build_and_test(test-mtmd-impl.cpp)
+target_link_libraries(test-mtmd-impl PRIVATE mtmd)
+
 # GGUF model data fetcher library for tests that need real model metadata
 # Only compile when cpp-httplib has SSL support (CPPHTTPLIB_OPENSSL_SUPPORT)
 if (TARGET cpp-httplib)

tests/test-mtmd-impl.cpp

@@ -0,0 +1,88 @@
+#include "testing.h"
+
+#include "mtmd-image.h"
+
+#include <iostream>
+#include <string>
+#include <utility>
+#include <vector>
+
+// this test file contains:
+// 1. test cases for mtmd helpers
+// 2. test cases for internal mtmd components
+// internal headers can be included here
+
+struct test_registry {
+    using fn_t = void (*)(testing &);
+
+    struct entry {
+        std::string name;
+        fn_t fn;
+    };
+
+    static std::vector<entry> & all() {
+        static std::vector<entry> entries;
+        return entries;
+    }
+
+    test_registry(const char * name, fn_t fn) {
+        all().push_back({ name, fn });
+    }
+};
+
+#define MAKE_TEST(name)                                               \
+    static void name(testing & t);                                    \
+    static const test_registry test_registry_ ## name(#name, &name);  \
+    static void name(testing & t)
+
+
+//
+// mtmd_image
+//
+
+MAKE_TEST(test_image_preprocessor_lfm2) {
+    clip_hparams hparams;
+    hparams.patch_size = 16;
+    hparams.n_merge = 2;
+    hparams.set_limit_image_tokens(64, 256);
+
+    // { image size, expected tiling }
+    const std::vector<std::pair<clip_image_size, bool>> cases = {
+        { {  704, 704 }, false },
+        // 720 / (patch_size * n_merge) is exactly 22.5, so this only matches HF
+        // if round_by_factor rounds half to even (22) instead of away from zero (23)
+        { {  720, 720 }, false },
+        { {  736, 736 }, true  },
+        { { 1024, 977 }, true  },
+        { { 1056, 384 }, false },
+    };
+
+    for (const auto & [size, expected] : cases) {
+        const bool actual = mtmd_image_preprocessor_lfm2::should_tile(hparams, size);
+
+        t.assert_equal(
+            "tiling for " + std::to_string(size.width) + "x" + std::to_string(size.height),
+            std::string(expected ? "tiled" : "single"),
+            std::string(actual   ? "tiled" : "single"));
+    }
+}
+
+//
+// main
+//
+
+int main(int argc, char ** argv) {
+    testing t(std::cout);
+    t.verbose = true;
+
+    // usage: test-mtmd-impl [filter_regex]
+    for (int i = 1; i < argc; i++) {
+        t.set_filter(argv[i]);
+    }
+
+    for (const auto & e : test_registry::all()) {
+        t.test(e.name, e.fn);
+    }
+
+    return t.summary();
+}

tools/mtmd/CMakeLists.txt

@@ -78,10 +78,8 @@ set_target_properties(mtmd PROPERTIES
 )
 
 target_link_libraries     (mtmd PUBLIC ggml llama)
-target_link_libraries     (mtmd PRIVATE Threads::Threads vendor-hash)
+target_link_libraries     (mtmd PRIVATE Threads::Threads vendor::hash vendor::miniaudio vendor::stb vendor::sheredom)
 target_include_directories(mtmd PUBLIC  .)
-target_include_directories(mtmd PRIVATE ../..)
-target_include_directories(mtmd PRIVATE ../../vendor)
 target_compile_features   (mtmd PRIVATE cxx_std_17)
 
 if (MTMD_VIDEO)
@@ -92,6 +90,9 @@ if (BUILD_SHARED_LIBS)
     set_target_properties     (mtmd PROPERTIES POSITION_INDEPENDENT_CODE ON)
     target_compile_definitions(mtmd PRIVATE LLAMA_BUILD)
     target_compile_definitions(mtmd PUBLIC  LLAMA_SHARED)
+
+    # export all symbols so that internal components can be tested by test-mtmd-impl
+    set_target_properties     (mtmd PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS ON)
 endif()
 
 set(MTMD_PUBLIC_HEADERS

tools/mtmd/clip-impl.h

@@ -858,6 +858,9 @@ static std::ifstream open_ifstream_binary(const std::string & fname) {
 }
 #endif
 
+// in test-mtmd-impl, we include woth common.h and this file, and these functions are duplicated
+// this is a quick fix to avoid compilation errors
+#ifndef DIRECTORY_SEPARATOR
 static std::string string_format(const char * fmt, ...) {
     va_list ap;
     va_list ap2;
@@ -915,6 +918,7 @@ inline bool string_ends_with(std::string_view str, std::string_view suffix) {
     return str.size() >= suffix.size() &&
            str.compare(str.size() - suffix.size(), suffix.size(), suffix) == 0;
 }
+#endif
 
 //
 // gguf utils

tools/mtmd/mtmd-helper.cpp

@@ -12,7 +12,7 @@
 #include "mtmd-helper-common.h"
 #include "llama.h"
 
-#include "hash.h"
+#include "hash/hash.h"
 
 #include <algorithm>
 #include <cinttypes>

tools/mtmd/mtmd-image.cpp

@@ -1013,14 +1013,31 @@ mtmd_image_preproc_out mtmd_image_preprocessor_lfm2::preprocess(const clip_image
     return output;
 }
 
+bool mtmd_image_preprocessor_lfm2::should_tile(
+        const clip_hparams & hparams,
+        const clip_image_size & original_size) {
+    const int align_size = hparams.patch_size * hparams.n_merge;
+
+    const auto round_by_factor = [align_size](float x) {
+        // see https://github.com/ggml-org/llama.cpp/pull/27057#discussion_r3796264887
+        return static_cast<int>(std::nearbyint(static_cast<double>(x) / align_size)) * align_size;
+    };
+
+    const int h_bar = std::max(hparams.patch_size, round_by_factor(original_size.height));
+    const int w_bar = std::max(hparams.patch_size, round_by_factor(original_size.width));
+
+    return static_cast<double>(h_bar) * static_cast<double>(w_bar) >
+           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 inst;
     const int align_size = hparams.patch_size * hparams.n_merge;
     inst.overview_size = img_tool::calc_size_preserved_ratio(
                             original_size,
                             { align_size, hparams.image_min_pixels, hparams.image_max_pixels, 0 });
-    // tile if either dimension exceeds tile_size with tolerance
-    const bool needs_tiling = original_size.width > tile_size * max_pixels_tolerance || original_size.height > tile_size * max_pixels_tolerance;
+
+    const bool needs_tiling = should_tile(hparams, original_size);
 
     if (!needs_tiling) {
         inst.refined_size = clip_image_size{0, 0};

tools/mtmd/mtmd-image.h

@@ -148,6 +148,8 @@ struct mtmd_image_preprocessor_lfm2 : 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;
 
+    static bool should_tile(const clip_hparams & hparams, const clip_image_size & original_size);
+
 private:
     clip_image_size find_closest_aspect_ratio(
             float aspect_ratio,

vendor/CMakeLists.txt

@@ -0,0 +1,11 @@
+# mtmd needs these even when common is not built
+add_subdirectory(hash)
+add_subdirectory(miniaudio)
+add_subdirectory(nlohmann)
+add_subdirectory(sheredom)
+add_subdirectory(stb)
+
+# only used by common
+if (LLAMA_BUILD_COMMON)
+    add_subdirectory(cpp-httplib)
+endif()

vendor/cpp-httplib/CMakeLists.txt

@@ -9,6 +9,8 @@ set(CMAKE_POSITION_INDEPENDENT_CODE ON)
 
 add_library(${TARGET} STATIC httplib.cpp httplib.h)
 
+add_library(vendor::cpp-httplib ALIAS ${TARGET})
+
 # disable warnings in 3rd party code
 if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
     target_compile_options(${TARGET} PRIVATE /w)

vendor/hash/CMakeLists.txt

@@ -16,6 +16,8 @@ add_library(${TARGET} STATIC
     ${VENDOR_SRCS}
 )
 
+add_library(vendor::hash ALIAS ${TARGET})
+
 target_compile_features(${TARGET} PRIVATE cxx_std_17)
 
 # disable warnings in 3rd party code, but keep them for hash.cpp
@@ -29,5 +31,8 @@ set_source_files_properties(${VENDOR_SRCS} PROPERTIES COMPILE_OPTIONS ${NO_WARN_
 # sha1 lives in a namespace to avoid a clash with boringssl, see scripts/sync_vendor.py
 set_source_files_properties(sha1/sha1.c PROPERTIES LANGUAGE CXX)
 
-# sha256.c includes "rotate-bits/rotate-bits.h", so consumers get this dir too
-target_include_directories(${TARGET} PUBLIC .)
+# expose the vendor/ root so consumers can include via "hash/hash.h"
+target_include_directories(${TARGET} PUBLIC ..)
+
+# internal includes of the vendored sources, e.g. sha256.c -> "rotate-bits/rotate-bits.h"
+target_include_directories(${TARGET} PRIVATE .)

vendor/miniaudio/CMakeLists.txt

@@ -0,0 +1,6 @@
+# header-only: interface target exposing the vendor/ root so consumers
+# can include via <miniaudio/miniaudio.h>
+add_library(miniaudio INTERFACE)
+add_library(vendor::miniaudio ALIAS miniaudio)
+
+target_include_directories(miniaudio INTERFACE ..)

vendor/nlohmann/CMakeLists.txt

@@ -0,0 +1,6 @@
+# header-only: interface target exposing the vendor/ root so consumers
+# can include via <nlohmann/json.hpp>
+add_library(nlohmann INTERFACE)
+add_library(vendor::nlohmann ALIAS nlohmann)
+
+target_include_directories(nlohmann INTERFACE ..)

vendor/sheredom/CMakeLists.txt

@@ -0,0 +1,6 @@
+# header-only: interface target exposing the vendor/ root so consumers
+# can include via <sheredom/subprocess.h>
+add_library(sheredom INTERFACE)
+add_library(vendor::sheredom ALIAS sheredom)
+
+target_include_directories(sheredom INTERFACE ..)

vendor/stb/CMakeLists.txt

@@ -0,0 +1,6 @@
+# header-only: interface target exposing the vendor/ root so consumers
+# can include via <stb/stb_image.h>
+add_library(stb INTERFACE)
+add_library(vendor::stb ALIAS stb)
+
+target_include_directories(stb INTERFACE ..)