ggml : allow MUL_MAT_ID -1 index to skip the computation (#26631 )

* ggml : allow MUL_MAT_ID -1 index to skip the computation * nits * update vulkan code * ggml-webgpu: ensure all threads reach workgroup barriers in mul_mat_id_vec (cherry picked from commit d8a4b9efb6e2b9abf3a866e847756b7e6bbc6b69) * fix cuda * openvino * exclude CUDA from this PR * add TODO tag --------- Co-authored-by: yomaytk <yoshimura.masashi.frbs@gmail.com>
tests : add speculative rejection draft sampling test (#30213 )
2026-10-24 05:27:38 -05:00 · 2026-10-24 12:15:36 +02:00 · 2026-10-24 12:47:49 +03:00 · 2026-10-24 09:21:24 +03:00 · 2026-10-24 06:05:05 +02:00 · 2026-10-24 12:04:05 +08:00
94 changed files with 2315 additions and 328 deletions
@@ -60,4 +60,4 @@ jobs:
      - name: Test
        id: ggml-ci
        run: |
-          GG_BUILD_METAL=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+          GG_BUILD_METAL=1 bash ./ci/run.sh ~/results/llama.cpp/$RUNNER_NAME ~/mnt/llama.cpp
@@ -150,7 +150,7 @@ jobs:
        id: ggml-ci
        run: |
          vulkaninfo --summary
-          GG_BUILD_VULKAN=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+          GG_BUILD_VULKAN=1 bash ./ci/run.sh ~/results/llama.cpp/$RUNNER_NAME ~/mnt/llama.cpp

  gpu-vulkan-intel-linux:
    runs-on: [self-hosted, Linux, Intel]
@@ -131,4 +131,4 @@ jobs:
        id: ggml-ci
        run: |
          GG_BUILD_WEBGPU=1 GG_BUILD_WEBGPU_DAWN_PREFIX="$GITHUB_WORKSPACE/dawn" \
-            bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+            bash ./ci/run.sh ~/results/llama.cpp/$RUNNER_NAME ~/mnt/llama.cpp
@@ -185,18 +185,28 @@ fi

 ## helpers

-# download a file if it does not exist or if it is outdated
+# download a file unless already cached in $MNT
 function gg_wget {
    local out=$1
    local url=$2

    local cwd=`pwd`
+    local name=`basename $url`
+    local tmp=.gg-wget.$$.tmp

    mkdir -p $out
    cd $out

-    # should not re-download if file is the same
-    wget -nv -c -N $url
+    # cache first: HF artifacts are static, HF sends no Last-Modified for -N
+    if [ ! -s $name ]; then
+        # publish with mv: concurrent runners never read partial files
+        rm -rf $tmp && mkdir $tmp
+        (cd $tmp && wget -nv $url)
+        if [ -s $tmp/$name ]; then
+            mv -f $tmp/$name $name
+        fi
+        rm -rf $tmp
+    fi

    cd $cwd
 }
@@ -312,8 +322,8 @@ function gg_run_test_scripts {

    set -e

-    (cd ./tools/gguf-split && time bash tests.sh "$SRC/build-ci-release/bin" "$MNT/models") 2>&1 | tee -a $OUT/${ci}-scripts.log
-    (cd ./tools/quantize   && time bash tests.sh "$SRC/build-ci-release/bin" "$MNT/models") 2>&1 | tee -a $OUT/${ci}-scripts.log
+    (cd ./tools/gguf-split && time bash tests.sh "$SRC/build-ci-release/bin" "$OUT/scratch") 2>&1 | tee -a $OUT/${ci}-scripts.log
+    (cd ./tools/quantize   && time bash tests.sh "$SRC/build-ci-release/bin" "$OUT/scratch") 2>&1 | tee -a $OUT/${ci}-scripts.log

    set +e
 }
@@ -374,10 +384,14 @@ function gg_run_qwen3_0_6b {


    gg_wget models-mnt/wikitext/ https://huggingface.co/datasets/ggml-org/ci/resolve/main/wikitext-2-raw-v1.zip
-    unzip -o models-mnt/wikitext/wikitext-2-raw-v1.zip -d models-mnt/wikitext/
+
+    rm -rf $OUT/scratch/qwen3-0.6B $OUT/scratch/wikitext
+    mkdir -p $OUT/scratch/qwen3-0.6B
+    unzip -o models-mnt/wikitext/wikitext-2-raw-v1.zip -d $OUT/scratch/wikitext/

    path_models="../models-mnt/qwen3/0.6B"
-    path_wiki="../models-mnt/wikitext/wikitext-2-raw"
+    path_local="$OUT/scratch/qwen3-0.6B"
+    path_wiki="$OUT/scratch/wikitext/wikitext-2-raw"

    rm -rf build-ci-release && mkdir build-ci-release && cd build-ci-release

@@ -386,76 +400,15 @@ function gg_run_qwen3_0_6b {
    (cmake -G "${CMAKE_GENERATOR}" -DCMAKE_BUILD_TYPE=Release ${CMAKE_EXTRA} .. ) 2>&1 | tee -a $OUT/${ci}-cmake.log
    (time cmake --build . --config Release -j$(nproc)) 2>&1 | tee -a $OUT/${ci}-make.log

-    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_models}/ggml-model-f16.gguf  --outtype f16
-    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_models}/ggml-model-bf16.gguf --outtype bf16
+    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_local}/ggml-model-f16.gguf  --outtype f16
+    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_local}/ggml-model-bf16.gguf --outtype bf16

-    model_f16="${path_models}/ggml-model-f16.gguf"
-    model_bf16="${path_models}/ggml-model-bf16.gguf"
-    model_q8_0="${path_models}/ggml-model-q8_0.gguf"
-    model_q4_0="${path_models}/ggml-model-q4_0.gguf"
-    model_q4_1="${path_models}/ggml-model-q4_1.gguf"
-    model_q5_0="${path_models}/ggml-model-q5_0.gguf"
-    model_q5_1="${path_models}/ggml-model-q5_1.gguf"
-    model_q2_k="${path_models}/ggml-model-q2_k.gguf"
-    model_q3_k="${path_models}/ggml-model-q3_k.gguf"
-    model_q4_k="${path_models}/ggml-model-q4_k.gguf"
-    model_q5_k="${path_models}/ggml-model-q5_k.gguf"
-    model_q6_k="${path_models}/ggml-model-q6_k.gguf"
+    model_f16="${path_local}/ggml-model-f16.gguf"
+    model_bf16="${path_local}/ggml-model-bf16.gguf"
+    model_q4_0="${path_local}/ggml-model-q4_0.gguf"

    wiki_test="${path_wiki}/wiki.test.raw"

-    ./bin/llama-quantize ${model_bf16} ${model_q8_0} q8_0 $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q4_0} q4_0 $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q4_1} q4_1 $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q5_0} q5_0 $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q5_1} q5_1 $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q2_k} q2_k $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q3_k} q3_k $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q4_k} q4_k $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q5_k} q5_k $(nproc)
-    ./bin/llama-quantize ${model_bf16} ${model_q6_k} q6_k $(nproc)
-
-    (time ./bin/llama-fit-params --model ${model_f16} 2>&1 | tee -a $OUT/${ci}-fp-f16.log)
-
-    (time ./bin/llama-completion -no-cnv --model ${model_f16}  -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-f16.log
-    (time ./bin/llama-completion -no-cnv --model ${model_bf16} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-bf16.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q8_0} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q8_0.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q4_0} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q4_0.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q4_1} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q4_1.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q5_0} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q5_0.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q5_1} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q5_1.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q2_k} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q2_k.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q3_k} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q3_k.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q4_k} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q4_k.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q5_k} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q5_k.log
-    (time ./bin/llama-completion -no-cnv --model ${model_q6_k} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-q6_k.log
-
-    (time ./bin/llama-perplexity --model ${model_f16}  -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-f16.log
-    if [ -z ${GG_BUILD_NO_BF16} ]; then
-        (time ./bin/llama-perplexity --model ${model_bf16} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-bf16.log
-    fi
-    (time ./bin/llama-perplexity --model ${model_q8_0} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q8_0.log
-    (time ./bin/llama-perplexity --model ${model_q4_0} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q4_0.log
-    (time ./bin/llama-perplexity --model ${model_q4_1} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q4_1.log
-    (time ./bin/llama-perplexity --model ${model_q5_0} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q5_0.log
-    (time ./bin/llama-perplexity --model ${model_q5_1} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q5_1.log
-    (time ./bin/llama-perplexity --model ${model_q2_k} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q2_k.log
-    (time ./bin/llama-perplexity --model ${model_q3_k} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q3_k.log
-    (time ./bin/llama-perplexity --model ${model_q4_k} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q4_k.log
-    (time ./bin/llama-perplexity --model ${model_q5_k} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q5_k.log
-    (time ./bin/llama-perplexity --model ${model_q6_k} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-q6_k.log
-
-    (time ./bin/llama-imatrix --model ${model_f16} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-imatrix.log
-
-    # TODO: these tests are extremely slow on WebGPU - re-enable when fixed
-    #       https://github.com/ggml-org/llama.cpp/actions/runs/39639985732
-    if [ -z ${GG_BUILD_WEBGPU} ]; then
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa off --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa off                ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-    fi
-
    function check_ppl {
        qnt="$1"
        ppl=$(echo "$2" | grep -oE "[0-9]+\.[0-9]+" | tail -n 1)
@@ -469,20 +422,56 @@ function gg_run_qwen3_0_6b {
        return 0
    }

+    for qnt in q8_0 q4_0 q4_1 q5_0 q5_1 q2_k q3_k q4_k q5_k q6_k; do
+        model="${path_local}/ggml-model-${qnt}.gguf"
+
+        ./bin/llama-quantize ${model_bf16} ${model} ${qnt} $(nproc)
+
+        (time ./bin/llama-completion -no-cnv --model ${model} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-${qnt}.log
+        (time ./bin/llama-perplexity --model ${model} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-${qnt}.log
+
+        # note: ppl > 20.0 for q2_k with this model
+        if [ "${qnt}" != "q2_k" ]; then
+            check_ppl "${qnt}" "$(cat $OUT/${ci}-tg-${qnt}.log | grep "^\[1\]")"
+        fi
+
+        if [ "${qnt}" = "q4_0" ]; then
+            # q4_0 is kept for other tests
+            q4_0="${path_models}/ggml-model-q4_0.gguf"
+            ln ${model} "$q4_0" 2>/dev/null || { cp ${model} $q4_0.tmp.$$ && mv -f $q4_0.tmp.$$ $q4_0; }
+        else
+            rm -f ${model}
+        fi
+    done
+
+    (time ./bin/llama-fit-params --model ${model_f16} 2>&1 | tee -a $OUT/${ci}-fp-f16.log)
+
+    (time ./bin/llama-completion -no-cnv --model ${model_f16}  -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-f16.log
+    (time ./bin/llama-completion -no-cnv --model ${model_bf16} -ngl 99 -c 1024 -s 1234 -n 64 --ignore-eos -p "I believe the meaning of life is" ) 2>&1 | tee -a $OUT/${ci}-tg-bf16.log
+
+    (time ./bin/llama-perplexity --model ${model_f16}  -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-f16.log
+    if [ -z ${GG_BUILD_NO_BF16} ]; then
+        (time ./bin/llama-perplexity --model ${model_bf16} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-tg-bf16.log
+    fi
+
+    (time ./bin/llama-imatrix --model ${model_f16} -f ${wiki_test} -ngl 99 -c 1024 -b 512 --chunks 2 ) 2>&1 | tee -a $OUT/${ci}-imatrix.log
+
+    # TODO: these tests are extremely slow on WebGPU - re-enable when fixed
+    #       https://github.com/ggml-org/llama.cpp/actions/runs/39639985732
+    if [ -z ${GG_BUILD_WEBGPU} ]; then
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa off --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa off                ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+    fi
+
    check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")"
    if [ -z ${GG_BUILD_NO_BF16} ]; then
        check_ppl "bf16" "$(cat $OUT/${ci}-tg-bf16.log | grep "^\[1\]")"
    fi
-    check_ppl "q8_0" "$(cat $OUT/${ci}-tg-q8_0.log | grep "^\[1\]")"
-    check_ppl "q4_0" "$(cat $OUT/${ci}-tg-q4_0.log | grep "^\[1\]")"
-    check_ppl "q4_1" "$(cat $OUT/${ci}-tg-q4_1.log | grep "^\[1\]")"
-    check_ppl "q5_0" "$(cat $OUT/${ci}-tg-q5_0.log | grep "^\[1\]")"
-    check_ppl "q5_1" "$(cat $OUT/${ci}-tg-q5_1.log | grep "^\[1\]")"
-   #check_ppl "q2_k" "$(cat $OUT/${ci}-tg-q2_k.log | grep "^\[1\]")" # note: ppl > 20.0 for this quant and model
-    check_ppl "q3_k" "$(cat $OUT/${ci}-tg-q3_k.log | grep "^\[1\]")"
-    check_ppl "q4_k" "$(cat $OUT/${ci}-tg-q4_k.log | grep "^\[1\]")"
-    check_ppl "q5_k" "$(cat $OUT/${ci}-tg-q5_k.log | grep "^\[1\]")"
-    check_ppl "q6_k" "$(cat $OUT/${ci}-tg-q6_k.log | grep "^\[1\]")"
+
+    # scratch is rebuilt on every run; q4_0 is already published to $MNT
+    rm -rf $OUT/scratch/qwen3-0.6B $OUT/scratch/wikitext

    set +e
 }
@@ -505,6 +494,8 @@ function gg_run_embd_bge_small {
    gg_wget models-mnt/bge-small/1_Pooling https://huggingface.co/BAAI/bge-small-en-v1.5/raw/main/1_Pooling/config.json

    path_models="../models-mnt/bge-small"
+    path_local="$OUT/scratch/bge-small"
+    rm -rf $path_local && mkdir -p $path_local

    rm -rf build-ci-release && mkdir build-ci-release && cd build-ci-release

@@ -513,10 +504,10 @@ function gg_run_embd_bge_small {
    (cmake -G "${CMAKE_GENERATOR}" -DCMAKE_BUILD_TYPE=Release ${CMAKE_EXTRA} .. ) 2>&1 | tee -a $OUT/${ci}-cmake.log
    (time cmake --build . --config Release -j$(nproc)) 2>&1 | tee -a $OUT/${ci}-make.log

-    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_models}/ggml-model-f16.gguf
+    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_local}/ggml-model-f16.gguf

-    model_f16="${path_models}/ggml-model-f16.gguf"
-    model_q8_0="${path_models}/ggml-model-q8_0.gguf"
+    model_f16="${path_local}/ggml-model-f16.gguf"
+    model_q8_0="${path_local}/ggml-model-q8_0.gguf"

    ./bin/llama-quantize ${model_f16} ${model_q8_0} q8_0

@@ -525,6 +516,8 @@ function gg_run_embd_bge_small {
    (time ./bin/llama-embedding --model ${model_f16}  -p "I believe the meaning of life is" -ngl 99 -c 0 --no-op-offload) 2>&1 | tee -a $OUT/${ci}-tg-f16.log
    (time ./bin/llama-embedding --model ${model_q8_0} -p "I believe the meaning of life is" -ngl 99 -c 0 --no-op-offload) 2>&1 | tee -a $OUT/${ci}-tg-q8_0.log

+    rm -rf $path_local
+
    set +e
 }

@@ -541,6 +534,8 @@ function gg_run_rerank_tiny {
    gg_wget models-mnt/rerank-tiny/ https://huggingface.co/jinaai/jina-reranker-v1-tiny-en/raw/main/vocab.json

    path_models="../models-mnt/rerank-tiny"
+    path_local="$OUT/scratch/rerank-tiny"
+    rm -rf $path_local && mkdir -p $path_local

    rm -rf build-ci-release && mkdir build-ci-release && cd build-ci-release

@@ -549,9 +544,9 @@ function gg_run_rerank_tiny {
    (cmake -G "${CMAKE_GENERATOR}" -DCMAKE_BUILD_TYPE=Release ${CMAKE_EXTRA} .. ) 2>&1 | tee -a $OUT/${ci}-cmake.log
    (time cmake --build . --config Release -j$(nproc)) 2>&1 | tee -a $OUT/${ci}-make.log

-    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_models}/ggml-model-f16.gguf
+    python3 ../convert_hf_to_gguf.py ${path_models} --outfile ${path_local}/ggml-model-f16.gguf

-    model_f16="${path_models}/ggml-model-f16.gguf"
+    model_f16="${path_local}/ggml-model-f16.gguf"

    (time ./bin/llama-fit-params --model ${model_f16} 2>&1 | tee -a $OUT/${ci}-fp-f16.log)

@@ -581,6 +576,8 @@ function gg_run_rerank_tiny {
    check_score "rerank score 1" "$(cat $OUT/${ci}-rk-f16.log | grep "rerank score 1")" "0.00" "0.05" | tee -a $OUT/${ci}-rk-f16.log
    check_score "rerank score 2" "$(cat $OUT/${ci}-rk-f16.log | grep "rerank score 2")" "0.10" "0.30" | tee -a $OUT/${ci}-rk-f16.log

+    rm -rf $path_local
+
    set +e
 }

@@ -687,11 +684,11 @@ if [ -z ${GG_BUILD_LOW_PERF} ]; then
    ln -sfn ${mnt_models} ${SRC}/models-mnt

    # Create a fresh python3 venv and enter it
-    if ! python3 -m venv "$MNT/venv"; then
-        echo "Error: Failed to create Python virtual environment at $MNT/venv."
+    if ! python3 -m venv "$OUT/venv"; then
+        echo "Error: Failed to create Python virtual environment at $OUT/venv."
        exit 1
    fi
-    source "$MNT/venv/bin/activate"
+    source "$OUT/venv/bin/activate"

    pip install -r ${SRC}/requirements.txt --disable-pip-version-check
    pip install --editable gguf-py --disable-pip-version-check
@@ -78,6 +78,8 @@ add_library(${TARGET}
    download.h
    fit.cpp
    fit.h
+    file-lock.cpp
+    file-lock.h
    hf-cache.cpp
    hf-cache.h
    http.h
@@ -184,3 +186,8 @@ if (LLAMA_LLGUIDANCE)
 endif()

 target_link_libraries(${TARGET} PUBLIC llama Threads::Threads)
+
+if (CMAKE_SYSTEM_NAME MATCHES "AIX")
+    # AIX's flock() function comes from libbsd.a
+    target_link_libraries(${TARGET} PRIVATE -lbsd)
+endif()
@@ -470,6 +470,14 @@ void common_models_handler_apply(common_models_handler & handler, common_params
    auto opts = handler.opts; // copy
    opts.callback = callback;

+    // another process may hold the cache lock during finalization:
+    // let the download callback cancel that wait too
+    auto finalize = [callback](const hf_cache::hf_file & f) {
+        return hf_cache::finalize_file(f, [callback]() {
+            return !(callback && callback->is_cancelled());
+        });
+    };
+
    // handle plain "url" if needed
    auto handle_url = [&](common_params_model & model) {
        if (!model.url.empty()) {
@@ -518,7 +526,7 @@ void common_models_handler_apply(common_models_handler & handler, common_params
    }

    // handle hf_plan tasks
-    auto add_tasks = [&opts, &tasks](const hf_cache::hf_files  & model_files,
+    auto add_tasks = [&opts, &tasks, &finalize](const hf_cache::hf_files  & model_files,
                                    const hf_cache::hf_file    & primary,
                                    common_params_model        & model) {
        for (size_t i = 0; i < model_files.size(); ++i) {
@@ -527,9 +535,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
            tasks.emplace_back(model_file, opts, [&, is_primary]() {
                if (is_primary) {
                    // the primary file is the first split (00001-of), use it as model path
-                    model.path = hf_cache::finalize_file(model_file);
+                    model.path = finalize(model_file);
                } else {
-                    hf_cache::finalize_file(model_file);
+                    finalize(model_file);
                }
            });
        }
@@ -581,9 +589,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan_spec.mtp, opts, [&]() {
            // only use the discovered MTP head when no draft path is set yet
            if (params.speculative.draft.mparams.path.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan_spec.mtp);
+                params.speculative.draft.mparams.path = finalize(plan_spec.mtp);
            } else {
-                hf_cache::finalize_file(plan_spec.mtp);
+                finalize(plan_spec.mtp);
            }
        });
    }
@@ -591,9 +599,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan_spec.dflash, opts, [&]() {
            // only use the discovered DFlash sidecar when no draft path is set yet
            if (params.speculative.draft.mparams.path.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan_spec.dflash);
+                params.speculative.draft.mparams.path = finalize(plan_spec.dflash);
            } else {
-                hf_cache::finalize_file(plan_spec.dflash);
+                finalize(plan_spec.dflash);
            }
        });
    }
@@ -601,9 +609,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan_spec.eagle3, opts, [&]() {
            // only use the discovered Eagle3 sidecar when no draft path is set yet
            if (params.speculative.draft.mparams.path.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan_spec.eagle3);
+                params.speculative.draft.mparams.path = finalize(plan_spec.eagle3);
            } else {
-                hf_cache::finalize_file(plan_spec.eagle3);
+                finalize(plan_spec.eagle3);
            }
        });
    }
@@ -611,9 +619,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan_spec.dspark, opts, [&]() {
            // only use the discovered DSpark sidecar when no draft path is set yet
            if (params.speculative.draft.mparams.path.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan_spec.dspark);
+                params.speculative.draft.mparams.path = finalize(plan_spec.dspark);
            } else {
-                hf_cache::finalize_file(plan_spec.dspark);
+                finalize(plan_spec.dspark);
            }
        });
    }
@@ -634,16 +642,16 @@ void common_models_handler_apply(common_models_handler & handler, common_params
    }
    if (!plan.mmproj.local_path.empty()) {
        tasks.emplace_back(plan.mmproj, opts, [&]() {
-            params.mmproj.path = hf_cache::finalize_file(plan.mmproj);
+            params.mmproj.path = finalize(plan.mmproj);
        });
    }
    if (!plan.mtp.local_path.empty() && !had_spec_url) {
        tasks.emplace_back(plan.mtp, opts, [&]() {
            // only fall back to the discovered MTP head when no draft was explicitly provided
            if (params.speculative.draft.mparams.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan.mtp);
+                params.speculative.draft.mparams.path = finalize(plan.mtp);
            } else {
-                hf_cache::finalize_file(plan.mtp);
+                finalize(plan.mtp);
            }
        });
    }
@@ -651,9 +659,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan.dflash, opts, [&]() {
            // only fall back to the discovered DFlash sidecar when no draft was explicitly provided
            if (params.speculative.draft.mparams.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan.dflash);
+                params.speculative.draft.mparams.path = finalize(plan.dflash);
            } else {
-                hf_cache::finalize_file(plan.dflash);
+                finalize(plan.dflash);
            }
        });
    }
@@ -661,9 +669,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan.eagle3, opts, [&]() {
            // only fall back to the discovered Eagle3 sidecar when no draft was explicitly provided
            if (params.speculative.draft.mparams.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan.eagle3);
+                params.speculative.draft.mparams.path = finalize(plan.eagle3);
            } else {
-                hf_cache::finalize_file(plan.eagle3);
+                finalize(plan.eagle3);
            }
        });
    }
@@ -671,9 +679,9 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan.dspark, opts, [&]() {
            // only fall back to the discovered DSpark sidecar when no draft was explicitly provided
            if (params.speculative.draft.mparams.empty()) {
-                params.speculative.draft.mparams.path = hf_cache::finalize_file(plan.dspark);
+                params.speculative.draft.mparams.path = finalize(plan.dspark);
            } else {
-                hf_cache::finalize_file(plan.dspark);
+                finalize(plan.dspark);
            }
        });
    }
@@ -681,7 +689,7 @@ void common_models_handler_apply(common_models_handler & handler, common_params
        tasks.emplace_back(plan.preset, opts, [&]() {
            // if HF repo is a preset repo, we simply run server in router mode with the preset.ini file
            params.models_preset_hf = params.model.hf_repo; // only for showing a warning
-            params.models_preset    = hf_cache::finalize_file(plan.preset);
+            params.models_preset    = finalize(plan.preset);
            // clear the model so the server starts in router mode
            params.model.path.clear();
            params.model.hf_repo.clear();
@@ -26,6 +26,7 @@
 #include <fstream>
 #include <iostream>
 #include <iterator>
+#include <random>
 #include <regex>
 #include <sstream>
 #include <string>
@@ -917,9 +918,13 @@ std::string fs_path_to_utf8(const std::filesystem::path & path) {
 }

 void fs_write_atomic(const std::filesystem::path & path, const std::string & data) {
+    // a random suffix gives every writer its own temporary file, across processes and containers
+    std::random_device rd;
+    const uint64_t id = ((uint64_t) rd() << 32) | rd();
+
    std::error_code ec;
    std::filesystem::path path_tmp = path;
-    path_tmp += ".tmp";
+    path_tmp += "." + std::to_string(id) + ".tmp";

    if (path.has_parent_path()) {
        std::filesystem::create_directories(path.parent_path(), ec);
@@ -4,6 +4,7 @@
 #include "common.h"
 #include "log.h"
 #include "download.h"
+#include "file-lock.h"
 #include "hf-cache.h"
 #include "json.h"

@@ -77,6 +78,18 @@ std::pair<std::string, std::string> common_download_split_repo_tag(const std::st
    return {hf_repo, tag};
 }

+// file name from url, without the query string
+static std::string url_get_filename(const std::string & url) {
+    std::string filename = url;
+    if (auto pos = filename.rfind('/'); pos != std::string::npos) {
+        filename = filename.substr(pos + 1);
+    }
+    if (auto pos = filename.find('?'); pos != std::string::npos) {
+        filename = filename.substr(0, pos);
+    }
+    return filename;
+}
+
 class ProgressBar : public common_download_callback {
    static inline std::mutex mutex;
    static inline std::map<const ProgressBar *, int> lines;
@@ -100,14 +113,8 @@ public:
    ProgressBar() = default;

    void on_start(const common_download_progress & p) override {
-        filename = p.url;
+        filename = url_get_filename(p.url);

-        if (auto pos = filename.rfind('/'); pos != std::string::npos) {
-            filename = filename.substr(pos + 1);
-        }
-        if (auto pos = filename.find('?'); pos != std::string::npos) {
-            filename = filename.substr(0, pos);
-        }
        for (size_t i = 0; i < filename.size(); ++i) {
            if ((filename[i] & 0xC0) != 0x80) {
                if (len++ == 39) {
@@ -257,15 +264,62 @@ static bool common_pull_file(httplib::Client & cli,
 static int common_download_file_single_online(const std::string & url,
                                              const std::string & path,
                                              const common_download_opts & opts,
-                                              bool skip_etag) {
+                                              bool skip_etag,
+                                              const std::string & lock_path,
+                                              const std::string & alt_path) {
    static const int max_attempts        = 3;
    static const int retry_delay_seconds = 2;

-    const bool file_exists = std::filesystem::exists(std::filesystem::u8path(path));
+    const std::string path_temporary = path + PARTIAL_SUFFIX;
+    const std::string path_lock = lock_path.empty() ? path + ".lock" : lock_path;
+
+    common_download_progress p;
+    p.url = url;
+
+    if (opts.callback) {
+        opts.callback->on_start(p);
+    }
+
+    // the completion callback may itself take the same lock (the HF cache
+    // finalizes blobs under it), so every exit goes through finish()
+    common_file_lock lock(path_lock);
+    auto finish = [&](int status) {
+        lock.close();
+        if (opts.callback) {
+            opts.callback->on_done(p, is_http_status_ok(status));
+        }
+        return status;
+    };
+
+    // alt_path is the HF snapshot: the blob may have been moved there when symlinks are unavailable
+    auto cache_hit = [&] {
+        return std::filesystem::exists(std::filesystem::u8path(path)) ||
+               (!alt_path.empty() && std::filesystem::exists(std::filesystem::u8path(alt_path)));
+    };
+
+    if (skip_etag && cache_hit()) {
+        LOG_DBG("%s: using cached file: %s\n", __func__, path.c_str());
+        return finish(304); // 304 Not Modified - fake cached response
+    }
+
+    // take the lock before touching the cache, so two processes cannot race on the same file
+    bool waited = false;
+    if (!lock.acquire([&] {
+            if (!waited) {
+                LOG_INF("%s: file '%s' is being downloaded by another process, waiting...\n", __func__, url_get_filename(url).c_str());
+                waited = true;
+            }
+            return !(opts.callback && opts.callback->is_cancelled());
+        })) {
+        return finish(-1);
+    }
+
+    // under the lock: another process may have finished while we waited
+    bool file_exists = cache_hit();

    if (file_exists && skip_etag) {
        LOG_DBG("%s: using cached file: %s\n", __func__, path.c_str());
-        return 304; // 304 Not Modified - fake cached response
+        return finish(304); // 304 Not Modified - fake cached response
    }

    auto [cli, parts] = common_http_client(url);
@@ -294,9 +348,9 @@ static int common_download_file_single_online(const std::string & url,
        LOG_TRC("%s: HEAD failed, status: %d\n", __func__, head ? head->status : -1);
        if (file_exists) {
            LOG_TRC("%s: using cached file (HEAD failed): %s\n", __func__, path.c_str());
-            return 304; // 304 Not Modified - fake cached response
+            return finish(304); // 304 Not Modified - fake cached response
        }
-        return head ? head->status : -1;
+        return finish(head ? head->status : -1);
    }

    std::string etag;
@@ -304,8 +358,6 @@ static int common_download_file_single_online(const std::string & url,
        etag = head->get_header_value("ETag");
    }

-    common_download_progress p;
-    p.url = url;
    if (head->has_header("Content-Length")) {
        try {
            p.total = std::stoull(head->get_header_value("Content-Length"));
@@ -322,16 +374,16 @@ static int common_download_file_single_online(const std::string & url,
    if (file_exists) {
        if (etag.empty()) {
            LOG_DBG("%s: using cached file (no server etag): %s\n", __func__, path.c_str());
-            return 304; // 304 Not Modified - fake cached response
+            return finish(304); // 304 Not Modified - fake cached response
        }
        if (!last_etag.empty() && last_etag == etag) {
            LOG_DBG("%s: using cached file (same etag): %s\n", __func__, path.c_str());
-            return 304; // 304 Not Modified - fake cached response
+            return finish(304); // 304 Not Modified - fake cached response
        }
        // pass this point, the file exists but is different from the server version, so we need to redownload it
        if (remove(path.c_str()) != 0) {
            LOG_ERR("%s: unable to delete file: %s\n", __func__, path.c_str());
-            return -1;
+            return finish(-1);
        }
    }

@@ -341,7 +393,6 @@ static int common_download_file_single_online(const std::string & url,
    }

    bool success = false;
-    const std::string path_temporary = path + PARTIAL_SUFFIX;
    int delay = retry_delay_seconds;

    if (!skip_etag && !etag.empty()) {
@@ -349,7 +400,7 @@ static int common_download_file_single_online(const std::string & url,
            LOG_WRN("%s: partial file does not match the server etag, discarding it: %s\n", __func__, path_temporary.c_str());
            if (remove(path_temporary.c_str()) != 0) {
                LOG_ERR("%s: unable to delete file: %s\n", __func__, path_temporary.c_str());
-                return -1;
+                return finish(-1);
            }
        }
        write_etag(path_temporary, etag);
@@ -359,7 +410,7 @@ static int common_download_file_single_online(const std::string & url,
        // report the resume offset, so a resumed file does not start at 0
        p.downloaded = supports_ranges && std::filesystem::exists(path_temporary) ?
                       std::filesystem::file_size(path_temporary) : 0;
-        opts.callback->on_start(p);
+        opts.callback->on_update(p);
    }

    for (int i = 0; i < max_attempts; ++i) {
@@ -403,16 +454,12 @@ static int common_download_file_single_online(const std::string & url,
        }
    }

-    if (opts.callback) {
-        opts.callback->on_done(p, success);
-    }
    // a stopped download keeps path_temporary, so the next attempt resumes from it
    if (!success) {
        LOG_ERR("%s: download failed after %d attempts\n", __func__, max_attempts);
-        return -1; // max attempts reached
    }

-    return head->status;
+    return finish(success ? head->status : -1);
 }

 std::pair<long, std::vector<char>> common_remote_get_content(const std::string          & url,
@@ -452,14 +499,16 @@ std::pair<long, std::vector<char>> common_remote_get_content(const std::string
 int common_download_file_single(const std::string & url,
                                const std::string & path,
                                const common_download_opts & opts,
-                                bool skip_etag) {
+                                bool skip_etag,
+                                const std::string & lock_path,
+                                const std::string & alt_path) {
    if (!opts.offline) {
        ProgressBar tty_cb;
        common_download_opts online_opts = opts;
        if (!online_opts.callback) {
            online_opts.callback = &tty_cb;
        }
-        return common_download_file_single_online(url, path, online_opts, skip_etag);
+        return common_download_file_single_online(url, path, online_opts, skip_etag, lock_path, alt_path);
    }

    if (!std::filesystem::exists(std::filesystem::u8path(path))) {
@@ -821,7 +870,7 @@ void common_download_run_tasks(const std::vector<common_download_task> & tasks)
    for (const auto & task : tasks) {
        futures.push_back(std::async(std::launch::async,
            [&task]() {
-                return common_download_file_single(task.url, task.local_path, task.opts, task.is_hf);
+                return common_download_file_single(task.url, task.local_path, task.opts, task.is_hf, task.lock_path, task.alt_path);
            }
        ));
    }
@@ -69,12 +69,15 @@ struct common_download_task {
    std::string local_path;
    std::function<void()> on_done;
    bool is_hf = false;
+    std::string lock_path;
+    std::string alt_path; // HF snapshot path, also treated as cached (the blob may be moved there)

    common_download_task() = default;
    common_download_task(hf_cache::hf_file f,
            const common_download_opts & opts,
            std::function<void()> on_done = nullptr)
-        : opts(opts), url(f.url), local_path(f.local_path), on_done(on_done), is_hf(true) {}
+        : opts(opts), url(f.url), local_path(f.local_path), on_done(on_done), is_hf(true),
+          lock_path(hf_cache::get_lock_path(f.repo_id, f.oid)), alt_path(f.final_path) {}
 };

 void common_download_run_tasks(const std::vector<common_download_task> & tasks);
@@ -92,10 +95,14 @@ std::string common_download_resolve_path(const std::string & hf_repo_with_tag, c
 // download single file from url to local path
 // returns status code or -1 on error
 // skip_etag: if true, don't read/write .etag files (for HF cache where filename is the hash)
+// lock_path: lock file used to serialize downloads of the same file; when empty, "<path>.lock" is used
+// alt_path: additional path treated as present in cache (see cache_hit in download.cpp)
 int common_download_file_single(const std::string & url,
                                const std::string & path,
                                const common_download_opts & opts = {},
-                                bool skip_etag = false);
+                                bool skip_etag = false,
+                                const std::string & lock_path = "",
+                                const std::string & alt_path = "");

 // resolve and download model from Docker registry
 // return local path to downloaded model file
@@ -0,0 +1,179 @@
+#include "file-lock.h"
+
+#include "log.h"
+
+#include <chrono>
+#include <filesystem>
+#include <thread>
+
+#include <fcntl.h>
+#include <sys/stat.h>
+
+#if defined(_WIN32)
+#define WIN32_LEAN_AND_MEAN
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#include <windows.h>
+#include <io.h>
+#else
+#include <sys/file.h>
+#include <unistd.h>
+#include <cerrno>
+#include <cstring>
+#endif
+
+common_file_lock::common_file_lock(const std::string & lock_path) {
+    this->lock_path = lock_path;
+}
+
+#if defined(_WIN32)
+// the \\?\ prefix (for paths longer than 255 chars) is only accepted on absolute paths
+// with backslash separators: normalize first, then check the length of the absolute form
+static std::wstring as_extended_path(const std::string & path) {
+    // convert from the narrow encoding std::filesystem uses on this system
+    int wlen = MultiByteToWideChar(CP_ACP, 0, path.c_str(), (int) path.size(), nullptr, 0);
+    if (wlen <= 0) {
+        return std::wstring(path.begin(), path.end()); // best effort for exotic encodings
+    }
+    std::wstring wide(wlen, L'\0');
+    MultiByteToWideChar(CP_ACP, 0, path.c_str(), (int) path.size(), wide.data(), wlen);
+
+    if (wide.rfind(L"\\\\?\\", 0) == 0) {
+        return wide;
+    }
+
+    // GetFullPathNameW resolves relative paths and normalizes separators
+    DWORD size = GetFullPathNameW(wide.c_str(), 0, nullptr, nullptr);
+    if (size == 0) {
+        return wide;
+    }
+    std::wstring full(size, L'\0');
+    DWORD written = GetFullPathNameW(wide.c_str(), size, full.data(), nullptr);
+    if (written == 0 || written >= size) {
+        return wide;
+    }
+    full.resize(written);
+
+    if (full.size() <= 255) {
+        return wide; // short enough: open the path exactly as given
+    }
+    if (full.rfind(L"\\\\", 0) == 0) {
+        return L"\\\\?\\UNC\\" + full.substr(2); // \\server\share -> \\?\UNC\server\share
+    }
+    return L"\\\\?\\" + full;
+}
+#endif
+
+common_file_lock::status common_file_lock::try_lock() {
+    if (lock_path.empty()) {
+        return status::acquired; // no lock path, nothing to lock
+    }
+#if defined(_WIN32)
+    if (fd < 0) {
+        // opened with a wide path: extended paths cannot go through the narrow CRT
+        std::wstring lock_path_w = as_extended_path(lock_path);
+        fd = _wopen(lock_path_w.c_str(), _O_RDWR | _O_CREAT | _O_BINARY | _O_NOINHERIT, _S_IREAD | _S_IWRITE);
+        if (fd < 0) {
+            LOG_ERR("%s: cannot open lock file: %s\n", __func__, lock_path.c_str());
+            return status::error;
+        }
+    }
+    // lock the first byte, the file itself stays empty
+    OVERLAPPED ov = {};
+    locked = LockFileEx((HANDLE) _get_osfhandle(fd),
+                        LOCKFILE_EXCLUSIVE_LOCK | LOCKFILE_FAIL_IMMEDIATELY, 0, 1, 0, &ov);
+    if (locked) {
+        return status::acquired;
+    }
+    if (GetLastError() == ERROR_LOCK_VIOLATION) {
+        return status::busy;
+    }
+    LOG_ERR("%s: file locking failed (%d)\n", __func__, (int) GetLastError());
+    return status::error;
+#else
+    if (fd < 0) {
+        int open_flags = O_RDWR | O_CREAT | O_CLOEXEC;
+#ifdef O_NOFOLLOW
+        open_flags |= O_NOFOLLOW;
+#endif
+        // group-writable, so users sharing a cache can wait on the lock
+        fd = open(lock_path.c_str(), open_flags, 0664);
+        if (fd < 0) {
+            LOG_ERR("%s: cannot open lock file: %s\n", __func__, lock_path.c_str());
+            return status::error;
+        }
+        // open() filters the mode through the umask
+        (void) fchmod(fd, 0664);
+    }
+
+    if (flock(fd, LOCK_EX | LOCK_NB) == 0) {
+        locked = true;
+        return status::acquired;
+    }
+    // EACCES is reported as contention on some filesystems
+    if (errno == EWOULDBLOCK || errno == EAGAIN || errno == EACCES) {
+        return status::busy;
+    }
+    // flock is not implemented (some NFS/FUSE mounts): report it instead of racing
+    LOG_ERR("%s: file locking failed (%s)\n", __func__, strerror(errno));
+    return status::error;
+#endif
+}
+
+bool common_file_lock::acquire(const std::function<bool()> & keep_waiting) {
+    // the lock file is opened by try_lock, its parent must exist by then
+    if (!lock_path.empty()) {
+        std::error_code ec;
+        auto lock_dir = std::filesystem::path(lock_path).parent_path();
+        if (!lock_dir.empty()) {
+            std::filesystem::create_directories(lock_dir, ec);
+            if (ec) {
+                LOG_ERR("%s: unable to create lock directory '%s': %s\n", __func__, lock_dir.string().c_str(), ec.message().c_str());
+                return false;
+            }
+        }
+    }
+
+    while (true) {
+        switch (try_lock()) {
+            case status::acquired:
+                return true;
+            case status::busy:
+                if (keep_waiting && !keep_waiting()) {
+                    return false;
+                }
+                std::this_thread::sleep_for(std::chrono::milliseconds(100));
+                continue;
+            case status::error:
+                return false;
+        }
+    }
+}
+
+// the lock file is never deleted: another process could then create a fresh
+// file at the same path and lock that
+void common_file_lock::close() {
+#if defined(_WIN32)
+    if (fd >= 0) {
+        if (locked) {
+            OVERLAPPED ov = {};
+            UnlockFileEx((HANDLE) _get_osfhandle(fd), 0, 1, 0, &ov);
+        }
+        _close(fd);
+    }
+#else
+    if (fd >= 0) {
+        if (locked) {
+            flock(fd, LOCK_UN);
+        }
+        ::close(fd);
+    }
+#endif
+    fd     = -1;
+    locked = false;
+}
+
+common_file_lock::~common_file_lock() {
+    close();
+}
@@ -0,0 +1,36 @@
+#pragma once
+
+#include <functional>
+#include <string>
+
+// Cross-process exclusive lock (flock on POSIX, LockFileEx on Windows).
+// Lazy: acquire() opens the lock file, so a read-only cache works if the lock is never acquired.
+class common_file_lock {
+public:
+    common_file_lock(const std::string & lock_path);
+    ~common_file_lock();
+
+    // take the lock, waiting while another process holds it;
+    // keep_waiting is polled before each retry, false cancels the wait.
+    // Also returns false when locking is unavailable.
+    bool acquire(const std::function<bool()> & keep_waiting = {});
+
+    void close();
+
+    common_file_lock(const common_file_lock &) = delete;
+    common_file_lock & operator=(const common_file_lock &) = delete;
+
+private:
+    enum class status {
+        acquired,
+        busy,  // another process holds the lock
+        error, // file locking is unavailable: no exclusion is possible
+    };
+
+    status try_lock();
+
+    std::string lock_path;
+
+    int  fd     = -1;
+    bool locked = false;
+};
@@ -2,6 +2,7 @@

 #include "build-info.h"
 #include "common.h"
+#include "file-lock.h"
 #include "log.h"
 #include "http.h"
 #include "json.h"
@@ -161,6 +162,14 @@ static bool is_valid_oid(const std::string & oid) {
    return is_hex_string(oid, 40) || is_hex_string(oid, 64);
 }

+std::string get_lock_path(const std::string & repo_id, const std::string & oid) {
+    if (!is_valid_repo_id(repo_id) || !is_valid_oid(oid)) {
+        return {};
+    }
+    fs::path lock_path = fs::path(get_cache_path()) / ".locks" / repo_to_folder_name(repo_id) / (oid + ".lock");
+    return lock_path.string();
+}
+
 static bool is_valid_subpath(const fs::path & path, const fs::path & subpath) {
    if (subpath.is_absolute()) {
        return false; // never do a / b with b absolute
@@ -437,7 +446,7 @@ hf_files get_cached_files(const std::string & repo_id) {
    return files;
 }

-std::string finalize_file(const hf_file & file) {
+std::string finalize_file(const hf_file & file, const std::function<bool()> & keep_waiting) {
    static std::atomic<bool> symlinks_disabled{false};

    std::error_code ec;
@@ -452,6 +461,24 @@ std::string finalize_file(const hf_file & file) {
        return file.final_path;
    }

+    // the lock file is shared with huggingface_hub: another process may be
+    // finalizing the same blob
+    common_file_lock lock(get_lock_path(file.repo_id, file.oid));
+    if (!lock.acquire(keep_waiting)) {
+        // without the lock, another process could move the blob while we check it
+        throw std::runtime_error("failed to lock cache for finalization: " + file.final_path);
+    }
+
+    // acquire() only polls keep_waiting while waiting for a busy lock
+    if (keep_waiting && !keep_waiting()) {
+        throw std::runtime_error("download cancelled before finalization: " + file.final_path);
+    }
+
+    // re-check under the lock: the winner may have created the link already
+    if (fs::exists(final_path, ec)) {
+        return file.final_path;
+    }
+
    fs::create_directories(final_path.parent_path(), ec);

    if (!symlinks_disabled) {
@@ -1,5 +1,6 @@
 #pragma once

+#include <functional>
 #include <string>
 #include <vector>

@@ -26,8 +27,14 @@ hf_files get_repo_files(

 hf_files get_cached_files(const std::string & repo_id = {});

-// Create snapshot path (link or move/copy) and return it
-std::string finalize_file(const hf_file & file);
+// Lock file path for a blob, shared with the hf CLI: <cache>/.locks/<repo folder>/<oid>.lock
+// Pure path computation: creates nothing. Returns an empty string on invalid repo or oid.
+std::string get_lock_path(const std::string & repo_id, const std::string & oid);
+
+// Create snapshot path (link or move/copy) and return it.
+// Throws when the lock cannot be taken or keep_waiting returns false.
+// A completed snapshot (final_path present) is returned without the lock, so a read-only cache works.
+std::string finalize_file(const hf_file & file, const std::function<bool()> & keep_waiting = {});

 // Remove the entire cached directory for a repo, returns true if removed
 bool remove_cached_repo(const std::string & repo_id);
@@ -286,6 +286,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
    "SeedOssForCausalLM": "olmo",
    "SmallThinkerForCausalLM": "smallthinker",
    "SmolLM3ForCausalLM": "llama",
+    "SopranoTextModel": "soprano",
    "Spark2_5ForCausalLM": "spark2_5",
    "SolarOpenForCausalLM": "glm",
    "StableLMEpochForCausalLM": "stablelm",
@@ -380,6 +381,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
    "Qwen3ASRForConditionalGeneration": "qwen3vl",
    "Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
    "PocketTTSModel": "pockettts",
+    "SopranoModel": "soprano",
    "Qwen3TTSForConditionalGeneration": "qwen3tts",
    "Qwen3VLForConditionalGeneration": "qwen3vl",
    "Qwen3VLMoeForConditionalGeneration": "qwen3vl",
@@ -0,0 +1,116 @@
+from __future__ import annotations
+
+import json
+from pathlib import Path
+from typing import Any
+
+import torch
+
+from .base import ModelBase, MmprojModel, gguf, logger
+from .qwen import Qwen3Model
+
+
+def _is_soprano_checkpoint(dir_model: Path) -> bool:
+    if not (dir_model / "decoder.pth").is_file() or not (dir_model / "config.json").is_file():
+        return False
+    with open(dir_model / "config.json", encoding="utf-8") as f:
+        config = json.load(f)
+    if config.get("model_type") != "qwen3" or config.get("bos_token_id") != 3 or config.get("eos_token_id") != 3:
+        return False
+    if not (dir_model / "tokenizer.json").is_file():
+        return False
+    with open(dir_model / "tokenizer.json", encoding="utf-8") as f:
+        added = {t["id"]: t["content"] for t in json.load(f).get("added_tokens", [])}
+    return [added.get(i) for i in range(4)] == ["[UNK]", "[TEXT]", "[START]", "[STOP]"]
+
+
+@ModelBase.register_hparams_loader(_is_soprano_checkpoint)
+def _load_soprano_hparams(dir_model: Path) -> dict[str, Any]:
+    logger.info("gguf: detected Soprano checkpoint")
+    hparams = ModelBase.load_hparams(dir_model, False, guess=False)
+    # the text model stays Qwen3, the mmproj holds the decoder
+    hparams["text_config"] = {"architectures": ["SopranoTextModel"]}
+    hparams["architectures"] = ["SopranoModel"]
+    return hparams
+
+
+@ModelBase.register("SopranoTextModel")
+class SopranoTextModel(Qwen3Model):
+    model_arch = gguf.MODEL_ARCH.QWEN3
+
+    def get_vocab_base_pre(self, tokenizer) -> str:
+        return "soprano"
+
+    def set_vocab(self):
+        super().set_vocab()
+        self.gguf_writer.add_unk_token_id(0)
+
+
+@ModelBase.register("SopranoModel")
+@ModelBase.example("ekwek/Soprano-80M", "ekwek/Soprano-1.1-80M")
+class SopranoModel(MmprojModel):
+    has_vision_encoder = False
+    has_audio_encoder = False
+
+    def _load_decoder(self) -> dict[str, torch.Tensor]:
+        state = torch.load(self.dir_model / "decoder.pth", map_location="cpu", mmap=True, weights_only=True)
+        return state.get("state_dict", state)
+
+    def get_audio_config(self):
+        state = self._load_decoder()
+        n_embd, n_inp, _ = state["decoder.embed.weight"].shape
+        n_ff = state["decoder.convnext.0.pwconv1.weight"].shape[0]
+        n_layer = sum(1 for name in state if name.endswith(".pwconv1.weight"))
+        if n_inp != self.n_embd_text:
+            raise ValueError(f"Soprano decoder input dim {n_inp} does not match the backbone dim {self.n_embd_text}")
+        return {"num_hidden_layers": n_layer, "hidden_size": n_embd, "intermediate_size": n_ff}
+
+    def set_gguf_parameters(self):
+        self.gguf_writer.add_file_type(self.ftype)
+        self.gguf_writer.add_clip_has_gen_audio_encoder(True)
+        self.gguf_writer.add_clip_gen_audio_projector_type(gguf.VisionProjectorType.SOPRANO)
+        self.gguf_writer.add_gen_audio_projection_dim(self.n_embd_text)
+        self.gguf_writer.add_gen_audio_embedding_length(self.hparams["hidden_size"])
+        self.gguf_writer.add_gen_audio_feed_forward_length(self.hparams["intermediate_size"])
+        self.gguf_writer.add_gen_audio_block_count(self.hparams["num_hidden_layers"])
+        self.gguf_writer.add_gen_audio_head_count(1)
+        self.gguf_writer.add_gen_audio_attention_layernorm_eps(1e-6)
+
+    def get_tensors(self):
+        yield from self._load_decoder().items()
+
+    def tensor_force_quant(self, name, new_name, bid, n_dims):
+        if "dwconv.weight" in new_name or (name == "decoder.embed.weight" and n_dims == 3):
+            return gguf.GGMLQuantizationType.F16
+        return super().tensor_force_quant(name, new_name, bid, n_dims)
+
+    def modify_tensors(self, data_torch, name, bid):
+        if name == "head.istft.window":
+            # the ISTFT in mtmd builds its own window
+            if not torch.equal(data_torch.float(), torch.hann_window(data_torch.shape[0])):
+                raise ValueError("Soprano requires a periodic Hann window")
+            return
+        if name == "decoder.embed.weight" and data_torch.shape[-1] == 1:
+            data_torch = data_torch.squeeze(-1)
+
+        T = gguf.MODEL_TENSOR
+        if name.startswith("decoder.convnext."):
+            _, _, blk, key, *suffix = name.split(".")
+            block_map = {
+                "dwconv":  T.A_GEN_WAV_UP_DWCONV,
+                "norm":    T.A_GEN_WAV_UP_NORM,
+                "pwconv1": T.A_GEN_WAV_UP_PW1,
+                "pwconv2": T.A_GEN_WAV_UP_PW2,
+                "gamma":   T.A_GEN_WAV_UP_GAMMA,
+            }
+            yield self.format_tensor_name(block_map[key], int(blk), suffix="".join("." + s for s in suffix)), data_torch
+            return
+
+        tensor_map = {
+            "decoder.embed":            T.A_GEN_WAV_INPUT,
+            "decoder.norm":             T.A_GEN_WAV_NORM,
+            "decoder.final_layer_norm": T.A_GEN_WAV_OUTPUT_NORM,
+            "head.out":                 T.A_GEN_WAV_OUTPUT,
+        }
+        key, suffix = name.rsplit(".", 1)
+        yield self.format_tensor_name(tensor_map[key], suffix="." + suffix), data_torch
@@ -926,6 +926,7 @@ extern "C" {
            enum   ggml_type      type);

    // dst[i0, i1, i2] = a[i0, i1, i2] + b[i0, ids[i1, i2]]
+    // an id of -1 adds nothing: dst[i0, i1, i2] = a[i0, i1, i2]
    GGML_API struct ggml_tensor * ggml_add_id(
            struct ggml_context * ctx,
            struct ggml_tensor  * a,
@@ -1477,6 +1478,16 @@ extern "C" {
            enum ggml_prec       prec,
            int                  idx);

+    // [TAG_GGML_PREC]
+    // get the accumulator / src[idx] precision set by the functions above
+    // return GGML_PREC_UNDEFINED when the op does not carry the requested precision
+    GGML_API enum ggml_prec ggml_prec_get_acc(
+            const struct ggml_tensor * a);
+
+    GGML_API enum ggml_prec ggml_prec_get_src(
+            const struct ggml_tensor * a,
+            int                        idx);
+
    // A: k columns, n rows => [ne03, ne02, n, k]
    // B: k columns, m rows  (i.e. we transpose it internally) => [ne03 * x, ne02 * y, m, k]
    // result is n columns, m rows => [ne03 * x, ne02 * y, m, n]
@@ -1498,6 +1509,7 @@ extern "C" {
            enum ggml_op_hint    hint);

    // indirect matrix multiplication
+    // an id of -1 skips the slot and zeroes the matching row of the result
    //
    // [TAG_MUL_MAT_ID_DUPLICATE_ID]
    // note: this op has undefined behaviour if a given `a` is referenced more than once for a given `b`
@@ -1652,7 +1652,12 @@ static void ggml_compute_forward_mul_mat_id(
            for (int id = 0; id < n_ids; ++id) {
                const int32_t i02 = *(const int32_t *) ((const char *) ids->data + iid1*ids->nb[1] + id*ids->nb[0]);

-                assert(i02 >= 0 && i02 < n_as);
+                assert(i02 == -1 || (i02 >= 0 && i02 < n_as));
+
+                if (i02 == -1) {
+                    memset((char *) dst->data + id*nb1 + iid1*nb2, 0, ne0*sizeof(float));
+                    continue;
+                }

                MMID_MATRIX_ROW(i02, matrix_row_counts[i02]) = (struct mmid_row_mapping) {id, iid1};
                matrix_row_counts[i02] += 1;
@@ -743,12 +743,17 @@ static void ggml_compute_forward_add_id_f32(
        // src1 indices
        const int i11 = *(int32_t *) ((char *) src2->data + i1*nb20 + i2*nb21);

-        GGML_ASSERT(i11 >= 0 && i11 < ne11);
+        GGML_ASSERT(i11 == -1 || (i11 >= 0 && i11 < ne11));

-        ggml_vec_add_f32(ne0,
-                (float *) ((char *) dst->data  + i3*nb3  + i2*nb2  + i1*nb1 ),
-                (float *) ((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01),
-                (float *) ((char *) src1->data + i11*nb11));
+        float       * dst_row  = (float *) ((char *) dst->data  + i3*nb3  + i2*nb2  + i1*nb1 );
+        const float * src0_row = (float *) ((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01);
+
+        if (i11 == -1) {
+            memcpy(dst_row, src0_row, ne0*sizeof(float));
+            continue;
+        }
+
+        ggml_vec_add_f32(ne0, dst_row, src0_row, (float *) ((char *) src1->data + i11*nb11));
    }
 }

@@ -4855,7 +4855,12 @@ template <typename BLOC_TYPE, int64_t INTER_SIZE, int64_t NB_COLS, ggml_type PAR
                    const int32_t i02 =
                        *(const int32_t *) ((const char *) ids->data + iid1 * ids->nb[1] + id * ids->nb[0]);

-                    GGML_ASSERT(i02 >= 0 && i02 < n_as);
+                    GGML_ASSERT(i02 == -1 || (i02 >= 0 && i02 < n_as));
+
+                    if (i02 == -1) {
+                        memset((char *) dst->data + id*nb1 + iid1*nb2, 0, ne0*sizeof(float));
+                        continue;
+                    }

                    MMID_MATRIX_ROW(i02, matrix_row_counts[i02]) = { id, iid1 };
                    matrix_row_counts[i02] += 1;
@@ -682,7 +682,12 @@ template <typename BLOC_TYPE, int64_t INTER_SIZE, int64_t NB_COLS> class tensor_
                    const int32_t i02 =
                        *(const int32_t *) ((const char *) ids->data + iid1 * ids->nb[1] + id * ids->nb[0]);

-                    GGML_ASSERT(i02 >= 0 && i02 < n_as);
+                    GGML_ASSERT(i02 == -1 || (i02 >= 0 && i02 < n_as));
+
+                    if (i02 == -1) {
+                        memset((char *) dst->data + id * nb1 + iid1 * nb2, 0, ne0 * sizeof(float));
+                        continue;
+                    }

                    MMID_MATRIX_ROW(i02, matrix_row_counts[i02]) = { id, iid1 };
                    matrix_row_counts[i02] += 1;
@@ -1409,14 +1409,14 @@ namespace ggml_cuda_mma {
            tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) {
 #if defined(AMD_WMMA_AVAILABLE)
 #if defined(RDNA4)
-        using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16;
+        using bf16x8_t = __attribute__((ext_vector_type(8))) short;
        using floatx8_t = __attribute__((ext_vector_type(8))) float;
        floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
        const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
        const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
        acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
 #elif defined(RDNA3)
-        using bf16x16_t = __attribute__((ext_vector_type(16))) __bf16;
+        using bf16x16_t = __attribute__((ext_vector_type(16))) short;
        using floatx8_t = __attribute__((ext_vector_type(8))) float;
        floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
        const bf16x16_t& a_frag = reinterpret_cast<const bf16x16_t&>(A.x[0]);
@@ -60,10 +60,10 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
            const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
            const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);

-            const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
-            const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
-            const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
-            const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
+            const int v0 = ggml_hip_perm(0x000001FF, 0x000001FF, q0_indices);
+            const int v1 = ggml_hip_perm(0x000001FF, 0x000001FF, q1_indices);
+            const int v2 = ggml_hip_perm(0x000001FF, 0x000001FF, q2_indices);
+            const int v3 = ggml_hip_perm(0x000001FF, 0x000001FF, q3_indices);
 #else
            // unpack crumbs into nibble indices
            const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8,  9, 12, 13]
@@ -162,8 +162,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
            const uint32_t qx_indices = (q & 0x03) | ((q & 0x0C) << 6) | ((q & 0x30) << 12) | ((q & 0xC0) << 18);
            const uint32_t qy_bits    = q >> 8;
            const uint32_t qy_indices = (qy_bits & 0x03) | ((qy_bits & 0x0C) << 6) | ((qy_bits & 0x30) << 12) | ((qy_bits & 0xC0) << 18);
-            const int qx = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qx_indices);
-            const int qy = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qy_indices);
+            const int      qx         = ggml_hip_perm(0x020100FF, 0x020100FF, qx_indices);
+            const int      qy         = ggml_hip_perm(0x020100FF, 0x020100FF, qy_indices);
 #else
            // unpack even and odd crumbs into byte values
            const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
@@ -435,6 +435,7 @@ template <ggml_type type, int J, bool fallback, ggml_prec prec_src1> static __de
    constexpr int warp_size = ggml_cuda_get_physical_warp_size();
    constexpr int nwarps    = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size;
    constexpr int I         = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1);
+    static_assert(I % warp_size == 0, "DP4A MMQ tile height must be a multiple of the physical warp size");

    const bool y_scale_used = y_scale != nullptr;

@@ -40,18 +40,18 @@ static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, con
    const uint32_t q_odd  = (q4 >> 4);

    // Perform lookups in the lower half of the table (indices 0-7).
-    uint32_t v_even_low = __builtin_amdgcn_perm(values[1], values[0], q_even & 0x07070707);
-    uint32_t v_odd_low = __builtin_amdgcn_perm(values[1], values[0], q_odd & 0x07070707);
+    uint32_t v_even_low = ggml_hip_perm(values[1], values[0], q_even & 0x07070707);
+    uint32_t v_odd_low  = ggml_hip_perm(values[1], values[0], q_odd & 0x07070707);

    // Perform lookups in the upper half of the table (indices 8-15).
-    uint32_t v_even_high = __builtin_amdgcn_perm(values[3], values[2], q_even & 0x07070707);
-    uint32_t v_odd_high = __builtin_amdgcn_perm(values[3], values[2], q_odd & 0x07070707);
+    uint32_t v_even_high = ggml_hip_perm(values[3], values[2], q_even & 0x07070707);
+    uint32_t v_odd_high  = ggml_hip_perm(values[3], values[2], q_odd & 0x07070707);

    // Select between the low and high results based on the MSB of each index nibble.
    uint32_t mask_even = 0x03020100 | ((q_even & 0x08080808) >> 1);
-    uint32_t res_x = __builtin_amdgcn_perm(v_even_high, v_even_low, mask_even);
+    uint32_t res_x     = ggml_hip_perm(v_even_high, v_even_low, mask_even);
    uint32_t mask_odd = 0x03020100 | ((q_odd & 0x08080808) >> 1);
-    uint32_t res_y = __builtin_amdgcn_perm(v_odd_high, v_odd_low, mask_odd);
+    uint32_t res_y     = ggml_hip_perm(v_odd_high, v_odd_low, mask_odd);

    return make_int2(res_x, res_y);
 #elif !defined(GGML_USE_MUSA)
@@ -720,10 +720,10 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
        const uint32_t q2_indices = (q2_bits & 0x01) | ((q2_bits & 0x02) << 7) | ((q2_bits & 0x04) << 14) | ((q2_bits & 0x08) << 21);
        const uint32_t q3_indices = (q3_bits & 0x01) | ((q3_bits & 0x02) << 7) | ((q3_bits & 0x04) << 14) | ((q3_bits & 0x08) << 21);

-        const int v0 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q0_indices);
-        const int v1 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q1_indices);
-        const int v2 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q2_indices);
-        const int v3 = __builtin_amdgcn_perm(0x000001FF, 0x000001FF, q3_indices);
+        const int v0 = ggml_hip_perm(0x000001FF, 0x000001FF, q0_indices);
+        const int v1 = ggml_hip_perm(0x000001FF, 0x000001FF, q1_indices);
+        const int v2 = ggml_hip_perm(0x000001FF, 0x000001FF, q2_indices);
+        const int v3 = ggml_hip_perm(0x000001FF, 0x000001FF, q3_indices);
 #else
        // unpack crumbs into nibble indices
        const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8,  9, 12, 13]
@@ -777,8 +777,8 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
        const uint32_t qx_indices = (q & 0x03) | ((q & 0x0C) << 6) | ((q & 0x30) << 12) | ((q & 0xC0) << 18);
        const uint32_t qy_bits    = q >> 8;
        const uint32_t qy_indices = (qy_bits & 0x03) | ((qy_bits & 0x0C) << 6) | ((qy_bits & 0x30) << 12) | ((qy_bits & 0xC0) << 18);
-        const int qx = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qx_indices);
-        const int qy = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qy_indices);
+        const int      qx         = ggml_hip_perm(0x020100FF, 0x020100FF, qx_indices);
+        const int      qy         = ggml_hip_perm(0x020100FF, 0x020100FF, qy_indices);
 #else
        // unpack even and odd crumbs into byte values
        const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
@@ -311,3 +311,22 @@ static __device__ __forceinline__ unsigned int __vcmpne4(unsigned int a, unsigne

    return (ne_any_bit >> 7) * 0xff;
 }
+
+// v_perm_b32 with selector bytes in [0, 7]: 0-3 pick bytes of lo, 4-7 pick bytes of hi.
+static __device__ __forceinline__ uint32_t ggml_hip_perm(const uint32_t hi, const uint32_t lo, const uint32_t sel) {
+#if __has_builtin(__builtin_amdgcn_is_invocable)
+    if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_perm)) {
+        return __builtin_amdgcn_perm(hi, lo, sel);
+    }
+    const uint64_t src = ((uint64_t) hi << 32) | lo;
+    uint32_t res = 0;
+#pragma unroll
+    for (int i = 0; i < 4; ++i) {
+        const uint32_t byte_sel = (sel >> (8 * i)) & 0x07;
+        res |= (uint32_t) ((src >> (8 * byte_sel)) & 0xFF) << (8 * i);
+    }
+    return res;
+#else
+    return __builtin_amdgcn_perm(hi, lo, sel);
+#endif // __has_builtin(__builtin_amdgcn_is_invocable)
+}
@@ -73,6 +73,7 @@ htp_op_code htp_op_remap(const struct ggml_tensor * t) {
                case GGML_UNARY_OP_RELU:       return HTP_OP_UNARY_RELU;
                case GGML_UNARY_OP_STEP:       return HTP_OP_UNARY_STEP;
                case GGML_UNARY_OP_CEIL:       return HTP_OP_UNARY_CEIL;
+                case GGML_UNARY_OP_SGN:        return HTP_OP_UNARY_SGN;
                default:                       break;
            }
            break;
@@ -1392,6 +1393,52 @@ static bool htp_kparams_solve_binary(
    return true;
 }

+// collapse contiguous dims of elementwise ops so that rows fill at least one HVX vector
+static void htp_unary_collapse_dims(struct ggml_tensor * src, struct ggml_tensor * dst, size_t elem_size) {
+    if (src->nb[0] != elem_size || dst->nb[0] != elem_size) {
+        return;
+    }
+
+    while (src->ne[1] > 1 || src->ne[2] > 1 || src->ne[3] > 1) {
+        const size_t row_size    = (size_t) src->ne[0] * elem_size;
+        const bool   rows_contig = src->nb[1] == row_size && dst->nb[1] == (size_t) dst->ne[0] * elem_size;
+        const bool   can_merge   = (src->ne[1] == 1 && dst->ne[1] == 1) ||
+                                   (rows_contig && (src->ne[0] == 1 || (row_size < 128 && row_size * src->ne[1] <= 2048)));
+        if (!can_merge) {
+            break;
+        }
+        src->ne[0] *= src->ne[1];
+        src->ne[1]  = src->ne[2];
+        src->ne[2]  = src->ne[3];
+        src->ne[3]  = 1;
+        dst->ne[0] *= dst->ne[1];
+        dst->ne[1]  = dst->ne[2];
+        dst->ne[2]  = dst->ne[3];
+        dst->ne[3]  = 1;
+        src->nb[1]  = src->nb[2];
+        src->nb[2]  = src->nb[3];
+        dst->nb[1]  = dst->nb[2];
+        dst->nb[2]  = dst->nb[3];
+    }
+
+    // ne1 too big to merge in full: merge it in chunks of k rows (k divides ne1)
+    const size_t row_size = (size_t) src->ne[0] * elem_size;
+    if (src->ne[1] > 1 && row_size < 128 && src->nb[1] == row_size && dst->nb[1] == row_size) {
+        int64_t k = 2048 / row_size;
+        while (k > 1 && src->ne[1] % k != 0) {
+            k--;
+        }
+        if (k > 1) {
+            src->ne[0] *= k;
+            src->ne[1] /= k;
+            src->nb[1] *= k;
+            dst->ne[0] *= k;
+            dst->ne[1] /= k;
+            dst->nb[1] *= k;
+        }
+    }
+}
+
 static bool htp_kparams_solve_unary_impl(
    const struct ggml_hexagon_session * sess,
    uint32_t op,
@@ -1435,6 +1482,7 @@ static bool htp_kparams_solve_unary_impl(
                case HTP_OP_UNARY_ABS:
                case HTP_OP_UNARY_STEP:
                case HTP_OP_UNARY_CEIL:
+                case HTP_OP_UNARY_SGN:
                    break;
                default:
                    return false;
@@ -1451,9 +1499,14 @@ static bool htp_kparams_solve_unary_impl(
    struct ggml_tensor src0_c = *src0;
    struct ggml_tensor dst_c  = *dst;
    if (htp_unary_is_elementwise(op)) {
-        HTP_UNARY_COLLAPSE_DIMS(src0_c.ne, src0_c.nb, dst_c.ne, dst_c.nb, elem_size);
+        htp_unary_collapse_dims(&src0_c, &dst_c, elem_size);
        src0 = &src0_c;
        dst  = &dst_c;
+        for (int i = 0; i < HTP_OP_MAX_DIMS; i++) {
+            kparams->ne[i]      = (uint32_t) src0->ne[i];
+            kparams->src0_nb[i] = (uint32_t) src0->nb[i];
+            kparams->dst_nb[i]  = (uint32_t) dst->nb[i];
+        }
    }

    const size_t dst_data_row_size  = dst->ne[0]  * ggml_type_size(dst->type);
@@ -113,6 +113,7 @@ enum htp_op_code {
    HTP_OP_GLU_REGLU,
    HTP_OP_GLU_SIGMOID,
    HTP_OP_UNARY_CEIL,
+    HTP_OP_UNARY_SGN,

    HTP_OP_INVALID
 };
@@ -677,6 +677,56 @@ static inline void hvx_step_f32_aa(uint8_t * restrict dst, const uint8_t * restr
    }
 }

+static inline void hvx_sgn_f32_aa(uint8_t * restrict dst, const uint8_t * restrict src, uint32_t n) {
+    assert((unsigned long) dst % 128 == 0);
+    assert((unsigned long) src % 128 == 0);
+
+    HVX_Vector * restrict vdst = (HVX_Vector *) dst;
+    HVX_Vector * restrict vsrc = (HVX_Vector *) src;
+
+    const uint32_t elem_size = sizeof(float);
+    const uint32_t epv       = 128 / elem_size;
+    const uint32_t nvec      = n / epv;
+    const uint32_t nloe      = n % epv;
+
+    uint32_t i = 0;
+
+    _Pragma("unroll(4)")
+    for (; i < nvec; i++) {
+        vdst[i] = hvx_vec_sgn_f32(vsrc[i]);
+    }
+    if (nloe) {
+        HVX_Vector v = hvx_vec_sgn_f32(vsrc[i]);
+        hvx_vec_store_a((void *) &vdst[i], nloe * elem_size, v);
+    }
+}
+
+static inline void hvx_sgn_f16_aa(uint8_t * restrict dst, const uint8_t * restrict src, uint32_t n) {
+    assert((unsigned long) dst % 128 == 0);
+    assert((unsigned long) src % 128 == 0);
+
+    HVX_Vector * restrict vdst = (HVX_Vector *) dst;
+    HVX_Vector * restrict vsrc = (HVX_Vector *) src;
+
+    const uint32_t elem_size = sizeof(_Float16);
+    const uint32_t epv       = 128 / elem_size;
+    const uint32_t nvec      = n / epv;
+    const uint32_t nloe      = n % epv;
+
+    uint32_t i = 0;
+
+    _Pragma("unroll(4)")
+    for (; i < nvec; i++) {
+        HVX_VectorPair p = hvx_vec_f16_to_f32(vsrc[i]);
+        vdst[i] = hvx_vec_f32_to_f16(hvx_vec_sgn_f32(Q6_V_lo_W(p)), hvx_vec_sgn_f32(Q6_V_hi_W(p)));
+    }
+    if (nloe) {
+        HVX_VectorPair p = hvx_vec_f16_to_f32(vsrc[i]);
+        HVX_Vector v = hvx_vec_f32_to_f16(hvx_vec_sgn_f32(Q6_V_lo_W(p)), hvx_vec_sgn_f32(Q6_V_hi_W(p)));
+        hvx_vec_store_a((void *) &vdst[i], nloe * elem_size, v);
+    }
+}
+
 static inline void hvx_step_f16_aa(uint8_t * restrict dst, const uint8_t * restrict src, uint32_t n) {
    assert((unsigned long) dst % 128 == 0);
    assert((unsigned long) src % 128 == 0);
@@ -118,6 +118,17 @@ static inline HVX_Vector hvx_vec_step_f32(HVX_Vector v) {
    return Q6_V_vmux_QVV(q, one, zero);
 }

+static inline HVX_Vector hvx_vec_sgn_f32(HVX_Vector v) {
+    const HVX_Vector zero       = Q6_V_vzero();
+    const HVX_Vector one        = hvx_vec_splat_f32(1.0f);
+    const HVX_Vector sign_mask  = Q6_V_vsplat_R(0x80000000);
+    const HVX_Vector mag        = hvx_vec_abs_f32(v);
+    const HVX_Vector signed_one = Q6_V_vor_VV(one, Q6_V_vand_VV(v, sign_mask));
+    // The ordered magnitude comparison maps NaN and both signed zeros to zero.
+    const HVX_VectorPred nonzero = Q6_Q_vcmp_gt_VsfVsf(mag, zero);
+    return Q6_V_vmux_QVV(nonzero, signed_one, zero);
+}
+
 static inline HVX_Vector hvx_vec_step_f16(HVX_Vector v) {
    const HVX_Vector zero = Q6_V_vzero();
    const HVX_Vector one  = hvx_vec_splat_f16((_Float16) 1.0f);
@@ -946,6 +946,7 @@ static int execute_op(struct htp_ops_context * octx) {
        case HTP_OP_UNARY_RELU:
        case HTP_OP_UNARY_STEP:
        case HTP_OP_UNARY_CEIL:
+        case HTP_OP_UNARY_SGN:
        case HTP_OP_L2_NORM:
        case HTP_OP_HADAMARD:
            return op_unary(octx);
@@ -1413,10 +1413,10 @@ static void hvx_mv_id_tiled(unsigned int nth, unsigned int ith, void * data) {

    for (uint32_t ie1 = 0; ie1 < n_aids; ++ie1) {  // for each expert
        const int32_t eid = *(const int32_t *) ((const uint8_t *) src2->data + ie1 * src2->nb[0]);
-        if (eid < 0) {
+        if (eid == -1) {
            continue;
        }
-        assert(eid < (int32_t) n_ids);
+        assert(eid >= 0 && eid < (int32_t) n_ids);

        const dma_addr_t src0_row = src0->data + eid * nb02;
        const uint8_t * restrict act_col = (const uint8_t *) act_data;
@@ -4166,10 +4166,10 @@ static inline void scan_expert_ids_n(
        const int32_t * row_ptr = (const int32_t *) (ids_data + iid1 * ids_nb1);
        for (uint32_t id = 0; id < n_ids; ++id) {
            const int32_t i02 = row_ptr[id];
-            if (i02 < 0) {
+            if (i02 == -1) {
                continue;
            }
-            assert(i02 < n_as);
+            assert(i02 >= 0 && i02 < n_as);

            if (matrix_rows) {
                matrix_rows[i02 * mapping_stride + counts[i02]] = (struct mmid_row_mapping) { id, iid1 };
@@ -4179,6 +4179,28 @@ static inline void scan_expert_ids_n(
    }
 }

+// the matmuls only write the rows that an expert owns, so the rows of the skipped slots are zeroed here
+static inline void zero_skipped_dst_rows(
+    const struct htp_tensor * ids,
+    uint32_t n_ids,
+    const struct htp_tensor * dst,
+    uint32_t ne0,
+    size_t nb1,
+    size_t nb2
+) {
+    const uint8_t * ids_data = (const uint8_t *) ids->data;
+
+    for (uint32_t iid1 = 0; iid1 < ids->ne[1]; ++iid1) {
+        for (uint32_t id = 0; id < n_ids; ++id) {
+            const int32_t i02 = *(const int32_t *) (ids_data + iid1 * ids->nb[1] + id * ids->nb[0]);
+            if (i02 != -1) {
+                continue;
+            }
+            memset((uint8_t *) (dst->data + id * nb1 + iid1 * nb2), 0, ne0 * sizeof(float));
+        }
+    }
+}
+
 static inline void scan_expert_ids(
    const struct htp_tensor * ids,
    uint32_t n_ids,
@@ -4204,10 +4226,10 @@ static inline void scan_expert_ids(
            const int32_t * row_ptr = (const int32_t *) (ids_data + iid1 * ids_nb1);
            for (uint32_t id = 0; id < n_ids; ++id) {
                const int32_t i02 = *(const int32_t *) ((const uint8_t *) row_ptr + id * ids_nb0);
-                if (i02 < 0) {
+                if (i02 == -1) {
                    continue;
                }
-                assert(i02 < n_as);
+                assert(i02 >= 0 && i02 < n_as);

                if (matrix_rows) {
                    matrix_rows[i02 * mapping_stride + counts[i02]] = (struct mmid_row_mapping) { id, iid1 };
@@ -4287,6 +4309,8 @@ int op_matmul_id(struct htp_ops_context * octx) {
        scan_expert_ids(ids, n_ids, n_as, matrix_row_counts, matrix_rows, mapping_stride);
    }

+    zero_skipped_dst_rows(ids, n_ids, dst, ne0, nb1, nb2);
+
    mmctx->matrix_row_counts    = matrix_row_counts;
    mmctx->matrix_rows          = matrix_rows;
    mmctx->mapping_stride       = mapping_stride;
@@ -440,6 +440,20 @@ static void ceil_f16(const void * restrict src,
    }
 }

+static void sgn_f16(const void * restrict src,
+                    void * restrict dst,
+                    const uint32_t num_rows,
+                    const struct htp_unary_context * uctx) {
+    htp_unary_op_preamble;
+
+    for (uint32_t ir = 0; ir < num_rows; ir++) {
+        const uint8_t * restrict src_local = (const uint8_t *)src + (ir * src0_row_size_aligned);
+        uint8_t * restrict dst_local       = (uint8_t *)dst + (ir * dst_row_size_aligned);
+
+        hvx_sgn_f16_aa((uint8_t *) dst_local, (const uint8_t *) src_local, ne0);
+    }
+}
+
 static void l2_norm_f16(const void * restrict src,
                        void * restrict dst,
                        const uint32_t num_rows,
@@ -976,6 +990,20 @@ static void ceil_f32(const void * restrict src,
    }
 }

+static void sgn_f32(const void * restrict src,
+                    void * restrict dst,
+                    const uint32_t num_rows,
+                    const struct htp_unary_context * uctx) {
+    htp_unary_op_preamble;
+
+    for (uint32_t ir = 0; ir < num_rows; ir++) {
+        const uint8_t * restrict src_local = (const uint8_t *)src + (ir * src0_row_size_aligned);
+        uint8_t * restrict dst_local       = (uint8_t *)dst + (ir * dst_row_size_aligned);
+
+        hvx_sgn_f32_aa(dst_local, src_local, ne0);
+    }
+}
+
 static void log_f32(const void * restrict src,
                    void * restrict dst,
                    const uint32_t num_rows,
@@ -1094,6 +1122,11 @@ static void tile_ceil_f32(void * restrict dst, const void * restrict src, uint32
    hvx_ceil_f32_aa((uint8_t *) dst, (const uint8_t *) src, tw);
 }

+static void tile_sgn_f32(void * restrict dst, const void * restrict src, uint32_t tw, const struct htp_unary_context * uctx) {
+    (void) uctx;
+    hvx_sgn_f32_aa((uint8_t *) dst, (const uint8_t *) src, tw);
+}
+
 static void tri_apply_tile_f32(const void * restrict src, void * restrict dst,
                               uint32_t tile_elems, uint32_t col_start, uint32_t i01,
                               uint32_t ne0, int32_t ttype) {
@@ -1738,6 +1771,7 @@ static int execute_op_unary(struct htp_ops_context * octx) {
        case HTP_OP_UNARY_RELU:      op_type = "relu-f32";                                   break;
        case HTP_OP_UNARY_STEP:      op_type = is_f16 ? "step-f16"     : "step-f32";         break;
        case HTP_OP_UNARY_CEIL:      op_type = is_f16 ? "ceil-f16"     : "ceil-f32";         break;
+        case HTP_OP_UNARY_SGN:       op_type = is_f16 ? "sgn-f16"      : "sgn-f32";          break;
        case HTP_OP_L2_NORM:         op_type = is_f16 ? "l2norm-f16"   : "l2norm-f32";       break;
        case HTP_OP_TRI:             op_type = "tri-f32";                                    break;
        case HTP_OP_HADAMARD:        op_type = is_f16 ? "hadamard-f16" : "hadamard-f32";     break;
@@ -1752,9 +1786,9 @@ static int execute_op_unary(struct htp_ops_context * octx) {
        return HTP_STATUS_INVAL_PARAMS;
    }

-    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
    const size_t src0_elem_size = is_f16 ? sizeof(_Float16) : sizeof(float);
    const size_t dst_elem_size  = (dst->type == HTP_TYPE_F16) ? sizeof(_Float16) : sizeof(float);
+    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
    const size_t src0_data_row_size = src0->ne[0] * src0_elem_size;
    const size_t dst_data_row_size  = dst->ne[0]  * dst_elem_size;

@@ -1872,6 +1906,7 @@ static int execute_op_unary(struct htp_ops_context * octx) {
            case HTP_OP_UNARY_RELU:      compute_func = (void *) tile_relu_f32;           break;
            case HTP_OP_UNARY_STEP:      compute_func = (void *) tile_step_f32;           break;
            case HTP_OP_UNARY_CEIL:      compute_func = (void *) tile_ceil_f32;           break;
+            case HTP_OP_UNARY_SGN:       compute_func = (void *) tile_sgn_f32;            break;
            case HTP_OP_TRI:
                task_func    = unary_thread_tiled_tri_f32;
                compute_func = (void *) tri_apply_tile_f32;
@@ -1892,6 +1927,7 @@ static int execute_op_unary(struct htp_ops_context * octx) {
            case HTP_OP_UNARY_LOG:       compute_func = (void *) log_f16;                 break;
            case HTP_OP_UNARY_STEP:      compute_func = (void *) step_f16;                break;
            case HTP_OP_UNARY_CEIL:      compute_func = (void *) ceil_f16;                break;
+            case HTP_OP_UNARY_SGN:       compute_func = (void *) sgn_f16;                 break;
            case HTP_OP_HADAMARD:        compute_func = (void *) hadamard_f16;            break;
            default:                     break;
        }
@@ -1923,6 +1959,7 @@ static int execute_op_unary(struct htp_ops_context * octx) {
            case HTP_OP_UNARY_RELU:      compute_func = (void *) relu_f32;                break;
            case HTP_OP_UNARY_STEP:      compute_func = (void *) step_f32;                break;
            case HTP_OP_UNARY_CEIL:      compute_func = (void *) ceil_f32;                break;
+            case HTP_OP_UNARY_SGN:       compute_func = (void *) sgn_f32;                 break;
            case HTP_OP_L2_NORM:         compute_func = (void *) l2_norm_f32;             break;
            case HTP_OP_HADAMARD:        compute_func = (void *) hadamard_f32;            break;
            case HTP_OP_TRI:
@@ -1952,10 +1989,14 @@ int op_unary(struct htp_ops_context * octx) {
                return execute_op_unary(octx);
            }

+            const struct htp_unary_kernel_params * kparams = (const struct htp_unary_kernel_params *) octx->kernel_params;
+
            struct htp_tensor src0_c = *octx->src[0];
            struct htp_tensor dst_c  = *octx->dst;
-            const size_t elem_size = (src0_c.type == HTP_TYPE_F16) ? sizeof(_Float16) : sizeof(float);
-            HTP_UNARY_COLLAPSE_DIMS(src0_c.ne, src0_c.nb, dst_c.ne, dst_c.nb, elem_size);
+            memcpy(src0_c.ne, kparams->ne,      sizeof(src0_c.ne));
+            memcpy(dst_c.ne,  kparams->ne,      sizeof(dst_c.ne));
+            memcpy(src0_c.nb, kparams->src0_nb, sizeof(src0_c.nb));
+            memcpy(dst_c.nb,  kparams->dst_nb,  sizeof(dst_c.nb));

            const struct htp_tensor * orig_src0 = octx->src[0];
            const struct htp_tensor * orig_dst  = octx->dst;
@@ -24,6 +24,11 @@ struct htp_unary_kernel_params {
    struct fastdiv_values div_ne02;
    struct fastdiv_values div_ne012;
    struct fastdiv_values div_tpr;
+
+    // collapsed shape of elementwise ops
+    uint32_t  ne[HTP_OP_MAX_DIMS];
+    uint32_t  src0_nb[HTP_OP_MAX_DIMS];
+    uint32_t  dst_nb[HTP_OP_MAX_DIMS];
 };

 #if defined(__cplusplus)
@@ -56,6 +61,7 @@ static inline bool htp_op_is_unary(uint32_t opcode) {
        case HTP_OP_UNARY_RELU:
        case HTP_OP_UNARY_STEP:
        case HTP_OP_UNARY_CEIL:
+        case HTP_OP_UNARY_SGN:
        case HTP_OP_L2_NORM:
        case HTP_OP_TRI:
        case HTP_OP_HADAMARD:
@@ -79,31 +85,6 @@ static inline bool htp_unary_is_elementwise(uint32_t opcode) {
    }
 }

-#define HTP_UNARY_COLLAPSE_DIMS(src_ne, src_nb, dst_ne, dst_nb, elem_size)                   \
-    do {                                                                                     \
-        while (((src_ne)[1] > 1 || (src_ne)[2] > 1 || (src_ne)[3] > 1) &&                    \
-               (size_t) (src_nb)[0] == (elem_size) &&                                        \
-               (size_t) (dst_nb)[0] == (elem_size) &&                                        \
-               (((src_ne)[1] == 1 && (dst_ne)[1] == 1) ||                                    \
-                ((size_t) (src_nb)[1] == (size_t) (src_ne)[0] * (size_t) (src_nb)[0] &&      \
-                 (size_t) (dst_nb)[1] == (size_t) (dst_ne)[0] * (size_t) (dst_nb)[0] &&      \
-                 ((src_ne)[0] == 1 ||                                                        \
-                  (uint64_t) (src_ne)[0] * (uint64_t) (src_ne)[1] * (elem_size) <= 128)))) { \
-            (src_ne)[0] *= (src_ne)[1];                                                      \
-            (src_ne)[1]  = (src_ne)[2];                                                      \
-            (src_ne)[2]  = (src_ne)[3];                                                      \
-            (src_ne)[3]  = 1;                                                                \
-            (dst_ne)[0] *= (dst_ne)[1];                                                      \
-            (dst_ne)[1]  = (dst_ne)[2];                                                      \
-            (dst_ne)[2]  = (dst_ne)[3];                                                      \
-            (dst_ne)[3]  = 1;                                                                \
-            (src_nb)[1]  = (src_nb)[2];                                                      \
-            (src_nb)[2]  = (src_nb)[3];                                                      \
-            (dst_nb)[1]  = (dst_nb)[2];                                                      \
-            (dst_nb)[2]  = (dst_nb)[3];                                                      \
-        }                                                                                    \
-    } while (0)
-
 struct htp_unary_vtcm_layout {
    size_t total_bytes;
    size_t off_src0;
@@ -1283,6 +1283,15 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_map0(g
    return res;
 }

+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_zero(ggml_metal_library_t lib) {
+    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, "kernel_mul_mm_id_zero");
+    if (!res.pipeline) {
+        res = ggml_metal_library_compile_pipeline(lib, "kernel_mul_mm_id_zero", "kernel_mul_mm_id_zero", nullptr);
+    }
+
+    return res;
+}
+
 ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id(ggml_metal_library_t lib, const ggml_tensor * op) {
    char base[256];
    char name[256];
@@ -140,6 +140,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_mm
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm            (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv            (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_map0    (ggml_metal_library_t lib, int ne02, int ne20);
+struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_zero    (ggml_metal_library_t lib);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax(ggml_metal_library_t lib);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax_part(ggml_metal_library_t lib);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id         (ggml_metal_library_t lib, const struct ggml_tensor * op);
@@ -572,6 +572,13 @@ typedef struct {
    uint64_t nb21;
 } ggml_metal_kargs_mul_mm_id_map0;

+typedef struct {
+    int32_t  ne0;
+    uint64_t nb1;
+    uint64_t nb2;
+    uint64_t nb21;
+} ggml_metal_kargs_mul_mm_id_zero;
+
 typedef struct {
    int32_t  ne00;
    int32_t  ne01;
@@ -2935,6 +2935,25 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
            ggml_metal_encoder_dispatch_threadgroups(enc, 1, 1, 1, ne02, 1, 1);
        }

+        // zero the rows that no expert owns - disjoint from what kernel_mul_mm_id writes
+        {
+            ggml_metal_kargs_mul_mm_id_zero args = {
+                /*.ne0  =*/ ne0,
+                /*.nb1  =*/ nb1,
+                /*.nb2  =*/ nb2,
+                /*.nb21 =*/ nb21,
+            };
+
+            auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_zero(lib);
+
+            ggml_metal_encoder_set_pipeline(enc, pipeline);
+            ggml_metal_encoder_set_bytes   (enc, &args, sizeof(args), 0);
+            ggml_metal_encoder_set_buffer  (enc, bid_src2, 1);
+            ggml_metal_encoder_set_buffer  (enc, bid_dst,  2);
+
+            ggml_metal_encoder_dispatch_threadgroups(enc, ne20, ne21, 1, 32, 1, 1);
+        }
+
        ggml_metal_op_concurrency_reset(ctx);

        if (use_amax) {
@@ -185,6 +185,14 @@ kernel void kernel_add_id(

    device       float * dst_row  = (device       float *)((device char *)dst  +  i1*nb1       + i2*nb2);
    device const float * src0_row = (device const float *)((device char *)src0 +  i1*args.nb01 + i2*args.nb02);
+
+    if (i11 == -1) {
+        for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
+            dst_row[i0] = src0_row[i0];
+        }
+        return;
+    }
+
    device const float * src1_row = (device const float *)((device char *)src1 + i11*args.nb11);

    for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
@@ -391,8 +391,10 @@ kernel void kernel_mul_mm_id_map0(
        if (i21 + tpitg < args.ne21) {
            device const int32_t * src2_i32 = (device const int32_t *) (src2 + (i21 + tpitg)*args.nb21);

+            // 0xFFFF marks a skipped slot, it never matches an expert since ne02 fits in a threadgroup
            FOR_UNROLL (short i20 = 0; i20 < ne20; i20++) {
-                sids[tpitg*ne20 + i20] = src2_i32[i20];
+                const int32_t id = src2_i32[i20];
+                sids[tpitg*ne20 + i20] = id == -1 ? 0xFFFF : (uint16_t) id;
            }
        }

@@ -506,6 +508,30 @@ kernel void kernel_mul_mm_id_map0(
    }
 }

+// kernel_mul_mm_id only writes the rows that an expert owns, so the skipped rows are zeroed here
+kernel void kernel_mul_mm_id_zero(
+        constant ggml_metal_kargs_mul_mm_id_zero & args,
+        device const char * src2,
+        device       char * dst,
+        uint3   tgpig[[threadgroup_position_in_grid]],
+        ushort3 tpitg[[thread_position_in_threadgroup]],
+        ushort3   ntg[[threads_per_threadgroup]]) {
+    const int i20 = tgpig.x; // slot
+    const int i21 = tgpig.y; // token
+
+    const int32_t id = ((device const int32_t *) (src2 + i21*args.nb21))[i20];
+
+    if (id != -1) {
+        return;
+    }
+
+    device float * dst_row = (device float *) (dst + i20*args.nb1 + i21*args.nb2);
+
+    for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
+        dst_row[i0] = 0.0f;
+    }
+}
+
 kernel void kernel_mul_mm_id_amax_part_f32(
        constant ggml_metal_kargs_mul_mm_id_amax & args,
        device   const char * src1,
@@ -3302,7 +3302,8 @@ kernel void kernel_mul_mv_id(
        uint3  tgpig[[threadgroup_position_in_grid]],
        ushort tiitg[[thread_index_in_threadgroup]],
        ushort tiisg[[thread_index_in_simdgroup]],
-        ushort sgitg[[simdgroup_index_in_threadgroup]]) {
+        ushort sgitg[[simdgroup_index_in_threadgroup]],
+        ushort3 ntg [[threads_per_threadgroup]]) {
    const int iid1 = tgpig.z/args.nei0;
    const int idx  = tgpig.z%args.nei0;

@@ -3316,11 +3317,20 @@ kernel void kernel_mul_mv_id(
    const int64_t i1 = idx;
    const int64_t i2 = i12;

+    device char * dst_cur = dst + (i1*args.ne0 + i2*args.ne1*args.ne0)*sizeof(float);
+
+    if (i02 == -1) {
+        device float * dst_f32 = (device float *) dst_cur;
+        const ushort nth = ntg.x*ntg.y*ntg.z;
+        for (int i0 = tiitg; i0 < args.ne0; i0 += nth) {
+            dst_f32[i0] = 0.0f;
+        }
+        return;
+    }
+
    device const char * src0_cur = src0s + i02*args.nb02;
    device const char * src1_cur = src1  + i11*args.nb11 + i12*args.nb12;

-    device char * dst_cur = dst + (i1*args.ne0 + i2*args.ne1*args.ne0)*sizeof(float);
-
    ggml_metal_kargs_mul_mv args0 = {
        /*.ne00 =*/ args.ne00,
        /*.ne01 =*/ args.ne01,
@@ -1047,6 +1047,7 @@ struct ggml_backend_opencl_context {
    cl_kernel kernel_moe_reorder_b;
    cl_kernel kernel_moe_histogram, kernel_moe_scan, kernel_moe_fill, kernel_moe_scatter;
    cl_kernel kernel_moe_scatter_stable = nullptr;   // deterministic slot assignment
+    cl_kernel kernel_moe_zero_dst;
    cl_kernel kernel_moe_combine_f32 = nullptr;   // fused router-weight mul + cross-expert sum
    cl_kernel kernel_moe_combine_bias_f32 = nullptr;  // same, with the down-projection bias add folded in
    cl_kernel kernel_mul_mv_id_q4_0_f32_8x_flat;
@@ -5547,6 +5548,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
            build_program_from_source(backend_ctx, kernel_src.c_str(), CL_moe_compile_opts);

        CL_CHECK((backend_ctx->kernel_moe_histogram = clCreateKernel(prog, "kernel_moe_histogram", &err), err));
+        CL_CHECK((backend_ctx->kernel_moe_zero_dst = clCreateKernel(prog, "kernel_moe_zero_dst", &err), err));
        CL_CHECK((backend_ctx->kernel_moe_scan = clCreateKernel(prog, "kernel_moe_scan", &err), err));
        CL_CHECK((backend_ctx->kernel_moe_fill = clCreateKernel(prog, "kernel_moe_fill", &err), err));
        CL_CHECK((backend_ctx->kernel_moe_scatter = clCreateKernel(prog, "kernel_moe_scatter", &err), err));
@@ -26961,6 +26963,39 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,

    GGML_ASSERT(ne00 == ne10);

+    // the matmul kernels only write the rows an expert owns, so zero the skipped rows here
+    {
+        cl_int err;
+
+        cl_buffer_region region;
+        region.origin = offset2;
+        region.size   = nb21 * ne21;
+        cl_mem buf_ids = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
+        CL_CHECK(err);
+
+        region.origin = offsetd;
+        region.size   = (size_t) ne0 * ne1 * ne2 * sizeof(float);
+        cl_mem buf_dst = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
+        CL_CHECK(err);
+
+        const cl_uint ids_stride = nb21 / sizeof(int);
+
+        cl_kernel kernel_zero = backend_ctx->kernel_moe_zero_dst;
+        CL_CHECK(clSetKernelArg(kernel_zero, 0, sizeof(cl_mem),  &buf_ids));
+        CL_CHECK(clSetKernelArg(kernel_zero, 1, sizeof(cl_mem),  &buf_dst));
+        CL_CHECK(clSetKernelArg(kernel_zero, 2, sizeof(cl_uint), &ne21));
+        CL_CHECK(clSetKernelArg(kernel_zero, 3, sizeof(cl_uint), &ne20));
+        CL_CHECK(clSetKernelArg(kernel_zero, 4, sizeof(cl_uint), &ids_stride));
+        CL_CHECK(clSetKernelArg(kernel_zero, 5, sizeof(cl_uint), &ne0));
+
+        size_t zero_local[3]  = {1, 1, 64};
+        size_t zero_global[3] = {(size_t) ne21, (size_t) ne20, 64};
+        backend_ctx->enqueue_ndrange_kernel(kernel_zero, 3, zero_global, zero_local, dst);
+
+        CL_CHECK(clReleaseMemObject(buf_ids));
+        CL_CHECK(clReleaseMemObject(buf_dst));
+    }
+
    int sgs   = 32; // subgroup size
    int nsg   = 1;  // number of subgroups
    int nrows = 1;  // number of row in src1
@@ -34,6 +34,14 @@ kernel void kernel_add_id(

    global float * dst_row  = (global float *)((global char *)dst  + i1*nb1 + i2*nb2);
    global float * src0_row = (global float *)((global char *)src0 + i1*nb01 + i2*nb02);
+
+    if (i11 == -1) { // skipped slot, add nothing
+        for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
+            dst_row[i0] = src0_row[i0];
+        }
+        return;
+    }
+
    global float * src1_row = (global float *)((global char *)src1 + i11*nb11);

    for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
@@ -89,6 +89,9 @@ __kernel void kernel_gemv_moe_mxfp4_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -188,6 +191,9 @@ __kernel void kernel_gemv_moe_mxfp4_f32_ns_wimg(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f;
@@ -44,6 +44,9 @@ __kernel void kernel_gemv_moe_q4_0_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q4_1_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -61,6 +61,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int num_superblocks = ne00 / QK_K;
    int num_subblocks = ne00 / 32;
@@ -180,6 +183,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns_wimg(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int num_superblocks = ne00 / QK_K;
    int num_subblocks = ne00 / 32;
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q5_0_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -46,6 +46,9 @@ __kernel void kernel_gemv_moe_q5_1_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }
    uint expert_offset = expert_id * ne00 * ne01 / 32;

    __private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -62,6 +62,9 @@ __kernel void kernel_gemv_moe_q5_k_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int num_superblocks = ne00 / QK_K;
    int num_subblocks = ne00 / 32;
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q6_k_f32_ns(
    uint i11 = i20 % ne11;

    uint expert_id = src2[i20];
+    if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int num_superblocks = ne00 / QK_K;
    int num_subblocks = ne00 / 32;  // 8 sub-blocks of 32 per super-block
@@ -18,6 +18,9 @@ __kernel void kernel_moe_histogram(
    }

    int expert_id = input[n * n_experts + k];
+    if (expert_id == -1) { // skipped slot, owned by no expert
+        return;
+    }
    atomic_inc(&hist[expert_id]);
 }

@@ -63,6 +66,9 @@ __kernel void kernel_moe_scatter(
    }

    int val = input[n * n_experts + k];
+    if (val == -1) { // skipped slot, owned by no expert
+        return;
+    }

    int local_slot = atomic_inc(&slot_counter[val]);

@@ -74,6 +80,32 @@ __kernel void kernel_moe_scatter(
    emap[tile_idx] = val;
 }

+// the gemm only writes the rows an expert owns, so the skipped rows (negative ids) are zeroed here
+__kernel void kernel_moe_zero_dst(
+    __global const int * input,
+    __global float * dst,
+    uint N,
+    uint topK,
+    uint n_experts,
+    uint ne0
+) {
+    uint n = get_global_id(0);
+    uint k = get_global_id(1);
+
+    if (n >= N || k >= topK) {
+        return;
+    }
+
+    if (input[n * n_experts + k] != -1) {
+        return;
+    }
+
+    __global float * dst_row = dst + ((ulong)n * topK + k) * ne0;
+    for (uint i = get_local_id(2); i < ne0; i += get_local_size(2)) {
+        dst_row[i] = 0.0f;
+    }
+}
+
 // Deterministic replacement for kernel_moe_scatter.
 //
 // kernel_moe_scatter takes each token's slot from atomic_inc(slot_counter[expert]),
@@ -172,6 +172,9 @@ kernel void kernel_mul_mv_id_mxfp4_f32(
    const int idx  = get_group_id(2)%ne20;

    int i02 = ((global int *) (src2 + iid1*nb21))[idx];
+    if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int i11 = idx % ne11;
    int i12 = iid1;
@@ -101,6 +101,9 @@ kernel void kernel_mul_mv_id_mxfp4_f32_flat(
    const int idx  = get_group_id(2) % ne20;

    uint i02 = ((global uint *) (src2 + offset2 + iid1 * nb21))[idx];
+    if ((int)i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int i11 = idx % ne11;

@@ -267,6 +267,9 @@ kernel void kernel_mul_mv_id_q4_0_f32_8x_flat(
    const int idx  = get_group_id(2)%ne20;

    const int i02 = ((global int *)(src2 + iid1*nb21))[idx];
+    if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    const int i11 = idx%ne11;
    const int i12 = iid1;
@@ -73,6 +73,9 @@ kernel void kernel_mul_mv_id_q8_0_f32(
    int idx  = get_group_id(2)%ne20;

    int i02 = ((global int *) (src2 + iid1*nb21))[idx];
+    if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int i11_ = idx % ne11;
    int i12_ = iid1;
@@ -72,6 +72,9 @@ kernel void kernel_mul_mv_id_q8_0_f32_flat(
    int idx  = (int)get_group_id(2)%ne20;

    int i02 = ((global int *) (src2 + iid1*nb21))[idx];
+    if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
+        return;
+    }

    int i11_ = idx % ne11;
    int i12_ = iid1;
@@ -9,8 +9,12 @@
 #include <openvino/op/constant.hpp>
 #include <openvino/op/convert.hpp>
 #include <openvino/op/gather.hpp>
+#include <openvino/op/greater_eq.hpp>
+#include <openvino/op/maximum.hpp>
+#include <openvino/op/multiply.hpp>
 #include <openvino/op/reshape.hpp>
 #include <openvino/op/shape_of.hpp>
+#include <openvino/op/unsqueeze.hpp>

 namespace ov {
 namespace frontend {
@@ -52,8 +56,16 @@ OutputVector translate_add_id(const NodeContext & context) {
        ids = std::make_shared<ov::op::v0::Convert>(ids, ov::element::i32);
    }

+    // an id of -1 marks a skipped slot: gather row 0 for it and zero the bias with a mask
+    auto zero = ov::op::v0::Constant::create(ids.get_element_type(), ov::Shape{}, {0});
+    ov::Output<ov::Node> keep = std::make_shared<ov::op::v1::GreaterEqual>(ids, zero);
+    keep = std::make_shared<ov::op::v0::Convert>(keep, bias.get_element_type());
+    keep = std::make_shared<ov::op::v0::Unsqueeze>(keep, ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, {-1}));
+    ids = std::make_shared<ov::op::v1::Maximum>(ids, zero);
+
    auto gather_axis = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{}, {0});
    ov::Output<ov::Node> selected_bias = std::make_shared<ov::op::v8::Gather>(bias, ids, gather_axis);
+    selected_bias = std::make_shared<ov::op::v1::Multiply>(selected_bias, keep);
    selected_bias = std::make_shared<ov::op::v1::Reshape>(
        selected_bias, std::make_shared<ov::op::v3::ShapeOf>(input, ov::element::i64), false);

@@ -26,6 +26,15 @@ static void add_id_kernel(
  float* dst_row = (float*)((char*)dst + i1 * nb1 + i2 * nb2);
  const float* src0_row =
      (const float*)((const char*)src0 + i1 * nb01 + i2 * nb02);
+
+  if (i11 == -1) { // skipped slot, add nothing
+    for (int64_t i0 = item_ct1.get_local_id(2); i0 < ne0;
+         i0 += item_ct1.get_local_range(2)) {
+      dst_row[i0] = src0_row[i0];
+    }
+    return;
+  }
+
  const float* src1_row = (const float*)((const char*)src1 + i11 * nb11);

  for (int64_t i0 = item_ct1.get_local_id(2); i0 < ne0;
@@ -5700,6 +5700,23 @@ __dpct_inline__ static void k_copy_dst_from_contiguous(
    }
 }

+__dpct_inline__ static void k_zero_dst_rows(
+    char *__restrict__ dst_original,
+    const mmid_row_mapping *__restrict__ row_mapping, int64_t ne0, size_t nb1,
+    size_t nb2, const sycl::nd_item<3> &item_ct1) {
+    int32_t i = item_ct1.get_group(2);
+
+    const int32_t i1 = row_mapping[i].i1;
+    const int32_t i2 = row_mapping[i].i2;
+
+    float * dst_row_original = (float *)(dst_original + i1*nb1 + i2*nb2);
+
+    for (int j = item_ct1.get_local_id(2); j < ne0;
+         j += item_ct1.get_local_range(2)) {
+        dst_row_original[j] = 0.0f;
+    }
+}
+
 // returns true when ggml_sycl_mul_mat_id takes the fallback path that requires stream synchronization
 static bool ggml_sycl_mul_mat_id_needs_sync(const ggml_tensor * dst) {
    const ggml_tensor * src0 = dst->src[0];
@@ -5785,14 +5802,18 @@ static void mmid_counting_sort_rows(
        int64_t n_ids, int64_t n_as, int64_t n_routed_rows,
        std::vector<int64_t> & expert_counts,
        std::vector<int64_t> & expert_row_offsets,
-        std::vector<mmid_row_mapping> & routed_row_src) {
+        std::vector<mmid_row_mapping> & routed_row_src,
+        std::vector<mmid_row_mapping> & skipped_row_dst) {

    // frequencies: how many routed rows each expert "owns"
    expert_counts.assign(n_as, 0);
    for (int64_t iid1 = 0; iid1 < ids->ne[1]; iid1++) {
        for (int64_t id = 0; id < n_ids; id++) {
            const int32_t row_id_i = *(const int32_t *) (ids_host + iid1*ids->nb[1] + id*ids->nb[0]);
-            GGML_ASSERT(row_id_i >= 0 && row_id_i < n_as);
+            GGML_ASSERT(row_id_i == -1 || (row_id_i >= 0 && row_id_i < n_as));
+            if (row_id_i == -1) {
+                continue;
+            }
            expert_counts[row_id_i]++;
        }
    }
@@ -5805,10 +5826,17 @@ static void mmid_counting_sort_rows(

    std::vector<int64_t> expert_row_next = expert_row_offsets;
    routed_row_src.resize(n_routed_rows);
+    skipped_row_dst.clear();
    for (int64_t iid1 = 0; iid1 < ids->ne[1]; iid1++) {
        for (int64_t id = 0; id < n_ids; id++) {
            const int32_t row_id_i = *(const int32_t *) (ids_host + iid1*ids->nb[1] + id*ids->nb[0]);
-            GGML_ASSERT(row_id_i >= 0 && row_id_i < n_as);
+            GGML_ASSERT(row_id_i == -1 || (row_id_i >= 0 && row_id_i < n_as));
+
+            // a skipped slot has no expert, its dst row is zeroed instead
+            if (row_id_i == -1) {
+                skipped_row_dst.push_back({(int32_t) id, (int32_t) iid1});
+                continue;
+            }

            // find and validate the next free row for a given expert (row_id_i)
            const int64_t routed_row = expert_row_next[row_id_i]++;
@@ -5817,6 +5845,7 @@ static void mmid_counting_sort_rows(
            routed_row_src[routed_row] = {(int32_t) id, (int32_t) iid1};
        }
    }
+    routed_row_src.resize(expert_row_offsets[n_as]);
 }

 static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
@@ -5884,7 +5913,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
        for (int64_t iid1 = 0; iid1 < ids->ne[1]; iid1++) {
            for (int64_t id = 0; id < n_ids; id++) {
                const int32_t i02 = *(const int32_t *) (ids_host.data() + iid1*ids->nb[1] + id*ids->nb[0]);
-                GGML_ASSERT(i02 >= 0 && i02 < n_as);
+                GGML_ASSERT(i02 == -1 || (i02 >= 0 && i02 < n_as));

                const int64_t i11 = id % ne11;
                const int64_t i12 = iid1;
@@ -5892,6 +5921,12 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
                const int64_t i1 = id;
                const int64_t i2 = i12;

+            if (i02 == -1) { // skipped slot
+                SYCL_CHECK(CHECK_TRY_ERROR(
+                        stream->memset(dst_original + i1*nb1 + i2*nb2, 0, ne0*sizeof(float))));
+                continue;
+            }
+
            src0_row.data = src0_original + i02*nb02;
            src1_row.data = src1_original + i11*nb11 + i12*nb12;
            dst_row.data = dst_original + i1*nb1 + i2*nb2;
@@ -5913,20 +5948,43 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
        std::vector<int64_t> expert_row_offsets;
        // the sources (slot/token pairs) of contiguous rows to guide k_copy_src1_to_contiguous
        std::vector<mmid_row_mapping> & routed_row_src = ctx.mmid_row_mapping_host;
+        // the slot/token pairs of the skipped slots, whose dst rows are zeroed
+        std::vector<mmid_row_mapping> skipped_row_dst;

        mmid_counting_sort_rows(ids, ids_host.data(), n_ids, n_as, n_routed_rows,
-                                expert_row_counts, expert_row_offsets, routed_row_src);
+                                expert_row_counts, expert_row_offsets, routed_row_src, skipped_row_dst);
+
+        const int64_t n_valid_rows = (int64_t) routed_row_src.size();

        ggml_sycl_pool_alloc<mmid_row_mapping> dev_row_mapping(ctx.pool(), n_routed_rows);
        SYCL_CHECK(CHECK_TRY_ERROR(
-                stream->memcpy(dev_row_mapping.get(), routed_row_src.data(), n_routed_rows*sizeof(mmid_row_mapping))));
+                stream->memcpy(dev_row_mapping.get(), routed_row_src.data(), n_valid_rows*sizeof(mmid_row_mapping))));

        const unsigned int max_work_group_size = ggml_sycl_info().max_work_group_sizes[ctx.device];
        assert(max_work_group_size % (WARP_SIZE * WARP_SIZE) == 0);

-        {
+        if (!skipped_row_dst.empty()) {
+            ggml_sycl_pool_alloc<mmid_row_mapping> dev_skipped(ctx.pool(), skipped_row_dst.size());
+            SYCL_CHECK(CHECK_TRY_ERROR(
+                    stream->memcpy(dev_skipped.get(), skipped_row_dst.data(), skipped_row_dst.size()*sizeof(mmid_row_mapping))));
+
+            sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, max_work_group_size));
+            sycl::range<3> grid_dims(1, 1, skipped_row_dst.size());
+            stream->submit([&](sycl::handler &cgh) {
+                mmid_row_mapping *__restrict dev_skipped_get = dev_skipped.get();
+
+                cgh.parallel_for(
+                    sycl::nd_range<3>(grid_dims * block_dims, block_dims),
+                    [=](sycl::nd_item<3> item_ct1) {
+                        k_zero_dst_rows(dst_original, dev_skipped_get, ne0, nb1, nb2, item_ct1);
+                    });
+            });
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->wait()));
+        }
+
+        if (n_valid_rows > 0) {
            sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne10, max_work_group_size));
-            sycl::range<3> grid_dims(1, 1, n_routed_rows);
+            sycl::range<3> grid_dims(1, 1, n_valid_rows);
            stream->submit([&](sycl::handler &cgh) {
                char *__restrict src1_contiguous_get =
                    src1_contiguous.get();
@@ -5992,9 +6050,9 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
            ggml_sycl_mul_mat(ctx, &src0_row, &src1_row, &dst_row);
        }

-        {
+        if (n_valid_rows > 0) {
            sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, max_work_group_size));
-            sycl::range<3> grid_dims(1, 1, n_routed_rows);
+            sycl::range<3> grid_dims(1, 1, n_valid_rows);
            stream->submit([&](sycl::handler &cgh) {
                const char *__restrict dst_contiguous_get =
                    dst_contiguous.get();
@@ -3149,7 +3149,6 @@ static void mul_mat_vec_q_moe(
    const int expert_idx = item_ct1.get_group(1);
    const int i02        = ids_dev[expert_idx];

-    const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
    const char * vy = (const char *) vy_base + (size_t) expert_idx * src1_row_stride;
    float *      dst = (float *) ((char *) dst_base + (size_t) expert_idx * dst_row_stride);

@@ -3159,6 +3158,15 @@ static void mul_mat_vec_q_moe(
        return;
    }

+    if (i02 == -1) { // skipped slot
+        if (item_ct1.get_local_id(2) == 0) {
+            dst[row] = 0.0f;
+        }
+        return;
+    }
+
+    const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
+
    const int     blocks_per_row  = ncols / qk;
    constexpr int blocks_per_warp = (vdr * WARP_SIZE + qi - 1) / qi;

@@ -3355,7 +3363,6 @@ static void mul_mat_vec_q_moe_reorder(
    const int expert_idx = item_ct1.get_group(1);
    const int i02        = ids_dev[expert_idx];

-    const char * vx  = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
    const char * vy  = (const char *) vy_base + (size_t) expert_idx * src1_row_stride;
    float *      dst = (float *) ((char *) dst_base + (size_t) expert_idx * dst_row_stride);

@@ -3366,6 +3373,15 @@ static void mul_mat_vec_q_moe_reorder(

    const auto sg = item_ct1.get_sub_group();

+    if (i02 == -1) { // skipped slot
+        if (sg.leader()) {
+            dst[row] = 0.0f;
+        }
+        return;
+    }
+
+    const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
+
    constexpr int vdr_mmvq = reorder_vec_dot_q_sycl::gtype == GGML_TYPE_Q8_0 ?
                            VDR_Q8_0_Q8_1_MMVQ : block_traits::vdr_mmvq;
    const int     blocks_per_row              = ncols / block_traits::qk;
@@ -222,6 +222,9 @@ struct vk_op_count_experts_push_constants {
    uint32_t ne00L;
    uint32_t row_tile_size;
    uint32_t row_ids_offset;
+    uint32_t ne0_dst;
+    uint32_t stride_dst;
+    uint32_t batch_stride_dst;
 };

 struct vk_op_glu_push_constants {
@@ -3585,9 +3585,9 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
    ggml_vk_create_pipeline(device, device->pipeline_count_equal_i32, "count_equal_i32", count_equal_i32_len, count_equal_i32_data, "main", 3, sizeof(vk_op_push_constants), {512, 1, 1}, { device->subgroup_size }, 1);

    if (device->subgroup_arithmetic && device->subgroup_vote && device->subgroup_require_full_support) {
-        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_subgroup_len, count_experts_subgroup_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true, true);
+        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_subgroup_len, count_experts_subgroup_data, "main", 3, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true, true);
    } else {
-        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
+        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 3, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
    }

    // comb holds a token's 4x4 matrix in one 16-lane slice of a subgroup, so it
@@ -6005,20 +6005,49 @@ void ggml_vk_matmul(
    ctx->prealloc_split_k_need_sync = true;
 }

-static bool ggml_vk_get_mul_mat_mat_f16acc(ggml_backend_vk_context * ctx, ggml_type src0_type, ggml_type src1_type, ggml_prec prec) {
+static bool ggml_vk_src1_prec_allows_q8(const ggml_tensor * dst) {
+    switch (ggml_prec_get_src(dst, 1)) {
+        case GGML_PREC_UNDEFINED:
+        case GGML_PREC_Q8:
+        case GGML_PREC_Q4:
+            return true;
+        case GGML_PREC_F32:
+        case GGML_PREC_BF16:
+        case GGML_PREC_F16:
+            return false;
+    }
+    return false;
+}
+
+static bool ggml_vk_prec_allows_f16acc(ggml_prec prec) {
+    switch (prec) {
+        case GGML_PREC_F32:
+        case GGML_PREC_BF16:
+            return false;
+        case GGML_PREC_UNDEFINED:
+        case GGML_PREC_F16:
+        case GGML_PREC_Q8:
+        case GGML_PREC_Q4:
+            return true;
+    }
+    return false;
+}
+
+static bool ggml_vk_get_mul_mat_mat_f16acc(ggml_backend_vk_context * ctx, ggml_type src0_type, ggml_type src1_type, ggml_prec acc_prec) {
    if (src0_type == GGML_TYPE_F32 || src0_type == GGML_TYPE_BF16) return false;
    if (src1_type == GGML_TYPE_Q8_1) return false;
+    const bool prec_allows_f16 = ggml_vk_prec_allows_f16acc(acc_prec);
    if (src0_type == GGML_TYPE_F16) {
-        return prec == GGML_PREC_DEFAULT && ctx->device->fp16 && !(ctx->device->coopmat_support && !ctx->device->coopmat_acc_f16_support);
+        return prec_allows_f16 && ctx->device->fp16 && !(ctx->device->coopmat_support && !ctx->device->coopmat_acc_f16_support);
    }
    // quant types
    if (ctx->device->coopmat2) {
-        return prec == GGML_PREC_DEFAULT;
+        return prec_allows_f16;
    }
    if (ctx->device->coopmat_support) {
-        return ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec == GGML_PREC_DEFAULT;
+        return ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec_allows_f16;
    }
-    return ctx->device->fp16 && prec == GGML_PREC_DEFAULT;
+    return ctx->device->fp16 && prec_allows_f16;
 }

 static const std::vector<vk_matmul_pipeline_pair>* ggml_vk_get_mul_mat_mat_pipeline_map(
@@ -6386,26 +6415,26 @@ static void ggml_vk_mul_mat_q_f16(ggml_backend_vk_context * ctx, vk_context& sub

    const bool is_coopmat = ctx->device->coopmat_support || ctx->device->coopmat2;
    const bool can_quantize_y = (ctx->device->integer_dot_product || ctx->device->coopmat_int_support) &&
-                                src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0;
+                                src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst);
    // coopmat shaders are fast enough that memory bandwidth reduction from f16 activations helps
    const bool prefer_f16_b = is_coopmat && ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32;

    const ggml_type y_kernel_type = ggml_vk_mul_mat_b_type(ctx, src0->type, src1, f16_type,
-                                        (ggml_prec)dst->op_params[0], false, can_quantize_y, prefer_f16_b);
+                                        ggml_prec_get_acc(dst), false, can_quantize_y, prefer_f16_b);

    const bool quantize_y       = y_kernel_type == GGML_TYPE_Q8_1;
    const bool y_f32_kernel     = y_kernel_type == GGML_TYPE_F32;
    const bool y_needs_reformat = !quantize_y && (!ggml_vk_dim01_contiguous(src1) || y_kernel_type != src1->type);

    const std::vector<vk_matmul_pipeline_pair>* mmp_map =
-        ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, (ggml_prec)dst->op_params[0]);
+        ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, ggml_prec_get_acc(dst));

    const bool qx_needs_dequant = mmp_map == nullptr || x_non_contig;
    const bool qy_needs_dequant = y_needs_reformat;

    if (qx_needs_dequant) {
        // dequant src0 to f16
-        mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, (ggml_prec)dst->op_params[0]);
+        mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, ggml_prec_get_acc(dst));
    }

    GGML_ASSERT(mmp_map != nullptr);
@@ -6736,7 +6765,7 @@ static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context&

    const ggml_type y_kernel_type = widen_bf16 ? GGML_TYPE_F32 : src1->type;
    const bool f16_f32_kernel = y_kernel_type == GGML_TYPE_F32;
-    bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_should_use_mmvq(ctx->device, ne01, ne11, ne10, src0->type);
+    bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && (swap_inputs || ggml_vk_src1_prec_allows_q8(dst)) && ggml_vk_should_use_mmvq(ctx->device, ne01, ne11, ne10, src0->type);

    vk_pipeline to_fp16_vk_0 = nullptr;
    vk_pipeline to_fp16_vk_1 = nullptr;
@@ -7435,26 +7464,26 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    const bool y_decode_vector_staging = false;
 #endif
    const bool can_quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 &&
-                                ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0;
+                                ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst);
    // Intel coopmat1 matrix cores prefer f16 activations for quantized weights
    const bool prefer_f16_b = ctx->device->coopmat_support && !ctx->device->coopmat2 &&
                              ctx->device->vendor_id == VK_VENDOR_ID_INTEL &&
                              ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32;

    const ggml_type y_kernel_type = ggml_vk_mul_mat_b_type(ctx, src0->type, src1, f16_type,
-                                        (ggml_prec)dst->op_params[0], true, can_quantize_y, prefer_f16_b);
+                                        ggml_prec_get_acc(dst), true, can_quantize_y, prefer_f16_b);

    const bool quantize_y   = y_kernel_type == GGML_TYPE_Q8_1;
    const bool y_f32_kernel = y_kernel_type == GGML_TYPE_F32;

    const std::vector<vk_matmul_pipeline_pair>* mmp_map =
-        ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, (ggml_prec)dst->op_params[0], true);
+        ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, ggml_prec_get_acc(dst), true);

    const bool qx_needs_dequant = mmp_map == nullptr || x_non_contig;

    if (qx_needs_dequant) {
        // dequant src0 to f16
-        mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, (ggml_prec)dst->op_params[0], true);
+        mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, ggml_prec_get_acc(dst), true);
    }

    GGML_ASSERT(mmp_map != nullptr);
@@ -7614,16 +7643,20 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
        ggml_vk_sync_buffers(ctx, subctx);
    }
    {
+        // this also zeroes the dst rows of the skipped slots (negative ids)
        vk_op_count_experts_push_constants pc = { (uint32_t)nei0,
                                           (uint32_t)nei1,
                                           (uint32_t)(nbi0 / ggml_type_size(ids->type)),
                                           (uint32_t)(nbi1 / ggml_type_size(ids->type)),
                                           (uint32_t)(get_misalign_bytes(ctx, ids) / ggml_type_size(ids->type)),
                                           (uint32_t)n_as,
-                                           0, 0, bn, (uint32_t)row_ids_offset };
+                                           0, 0, bn, (uint32_t)row_ids_offset,
+                                           (uint32_t)ne01,
+                                           (uint32_t)(dst->nb[1] / sizeof(float)),
+                                           (uint32_t)(dst->nb[2] / sizeof(float)) };
        init_pushconst_fastdiv(pc);
        ggml_vk_dispatch_pipeline(ctx, subctx, count_experts,
-            { d_ids, expert_count_buf }, pc,
+            { d_ids, expert_count_buf, vk_subbuffer{ d_D, d_buf_offset, d_sz } }, pc,
            { 1, 1, 1 });
    }

@@ -7749,7 +7782,7 @@ static void ggml_vk_mul_mat_vec_id_q_f16(ggml_backend_vk_context * ctx, vk_conte
    const bool y_non_contig = !ggml_vk_dim01_contiguous(src1);

    const bool f16_f32_kernel = src1->type == GGML_TYPE_F32;
-    bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_should_use_mmvq(ctx->device, ne01, ne12, ne10, src0->type);
+    bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst) && ggml_vk_should_use_mmvq(ctx->device, ne01, ne12, ne10, src0->type);

    vk_pipeline to_fp16_vk_0 = nullptr;
    vk_pipeline to_fp16_vk_1 = nullptr;
@@ -8128,7 +8161,7 @@ void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx, const
    uint32_t workgroups_y = (uint32_t)neq2;
    uint32_t workgroups_z = (uint32_t)neq3;

-    const bool f32acc = !ctx->device->fp16 || dst->op_params[3] == GGML_PREC_F32 || k->type == GGML_TYPE_BF16;
+    const bool f32acc = !ctx->device->fp16 || !ggml_vk_prec_allows_f16acc(ggml_prec_get_acc(dst)) || k->type == GGML_TYPE_BF16;

    // dequant K/V once into an f16 scratch, reordered KV layout so FA can read without a stride
    auto is_dense_kv_cache = [](const ggml_tensor * t) {
@@ -15516,9 +15549,9 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                        // If there's not enough shared memory for row_ids and the result tile, fallback to CPU
                        return false;
                    }
-                    if (ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) {
-                        return false;
-                    }
+                }
+                if (ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+                    return false;
                }
                switch (src0_type) {
                    case GGML_TYPE_F32:
@@ -27,14 +27,22 @@ void main() {
    const uint i1 = gl_WorkGroupID.x;
    const uint i2 = gl_WorkGroupID.y;

-    const uint i11 = data_c[i1 + i2 * p.s21];
+    const int i11 = data_c[i1 + i2 * p.s21];

    const uint s1 = p.ne0;
    const uint s2 = p.ne0 * p.ne1;

    const uint d0 = i1 * s1 + i2 * s2;
    const uint a0 = i1 * p.s01 + i2 * p.s02;
-    const uint b0 = i11 * p.s11;
+
+    if (i11 == -1) { // skipped slot, add nothing
+        for (uint i0 = gl_LocalInvocationID.x; i0 < p.ne0; i0 += BLOCK_SIZE) {
+            data_d[d0 + i0] = data_a[a0 + i0];
+        }
+        return;
+    }
+
+    const uint b0 = uint(i11) * p.s11;

    for (uint i0 = gl_LocalInvocationID.x; i0 < p.ne0; i0 += BLOCK_SIZE) {
        data_d[d0 + i0] = data_a[a0 + i0] + data_b[b0 + i0];
@@ -23,6 +23,9 @@ layout (push_constant) uniform parameter
    uint32_t ne00L;
    uint32_t row_tile_size;
    uint32_t row_ids_offset;
+    uint32_t ne0_dst;
+    uint32_t stride_dst;
+    uint32_t batch_stride_dst;
 } p;

 #define BLOCK_SIZE 256
@@ -32,6 +35,7 @@ layout(local_size_x = BLOCK_SIZE, local_size_y = 1, local_size_z = 1) in;

 layout (binding = 0) readonly buffer A {uint data_a[];};
 layout (binding = 1) writeonly buffer D {uint data_d[];};
+layout (binding = 2) writeonly buffer DST {float data_dst[];};

 shared uint vals[EXPERTS_PER_CHUNK];
 shared uint offsets[EXPERTS_PER_CHUNK];
@@ -47,6 +51,26 @@ void write_tiles(uint expert, uint count, uint tile_begin, uint n_tiles) {
    }
 }

+// the matmul writes only the rows that an expert owns, so zero the rows of the skipped slots
+void zero_skipped_rows() {
+    const uint num_elements = p.ne00 * p.ne01;
+    const uint tid = gl_LocalInvocationID.x;
+
+    for (uint idx = 0; idx < num_elements; idx++) {
+        const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
+        const uint i00 = idx - i01 * p.ne00;
+
+        if (int(data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00]) != -1) {
+            continue;
+        }
+
+        const uint d_off = i00 * p.stride_dst + i01 * p.batch_stride_dst;
+        for (uint i = tid; i < p.ne0_dst; i += BLOCK_SIZE) {
+            data_dst[d_off + i] = 0.0f;
+        }
+    }
+}
+
 void main() {
    const uint num_elements = p.ne00 * p.ne01;
    const uint tid = gl_LocalInvocationID.x;
@@ -123,4 +147,6 @@ void main() {
    if (tid == 0) {
        data_d[0] = total_tiles;
    }
+
+    zero_skipped_rows();
 }
@@ -42,6 +42,7 @@ layout (push_constant) uniform parameter

 #ifdef MUL_MAT_ID
 uint expert_id;
+bool expert_skip; // an id of -1 skips the slot, the result row is zeroed
 #endif

 void get_offsets(out uint a_offset, out uint b_offset, out uint d_offset) {
@@ -63,7 +64,9 @@ void get_offsets(out uint a_offset, out uint b_offset, out uint d_offset) {
        batch_idx_a = i03 * p.ne02 + i02;
    }
 #else
-    expert_id = data_ids[expert_i0 + p.expert_i1 * p.nbi1];
+    const int expert = data_ids[expert_i0 + p.expert_i1 * p.nbi1];
+    expert_skip = expert == -1;
+    expert_id   = expert_skip ? 0 : uint(expert); // clamp, so that the offsets below stay in range
 #endif

    a_offset =
@@ -113,6 +116,9 @@ void reduce_result(inout FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t
                    const uint expert_i0 = gl_GlobalInvocationID.y;
                    temp[j][n] *= FLOAT_TYPE(data_fuse1[expert_i0]);
                }
+                if (expert_skip) {
+                    temp[j][n] = FLOAT_TYPE(0);
+                }
 #else
                if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
                    temp[j][n] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
@@ -168,6 +174,9 @@ void reduce_result(FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t d_offs
                    const uint expert_i0 = gl_GlobalInvocationID.y;
                    temp[j][n] *= FLOAT_TYPE(data_fuse1[expert_i0]);
                }
+                if (expert_skip) {
+                    temp[j][n] = FLOAT_TYPE(0);
+                }
 #else
                if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
                    temp[j][n] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
@@ -213,6 +222,9 @@ void reduce_result(FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t d_offs
                    const uint expert_i0 = gl_GlobalInvocationID.y;
                    tmpsh[j][n][0] *= FLOAT_TYPE(data_fuse1[expert_i0]);
                }
+                if (expert_skip) {
+                    tmpsh[j][n][0] = FLOAT_TYPE(0);
+                }
 #else
                if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
                    tmpsh[j][n][0] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
@@ -1863,10 +1863,12 @@ static webgpu_encoded_op ggml_webgpu_mul_mat_id(webgpu_context & ctx,
    // params for mul_mat_id_gather.wgsl
    std::vector<uint32_t> gather_params = {
        (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src2) / ggml_type_size(src2->type)),
+        (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
        param_n_expert,
        param_n_expert_used,
        param_n_tokens,
        (uint32_t) (src2->nb[1] / ggml_type_size(src2->type)),
+        (uint32_t) dst->ne[0],
    };

    const size_t dst_offset       = ggml_webgpu_tensor_offset(dst);
@@ -1888,6 +1890,8 @@ static webgpu_encoded_op ggml_webgpu_mul_mat_id(webgpu_context & ctx,
                                          ggml_webgpu_tensor_binding_size(ctx, src2)),
        ggml_webgpu_make_bind_group_entry(1, ggml_webgpu_tensor_buf(dst), ids_cache_align_offset, ids_cache_bind_size),
        ggml_webgpu_make_bind_group_entry(2, ggml_webgpu_tensor_buf(dst), ids_tiles_align_offset, ids_tiles_bind_size),
+        ggml_webgpu_make_bind_group_entry(3, ggml_webgpu_tensor_buf(dst), ggml_webgpu_tensor_align_offset(ctx, dst),
+                                          ggml_webgpu_tensor_binding_size(ctx, dst)),
    };

    dispatches.push_back({
@@ -46,12 +46,24 @@ fn main(@builtin(workgroup_id) wg_id: vec3<u32>,
        let i2 = wg_linear / params.ne1;
        let i1 = wg_linear % params.ne1;

-        let i11 = u32(ids[params.offset_ids + i1 * params.nb20 + i2 * params.nb21]);
+        let i11 = i32(ids[params.offset_ids + i1 * params.nb20 + i2 * params.nb21]);

        let src0_row = params.offset_src0 + i1 * params.nb01 + i2 * params.nb02;
-        let src1_row = params.offset_src1 + i11 * params.nb11;
        let dst_row = params.offset_dst + i1 * params.ne0 + i2 * (params.ne0 * params.ne1);

+        if (i11 == -1) { // skipped slot, add nothing
+#ifdef INPLACE
+            // src0 already holds the result
+#else
+            for (var i = thread_id;i < params.ne0; i += WG_SIZE) {
+                dst[dst_row + i] = src0[src0_row + i];
+            }
+#endif
+            return;
+        }
+
+        let src1_row = params.offset_src1 + u32(i11) * params.nb11;
+
        for (var i = thread_id;i < params.ne0; i += WG_SIZE) {
 #ifdef INPLACE
            src0[src0_row + i] = src0[src0_row + i] + src1[src1_row + i];
@@ -2,19 +2,24 @@ enable f16;

 struct MulMatIdGatherParams {
    offset_ids: u32,
+    offset_dst: u32,

    n_expert: u32,
    n_expert_used: u32,
    n_tokens: u32,

    stride_ids_1: u32,
+
+    m: u32,
 };

@group(0) @binding(0) var<storage, read_write> ids: array<i32>;        // [n_expert_used, n_tokens]
@group(0) @binding(1) var<storage, read_write> ids_slot_pos: array<u32>; // [n_expert_used * n_tokens]
@group(0) @binding(2) var<storage, read_write> ids_tiles: array<u32>; // [n_tiles, (expert, first slot, n slots) * n_tiles]

-@group(0) @binding(3) var<uniform> params: MulMatIdGatherParams;
+@group(0) @binding(3) var<storage, read_write> dst: array<f32>; // [m, n_expert_used, n_tokens]
+
+@group(0) @binding(4) var<uniform> params: MulMatIdGatherParams;

 // per expert: the slot count, replaced by the first slot of the expert in ids_slot_pos by the prefix sum, then advanced as its write cursor
 var<workgroup> ids_count: array<atomic<u32>, N_EXPERT>;
@@ -36,8 +41,10 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
    for (var i = thread_id; i < ids_slots_num; i += WG_SIZE) {
        let row = i / params.n_expert_used;
        let col = i % params.n_expert_used;
-        let expert = u32(ids[params.offset_ids + row * params.stride_ids_1 + col]);
-        atomicAdd(&ids_count[expert], 1u);
+        let expert = ids[params.offset_ids + row * params.stride_ids_1 + col];
+        if (expert != -1) { // skipped slot, owned by no expert
+            atomicAdd(&ids_count[u32(expert)], 1u);
+        }
    }

    workgroupBarrier();
@@ -71,8 +78,23 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
    for (var i = thread_id; i < ids_slots_num; i += WG_SIZE) {
        let row = i / params.n_expert_used;
        let col = i % params.n_expert_used;
-        let expert = u32(ids[params.offset_ids + row * params.stride_ids_1 + col]);
-        let pos = atomicAdd(&ids_count[expert], 1u);
-        ids_slot_pos[pos] = i;
+        let expert = ids[params.offset_ids + row * params.stride_ids_1 + col];
+        if (expert != -1) {
+            let pos = atomicAdd(&ids_count[u32(expert)], 1u);
+            ids_slot_pos[pos] = i;
+        }
+    }
+
+    // the matmul writes only the rows that an expert owns, so zero the rows of the skipped slots
+    for (var i: u32 = 0; i < ids_slots_num; i += 1) {
+        let row = i / params.n_expert_used;
+        let col = i % params.n_expert_used;
+        if (ids[params.offset_ids + row * params.stride_ids_1 + col] != -1) {
+            continue;
+        }
+        let dst_row = params.offset_dst + i * params.m;
+        for (var r = thread_id; r < params.m; r += WG_SIZE) {
+            dst[dst_row + r] = 0.0f;
+        }
    }
 }
@@ -70,11 +70,18 @@ fn main(
    let own_expert = ids[params.offset_ids + slot];
    let row_base   = (wg_linear % output_groups) * OUTPUTS_PER_WG;

-    let src0_batch_offset = params.offset_src0 + own_expert * params.stride_02;
+    // a skipped slot (id -1) keeps the zero accumulator, so the reduction below writes zeros
+    // no early return: the barriers below must stay in uniform control flow
+    let skipped = i32(own_expert) == -1;
+
+    let src0_batch_offset = params.offset_src0 + select(own_expert, 0u, skipped) * params.stride_02;
    let src1_idx_base     = params.offset_src1 + (slot % params.b_ne1) * params.stride_11;
    let dst_idx_base      = params.offset_dst + slot * params.m + row_base;

-    let acc = accumulate_vec_dot(thread_id, row_base, src0_batch_offset, src1_idx_base);
+    var acc: array<array<f32, OUTPUTS_PER_WG>, NUM_COLS>;
+    if (!skipped) {
+        acc = accumulate_vec_dot(thread_id, row_base, src0_batch_offset, src1_idx_base);
+    }

 #ifdef USE_SUBGROUP_REDUCTION
    for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
@@ -375,7 +375,12 @@ static void ggml_zendnn_compute_forward_mul_mat_id(
        for (int id = 0; id < n_ids; ++id) {
            const int32_t i02 = *(const int32_t *)((const char *)ids->data + iid1*ids->nb[1] + id*ids->nb[0]);

-            GGML_ASSERT(i02 >= 0 && i02 < n_as);
+            GGML_ASSERT(i02 == -1 || (i02 >= 0 && i02 < n_as));
+
+            if (i02 == -1) {
+                memset((char *) dst->data + id*nb1 + iid1*nb2, 0, ne0*sizeof(float));
+                continue;
+            }

            if (matrix_row_counts[i02] == 0) {
                n_active_experts++;
@@ -3348,6 +3348,36 @@ bool ggml_prec_set_src(
    return true;
 }

+enum ggml_prec ggml_prec_get_acc(
+        const struct ggml_tensor * a) {
+    switch (a->op) {
+        case GGML_OP_MUL_MAT:
+        case GGML_OP_MUL_MAT_ID:
+            return (enum ggml_prec) ggml_get_op_params_i32(a, 0);
+        case GGML_OP_FLASH_ATTN_EXT:
+            return (enum ggml_prec) ggml_get_op_params_i32(a, 3);
+        default:
+            return GGML_PREC_UNDEFINED;
+    };
+}
+
+enum ggml_prec ggml_prec_get_src(
+        const struct ggml_tensor * a,
+        int                        idx) {
+    GGML_ASSERT(idx >= 0 && idx < GGML_MAX_SRC);
+
+    switch (a->op) {
+        case GGML_OP_MUL_MAT:
+        case GGML_OP_MUL_MAT_ID:
+            if (idx != 1) {
+                return GGML_PREC_UNDEFINED;
+            }
+            return (enum ggml_prec) ggml_get_op_params_i32(a, 2 + idx);
+        default:
+            return GGML_PREC_UNDEFINED;
+    };
+}
+
 // ggml_mul_mat

 static inline bool ggml_can_mul_mat(const struct ggml_tensor * t0, const struct ggml_tensor * t1) {
@@ -3408,6 +3438,8 @@ void ggml_mul_mat_set_hint(
    in b, n_expert_used can be broadcasted to match the n_expert_used of ids

    c ~= as[:,:,i] @ b[:,i%r,t], i = ids[e,t] for all e,t in ids
+
+    an id of -1 skips the slot: no matrix is selected, c[:,e,t] is set to zero
 */
 struct ggml_tensor * ggml_mul_mat_id(
        struct ggml_context * ctx,
@@ -1272,6 +1272,11 @@ class MODEL_TENSOR(IntEnum):
    A_GEN_WAV_SEANET_RES_CONV1 = auto()
    A_GEN_WAV_SEANET_RES_CONV2 = auto()
    A_GEN_WAV_SEANET_SCALE_CONV = auto() # strided upsample convtr
+    # soprano: Vocos decoder, the ConvNeXt blocks reuse A_GEN_WAV_UP_*
+    A_GEN_WAV_INPUT                = auto()
+    A_GEN_WAV_NORM                 = auto()
+    A_GEN_WAV_OUTPUT_NORM          = auto()
+    A_GEN_WAV_OUTPUT               = auto()
    A_MMPROJ              = auto()
    A_MMPROJ_FC           = auto()
    A_MM_NORM_PRE         = auto()
@@ -2084,6 +2089,10 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
    MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1: "a.gen.wav.seanet.blk.{bid}.res_conv1",
    MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2: "a.gen.wav.seanet.blk.{bid}.res_conv2",
    MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV: "a.gen.wav.seanet.blk.{bid}.scale_conv",
+    MODEL_TENSOR.A_GEN_WAV_INPUT:           "a.gen.wav.input",
+    MODEL_TENSOR.A_GEN_WAV_NORM:            "a.gen.wav.norm",
+    MODEL_TENSOR.A_GEN_WAV_OUTPUT_NORM:     "a.gen.wav.output_norm",
+    MODEL_TENSOR.A_GEN_WAV_OUTPUT:          "a.gen.wav.output",
    MODEL_TENSOR.A_MMPROJ:                  "mm.a.mlp.{bid}",
    MODEL_TENSOR.A_MMPROJ_FC:               "mm.a.fc",
    MODEL_TENSOR.A_MM_NORM_PRE:             "mm.a.norm_pre",
@@ -2462,6 +2471,10 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1,
        MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2,
        MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV,
+        MODEL_TENSOR.A_GEN_WAV_INPUT,
+        MODEL_TENSOR.A_GEN_WAV_NORM,
+        MODEL_TENSOR.A_GEN_WAV_OUTPUT_NORM,
+        MODEL_TENSOR.A_GEN_WAV_OUTPUT,
        MODEL_TENSOR.A_ENC_CONV_NORM_MEAN,
        MODEL_TENSOR.A_ENC_CONV_NORM_VAR,
        MODEL_TENSOR.A_ENC_MEL_FILTERS,
@@ -6447,6 +6460,7 @@ class VisionProjectorType:
    QWEN3TTS_GEN = "qwen3tts_gen" # audio generation: code_predictor
    POCKETTTS_SPKENC = "pockettts_spkenc" # audio: mimi encoder as voice-prompt encoder
    POCKETTTS_GEN = "pockettts_gen" # audio generation: flow-matching decoder + mimi decoder
+    SOPRANO = "soprano" # audio generation: Vocos decoder
    HUNYUANVL      = "hunyuanvl"
    PARAKEET       = "parakeet"  # audio
    MINIMAXM3      = "minimax_m3"
@@ -572,6 +572,10 @@ struct llm_tokenizer_bpe : llm_tokenizer {
                    "(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])|[^\\r\\n\\p{L}\\p{N}]?(?:\\p{L}|\\p{M}|\\u200C|\\u200D)+|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+",
                };
                break;
+            case LLAMA_VOCAB_PRE_TYPE_SOPRANO:
+                regex_exprs = { "\\p{N}", "\\s+|[\\p{L}\\p{N}_]+|[^\\p{L}\\p{N}_\\s]+" };
+                byte_encode = false;
+                break;
            case LLAMA_VOCAB_PRE_TYPE_WHITESPACE:
                // whitespace pre-tokenizer (jinaai/jina-embeddings-v2-base-zh)
                regex_exprs = {
@@ -647,6 +651,22 @@ struct llm_tokenizer_bpe_session {
            replace_all(normalized_text, "[" + zwnj, "[");
            replace_all(normalized_text, "<" + zwnj, "<");
            word_collection = unicode_regex_split(normalized_text, tokenizer.regex_exprs, tokenizer.byte_encode);
+        } else if (tok_pre == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
+            std::string normalized_text;
+            bool space = false;
+            for (uint32_t cpt : unicode_cpts_from_utf8(text)) {
+                if (unicode_cpt_flags_from_cpt(cpt).is_whitespace) {
+                    if (!space) {
+                        normalized_text += ' ';
+                    }
+                    space = true;
+                } else {
+                    // Unicode full lowercase expands LATIN CAPITAL LETTER I WITH DOT ABOVE.
+                    normalized_text += cpt == 0x0130 ? "i\xcc\x87" : unicode_cpt_to_utf8(unicode_tolower(cpt));
+                    space = false;
+                }
+            }
+            word_collection = unicode_regex_split(normalized_text, tokenizer.regex_exprs, tokenizer.byte_encode);
        } else {
            word_collection = unicode_regex_split(text, tokenizer.regex_exprs, tokenizer.byte_encode);
        }
@@ -769,7 +789,9 @@ struct llm_tokenizer_bpe_session {
                const std::string str = std::string(symbol.text, symbol.n);
                const auto token = vocab.text_to_token(str);

-                if (token == LLAMA_TOKEN_NULL) {
+                if (token == LLAMA_TOKEN_NULL && tok_pre == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
+                    output.push_back(vocab.token_unk());
+                } else if (token == LLAMA_TOKEN_NULL) {
                    for (auto j = str.begin(); j != str.end(); ++j) {
                        llama_token token_multibyte = LLAMA_TOKEN_NULL;
                        if (tokenizer.byte_encode) {
@@ -2730,6 +2752,9 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
                pre_type = LLAMA_VOCAB_PRE_TYPE_NEMOTRON_3_5_EMBD;
                ignore_merges = true;
                add_bos = true;
+            } else if (
+                tokenizer_pre == "soprano") {
+                pre_type = LLAMA_VOCAB_PRE_TYPE_SOPRANO;
            } else {
                throw std::runtime_error(format("unknown pre-tokenizer type: '%s'", tokenizer_pre.c_str()));
            }
@@ -3283,6 +3308,10 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
            }
        }

+        if (pre_type == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
+            special_eog_ids.insert(special_eos_id);
+        }
+
        // sanity checks
        if (special_eos_id != LLAMA_TOKEN_NULL && special_eog_ids.count(special_eos_id) == 0) {
            special_eog_ids.insert(special_eos_id);
@@ -4096,6 +4125,9 @@ int32_t llama_vocab::impl::token_to_piece(llama_token token, char * buf, int32_t
                    return _try_copy(token_text.data(), token_text.size());
                }
                if (attr & LLAMA_TOKEN_ATTR_NORMAL) {
+                    if (pre_type == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
+                        return _try_copy(token_text.data(), token_text.size());
+                    }
                    if (escape_whitespaces) {
                        // SPM-style BPE: tokens contain ▁ for spaces
                        std::string result = token_text;
@@ -73,6 +73,7 @@ enum llama_vocab_pre_type {
    LLAMA_VOCAB_PRE_TYPE_FLINT             = 62,
    LLAMA_VOCAB_PRE_TYPE_DOMYN_SMALL       = 63,
    LLAMA_VOCAB_PRE_TYPE_NEMOTRON_3_5_EMBD = 64,
+    LLAMA_VOCAB_PRE_TYPE_SOPRANO           = 65,
 };

 // HF Metaspace pre-tokenizer prepend_scheme
@@ -240,6 +240,26 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
        FIXTURES_SETUP generate-models
    )

+    # Target model for test-speculative-rejection-sampling
+    set(SPEC_TARGET_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-speculative-models/target")
+    file(MAKE_DIRECTORY "${SPEC_TARGET_DIR}")
+    llama_test(
+        test-llama-archs
+        NAME test-generate-speculative-target
+        LABEL main
+        ARGS -a llama -s 999 -o "${SPEC_TARGET_DIR}"
+    )
+    set_tests_properties(
+        test-generate-speculative-target
+        PROPERTIES FIXTURES_SETUP speculative-model
+    )
+
+    # To avoid recreating the model use the fixture exclude any option:
+    # -FA "^speculative-model$"
+    llama_build_and_test(test-speculative-rejection-sampling.cpp ARGS
+        -m "${SPEC_TARGET_DIR}/llama-dense.gguf")
+    set_tests_properties(test-speculative-rejection-sampling PROPERTIES FIXTURES_REQUIRED speculative-model)
+
    # Test context state handling across all architectures, using the generated dummy models
    llama_test(
        test-llama-context
@@ -1194,6 +1194,29 @@ static bool op_filter_entry_matches(std::string_view entry, std::string_view op_
    return std::regex_search(op_name.data(), op_name.data() + op_name.size(), re);
 }

+// TODO: [TAG_MUL_MAT_ID_SKIP] see https://github.com/ggml-org/llama.cpp/pull/26631
+static bool backend_lacks_skipped_ids(ggml_backend_t backend, ggml_cgraph * gf) {
+    const char * name = ggml_backend_reg_name(ggml_backend_dev_backend_reg(ggml_backend_get_device(backend)));
+    if (strcmp(name, "CUDA") != 0 && strcmp(name, "ROCm") != 0 && strcmp(name, "MUSA") != 0) {
+        return false;
+    }
+    for (int i = 0; i < ggml_graph_n_nodes(gf); i++) {
+        const ggml_tensor * node = ggml_graph_node(gf, i);
+        if (node->op != GGML_OP_MUL_MAT_ID && node->op != GGML_OP_ADD_ID) {
+            continue;
+        }
+        const ggml_tensor * ids = node->src[2];
+        std::vector<int32_t> row(ids->ne[0]);
+        for (int64_t i1 = 0; i1 < ids->ne[1]; i1++) {
+            ggml_backend_tensor_get(ids, row.data(), i1*ids->nb[1], ids->ne[0]*sizeof(int32_t));
+            if (std::find(row.begin(), row.end(), -1) != row.end()) {
+                return true;
+            }
+        }
+    }
+    return false;
+}
+
 struct test_case {
    virtual ~test_case() {}

@@ -1475,6 +1498,13 @@ struct test_case {
            initialize_tensors(ctx_weights.get());
        }

+        // TODO: [TAG_MUL_MAT_ID_SKIP] @ngxson : remove this once CUDA implements mul_mat_id with -1 = skipped
+        if (backend_lacks_skipped_ids(backend1, gf)) {
+            test_result result(ggml_backend_name(backend1), current_op_name, vars(), "test", false, false, "not supported");
+            print_test_result_locked(output_printer, result);
+            return test_status_t::NOT_SUPPORTED;
+        }
+
        // compare
        struct callback_userdata {
            bool   ok;
@@ -3564,6 +3594,16 @@ struct test_bin_bcast : public test_case {
    }
 };

+// set some ids to -1, so that MUL_MAT_ID and ADD_ID skip those slots
+// every 4th row skips all of its slots, the other rows skip only some of them
+static void mask_skipped_ids(std::vector<int32_t> & ids, int64_t row) {
+    for (size_t i = 0; i < ids.size(); i++) {
+        if (row % 4 == 1 || (row + (int64_t) i) % 3 == 0) {
+            ids[i] = -1;
+        }
+    }
+}
+
 // GGML_OP_ADD_ID
 struct test_add_id : public test_case {
    const ggml_type type_a;
@@ -3572,9 +3612,10 @@ struct test_add_id : public test_case {
    const int64_t n_experts;
    const int64_t n_experts_used;
    const int64_t n_token;
+    const bool skip; // set some ids to a negative value

    std::string vars() override {
-        return VARS_TO_STR6(type_a, type_b, n_embd, n_experts, n_experts_used, n_token);
+        return VARS_TO_STR7(type_a, type_b, n_embd, n_experts, n_experts_used, n_token, skip);
    }

    size_t op_size(ggml_tensor * t) override {
@@ -3586,9 +3627,10 @@ struct test_add_id : public test_case {
            int64_t n_embd = 128,
            int64_t n_experts = 16,
            int64_t n_experts_used = 8,
-            int64_t n_token = 10)
+            int64_t n_token = 10,
+            bool skip = false)
        : type_a(type_a), type_b(type_b), n_embd(n_embd),
-          n_experts(n_experts), n_experts_used(n_experts_used), n_token(n_token) {}
+          n_experts(n_experts), n_experts_used(n_experts_used), n_token(n_token), skip(skip) {}

    ggml_tensor * build_graph(ggml_context * ctx) override {
        ggml_tensor * a = ggml_new_tensor_3d(ctx, type_a, n_embd, n_experts_used, n_token);
@@ -3617,6 +3659,9 @@ struct test_add_id : public test_case {
                        data[i] = i % n_experts;
                    }
                    std::shuffle(data.begin(), data.end(), rng);
+                    if (skip) {
+                        mask_skipped_ids(data, r);
+                    }
                    ggml_backend_tensor_set(t, data.data(), r * t->nb[1], t->ne[0] * sizeof(int32_t));
                }
            } else {
@@ -5550,7 +5595,8 @@ struct test_mul_mat_w4a4 : public test_mul_mat {
    }
 };

-static void init_mul_mat_id_ids(ggml_context * ctx, int n_mats) {
+static void init_mul_mat_id_ids(ggml_context * ctx, int n_mats, bool skip = false) {
+
    std::random_device rd;
    std::default_random_engine rng(rd());
    for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
@@ -5563,12 +5609,15 @@ static void init_mul_mat_id_ids(ggml_context * ctx, int n_mats) {
                data[i] = i % n_mats;
            }
            std::shuffle(data.begin(), data.end(), rng);
+            if (skip) {
+                mask_skipped_ids(data, r);
+            }
            ggml_backend_tensor_set(t, data.data(), r * t->nb[1], t->ne[0] * sizeof(int32_t));
        }
    }
 }

-static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats, float amax = 1.0f) {
+static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats, bool skip = false, float amax = 1.0f) {
    for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
        if (t->type == GGML_TYPE_I32) {
            continue;
@@ -5579,7 +5628,7 @@ static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats, float amax =
            init_tensor_uniform(t);
        }
    }
-    init_mul_mat_id_ids(ctx, n_mats);
+    init_mul_mat_id_ids(ctx, n_mats, skip);
 }

 // GGML_OP_MUL_MAT_ID
@@ -5595,9 +5644,10 @@ struct test_mul_mat_id : public test_case {
    const float amax; // magnitude of src1
    const int64_t m_v; // rows of as in memory, the experts of as are strided for m_v > m, no view for m_v == 0
    const bool ids_offset;
+    const bool skip; // set some ids to a negative value

    std::string vars() override {
-        return VARS_TO_STR11(type_a, type_b, n_mats, n_used, b, m, n, k, amax, m_v, ids_offset);
+        return VARS_TO_STR12(type_a, type_b, n_mats, n_used, b, m, n, k, amax, m_v, ids_offset, skip);
    }

    double max_nmse_err() override {
@@ -5622,9 +5672,9 @@ struct test_mul_mat_id : public test_case {
    test_mul_mat_id(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
            int n_mats = 8, int n_used = 2, bool b = false,
            int64_t m = 32, int64_t n = 32, int64_t k = 32,
-            float amax = 1.0f, int64_t m_v = 0, bool ids_offset = false)
+            float amax = 1.0f, int64_t m_v = 0, bool ids_offset = false, bool skip = false)
        : type_a(type_a), type_b(type_b), n_mats(n_mats), n_used(n_used), b(b),
-            m(m), n(n), k(k), amax(amax), m_v(m_v), ids_offset(ids_offset) {
+            m(m), n(n), k(k), amax(amax), m_v(m_v), ids_offset(ids_offset), skip(skip) {
            GGML_ASSERT(n_used <= n_mats);
            GGML_ASSERT(m_v == 0 || m_v > m);
        }
@@ -5659,11 +5709,11 @@ struct test_mul_mat_id : public test_case {
    }

    void initialize_tensors(ggml_context * ctx) override {
-        init_mul_mat_id_tensors(ctx, n_mats, amax);
+        init_mul_mat_id_tensors(ctx, n_mats, skip, amax);
    }

    void reinit_perf_iter(ggml_context * ctx) override {
-        init_mul_mat_id_ids(ctx, n_mats);
+        init_mul_mat_id_ids(ctx, n_mats, skip);
    }
 };

@@ -5738,9 +5788,10 @@ struct test_mul_mat_id_fusion : public test_case {
    const int64_t k;
    const uint32_t o; // number of outputs
    const bool mul;
+    const bool skip; // set some ids to a negative value

    std::string vars() override {
-        return VARS_TO_STR10(type_a, type_b, n_mats, n_used, b, m, n, k, o, mul);
+        return VARS_TO_STR11(type_a, type_b, n_mats, n_used, b, m, n, k, o, mul, skip);
    }

    double max_nmse_err() override {
@@ -5754,9 +5805,9 @@ struct test_mul_mat_id_fusion : public test_case {

    test_mul_mat_id_fusion(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
            int n_mats = 8, int n_used = 2, bool b = false,
-            int64_t m = 32, int64_t n = 32, int64_t k = 32, uint32_t o = 1, bool mul = false)
+            int64_t m = 32, int64_t n = 32, int64_t k = 32, uint32_t o = 1, bool mul = false, bool skip = false)
        : type_a(type_a), type_b(type_b), n_mats(n_mats), n_used(n_used), b(b),
-            m(m), n(n), k(k), o(o), mul(mul) {
+            m(m), n(n), k(k), o(o), mul(mul), skip(skip) {
            GGML_ASSERT(n_used <= n_mats);
        }

@@ -5796,7 +5847,7 @@ struct test_mul_mat_id_fusion : public test_case {
    }

    void initialize_tensors(ggml_context * ctx) override {
-        init_mul_mat_id_tensors(ctx, n_mats);
+        init_mul_mat_id_tensors(ctx, n_mats, skip);
    }

    bool run_whole_graph() override { return true; }
@@ -7754,19 +7805,20 @@ struct test_mul_mat_vec_fusion : public test_case {
    const bool with_gate;
    const bool with_lane_scale;
    std::array<int64_t, 2> batch_dims;
+    const bool skip; // set some ids to a negative value (only for use_id)

    test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k,
                        bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true,
-                        bool with_lane_scale = false, std::array<int64_t, 2> batch_dims = {4, 2})
+                        bool with_lane_scale = false, std::array<int64_t, 2> batch_dims = {4, 2}, bool skip = false)
    : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias),
-        with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims) {
+        with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims), skip(skip) {
        if (use_id) {
            GGML_ASSERT(n_used <= n_mats);
        }
    }

    std::string vars() override {
-        return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims);
+        return VARS_TO_STR14(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims, skip);
    }

    std::string op_desc(ggml_tensor * t) override {
@@ -7917,7 +7969,7 @@ struct test_mul_mat_vec_fusion : public test_case {
                init_tensor_uniform(t);
            }
        } else {
-            init_mul_mat_id_tensors(ctx, n_mats);
+            init_mul_mat_id_tensors(ctx, n_mats, skip);
        }
    }

@@ -11453,6 +11505,21 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
        }
    }

+    // skipped slots (negative ids)
+    for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16, GGML_TYPE_Q4_0, GGML_TYPE_Q4_K, GGML_TYPE_Q8_0, GGML_TYPE_MXFP4}) {
+        for (int n_used : {1, 2, 4}) {
+            for (bool b : {false, true}) {
+                for (int n : {1, 4, 33, 129}) {
+                    test_cases.emplace_back(new test_mul_mat_id(type_a, GGML_TYPE_F32, 8, n_used, b, 512, n, 256, 1.0f, 0, false, true));
+                }
+            }
+        }
+    }
+    // n_used == n_mats, so every expert can be skipped
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 4, false, 512, 64, 256, 1.0f, 0, false, true));
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, 1, 1.0f, 0, false, true));
+    test_cases.emplace_back(new test_mul_mat_id_fusion(GGML_TYPE_F16, GGML_TYPE_F32, 16, 16, false, 32, 32, 32, 3, false, true));
+
    for (int bs : {1, 4, 512}) {
        for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q4_K}) {
            for (ggml_type type_b : {GGML_TYPE_F32}) {
@@ -11500,6 +11567,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
                    for (int n_embd : {32, 129}) {
                        for (int n_token : {1, 32, 129}) {
                            test_cases.emplace_back(new test_add_id(type_a, type_b, n_embd, n_mats, n_used, n_token));
+                            test_cases.emplace_back(new test_add_id(type_a, type_b, n_embd, n_mats, n_used, n_token, true));
                        }
                    }
                }
@@ -12238,6 +12306,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
                                    test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, m_batch, 32, 256,
                                        use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}));
                                }
+                                // the lane scale does a GET_ROWS on the ids, which does not accept -1
+                                if (use_id && !with_lane_scale) {
+                                    test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 8, 32, 256,
+                                        use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}, true));
+                                }
                            }
                        }
                    }
@@ -5,6 +5,7 @@
 #include "mtmd-internal.h"

 #include <cmath>
+#include <cstdint>
 #include <iostream>
 #include <stdexcept>
 #include <string>
@@ -95,6 +96,45 @@ static std::string image_size_to_str(const clip_image_size & size) {
    return std::to_string(size.width) + "x" + std::to_string(size.height);
 }

+static std::string resize_algo_to_str(resize_algo algo) {
+    std::string str;
+    switch (algo & ~RESIZE_ALGO_FLAG_TORCH) {
+        case RESIZE_ALGO_BILINEAR: str = "bilinear"; break;
+        case RESIZE_ALGO_BICUBIC:  str = "bicubic";  break;
+        case RESIZE_ALGO_LANCZOS:  str = "lanczos";  break;
+        default:                   str = "unknown";  break;
+    }
+    if (algo & RESIZE_ALGO_FLAG_TORCH) {
+        str += "_torch";
+    }
+    return str;
+}
+
+// pseudo-random pixels: s = s * 1103515245 + 12345 from s = 1, one byte (s >> 16) & 255 per channel
+static clip_image_u8 make_test_image(const clip_image_size & size) {
+    clip_image_u8 img;
+    img.set_size(size, false);
+    uint32_t s = 1;
+    auto next = [&s]() {
+        s = s * 1103515245u + 12345u;
+        return static_cast<uint8_t>(s >> 16);
+    };
+    for (int y = 0; y < size.height; y++) {
+        for (int x = 0; x < size.width; x++) {
+            img.set_pixel(x, y, { next(), next(), next() });
+        }
+    }
+    return img;
+}
+
+// FNV-1a of the pixels
+static uint32_t hash_pixels(const clip_image_u8 & img, uint32_t hash = 2166136261u) {
+    for (uint8_t v : img.get_ro_buf()) {
+        hash = (hash ^ v) * 16777619u;
+    }
+    return hash;
+}
+
 MAKE_TEST(test_image_preprocessor_lfm2) {
    clip_hparams hparams;
    hparams.patch_size = 16;
@@ -158,6 +198,42 @@ MAKE_TEST(test_image_preprocessor_lfm2) {
            std::string(mtmd_image_preprocessor_lfm2::should_tile(hparams, { 1024, 977 }) ? "tiled" : "single"));
 }

+MAKE_TEST(test_image_preprocessor_lfm2_torch_resize) {
+    clip_hparams hparams;
+    hparams.patch_size = 16;
+    hparams.n_merge = 2;
+    hparams.set_limit_image_tokens(64, 256);
+    hparams.preproc_min_tiles            = 2;
+    hparams.preproc_max_tiles            = 10;
+    hparams.preproc_tile_size            = 512;
+    hparams.preproc_max_pixels_tolerance = 2.0f;
+    // Lfm2VlImageProcessor stretches the image to the tile grid
+    hparams.image_pad_rf = PAD_NONE;
+
+    // { resize algo, image size, expected hash of the tiles and the thumbnail from Lfm2VlImageProcessor (torchvision 0.29, native uint8 resize: arm64, or x86 with AVX2/AVX512) }
+    const std::vector<std::tuple<resize_algo, clip_image_size, uint32_t>> cases = {
+        { RESIZE_ALGO_BILINEAR, {  620,  410 }, 0xf5e9c941 }, // 608x416
+        { RESIZE_ALGO_BILINEAR, { 2000, 1500 }, 0x3331f56c }, // 3x2 tiles, 576x416 thumbnail
+        { RESIZE_ALGO_BICUBIC,  {  620,  410 }, 0x2b34eac9 },
+        { RESIZE_ALGO_BICUBIC,  { 2000, 1500 }, 0xb035d942 },
+    };
+
+    for (const auto & [base_algo, size, expected] : cases) {
+        const resize_algo algo = static_cast<resize_algo>(base_algo | RESIZE_ALGO_FLAG_TORCH);
+        hparams.image_resize_algo_rf = algo;
+        hparams.image_resize_algo_ov = algo;
+        const mtmd_image_preprocessor_lfm2 preproc(hparams);
+        const auto sliced = preproc.slice_image(make_test_image(size), preproc.get_slice_instructions(size));
+
+        uint32_t hash = 2166136261u;
+        for (const auto & slice : sliced.slices) {
+            hash = hash_pixels(slice, hash);
+        }
+        hash = hash_pixels(sliced.overview, hash);
+        t.assert_equal("pixels of " + std::to_string(size.width) + "x" + std::to_string(size.height) + " with resize algo " + resize_algo_to_str(algo), expected, hash);
+    }
+}
+
 //
 // mtmd temporal merge
 //
@@ -0,0 +1,229 @@
+#include "sampling.h"
+#include "llama.h"
+#include "common.h"
+#include "arg.h"
+#include "log.h"
+
+#include <cmath>
+#include <cstdlib>
+#include <cstring>
+#include <optional>
+#include <random>
+#include <unordered_map>
+
+// Tests that common_sampler_sample_and_accept_n_rejection preserves the target
+// distribution regardless of what the draft distribution Q looks like.
+
+static int test_speculative_rejection_sampling(common_params & params, int n_trials) {
+    params.sampling.backend_sampling = false;
+    auto llama_init = common_init_from_params(params);
+
+    auto * model_tgt = llama_init->model();
+    auto * ctx_tgt   = llama_init->context();
+    if (!model_tgt || !ctx_tgt) {
+        LOG_ERR("failed to initialize target model and context\n");
+        return 1;
+    }
+
+    const llama_tokens prompt = { 1, 2, 3, 1, 2 };
+
+    // Process the prompt with the target model to get its true distribution for
+    // the position right after the prompt.
+    common_batch batch = common_batch_get_one(ctx_tgt, prompt);
+    if (llama_process(ctx_tgt, LLAMA_PROCESS_TYPE_DECODE, batch.get()) != 0) {
+        LOG_ERR("failed to decode the test prompt\n");
+        return 1;
+    }
+
+    common_params_sampling sampling;
+    sampling.seed     = 18;
+    sampling.samplers = { COMMON_SAMPLER_TYPE_TEMPERATURE };
+    sampling.temp     = 1.0f;
+
+    // Initialize the sampler for the target model.
+    common_sampler_ptr smpl(common_sampler_init(model_tgt, sampling));
+    if (!smpl) {
+        LOG_ERR("failed to initialize target sampler\n");
+        return 1;
+    }
+
+    // Sample once so the sampler chain runs and populates cur_p. We don't need
+    // the returned token id itself so the return value is ignored.
+    common_sampler_sample(smpl.get(), ctx_tgt, -1);
+
+    // Get the target candidates distribution in sorted order.
+    const llama_token_data_array * dist_tgt = common_sampler_get_candidates(smpl.get(), true);
+
+    // Clone smpl before the trial loop starts. We need the dist_tgt pointer to stay valid.
+    common_sampler_ptr smpl_verify(common_sampler_clone(smpl.get()));
+    if (!smpl_verify) {
+        LOG_ERR("failed to clone target sampler\n");
+        return 1;
+    }
+
+    const llama_vocab * vocab = llama_model_get_vocab(model_tgt);
+    const int n_vocab = llama_vocab_n_tokens(vocab);
+
+    std::mt19937 rng(18);
+    std::uniform_real_distribution<float> uni(0.0f, 1.0f);
+
+    std::unordered_map<llama_token, int> token_counts;
+
+    std::vector<llama_token_data> q(n_vocab);
+    for (int i = 0; i < n_vocab; ++i) {
+        q[i].id = i;
+        q[i].logit = 0.0f;
+    }
+
+    for (int trial = 0; trial < n_trials; ++trial) {
+        float q_sum = 0.0f;
+
+        for (auto & td : q) {
+            td.p = uni(rng);
+            q_sum += td.p;
+        }
+
+        const float inv_sum = 1.0f / q_sum;
+        for (auto & td : q) {
+            td.p *= inv_sum;
+        }
+
+        float u = uni(rng);
+
+        // Fallback to the last bucket to absorb residual probability from FP precision loss
+        llama_token draft_tok = q.back().id;
+        for (const auto & e : q) {
+            u -= e.p;
+            if (u <= 0.0f) {
+                draft_tok = e.id;
+                break;
+            }
+        }
+
+        // Both indices reuse the prompt's last output logits. The second entry
+        // is the bonus draw slot; it is always stale and its result is ignored.
+        const llama_tokens draft = { draft_tok };
+        const std::vector<std::vector<llama_token_data>> draft_q = { q };
+        const std::vector<int> idxs = { -1, -1 };
+
+        const std::vector<llama_token> result =
+            common_sampler_sample_and_accept_n_rejection(smpl_verify.get(), ctx_tgt, idxs, draft, draft_q, false);
+
+        token_counts[result[0]]++;
+    }
+
+    LOG("\n");
+    LOG("Empirical vs target distribution after %d trials:\n", n_trials);
+
+    bool   ok      = true;
+    size_t covered = 0;
+
+    auto get_count_for = [&](llama_token id) -> int {
+        const auto it = token_counts.find(id);
+        return it != token_counts.end() ? it->second : 0;
+    };
+
+    // Iterate over the target distribution
+    for (size_t i = 0; i < dist_tgt->size; ++i) {
+        const llama_token id  = dist_tgt->data[i].id;
+        const float       p_t = dist_tgt->data[i].p;
+
+        const int   count = get_count_for(id);
+        const float p_e   = (float) count / n_trials;
+        const float diff  = std::fabs(p_t - p_e);
+
+        // We only deal with ids in dist_tgt, and if each trial result was one
+        // of dist_tgt's candidates then covered ends up equal to the number of
+        // trials. If not this inconsistency tells us how many trials produced
+        // a token outside of the target's distribution.
+        covered += count;
+
+        const float standard_error = std::sqrt(p_t * (1.0f - p_t) / n_trials);
+        const float tolerance      = 6.0f * standard_error;
+
+        LOG("token %3d : target = %f, empirical = %f, diff = %f, tol = %f\n", id, p_t, p_e, diff, tolerance);
+
+        if (diff > tolerance) {
+            LOG_ERR("token %d: empirical probability %f too far from target %f (diff %f > tol %f)\n",
+                    id, p_e, p_t, diff, tolerance);
+            ok = false;
+        }
+    }
+
+    if (covered != (size_t) n_trials) {
+        LOG_ERR("%zu of %d trials produced a token outside dist_tgt's candidates\n",
+                (size_t) n_trials - covered, n_trials);
+        ok = false;
+    }
+
+    if (!ok) {
+        return 1;
+    }
+
+    return 0;
+}
+
+struct parsed_args {
+    int n_trials;
+    std::vector<char *> argv;
+};
+
+static std::optional<parsed_args> parse_n_trials_arg(int argc, char ** argv, int default_n_trials) {
+    parsed_args result;
+    result.n_trials = default_n_trials;
+    result.argv.push_back(argv[0]);
+
+    for (int i = 1; i < argc; ++i) {
+        if (strcmp(argv[i], "--n-trials") == 0) {
+            if (i + 1 >= argc) {
+                LOG_ERR("%s: --n-trials requires a number argument\n", __func__);
+                return std::nullopt;
+            }
+            int value = std::atoi(argv[i + 1]);
+            if (value <= 0) {
+                LOG_ERR("%s: --n-trials requires a positive number argument\n", __func__);
+                return std::nullopt;
+            }
+            result.n_trials = value;
+            i++;
+        } else {
+            result.argv.push_back(argv[i]);
+        }
+    }
+    result.argv.push_back(nullptr);
+
+    return result;
+}
+
+int main(int argc, char ** argv) {
+    std::setlocale(LC_NUMERIC, "C");
+
+    common_params params;
+    params.n_ctx = 256;
+
+    common_init();
+
+    auto parsed = parse_n_trials_arg(argc, argv, 5000);
+    if (!parsed) {
+        return 1;
+    }
+    const int fargc = (int) parsed->argv.size() - 1;
+
+    if (!common_params_parse(fargc, parsed->argv.data(), params, LLAMA_EXAMPLE_COMMON)) {
+        return 1;
+    }
+
+    LOG("Using target model '%s'\n", params.model.path.c_str());
+
+    llama_backend_init();
+
+    const int result = test_speculative_rejection_sampling(params, parsed->n_trials);
+    if (result == 0) {
+        LOG("\nAll tests passed.\n");
+    } else {
+        LOG("\nSome tests failed.\n");
+    }
+
+    llama_backend_free();
+    return result;
+}
@@ -67,6 +67,7 @@ add_library(mtmd
            models/pockettts-seanet.cpp
            models/pockettts-spkenc.cpp
            models/pockettts-gen.cpp
+            models/soprano.cpp
            models/step3vl.cpp
            models/siglip.cpp
            models/whisper-enc.cpp
@@ -101,6 +101,10 @@ struct clip_graph {

    ggml_tensor * build_inp_raw(int channels = 3);

+    // qwen-vl style temporal merge, handles n_batch == 1 (image) or n_batch == 2 (2 video frames)
+    // requires patch_embeddings_0 and patch_embeddings_1; returns tensor with shape [w, h, n_embd, 1]
+    ggml_tensor * build_inp_temporal_merge();
+
    // f16 if flash attn is enabled, set it with set_input_attn_mask()
    // idx is only needed when the graph has more than one mask
    ggml_tensor * build_inp_attn_mask(int64_t n_kv, int64_t n_q, int idx = 0);
@@ -300,6 +300,12 @@
 #define TN_A_GEN_WAV_SEANET_RES_CONV2  "a.gen.wav.seanet.blk.%d.res_conv2.%s"
 #define TN_A_GEN_WAV_SEANET_SCALE_CONV "a.gen.wav.seanet.blk.%d.scale_conv.%s"

+// soprano: Vocos decoder, the ConvNeXt blocks use TN_A_GEN_WAV_UP_*
+#define TN_A_GEN_WAV_INPUT           "a.gen.wav.input.%s"
+#define TN_A_GEN_WAV_NORM            "a.gen.wav.norm.%s"
+#define TN_A_GEN_WAV_OUTPUT_NORM     "a.gen.wav.output_norm.%s"
+#define TN_A_GEN_WAV_OUTPUT          "a.gen.wav.output.%s"
+
 // cogvlm
 #define TN_MM_POST_FC_NORM "mm.post_fc_norm.%s"
 #define TN_MM_H_TO_4H      "mm.up.%s"
@@ -547,6 +553,7 @@ enum projector_type {
    PROJECTOR_TYPE_QWEN3TTS_GEN,
    PROJECTOR_TYPE_POCKETTTS_SPKENC,
    PROJECTOR_TYPE_POCKETTTS_GEN,
+    PROJECTOR_TYPE_SOPRANO,
    PROJECTOR_TYPE_MUSE_GLIMMER,
    PROJECTOR_TYPE_SENSENOVA_U1,
    PROJECTOR_TYPE_COHERE2V,
@@ -623,6 +630,7 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
    { PROJECTOR_TYPE_QWEN3TTS_GEN,      "qwen3tts_gen"},
    { PROJECTOR_TYPE_POCKETTTS_SPKENC,  "pockettts_spkenc"},
    { PROJECTOR_TYPE_POCKETTTS_GEN,     "pockettts_gen"},
+    { PROJECTOR_TYPE_SOPRANO,           "soprano"},
    { PROJECTOR_TYPE_MUSE_GLIMMER,      "muse-glimmer"},
    { PROJECTOR_TYPE_SENSENOVA_U1,      "sensenova_u1"},
    { PROJECTOR_TYPE_COHERE2V,          "cohere2v"},
@@ -29,11 +29,14 @@ enum patch_merge_type {
    PATCH_MERGE_SPATIAL_UNPAD,
 };

-// all algos are Pillow-compatible (matching PIL.Image.resize output)
+// the low bits select the algo, Pillow-compatible (matching PIL.Image.resize output)
+// RESIZE_ALGO_FLAG_TORCH matches torchvision's antialiased resize of uint8 images on the CPU (int16 weights), bilinear and bicubic only
 enum resize_algo {
    RESIZE_ALGO_BILINEAR,
    RESIZE_ALGO_BICUBIC,
    RESIZE_ALGO_LANCZOS,
+
+    RESIZE_ALGO_FLAG_TORCH = 1 << 8,
 };

 // Padding style for img_tool::resize
@@ -207,6 +210,11 @@ struct clip_hparams {
    int32_t custom_image_min_tokens = -1;
    int32_t custom_image_max_tokens = -1;

+    void set_image_resize_algo(resize_algo algo, bool torch = false) {
+        GGML_ASSERT(!(torch && algo == RESIZE_ALGO_LANCZOS) && "torchvision has no lanczos");
+        image_resize_algo = static_cast<resize_algo>(torch ? algo | RESIZE_ALGO_FLAG_TORCH : algo);
+    }
+
    void set_limit_image_tokens(int n_tokens_min, int n_tokens_max) {
        const int patch_area = patch_size * patch_size * n_merge * n_merge;
        image_min_pixels = (custom_image_min_tokens > 0 ? custom_image_min_tokens : n_tokens_min) * patch_area;
@@ -579,6 +587,18 @@ struct clip_code2wav {
    ggml_tensor * dac_post_conv_b      = nullptr;
 };

+struct clip_vocos {
+    ggml_tensor * input_w        = nullptr;
+    ggml_tensor * input_b        = nullptr;
+    ggml_tensor * norm_w         = nullptr;
+    ggml_tensor * norm_b         = nullptr;
+    ggml_tensor * output_norm_w  = nullptr;
+    ggml_tensor * output_norm_b  = nullptr;
+    ggml_tensor * output_w       = nullptr;
+    ggml_tensor * output_b       = nullptr;
+    std::vector<clip_code2wav::upsample_block> blocks;
+};
+
 struct clip_model {
    clip_modality modality = CLIP_MODALITY_VISION;
    projector_type proj_type = PROJECTOR_TYPE_MLP;
@@ -845,6 +865,9 @@ struct clip_model {
    ggml_tensor * gen_upsample_w  = nullptr; // depthwise convtr, frame rate -> encoder frame rate
    std::vector<clip_layer> gen_tfm_layers; // mimi decoder_transformer

+    // soprano: Vocos decoder, hidden states -> STFT coefficients
+    clip_vocos vocos;
+
    // cogvlm
    ggml_tensor * mm_post_fc_norm_w = nullptr;
    ggml_tensor * mm_post_fc_norm_b = nullptr;
@@ -584,6 +584,36 @@ ggml_tensor * clip_graph::build_inp_raw(int channels) {
    return inp_raw;
 }

+// qwen-vl style temporal merge: Conv3D(kt=2) split into 2 Conv2D
+ggml_tensor * clip_graph::build_inp_temporal_merge() {
+    ggml_tensor * inp_raw = build_inp_raw();
+
+    GGML_ASSERT(img.nx() % (patch_size * 2) == 0);
+    GGML_ASSERT(img.ny() % (patch_size * 2) == 0);
+
+    const size_t nb1 = ggml_row_size(inp_raw->type, img.nx());
+    const size_t nb2 = ggml_row_size(inp_raw->type, img.nx() * img.ny());
+
+    if (n_batch == 1) {
+        // still image input
+        return ggml_add(ctx0,
+            ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_raw, patch_size, patch_size, 0, 0, 1, 1),
+            ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_raw, patch_size, patch_size, 0, 0, 1, 1));
+    } else if (n_batch == 2) {
+        // 2 frames input (video input)
+        ggml_tensor * inp_0 = ggml_view_3d(ctx0, inp_raw,
+                                    img.nx(), img.ny(), 3, nb1, nb2, 0);
+        ggml_tensor * inp_1 = ggml_view_3d(ctx0, inp_raw,
+                                    img.nx(), img.ny(), 3, nb1, nb2,
+                                    nb2 * 3); // move to the second frame
+        return ggml_add(ctx0,
+            ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_0, patch_size, patch_size, 0, 0, 1, 1),
+            ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_1, patch_size, patch_size, 0, 0, 1, 1));
+    } else {
+        GGML_ASSERT(false && "n_batch > 2 is not supported");
+    }
+}
+
 static std::string get_attn_mask_name(int idx) {
    return idx == 0 ? "attn_mask" : "attn_mask_" + std::to_string(idx);
 }
@@ -1168,6 +1198,11 @@ static std::unique_ptr<clip_graph> clip_get_graph_builder(clip_ctx * ctx, const
                const float top_p = params ? params->top_p : 1.0f;
                builder = std::make_unique<clip_graph_qwen3tts_gen>(ctx, img, gen_process, top_k, top_p);
            } break;
+        case PROJECTOR_TYPE_SOPRANO:
+            {
+                const int n_frames = params && params->feats ? (int) (params->feats->size() / ctx->model.hparams.projection_dim) : 2;
+                builder = std::make_unique<clip_graph_soprano>(ctx, img, n_frames);
+            } break;
        case PROJECTOR_TYPE_YOUTUVL:
            {
                builder = std::make_unique<clip_graph_youtuvl>(ctx, img);
@@ -1584,12 +1619,13 @@ struct clip_model_loader {
                        // default for older GGUFs
                        std::string resize_algo = "bilinear";
                        get_string(KEY_IMAGE_RESIZE_ALGO, resize_algo, false);
+                        // Lfm2VlImageProcessor and d1-omni vision.py resize bilinear and bicubic with torchvision
                        if (resize_algo == "bilinear") {
-                            hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                            hparams.set_image_resize_algo(RESIZE_ALGO_BILINEAR, true);
                        } else if (resize_algo == "bicubic") {
-                            hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                            hparams.set_image_resize_algo(RESIZE_ALGO_BICUBIC, true);
                        } else if (resize_algo == "lanczos") {
-                            hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
+                            hparams.set_image_resize_algo(RESIZE_ALGO_LANCZOS);
                        } else {
                            throw std::runtime_error(string_format("%s: unsupported image resize algo: %s\n", __func__, resize_algo.c_str()));
                        }
@@ -2073,6 +2109,12 @@ struct clip_model_loader {
                        hparams.flow_n_step       = 1;
                        hparams.gen_eos_threshold = -4.0f;
                    } break;
+                case PROJECTOR_TYPE_SOPRANO:
+                    {
+                        // SopranoDecoder defaults, n_fft comes from the head tensor
+                        hparams.audio_sample_rate = 32000;
+                        hparams.audio_hop_len     = 512;
+                    } break;
                case PROJECTOR_TYPE_PADDLEOCR:
                    {
                        hparams.n_merge = 2;
@@ -2452,6 +2494,7 @@ struct clip_model_loader {
            model.proj_type != PROJECTOR_TYPE_GEMMA3NV &&
            model.proj_type != PROJECTOR_TYPE_QWEN3TTS_SPKENC &&
            model.proj_type != PROJECTOR_TYPE_POCKETTTS_GEN &&
+            model.proj_type != PROJECTOR_TYPE_SOPRANO &&
            model.proj_type != PROJECTOR_TYPE_MUSE_SPARK_AUDIO);

        // layers
@@ -3391,6 +3434,31 @@ struct clip_model_loader {
                        c2w.dac_post_conv_b      = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV,  "bias"));
                    }
                } break;
+            case PROJECTOR_TYPE_SOPRANO:
+                {
+                    auto & v = model.vocos;
+                    v.input_w       = get_tensor(string_format(TN_A_GEN_WAV_INPUT, "weight"));
+                    v.input_b       = get_tensor(string_format(TN_A_GEN_WAV_INPUT, "bias"));
+                    v.norm_w        = get_tensor(string_format(TN_A_GEN_WAV_NORM, "weight"));
+                    v.norm_b        = get_tensor(string_format(TN_A_GEN_WAV_NORM, "bias"));
+                    v.output_norm_w = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT_NORM, "weight"));
+                    v.output_norm_b = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT_NORM, "bias"));
+                    v.output_w      = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT, "weight"));
+                    v.output_b      = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT, "bias"));
+                    v.blocks.resize(hparams.n_layer);
+                    for (int il = 0; il < hparams.n_layer; ++il) {
+                        auto & b = v.blocks[il];
+                        b.dwconv_w = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "weight"));
+                        b.dwconv_b = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "bias"));
+                        b.norm_w   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM, il, "weight"));
+                        b.norm_b   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM, il, "bias"));
+                        b.pw1_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1, il, "weight"));
+                        b.pw1_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1, il, "bias"));
+                        b.pw2_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2, il, "weight"));
+                        b.pw2_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2, il, "bias"));
+                        b.gamma    = get_tensor(string_format(TN_A_GEN_WAV_UP_GAMMA, il));
+                    }
+                } break;
            case PROJECTOR_TYPE_VOXTRAL:
            case PROJECTOR_TYPE_VOXTRAL_RT:
                {
@@ -5679,6 +5747,10 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
                    set_input_f32("inp_noise", noise);
                }
            } break;
+        case PROJECTOR_TYPE_SOPRANO:
+            {
+                set_input_f32("inp_feats", *params->feats);
+            } break;
        case PROJECTOR_TYPE_GEMMA4V:
        case PROJECTOR_TYPE_GEMMA4UV:
            {
@@ -6329,6 +6401,11 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
        out_audio.resize(ggml_nelements(audio));
        ggml_backend_tensor_get(audio, out_audio.data(), 0, ggml_nbytes(audio));

+        if (model.proj_type == PROJECTOR_TYPE_SOPRANO) {
+            std::vector<float> spectrum = std::move(out_audio);
+            clip_graph_soprano::decode_spectrum(model, spectrum, out_audio);
+        }
+
        // drop the tail audio that comes from the code-0 rear padding
        const int64_t n_codes    = params->codes ? model.gen_code_head_w->ne[2] + 1 : 0;
        const int64_t n_frames_w = hparams.wav_tfm_swa;
@@ -6536,6 +6613,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
            return ctx->model.spk_proj_w->ne[1];
        case PROJECTOR_TYPE_POCKETTTS_GEN:
            return ctx->model.gen_input_lin_w->ne[1];
+        case PROJECTOR_TYPE_SOPRANO:
+            return ctx->model.hparams.projection_dim;
        case PROJECTOR_TYPE_PARAKEET:
            return ctx->model.mm_1_w->ne[1];
        default:
@@ -14,7 +14,7 @@ ggml_cgraph * clip_graph_ling3vl::build() {
    // vision M-RoPE, same layout as qwen3vl: [row, col, row, col] quarters
    int mrope_sections[4] = {d_head/4, d_head/4, d_head/4, d_head/4};

-    ggml_tensor * inp = build_inp_with_temporal_merge();
+    ggml_tensor * inp = build_inp_temporal_merge();

    // spatial merge
    {
@@ -42,7 +42,6 @@ struct clip_graph_deepseek4v : clip_graph {
 struct clip_graph_qwen2vl : clip_graph {
    clip_graph_qwen2vl(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
    ggml_cgraph * build() override;
-    ggml_tensor * build_inp_with_temporal_merge();
 };

 struct clip_graph_qwen3vl : clip_graph_qwen2vl {
@@ -420,6 +419,15 @@ std::vector<c2w_state_slot> list_c2w_state_slots(const clip_hparams & hparams, c
 // same, for the streaming mimi decoder (pocket-tts GEN_WAV)
 std::vector<c2w_state_slot> list_pockettts_state_slots(const clip_hparams & hparams, const clip_model & model);

+struct clip_graph_soprano : clip_graph {
+    clip_graph_soprano(clip_ctx * ctx, const clip_image_f32 & img, int n_frames) : clip_graph(ctx, img), n_frames(n_frames) {}
+    ggml_cgraph * build() override;
+    static void decode_spectrum(const clip_model & model, const std::vector<float> & spectrum, std::vector<float> & pcm);
+    // SopranoDecoder default
+    static constexpr int upscale = 4;
+    int n_frames;
+};
+
 struct clip_graph_kimik25 : clip_graph {
    clip_graph_kimik25(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
    ggml_cgraph * build() override;
@@ -1,34 +1,5 @@
 #include "models.h"

-ggml_tensor * clip_graph_qwen2vl::build_inp_with_temporal_merge() {
-    ggml_tensor * inp_raw = build_inp_raw();
-
-    GGML_ASSERT(img.nx() % (patch_size * 2) == 0);
-    GGML_ASSERT(img.ny() % (patch_size * 2) == 0);
-
-    const size_t nb1 = ggml_row_size(inp_raw->type, img.nx());
-    const size_t nb2 = ggml_row_size(inp_raw->type, img.nx() * img.ny());
-
-    if (n_batch == 1) {
-        // still image input
-        return ggml_add(ctx0,
-            ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_raw, patch_size, patch_size, 0, 0, 1, 1),
-            ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_raw, patch_size, patch_size, 0, 0, 1, 1));
-    } else if (n_batch == 2) {
-        // 2 frames input (video input)
-        ggml_tensor * inp_0 = ggml_view_3d(ctx0, inp_raw,
-                                    img.nx(), img.ny(), 3, nb1, nb2, 0);
-        ggml_tensor * inp_1 = ggml_view_3d(ctx0, inp_raw,
-                                    img.nx(), img.ny(), 3, nb1, nb2,
-                                    nb2 * 3); // move to the second frame
-        return ggml_add(ctx0,
-            ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_0, patch_size, patch_size, 0, 0, 1, 1),
-            ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_1, patch_size, patch_size, 0, 0, 1, 1));
-    } else {
-        GGML_ASSERT(false && "n_batch > 2 is not supported");
-    }
-}
-
 ggml_cgraph * clip_graph_qwen2vl::build() {
    GGML_ASSERT(model.patch_bias == nullptr);
    GGML_ASSERT(model.class_embedding == nullptr);
@@ -45,7 +16,7 @@ ggml_cgraph * clip_graph_qwen2vl::build() {

    int mrope_sections[4] = {d_head/4, d_head/4, d_head/4, d_head/4};

-    ggml_tensor * inp = build_inp_with_temporal_merge();
+    ggml_tensor * inp = build_inp_temporal_merge();

    // second conv dimension
    {
@@ -13,7 +13,7 @@ ggml_cgraph * clip_graph_qwen3vl::build() {

    int mrope_sections[4] = {d_head/4, d_head/4, d_head/4, d_head/4};

-    ggml_tensor * inp = build_inp_with_temporal_merge();
+    ggml_tensor * inp = build_inp_temporal_merge();

    // spatial merge
    {
@@ -0,0 +1,78 @@
+#include "models.h"
+#include "../mtmd-audio.h"
+
+// Soprano Vocos: normalized Qwen3 hidden states -> STFT coefficients.
+ggml_cgraph * clip_graph_soprano::build() {
+    const auto & v = model.vocos;
+    const int n_inp = hparams.projection_dim;
+    const int n_up = upscale * (n_frames - 1) + 1;
+    ggml_tensor * cur = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_inp, n_frames);
+    ggml_set_name(cur, "inp_feats");
+    ggml_set_input(cur);
+
+    cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
+    cur = ggml_interpolate(ctx0, cur, n_up, n_inp, 1, 1, GGML_SCALE_MODE_BILINEAR | GGML_SCALE_FLAG_ALIGN_CORNERS);
+    if (ggml_n_dims(v.input_w) == 3) {
+        const int k = v.input_w->ne[0];
+        cur = ggml_conv_1d(ctx0, v.input_w, cur, 1, k / 2, 1);
+        cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
+    } else {
+        cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
+        cur = build_mm(v.input_w, cur);
+    }
+    cur = ggml_add(ctx0, cur, v.input_b);
+    cur = build_norm(cur, v.norm_w, v.norm_b, NORM_TYPE_NORMAL, eps, -1);
+
+    for (size_t il = 0; il < v.blocks.size(); ++il) {
+        const auto & b = v.blocks[il];
+        const int k = b.dwconv_w->ne[0];
+        ggml_tensor * residual = cur;
+        cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
+        cur = ggml_conv_1d_dw(ctx0, b.dwconv_w, cur, 1, k / 2, 1);
+        cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
+        cur = ggml_add(ctx0, cur, b.dwconv_b);
+        cur = build_norm(cur, b.norm_w, b.norm_b, NORM_TYPE_NORMAL, eps, il);
+        cur = ggml_add(ctx0, build_mm(b.pw1_w, cur), b.pw1_b);
+        cur = ggml_gelu_erf(ctx0, cur);
+        cur = ggml_add(ctx0, build_mm(b.pw2_w, cur), b.pw2_b);
+        cur = ggml_add(ctx0, residual, ggml_mul(ctx0, cur, b.gamma));
+        cb(cur, "vocos_block", il);
+    }
+    cur = build_norm(cur, v.output_norm_w, v.output_norm_b, NORM_TYPE_NORMAL, eps, -1);
+    cur = ggml_add(ctx0, build_mm(v.output_w, cur), v.output_b);
+    cb(cur, "vocos_head", -1);
+
+    // the head holds n_fft/2+1 log-magnitudes followed by n_fft/2+1 phases
+    const int n_bins = v.output_w->ne[1] / 2;
+    ggml_tensor * mag = ggml_view_2d(ctx0, cur, n_bins, n_up, cur->nb[1], 0);
+    ggml_tensor * phase = ggml_view_2d(ctx0, cur, n_bins, n_up, cur->nb[1], n_bins * sizeof(float));
+    mag = ggml_clamp(ctx0, ggml_exp(ctx0, ggml_cont(ctx0, mag)), 0.0f, 100.0f);
+    phase = ggml_cont(ctx0, phase);
+    ggml_tensor * re = ggml_mul(ctx0, mag, ggml_cos(ctx0, phase));
+    ggml_tensor * im = ggml_mul(ctx0, mag, ggml_sin(ctx0, phase));
+    re = ggml_reshape_3d(ctx0, re, 1, n_bins, n_up);
+    im = ggml_reshape_3d(ctx0, im, 1, n_bins, n_up);
+    cur = ggml_concat(ctx0, re, im, 0);
+    ggml_set_name(cur, "out_audio");
+    ggml_set_output(cur);
+    ggml_build_forward_expand(gf, cur);
+    return gf;
+}
+
+void clip_graph_soprano::decode_spectrum(const clip_model & model, const std::vector<float> & spectrum, std::vector<float> & pcm) {
+    const int n_frame_vals = model.vocos.output_w->ne[1];
+    const int n_fft = n_frame_vals - 2;
+    const int hop_length = model.hparams.audio_hop_len;
+    const size_t n_frames = spectrum.size() / n_frame_vals;
+    mtmd_audio_streaming_istft istft(n_fft, hop_length);
+    pcm.clear();
+    for (size_t i = 0; i < n_frames; ++i) {
+        auto frame = istft.process_frame(spectrum.data() + i * n_frame_vals);
+        pcm.insert(pcm.end(), frame.begin(), frame.end());
+    }
+    auto tail = istft.flush();
+    pcm.insert(pcm.end(), tail.begin(), tail.end());
+    // the shared ISTFT removes (n_fft - hop) / 2 samples, torch.istft(center=True) removes n_fft / 2
+    pcm.erase(pcm.begin(), pcm.begin() + hop_length / 2);
+    pcm.resize((n_frames - 1) * hop_length);
+}
@@ -1076,12 +1076,199 @@ private:
    bool   wav_header_sent = false;
 };

+// Soprano uses the hidden state that predicts each non-EOS token, including the prompt's last state.
+class soprano_gen_audio_pipeline : public mtmd_gen_audio_pipeline {
+public:
+    using mtmd_gen_audio_pipeline::mtmd_gen_audio_pipeline;
+
+    void reset() override {
+        prompt.clear();
+        features.clear();
+        h_next.clear();
+        audio.clear();
+        wav.clear();
+        pos = 0;
+        prompt_pos = 0;
+        decoded = false;
+        sent = false;
+    }
+
+    int32_t set_input(const mtmd_helper_gen_audio_inp * inp) override {
+        reset();
+        if (inp->speaker_ref || !inp->prompt || inp->prompt_len == 0 || inp->prompt_len > INT32_MAX ||
+            inp->seq_id < 0 || (uint32_t) inp->seq_id >= llama_n_seq_max(lctx)) {
+            LOG_ERR("mtmd_helper_gen_audio: soprano requires text, without a speaker reference\n");
+            return 1;
+        }
+        seq_id = inp->seq_id;
+        out_type = inp->out_type;
+        stream = inp->stream;
+        const llama_token text = find_special_token(vocab, "[TEXT]");
+        const llama_token start = find_special_token(vocab, "[START]");
+        const llama_token stop = find_special_token(vocab, "[STOP]");
+        if (text == LLAMA_TOKEN_NULL || start == LLAMA_TOKEN_NULL || stop == LLAMA_TOKEN_NULL) {
+            LOG_ERR("mtmd_helper_gen_audio: incompatible soprano vocabulary\n");
+            return 1;
+        }
+        const int required = llama_tokenize(vocab, inp->prompt, inp->prompt_len, nullptr, 0, false, false);
+        if (required >= 0 || (uint32_t) (-required + 3) >= llama_n_ctx_seq(lctx)) {
+            LOG_ERR("mtmd_helper_gen_audio: soprano prompt does not fit in the context\n");
+            return 1;
+        }
+        const int n = -required;
+        prompt.resize(n + 3);
+        prompt[0] = stop;
+        prompt[1] = text;
+        if (llama_tokenize(vocab, inp->prompt, inp->prompt_len, prompt.data() + 2, n, false, false) != n) {
+            return 1;
+        }
+        prompt.back() = start;
+        if (!llama_memory_seq_rm(llama_get_memory(lctx), seq_id, 0, -1)) {
+            LOG_ERR("mtmd_helper_gen_audio: cannot reset soprano sequence\n");
+            return 1;
+        }
+        if (!init_gen_state(inp->seed)) {
+            return 1;
+        }
+        return 0;
+    }
+
+    int32_t step_prompt(int32_t n_batch) override {
+        if (n_batch <= 0 || prompt.empty()) {
+            return -1;
+        }
+        const int n = std::min(n_batch, (int) prompt.size() - prompt_pos);
+        if (n == 0) {
+            return 0;
+        }
+        if (!decode(prompt.data() + prompt_pos, n)) {
+            return -1;
+        }
+        prompt_pos += n;
+        return (int) prompt.size() - prompt_pos;
+    }
+
+    int32_t step_gen(llama_token sampled, const float * h_state_in, const float ** h_state_out, bool * out_stop) override {
+        *h_state_out = nullptr;
+        *out_stop = false;
+        if (sampled < 0 || sampled >= llama_vocab_n_tokens(vocab) || !h_state_in || prompt_pos != (int) prompt.size() || prompt.empty()) {
+            return 1;
+        }
+        if (llama_vocab_is_eog(vocab, sampled) || (uint32_t) pos >= llama_n_ctx_seq(lctx)) {
+            *out_stop = true;
+            return 0;
+        }
+        features.insert(features.end(), h_state_in, h_state_in + n_embd);
+        if (!decode(&sampled, 1)) {
+            features.resize(features.size() - n_embd);
+            return 1;
+        }
+        const float * h = llama_get_embeddings_ith(lctx, -1);
+        if (!h) {
+            return 1;
+        }
+        h_next.assign(h, h + n_embd);
+        *h_state_out = h_next.data();
+        return 0;
+    }
+
+    int32_t get_output(int32_t * sample_rate, const char ** data, size_t * data_len, int64_t * n_samples) override {
+        *sample_rate = info.sample_rate;
+        *data_len = 0;
+        if (n_samples) { *n_samples = 0; }
+        if (stream) {
+            // the decoder needs the whole utterance, so the audio comes out in one chunk after flush()
+            if (!decoded || sent) {
+                return 0;
+            }
+            sent = true;
+        } else if (!decode_audio()) {
+            return 1;
+        }
+        if (n_samples) { *n_samples = audio.size(); }
+        if (out_type == MTMD_HELPER_GEN_AUDIO_OUTTYPE_PCM) {
+            *data = (const char *) audio.data();
+            *data_len = audio.size() * sizeof(float);
+        } else {
+            wav.clear();
+            if (!write_wav16(wav, audio, info.sample_rate)) {
+                return 1;
+            }
+            *data = wav.data();
+            *data_len = wav.size();
+        }
+        return 0;
+    }
+
+    // the whole utterance is decoded at once
+    int32_t flush() override {
+        return decode_audio() ? 0 : 1;
+    }
+
+private:
+    bool decode_audio() {
+        if (decoded) {
+            return true;
+        }
+        if (features.size() < 2 * (size_t) n_embd) {
+            LOG_ERR("mtmd_helper_gen_audio: soprano needs at least two generated frames\n");
+            return false;
+        }
+        mtmd_gen_inp inp = mtmd_gen_inp_default(mctx);
+        inp.type = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
+        inp.feats = features.data();
+        inp.n_feats = features.size();
+        mtmd_gen_out out{};
+        if (mtmd_gen_audio_process(mctx, gen_state.get(), &inp, &out) != 0) {
+            return false;
+        }
+        audio.assign(out.audio, out.audio + out.n_samples);
+        decoded = true;
+        return true;
+    }
+
+    bool decode(llama_token * tokens, int n) {
+        llama_batch batch = llama_batch_init(n, 0, 1);
+        batch.n_tokens = n;
+        for (int i = 0; i < n; ++i) {
+            batch.token[i] = tokens[i];
+            batch.pos[i] = pos + i;
+            batch.n_seq_id[i] = 1;
+            batch.seq_id[i][0] = seq_id;
+            batch.logits[i] = i == n - 1;
+        }
+        const int ret = llama_decode(lctx, batch);
+        llama_batch_free(batch);
+        if (ret != 0) {
+            LOG_ERR("mtmd_helper_gen_audio: soprano backbone decode failed\n");
+            return false;
+        }
+        pos += n;
+        return true;
+    }
+
+    llama_seq_id seq_id = 0;
+    llama_pos pos = 0;
+    int prompt_pos = 0;
+    std::vector<llama_token> prompt;
+    std::vector<float> features;
+    std::vector<float> h_next;
+    std::vector<float> audio;
+    std::vector<char> wav;
+    mtmd_helper_gen_audio_outtype out_type = MTMD_HELPER_GEN_AUDIO_OUTTYPE_WAV;
+    bool stream  = false;
+    bool decoded = false;
+    bool sent    = false;
+};
+
 static std::unique_ptr<mtmd_gen_audio_pipeline> make_pipeline(llama_context * lctx, mtmd_context * mctx) {
    switch (mtmd_gen_audio_get_info(mctx).type) {
        case MTMD_GEN_AUDIO_TYPE_QWEN3TTS:
            return std::unique_ptr<mtmd_gen_audio_pipeline>(new qwen3tts_gen_audio_pipeline(lctx, mctx));
        case MTMD_GEN_AUDIO_TYPE_POCKETTTS:
            return std::unique_ptr<mtmd_gen_audio_pipeline>(new pockettts_gen_audio_pipeline(lctx, mctx));
+        case MTMD_GEN_AUDIO_TYPE_SOPRANO:
+            return std::unique_ptr<mtmd_gen_audio_pipeline>(new soprano_gen_audio_pipeline(lctx, mctx));
        default:
            return nullptr;
    }
@@ -205,6 +205,8 @@ private:
    // 1. Separable filtering: horizontal pass followed by vertical pass
    // 2. Pre-computes normalized filter coefficients for each output pixel
    // 3. Fixed-point integer arithmetic (22 fractional bits) for speed and determinism
+    //
+    // RESIZE_ALGO_FLAG_TORCH is the same, except for the int16 weights of torchvision (aten/src/ATen/native/cpu/UpSampleKernel.cpp)
    static bool resize_pillow(
            const clip_image_u8 & img,
            clip_image_u8 & dst,
@@ -215,6 +217,13 @@ private:
        // This allows encoding fractional weights as integers: weight * 2^22
        const int PRECISION_BITS = 32 - 8 - 2;

+        const bool int16_weights = algo & RESIZE_ALGO_FLAG_TORCH;
+        algo = static_cast<resize_algo>(algo & ~RESIZE_ALGO_FLAG_TORCH);
+
+        if (int16_weights && algo == RESIZE_ALGO_LANCZOS) {
+            GGML_ABORT("unsupported resize algo: torchvision has no lanczos");
+        }
+
        // Filter support radius
        double filter_support;
        switch (algo) {
@@ -222,11 +231,11 @@ private:
            case RESIZE_ALGO_BICUBIC:  filter_support = 2.0; break;
            case RESIZE_ALGO_LANCZOS:  filter_support = 3.0; break;
            default:
-                throw std::runtime_error("Unsupported resize algorithm");
+                GGML_ABORT("unsupported resize algo: %d", (int) algo);
        }

        // Returns filter weight for distance x from pixel center
-        // Note: for bicubic, Pillow uses a = -0.5 while GGML/PyTorch use a = -0.75
+        // Note: for bicubic, Pillow (and PyTorch with antialias) uses a = -0.5 while GGML (and PyTorch without antialias) uses a = -0.75
        auto resample_filter = [algo](double x) -> double {
            if (algo == RESIZE_ALGO_LANCZOS) {
                if (-3.0 <= x && x < 3.0) {
@@ -351,12 +360,24 @@ private:
            // Formula: int32 = round(float * 2^PRECISION_BITS)
            weights.resize((size_t) outSize * ksize);

-            const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS
+            // torchvision rounds the weights to int16, with as many fractional bits as the largest weight allows;
+            // scaled back to PRECISION_BITS, they give the same output in the passes below
+            int precision = PRECISION_BITS;
+            if (int16_weights) {
+                const double weight_max = *std::max_element(pre_weights.begin(), pre_weights.end());
+                for (precision = 0; precision < PRECISION_BITS; precision++) {
+                    if (static_cast<int>(0.5 + weight_max * (1 << (precision + 1))) >= (1 << 15)) {
+                        break;
+                    }
+                }
+            }
+
+            const double fxp_scale = std::ldexp(1.0, precision); // 1.0 * 2^precision

            for (size_t i = 0; i < (size_t) outSize * ksize; i++) {
                // Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
                const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
-                weights[i] = static_cast<int32_t>(rounded);
+                weights[i] = static_cast<int32_t>(rounded) * (1 << (PRECISION_BITS - precision));
            }

            return ksize;
@@ -31,6 +31,7 @@ struct mtmd_image_preprocessor {
    const clip_hparams & hparams;

    mtmd_image_preprocessor(const clip_ctx * ctx): hparams(*clip_get_hparams(ctx)) {}
+    mtmd_image_preprocessor(const clip_hparams & hparams): hparams(hparams) {}

    virtual ~mtmd_image_preprocessor() = default;
    virtual mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const = 0;
@@ -59,6 +60,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_preprocessor_llava_uhd(const clip_hparams & hparams) : mtmd_image_preprocessor(hparams) {}
    mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;

    struct slice_coordinates {
@@ -624,7 +624,9 @@ struct mtmd_context {
        ctx_v = res.ctx_v;
        ctx_a = res.ctx_a;
        ctx_gen_a = res.ctx_gen_a;
-        if (!ctx_v && !ctx_a) {
+        // soprano has no speaker encoder, its mmproj only holds the decoder
+        const bool gen_only = ctx_gen_a && clip_get_projector_type(ctx_gen_a) == PROJECTOR_TYPE_SOPRANO;
+        if (!ctx_v && !ctx_a && !gen_only) {
            throw std::runtime_error(string_format("Failed to load CLIP model from %s\n", mmproj_fname));
        }

@@ -641,7 +643,8 @@ struct mtmd_context {

        // since we already validate n_embd of vision and audio mmproj,
        // we can safely assume that they are the same
-        int n_embd_clip = clip_n_mmproj_embd(ctx_v ? ctx_v : ctx_a);
+        // a gen-only mmproj is validated below
+        int n_embd_clip = ctx_v || ctx_a ? clip_n_mmproj_embd(ctx_v ? ctx_v : ctx_a) : n_embd_text;
        if (n_embd_text > 0 && n_embd_text != n_embd_clip) {
            throw std::runtime_error(string_format(
                "mismatch between text model (n_embd = %d) and mmproj (n_embd = %d)\n"
@@ -2082,6 +2085,10 @@ mtmd_gen_audio_info mtmd_gen_audio_get_info(const mtmd_context * ctx) {
            info.type = MTMD_GEN_AUDIO_TYPE_POCKETTTS;
            info.sample_rate = 24000;
            break;
+        case PROJECTOR_TYPE_SOPRANO:
+            info.type = MTMD_GEN_AUDIO_TYPE_SOPRANO;
+            info.sample_rate = 32000;
+            break;
        default:
            info.type = MTMD_GEN_AUDIO_TYPE_NONE;
            break;
@@ -2109,6 +2116,9 @@ mtmd_gen_inp mtmd_gen_inp_default(const mtmd_context * ctx) {
            inp.top_p = 1.0f;
            inp.temp  = 0.7f;
            break;
+        case PROJECTOR_TYPE_SOPRANO:
+            inp.type = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
+            break;
        default:
            break;
    }
@@ -2170,6 +2180,14 @@ static int32_t mtmd_gen_audio_process_impl(mtmd_context * ctx, mtmd_gen_state *

    *out = {};

+    if (clip_get_projector_type(ctx_clip) == PROJECTOR_TYPE_SOPRANO) {
+        const size_t n_embd = (size_t) clip_n_mmproj_embd(ctx_clip);
+        if (inp->type != MTMD_GEN_PROCESS_TYPE_GEN_WAV || !inp->feats || inp->n_feats < 2 * n_embd || inp->n_feats % n_embd != 0 || inp->codes) {
+            LOG_ERR("%s: soprano requires at least 2 frames of features, without codes\n", __func__);
+            return 1;
+        }
+    }
+
    if (inp->type == MTMD_GEN_PROCESS_TYPE_GEN_CODE) {
        const size_t n_embd = (size_t) clip_n_mmproj_embd(ctx_clip);

@@ -3046,6 +3064,9 @@ mtmd_memory_usage mtmd_get_memory_usage(const char * mmproj_fname,
        if (ctx->ctx_a) {
            merge(ctx->ctx_a);
        }
+        if (ctx->ctx_gen_a) {
+            merge(ctx->ctx_gen_a);
+        }
        mtmd_memory_usage res;
        res.backend_mem_usage = std::move(total_mem);
        res.image_max_tokens  = ctx->ctx_v ? clip_get_image_max_tokens(ctx->ctx_v) : -1;
@@ -427,6 +427,7 @@ enum mtmd_gen_audio_type {
    MTMD_GEN_AUDIO_TYPE_NONE, // not supported
    MTMD_GEN_AUDIO_TYPE_QWEN3TTS,
    MTMD_GEN_AUDIO_TYPE_POCKETTTS,
+    MTMD_GEN_AUDIO_TYPE_SOPRANO,
 };

 struct mtmd_gen_audio_info {
@@ -59,3 +59,28 @@ The [upstream repository](https://huggingface.co/kyutai/pocket-tts) holds one co
 python convert_hf_to_gguf.py path/to/pocket-tts/languages/english --outfile pocket-tts.gguf
 python convert_hf_to_gguf.py path/to/pocket-tts/languages/english --mmproj --outfile mmproj-pocket-tts.gguf
 ```
+
+## Soprano
+
+Available params:
+- `--tts-speaker-file` is not supported; the model uses a fixed voice
+- Note: `lang` is not used. Supply normalized English text
+
+Example usage:
+
+```sh
+llama-tts -m soprano.gguf \
+    -mm mmproj-soprano.gguf \
+    -p "Hello world" \
+    --temp 0 \
+    --output out.wav
+```
+
+**Note for GGUF conversion:**
+
+Both [Soprano-80M](https://huggingface.co/ekwek/Soprano-80M) and [Soprano-1.1-80M](https://huggingface.co/ekwek/Soprano-1.1-80M) are supported. Convert both files from the same model directory, including `decoder.pth`:
+
+```sh
+python convert_hf_to_gguf.py path/to/Soprano-1.1-80M --outtype f32 --outfile soprano.gguf
+python convert_hf_to_gguf.py path/to/Soprano-1.1-80M --mmproj --outtype f16 --outfile mmproj-soprano.gguf
+```
Author	SHA1	Message	Date
Xuan-Son Nguyenandyomaytk	701fd25b3a	ggml : allow MUL_MAT_ID -1 index to skip the computation (#26631 ) * ggml : allow MUL_MAT_ID -1 index to skip the computation * nits * update vulkan code * ggml-webgpu: ensure all threads reach workgroup barriers in mul_mat_id_vec (cherry picked from commit d8a4b9efb6e2b9abf3a866e847756b7e6bbc6b69) * fix cuda * openvino * exclude CUDA from this PR * add TODO tag --------- Co-authored-by: yomaytk <yoshimura.masashi.frbs@gmail.com>	2026-10-24 12:15:36 +02:00
Daniel Bevenius	3642238cc1	tests : add speculative rejection draft sampling test (#30213 ) * tests : add speculative rejection draft sampling test This commit adds a test for speculative rejection sampling test that excercises the speculative rejection sampling mechanism in common/speculative.cpp. The motivation for this is have a test that verifies that the speculative draft models emitts tokens that match the target models distribution. So when rejection sampling occurs it should still produce the same probability distribution as if we directly sampled from the target model. Running the test: ```console $ ctest --test-dir build -R "^test-speculative-rejection-sampling$" \ --n_trials 5000 -V ``` This will generate two models, one target and one draft model. For subsequent runs use the "fixture exclude any" option to avoid recreating the models: ```console $ ctest --test-dir build-cpu-debug -R \ "^test-speculative-rejection-sampling$" \ -FA "^speculative-models$" --n_trails 5000 -V ``` * remove draft model	2026-10-24 12:47:49 +03:00
Georgi Gerganov	6b9ba4f2e8	ci : fix data race between runners sharing the model cache (#31046 ) * ci : fix data race between runners sharing the model cache Two runners per mac mini share HOME, so concurrent jobs corrupt each other downloads, converted models and venv under the shared MNT cache. - gg_wget: download in a temp dir, publish with mv, so readers only ever see complete files (-N cannot be combined with -O) - convert, quantize and wikitext unzip now produce under a per-runner scratch dir in OUT, published to MNT via hardlink when missing - gguf-split and quantize tests get a per-runner TMP_DIR - python venv moved from MNT to OUT - metal, vulkan and webgpu mac jobs pass a per-runner OUT path Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine) * ci : clean up scratch models after tests Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine) * ci : fix q4_0 publish when a cached model exists cp -n skips an existing file with exit code 1, so the ln fallback aborted the qwen3 job under set -e. Publish through a unique temp file with an atomic mv: readers never see a partial file and a stale cached q4_0 gets replaced. Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine) * ci : cache-first gg_wget to avoid re-downloads HF sends no Last-Modified, so wget -N re-downloads the full files on every run. Skip the download when a non-empty cache file exists: the HF CI artifacts are static. Downloads still go through a temp dir and an atomic mv, so concurrent runners never read partial files. Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)	2026-10-24 09:21:24 +03:00
Ruben Ortlam	8e5a58df87	vulkan: ggml src1 precision support (#30656 ) * vulkan: update accumulator precision selection * handle fp32 overflows through scaling * handle overflows in copy-to-fp16 activations * add missing cases * simplify code * add tests * add ggml prec get helpers, use in Vulkan backend * clean up comments * revert comment edit * trim tests * also handle fp32 src0, fix test scaling src0 values too * don't scale if disabled * remove fp16 scaling * add check to mul_mat * fix inconsistency * rename to acc_prec	2026-10-24 06:05:05 +02:00
Arseniy Obolenskiy	1ddffe9f70	HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable (#30915 ) * HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable amdgcnspirv rejects target-specific builtins unless they are guarded by `__builtin_amdgcn_is_invocable`, so 66 of 145 HIP files failed to build Wrap the perm calls in a helper that keeps v_perm_b32 when the target has it and falls back to a portable byte select otherwise * move perm helper to vendors/hip.h	2026-10-24 12:04:05 +08:00
Xuan-Son Nguyen	7894e8b924	mtmd: move build_inp_temporal_merge to base class (#31065 )	2026-10-24 03:38:25 +02:00
Yuri KhrustalevandXuan Son Nguyen	edc353842c	mtmd : lfm2 resize matches torchvision (#30943 ) * mtmd : lfm2 bilinear resize matches torchvision Lfm2VlImageProcessor and d1-omni vision.py resize uint8 images with torchvision (antialias=True). That is Pillow's algorithm, except that the weights are int16 with as many fractional bits as the largest weight allows, where Pillow uses 22 bits. The Pillow port was 1-2 levels off on about 1% of the values of the d1-omni and LFM2.5-VL crops. RESIZE_ALGO_BILINEAR_TORCH rounds the weights of the Pillow port to that precision and scales them back to 22 bits, which gives the output of torchvision with the same passes, and "bilinear" maps to it for LFM2 and d1-omni. Other models keep the Pillow algorithms. * mtmd : lfm2 bicubic resize matches torchvision Lfm2VlImageProcessor resizes with torchvision for the bicubic checkpoints too (d1-3B, LFM2-VL), so "bicubic" maps to RESIZE_ALGO_BICUBIC_TORCH for LFM2: the Pillow bicubic filter (a = -0.5, as torchvision with antialias) with the int16 weights of torchvision. The Pillow port was 1-2 levels off on about 0.25% of the values of the d1-3B crops. * use bitfield, add hparams setter --------- Co-authored-by: Xuan Son Nguyen <son@huggingface.co>	2026-10-24 02:21:06 +02:00
Hans	9ae73e68f5	mtmd : add Soprano TTS support (#28607 ) Add Soprano-80M and Soprano-1.1-80M (https://huggingface.co/ekwek/Soprano-1.1-80M) as an audio generation model: a Qwen3 backbone plus a Vocos decoder that runs as a gen-only mmproj. - convert: detect Soprano checkpoints and export the backbone as qwen3 and the decoder.pth Vocos weights as the mmproj - vocab: add the soprano pre-tokenizer (lowercase, whitespace collapse, per-digit split, [UNK] fallback) - mtmd: add the Vocos graph and ISTFT, allow loading a gen-only mmproj for Soprano and include the gen-audio context in mtmd_get_memory_usage - mtmd-helper: add the Soprano pipeline; the decoder runs over the whole utterance, so stream mode sends the audio in one chunk after flush() - tts: document Soprano usage Assisted-by: Codex Assisted-by: Claude Opus 5.5	2026-10-24 01:09:50 +02:00
kurquharandMax Krasnyansky	0ae0e19766	hexagon: add SGN op (#30346 ) * hexagon : add SGN support Assisted-by: OpenCode * hexagon : optimize SGN with sign bit Assisted-by: OpenCode * hex-sgn: add f16 * hex-unary: update SGN to use generic unary_colapse and relax colapse checks to apply to more conditions * hex-unary: improve dims colapsing and run it in the host-side solver --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>	2026-10-23 14:58:44 -07:00
Adrien Gallouët	21ec307146	common : allow concurrent downloads across processes (#29312 ) Signed-off-by: Adrien Gallouët <angt@huggingface.co>	2026-10-23 23:52:56 +02:00
WalterCool	9145ddb955	hip: fix bf16 WMMA operand vector type for gfx11/gfx12 (#30885 ) The __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 (and _gfx12) builtins take their A/B operands as vectors of short (_ExtVector<16, short> / _ExtVector<8, short>), not __bf16. Declaring the fragments as ext_vector_type(N) __bf16 fails to compile: error: cannot initialize a parameter of type 'short __attribute__((ext_vector_type(16)))' with an lvalue of type 'const bf16x16_t' (vector of 16 '__bf16' values) short and __bf16 are both 16-bit and the value is only ever passed to the builtin via reinterpret_cast over raw bytes, so changing the element type is bit-identical. The operand type has been short since the builtin was introduced, so this works across LLVM/ROCm versions without a version gate. Fixes HIP builds for RDNA3 (gfx1150/1151) and RDNA4 (gfx1200/1201). Signed-off-by: Pablo Cholaky <waltercool@slash.cl>	2026-10-23 23:31:28 +02:00
thelittlefireman	1a24df1639	ggm-cuda: Enforce DP4A MMQ tile height multiple of warp size (#30837 ) Added static assertion to enforce DP4A MMQ tile height requirement.	2026-10-23 21:37:20 +02:00