server: refactor server entrypoint (#31113)

* server_role

* cont

.github/workflows/build-apple.yml

@@ -105,7 +105,7 @@ jobs:
           cd build
           # Metal Paravirtual devices are difficult to support -> disable
           # ref: https://github.com/ggml-org/llama.cpp/pull/19802#issuecomment-4013704023
-          ctest -L main -E "test-llama-archs|test-llama-context" --verbose --timeout 900
+          ctest -L main -E "test-llama-archs|test-llama-context|test-speculative-rejection-sampling" --verbose --timeout 900
 
   macos-latest-x64:
     runs-on: macos-15-intel

.github/workflows/ci-self-hosted-cpu.yml

@@ -56,7 +56,7 @@ jobs:
       - name: Test
         id: ggml-ci
         run: |
-          LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+          LLAMA_ARG_THREADS=$(nproc) GG_BUILD_HIGH_PERF=1 GG_BUILD_EXTRA_TESTS_0=1 bash ./ci/run.sh ~/results/llama.cpp/$RUNNER_NAME ~/mnt/llama.cpp
 
   cpu-arm64-high-perf-graviton4:
     runs-on: ah-ubuntu_24_04-c8g_8x

.github/workflows/ci-self-hosted-metal.yml

@@ -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

.github/workflows/ci-self-hosted-vulkan.yml

@@ -79,7 +79,7 @@ jobs:
         id: ggml-ci
         run: |
           vulkaninfo --summary
-          GG_BUILD_VULKAN=1 GGML_VK_DISABLE_COOPMAT2=1 bash ./ci/run.sh ~/results/llama.cpp ~/mnt/llama.cpp
+          GG_BUILD_VULKAN=1 GGML_VK_DISABLE_COOPMAT2=1 bash ./ci/run.sh ~/results/llama.cpp/$RUNNER_NAME ~/mnt/llama.cpp
 
       # - name: ccache-buckets-save
       #   if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
@@ -124,7 +124,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
 
       # - name: ccache-buckets-save
       #   if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
@@ -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]

.github/workflows/ci-self-hosted-webgpu.yml

@@ -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

ci/run.sh

@@ -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

ggml/include/ggml.h

@@ -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`

ggml/src/ggml-backend-meta.cpp

@@ -774,13 +774,6 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
         return src_ss[0];
     };
 
-    auto handle_pool_1d = [&](const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
-        if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0) {
-            return {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
-        }
-        return src_ss[0];
-    };
-
     auto handle_flash_attn_ext = [&](const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
         GGML_ASSERT(tensor->src[3] == nullptr || src_ss[3].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
 
@@ -997,14 +990,12 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
             case GGML_OP_CONV_3D:
             case GGML_OP_CONV_2D_DW:
             case GGML_OP_CONV_TRANSPOSE_2D:
+            case GGML_OP_POOL_1D:
             case GGML_OP_POOL_2D:
             case GGML_OP_POOL_2D_BACK:
             case GGML_OP_UPSCALE: {
                 split_state = handle_generic(src_ss, /*scalar_only =*/ true);
             } break;
-            case GGML_OP_POOL_1D: {
-                split_state = handle_pool_1d(src_ss);
-            } break;
             case GGML_OP_PAD: {
                 split_state = handle_pad(src_ss);
             } break;

ggml/src/ggml-cpu/ggml-cpu.c

@@ -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;

ggml/src/ggml-cpu/ops.cpp

@@ -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));
     }
 }
 

ggml/src/ggml-cpu/repack.cpp

@@ -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;

ggml/src/ggml-cpu/spacemit/ime.cpp

@@ -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;

ggml/src/ggml-cuda/mmq-load-tiles.cuh

@@ -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);

ggml/src/ggml-cuda/vecdotq.cuh

@@ -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);

ggml/src/ggml-cuda/vendors/hip.h

@@ -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)
+}

ggml/src/ggml-hexagon/htp/matmul-ops.c

@@ -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;

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

@@ -1225,40 +1225,6 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
     return res;
 }
 
-ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax_part(ggml_metal_library_t lib) {
-    char base[256];
-    char name[256];
-
-    snprintf(base, 256, "kernel_mul_mm_id_amax_part_f32");
-    snprintf(name, 256, "%s", base);
-
-    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
-    if (!res.pipeline) {
-        res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr);
-    }
-
-    res.smem = 32*sizeof(float);
-
-    return res;
-}
-
-ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax(ggml_metal_library_t lib) {
-    char base[256];
-    char name[256];
-
-    snprintf(base, 256, "kernel_mul_mm_id_amax_f32");
-    snprintf(name, 256, "%s", base);
-
-    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
-    if (!res.pipeline) {
-        res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr);
-    }
-
-    res.smem = 32*sizeof(float);
-
-    return res;
-}
-
 ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_map0(ggml_metal_library_t lib, int ne02, int ne20) {
     char base[256];
     char name[256];
@@ -1283,6 +1249,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];
@@ -1292,18 +1267,14 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id(ggml_m
 
     const bool bc_inp = op->src[0]->ne[0] % 32 != 0;
 
-    // src1 prec [TAG_GGML_PREC]
-    const bool amax = ggml_get_op_params_i32(op, 3) == GGML_PREC_F32;
-
     snprintf(base, 256, "kernel_mul_mm_id_%s_%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1));
-    snprintf(name, 256, "%s_bci=%d_amax=%d", base, bc_inp, amax);
+    snprintf(name, 256, "%s_bci=%d", base, bc_inp);
 
     ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
     if (!res.pipeline) {
         ggml_metal_cv_t cv = ggml_metal_cv_init();
 
         ggml_metal_cv_set_bool(cv, bc_inp, FC_MUL_MM + 0);
-        ggml_metal_cv_set_bool(cv, amax,   FC_MUL_MM + 6);
 
         res = ggml_metal_library_compile_pipeline(lib, base, name, cv);
 

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

@@ -140,8 +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_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_zero    (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);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id         (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argmax            (ggml_metal_library_t lib, const struct ggml_tensor * op);

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

@@ -14,8 +14,6 @@
 #define N_MM_SIMD_GROUP_X 2
 #define N_MM_SIMD_GROUP_Y 2
 
-#define N_MM_NPART_AMAX 256
-
 #define N_ARGMAX_NPART_MAX 256
 #define N_ARGMAX_CHUNK 4096
 
@@ -573,12 +571,11 @@ typedef struct {
 } ggml_metal_kargs_mul_mm_id_map0;
 
 typedef struct {
-    int32_t  ne00;
-    int32_t  ne01;
-    int32_t  ne02;
-    uint64_t nb01;
-    uint64_t nb02;
-} ggml_metal_kargs_mul_mm_id_amax;
+    int32_t  ne0;
+    uint64_t nb1;
+    uint64_t nb2;
+    uint64_t nb21;
+} ggml_metal_kargs_mul_mm_id_zero;
 
 typedef struct {
     int32_t  ne00;

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

@@ -2812,15 +2812,6 @@ size_t ggml_metal_op_mul_mat_id_extra_ids(const ggml_tensor * op) {
     return ggml_type_size(GGML_TYPE_I32)*(ne20*ne21 + ne02);
 }
 
-size_t ggml_metal_op_mul_mat_id_extra_amax(const ggml_tensor * op) {
-    assert(op->op == GGML_OP_MUL_MAT_ID);
-
-    GGML_UNUSED(op);
-
-    // 2 scaling factors (8 bytes) + N_MM_NPART_AMAX per-threadgroup scales for stage-1
-    return 8 + N_MM_NPART_AMAX*sizeof(float);
-}
-
 int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
     ggml_tensor * op = ctx->node(idx);
 
@@ -2872,39 +2863,6 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
         ggml_metal_buffer_id bid_ids = bid_tpe;
         bid_ids.offs += ggml_metal_op_mul_mat_id_extra_tpe(op);
 
-        ggml_metal_buffer_id bid_amax = bid_ids;
-        bid_amax.offs += ggml_metal_op_mul_mat_id_extra_ids(op);
-
-        // src1 prec [TAG_GGML_PREC]
-        const bool use_amax = ggml_get_op_params_i32(op, 3) == GGML_PREC_F32;
-
-        // src1 rescale factors, computed before the matmul
-        // ref: https://github.com/ggml-org/llama.cpp/pull/26223
-        if (use_amax) {
-            ggml_metal_kargs_mul_mm_id_amax args = {
-                /*.ne00 =*/ ne10,
-                /*.ne01 =*/ ne11,
-                /*.ne02 =*/ ne12,
-                /*.nb01 =*/ nb11,
-                /*.nb02 =*/ nb12,
-            };
-
-            auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_amax_part(lib);
-
-            const size_t smem = pipeline.smem;
-
-            GGML_ASSERT(smem <= props_dev->max_theadgroup_memory_size);
-
-            ggml_metal_encoder_set_pipeline(enc, pipeline);
-            ggml_metal_encoder_set_bytes   (enc, &args, sizeof(args), 0);
-            ggml_metal_encoder_set_buffer  (enc, bid_src1, 1);
-            ggml_metal_encoder_set_buffer  (enc, bid_amax, 2);
-
-            ggml_metal_encoder_set_threadgroup_memory_size(enc, smem, 0);
-
-            ggml_metal_encoder_dispatch_threadgroups(enc, N_MM_NPART_AMAX, 1, 1, 256, 1, 1);
-        }
-
         {
             ggml_metal_kargs_mul_mm_id_map0 args = {
                 ne10,
@@ -2935,20 +2893,28 @@ 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);
         }
 
-        ggml_metal_op_concurrency_reset(ctx);
+        // 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,
+            };
 
-        if (use_amax) {
-            auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_amax(lib);
+            auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_zero(lib);
 
             ggml_metal_encoder_set_pipeline(enc, pipeline);
-            ggml_metal_encoder_set_buffer  (enc, bid_amax, 0);
+            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, 1, 1, 1, 32, 1, 1);
-
-            // the next kernel has to wait for the amax data
-            ggml_metal_op_concurrency_reset(ctx);
+            ggml_metal_encoder_dispatch_threadgroups(enc, ne20, ne21, 1, 32, 1, 1);
         }
 
+        // this barrier is always needed because the next kernel has to wait for the id maps to be computed
+        ggml_metal_op_concurrency_reset(ctx);
+
         {
             auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id(lib, op);
 
@@ -2977,7 +2943,6 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
             ggml_metal_encoder_set_buffer  (enc, bid_tpe,  3);
             ggml_metal_encoder_set_buffer  (enc, bid_ids,  4);
             ggml_metal_encoder_set_buffer  (enc, bid_dst,  5);
-            ggml_metal_encoder_set_buffer  (enc, bid_amax, 6);
 
             const size_t smem = pipeline.smem;
 

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

@@ -36,7 +36,6 @@ size_t ggml_metal_op_mul_mat_id_extra_tpe(const struct ggml_tensor * op);
 
 // id map [n_tokens, n_expert]
 size_t ggml_metal_op_mul_mat_id_extra_ids(const struct ggml_tensor * op);
-size_t ggml_metal_op_mul_mat_id_extra_amax(const struct ggml_tensor * op);
 
 // return true if we should use the FA vector kernel for this op
 bool ggml_metal_op_flash_attn_ext_use_vec(const struct ggml_tensor * op);

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

@@ -226,7 +226,6 @@ static size_t ggml_backend_metal_buffer_type_get_alloc_size(ggml_backend_buffer_
             {
                 res += ggml_metal_op_mul_mat_id_extra_tpe(tensor);
                 res += ggml_metal_op_mul_mat_id_extra_ids(tensor);
-                res += ggml_metal_op_mul_mat_id_extra_amax(tensor);
             } break;
         case GGML_OP_FLASH_ATTN_EXT:
             {

ggml/src/ggml-metal/kernels/binbcast.metal

@@ -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) {

ggml/src/ggml-metal/kernels/mul_mm.metal

@@ -7,7 +7,6 @@ constant short FC_mul_mm_ne12  [[function_constant(FC_MUL_MM + 2)]];
 constant short FC_mul_mm_ne13  [[function_constant(FC_MUL_MM + 3)]];
 constant short FC_mul_mm_r2    [[function_constant(FC_MUL_MM + 4)]];
 constant short FC_mul_mm_r3    [[function_constant(FC_MUL_MM + 5)]];
-constant bool FC_mul_mm_id_amax [[function_constant(FC_MUL_MM + 6)]];
 constant int32_t FC_mul_mm_id_map0_ne02 [[function_constant(FC_MUL_MM + 7)]];
 
 // each block_q contains 16*nl weights
@@ -391,8 +390,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,82 +507,27 @@ kernel void kernel_mul_mm_id_map0(
     }
 }
 
-kernel void kernel_mul_mm_id_amax_part_f32(
-        constant ggml_metal_kargs_mul_mm_id_amax & args,
-        device   const char * src1,
-        device         char * dst,
-        threadgroup    char * shmem [[threadgroup(0)]],
-        uint  tgpig[[threadgroup_position_in_grid]],
-        ushort tiitg[[thread_index_in_threadgroup]],
-        ushort tiisg[[thread_index_in_simdgroup]],
-        ushort sgitg[[simdgroup_index_in_threadgroup]],
-        ushort   ntg[[threads_per_threadgroup]]) {
-    const int nrow = args.ne01*args.ne02;
+// 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
 
-    float lmax = 0.0f;
+    const int32_t id = ((device const int32_t *) (src2 + i21*args.nb21))[i20];
 
-    for (int ir = tgpig; ir < nrow; ir += N_MM_NPART_AMAX) {
-        const int i01 = ir % args.ne01;
-        const int i02 = ir / args.ne01;
-
-        device const float * row = (device const float *) (src1 + i02*args.nb02 + i01*args.nb01);
-
-        for (int i00 = tiitg; i00 < args.ne00; i00 += ntg) {
-            lmax = max(lmax, fabs(row[i00]));
-        }
+    if (id != -1) {
+        return;
     }
 
-    float amax = simd_max(lmax);
+    device float * dst_row = (device float *) (dst + i20*args.nb1 + i21*args.nb2);
 
-    threadgroup float * shared_amax = (threadgroup float *) shmem;
-
-    if (ntg > N_SIMDWIDTH) {
-        if (sgitg == 0) {
-            shared_amax[tiisg] = 0.0f;
-        }
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-
-        if (tiisg == 0) {
-            shared_amax[sgitg] = amax;
-        }
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-
-        amax = shared_amax[tiisg];
-        amax = simd_max(amax);
-    }
-
-    if (tiitg == 0) {
-        ((device float *) (dst + 8))[tgpig] = amax;
-    }
-}
-
-kernel void kernel_mul_mm_id_amax_f32(
-        device char * dst,
-        ushort tiitg[[thread_index_in_threadgroup]]) {
-    device const float * part = (device const float *) (dst + 8);
-
-    float amax = 0.0f;
-
-    for (int i = tiitg; i < N_MM_NPART_AMAX; i += N_SIMDWIDTH) {
-        amax = max(amax, part[i]);
-    }
-
-    amax = simd_max(amax);
-
-    if (tiitg == 0) {
-        // leave a comfortable margin below the f16 max of 65504
-        float scale = 1.0f;
-
-        // isfinite: src1 already inf/nan is not ours to fix - keep the
-        // scale at 1.0 instead of turning it into a different failure
-        if (isfinite(amax) && amax > 32768.0f) {
-            scale = exp2(ceil(log2(amax)) - 15.0f);
-        }
-
-        device float * d = (device float *) dst;
-
-        d[0] = 1.0f/scale; // exact: scale is a power of two
-        d[1] = scale;
+    for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
+        dst_row[i0] = 0.0f;
     }
 }
 
@@ -607,7 +553,6 @@ kernel void kernel_mul_mm_id(
         device const char * htpe,
         device const char * hids,
         device       char * dst,
-        device const char * amax,
         threadgroup  char * shmem [[threadgroup(0)]],
         uint3  tgpig[[threadgroup_position_in_grid]],
         ushort tiitg[[thread_index_in_threadgroup]],
@@ -663,10 +608,6 @@ kernel void kernel_mul_mm_id(
     // skip the upper half of the token tile when the expert did not fill it
     constexpr short NR1H = NR1/2;
 
-    // power-of-two rescaling
-    const float s1_inv   = FC_mul_mm_id_amax ? ((device const float *) amax)[0] : 1.0f;
-    const float s1_scale = FC_mul_mm_id_amax ? ((device const float *) amax)[1] : 1.0f;
-
 #ifndef GGML_METAL_HAS_TENSOR
     S0_8x8 ma[4];
     S1_8x8 mb[2];
@@ -772,7 +713,7 @@ kernel void kernel_mul_mm_id(
 
                 const short ib = 4*sx + sy;
 
-                *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) (*((device T1 *) y + i) * (T1) s1_inv) : 0;
+                *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0;
             }
         } else {
             const short sx = (tiitg%NL1);
@@ -785,7 +726,7 @@ kernel void kernel_mul_mm_id(
 
             const short ib = 4*sx + sy;
 
-            *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)((*((device T1_2x4 *) y)) * (T1) s1_inv);
+            *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)(*((device T1_2x4 *) y));
         }
 #else
         // load data and store to threadgroup memory
@@ -833,7 +774,7 @@ kernel void kernel_mul_mm_id(
                 //const short lx = (tiitg/NL1)%8;
                 //const short ly = i;
 
-                *(sb + NK*(8*sy + ly) + 8*sx + lx) = loop_k + iy + i < args.ne00 ? (S1) (*((device T1 *) y + i) * (T1) s1_inv) : 0;
+                *(sb + NK*(8*sy + ly) + 8*sx + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0;
             }
         } else {
             const short sx = (tiitg%NL1);
@@ -844,7 +785,7 @@ kernel void kernel_mul_mm_id(
             //const short lx = (tiitg/NL1)%8;
             //const short ly = i;
 
-            *(threadgroup S1_2x4 *)(sb + NK*(8*sy + ly) + 8*sx) = (S1_2x4)((*((device T1_2x4 *) y)) * (T1) s1_inv);
+            *(threadgroup S1_2x4 *)(sb + NK*(8*sy + ly) + 8*sx) = (S1_2x4)(*((device T1_2x4 *) y));
         }
 #endif
 
@@ -935,12 +876,12 @@ kernel void kernel_mul_mm_id(
 
         int i = tiisg;
         for (; i < nr0/4; i += 32) {
-            *(D4 + i) = *(C4 + i) * s1_scale;
+            *(D4 + i) = *(C4 + i);
         }
 
         i = (4*(nr0/4)) + tiisg;
         for (; i < nr0; i += 32) {
-            *(D + i) = *(C + i) * s1_scale;
+            *(D + i) = *(C + i);
         }
     }
 }

ggml/src/ggml-metal/kernels/mul_mv.metal

@@ -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,

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -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

ggml/src/ggml-opencl/kernels/add_id.cl

@@ -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)) {

ggml/src/ggml-opencl/kernels/gemv_moe_mxfp4_f32_ns.cl

@@ -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;

ggml/src/ggml-opencl/kernels/gemv_moe_q4_0_f32_ns.cl

@@ -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

ggml/src/ggml-opencl/kernels/gemv_moe_q4_1_f32_ns.cl

@@ -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

ggml/src/ggml-opencl/kernels/gemv_moe_q4_k_f32_ns.cl

@@ -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;

ggml/src/ggml-opencl/kernels/gemv_moe_q5_0_f32_ns.cl

@@ -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

ggml/src/ggml-opencl/kernels/gemv_moe_q5_1_f32_ns.cl

@@ -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

ggml/src/ggml-opencl/kernels/gemv_moe_q5_k_f32_ns.cl

@@ -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;

ggml/src/ggml-opencl/kernels/gemv_moe_q6_k_f32_ns.cl

@@ -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

ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl

@@ -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]),

ggml/src/ggml-opencl/kernels/mul_mv_id_mxfp4_f32.cl

@@ -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;

ggml/src/ggml-opencl/kernels/mul_mv_id_mxfp4_f32_flat.cl

@@ -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;
 

ggml/src/ggml-opencl/kernels/mul_mv_id_q4_0_f32_8x_flat.cl

@@ -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;

ggml/src/ggml-opencl/kernels/mul_mv_id_q8_0_f32.cl

@@ -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;

ggml/src/ggml-opencl/kernels/mul_mv_id_q8_0_f32_flat.cl

@@ -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;

ggml/src/ggml-openvino/openvino/op/add_id.cpp

@@ -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);
 

ggml/src/ggml-rpc/transport.cpp

@@ -517,6 +517,11 @@ bool socket_t::impl::rdma_recv(void * data, size_t size) {
 
         int slot = (int)wc.wr_id;
         size_t got = wc.byte_len;
+        if (got > rem) {
+            LOG_ERROR("RDMA recv overrun: got=%zu rem=%zu\n", got, rem);
+            c->post_rx(slot);
+            return false;
+        }
         memcpy(dst, c->rx_slot(slot), got);
 
         if (!c->post_rx(slot)) return false;

ggml/src/ggml-sycl/add-id.cpp

@@ -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;

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -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();

ggml/src/ggml-sycl/mmvq.cpp

@@ -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;

ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h

@@ -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 {

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

@@ -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:

ggml/src/ggml-vulkan/vulkan-shaders/add_id.comp

@@ -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];

ggml/src/ggml-vulkan/vulkan-shaders/count_experts.comp

@@ -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();
 }

ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_base.glsl

@@ -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]);

ggml/src/ggml-webgpu/ggml-webgpu.cpp

@@ -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({

ggml/src/ggml-webgpu/wgsl-shaders/add_id.wgsl

@@ -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];

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_id_gather.wgsl

@@ -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;
+        }
     }
 }

ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_id_vec.wgsl

@@ -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++) {

ggml/src/ggml-zendnn/ggml-zendnn.cpp

@@ -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++;

ggml/src/ggml.c

@@ -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,

src/llama-arch.cpp

@@ -1319,6 +1319,8 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) {
         case LLM_ARCH_MINIMAX_M2:
         case LLM_ARCH_MINIMAX_M3:
         case LLM_ARCH_MISTRAL4:
+        case LLM_ARCH_MELLUM: // [TAG_MOE_AMAX_SCALE]
+        case LLM_ARCH_FLINT:  // [TAG_MOE_AMAX_SCALE]
         case LLM_ARCH_KIMI_LINEAR:
         case LLM_ARCH_BAILINGMOE3:
         case LLM_ARCH_KIMI_K3:

src/llama-graph.cpp

@@ -2510,6 +2510,7 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
     {
         // TODO: this graph-level rescaling should only be applied when the corresponding backend does not support
         //       ggml_prec_set_src(experts, GGML_PREC_BF16, 1);
+        //       the amax op cannot be split across devices, so these architectures are denied -sm tensor [TAG_MOE_AMAX_SCALE]
         //       ref: https://github.com/ggml-org/llama.cpp/pull/30656#issuecomment-6269597922
         llm_graph_amax_src1 amax_src1 = { nullptr, nullptr };
 

src/models/minimax-m3.cpp

@@ -550,9 +550,9 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
         if ((uint32_t) il < hparams.n_layer_dense_lead) {
             // leading dense FFN (swigluoai)
             cur = build_ffn(cur,
-                    model.layers[il].ffn_up,   NULL, NULL,
-                    model.layers[il].ffn_gate, NULL, NULL,
-                    model.layers[il].ffn_down, NULL, NULL,
+                    model.layers[il].ffn_up,   NULL, model.layers[il].ffn_up_s,
+                    model.layers[il].ffn_gate, NULL, model.layers[il].ffn_gate_s,
+                    model.layers[il].ffn_down, NULL, model.layers[il].ffn_down_s,
                     NULL,
                     LLM_FFN_SWIGLU_OAI_MOE, LLM_FFN_PAR, il);
             cb(cur, "ffn_out", il);
@@ -568,14 +568,17 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
                     LLM_FFN_SWIGLU_OAI_MOE, hparams.expert_weights_norm,
                     hparams.expert_weights_scale,
                     (llama_expert_gating_func_type) hparams.expert_gating_func,
-                    il);
+                    il, nullptr, nullptr,
+                    model.layers[il].ffn_up_exps_s,
+                    model.layers[il].ffn_gate_exps_s,
+                    model.layers[il].ffn_down_exps_s);
             cb(moe_out, "ffn_moe_out", il);
 
             // shared expert (swigluoai)
             ggml_tensor * ffn_shexp = build_ffn(cur,
-                    model.layers[il].ffn_up_shexp,   NULL, NULL,
-                    model.layers[il].ffn_gate_shexp, NULL, NULL,
-                    model.layers[il].ffn_down_shexp, NULL, NULL,
+                    model.layers[il].ffn_up_shexp,   NULL, model.layers[il].ffn_up_shexp_s,
+                    model.layers[il].ffn_gate_shexp, NULL, model.layers[il].ffn_gate_shexp_s,
+                    model.layers[il].ffn_down_shexp, NULL, model.layers[il].ffn_down_shexp_s,
                     NULL,
                     LLM_FFN_SWIGLU_OAI_MOE, LLM_FFN_PAR, il);
             cb(ffn_shexp, "ffn_shexp", il);

tests/CMakeLists.txt

@@ -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

tests/test-backend-ops.cpp

@@ -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,23 +5609,21 @@ 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) {
     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;
-        }
-        if (amax != 1.0f && t->type == GGML_TYPE_F32) {
-            init_tensor_uniform(t, -amax, amax);
-        } else {
+        if (t->type != GGML_TYPE_I32) {
             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
@@ -5592,12 +5636,12 @@ struct test_mul_mat_id : public test_case {
     const int64_t m;
     const int64_t n;
     const int64_t k;
-    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_STR11(type_a, type_b, n_mats, n_used, b, m, n, k, m_v, ids_offset, skip);
     }
 
     double max_nmse_err() override {
@@ -5622,9 +5666,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)
+            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), m_v(m_v), ids_offset(ids_offset), skip(skip) {
             GGML_ASSERT(n_used <= n_mats);
             GGML_ASSERT(m_v == 0 || m_v > m);
         }
@@ -5650,20 +5694,15 @@ struct test_mul_mat_id : public test_case {
         ggml_tensor * out = ggml_mul_mat_id(ctx, as, b, ids);
         ggml_set_name(out, "out");
 
-        if (amax > 65504.0f) {
-            // src1 exceeds F16 range
-            ggml_prec_set_src(out, GGML_PREC_F32, 1);
-        }
-
         return out;
     }
 
     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);
     }
 
     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 +5777,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 +5794,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 +5836,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 +7794,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 +7958,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);
         }
     }
 
@@ -11285,8 +11326,8 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 8, 1, 64, {8, 1}, {1, 1}, {0, 1, 2, 3}, 0, 1, false, 16, 16));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 8, 16, 64, {8, 1}, {1, 1}, {0, 1, 2, 3}, 0, 1, false, 16, 16));
     // as is a view whose experts are strided by more rows than it uses
-    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 4, 2, false, 8,  1, 64, 1.0f, 64));
-    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 2, false, 8, 16, 64, 1.0f, 64));
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 4, 2, false, 8,  1, 64, 64));
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 2, false, 8, 16, 64, 64));
 
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 576, 512, 576, {1,1}, {1,1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 1, 2048, 8192, {1,  1}, {1, 1}));
@@ -11400,7 +11441,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     }
 
     for (ggml_type ta : {GGML_TYPE_F16, GGML_TYPE_Q4_0}) {
-        test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 1025, 10, false, 64, 33, 256, 1.0f, 0, true));
+        test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 1025, 10, false, 64, 33, 256, 0, true));
         test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 2050, 10, true, 64, 33, 256));
     }
 #endif
@@ -11415,12 +11456,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         test_cases.emplace_back(new test_mul_mat_id(type_a, GGML_TYPE_F32, 4, 4, false, 16, 10, 256));
     }
 
-    // test src1 f16 overflow
-    for (int n : {16, 32, 64}) {
-        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_K, GGML_TYPE_F32, 128, 4, false, 4096, n, 2048, 1e5f));
-        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8,   2, false, 512,  n, 256,  1e5f));
-    }
-
     for (ggml_type type_a : base_types) {
         for (ggml_type type_b : {GGML_TYPE_F32 /*, GGML_TYPE_F16 */}) {
             for (int n_mats : {4, 8}) {
@@ -11453,6 +11488,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, 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, 0, false, true));
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, 1, 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 +11550,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 +12289,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));
+                                }
                             }
                         }
                     }

tests/test-llama-context.cpp

@@ -60,10 +60,12 @@ static llama_tokens prepare_tokens(llama_model * model, const common_params & pa
     llama_tokens tokens;
 
     if (params.prompt.empty()) {
-        const int n_prompt = params.n_batch;
+        // the prompt must not exceed n_batch and n_ctx of cases
+        constexpr int n_prompt_max = 100;
+        const int n_prompt = std::min(params.n_batch, n_prompt_max);
 
         // this path is useful for model files that do not have a tokenizer
-        TLOG_INF("%s: no prompt provided, generating %d (n_batch) random tokens\n", __func__, n_prompt);
+        TLOG_INF("%s: no prompt provided, generating %d random tokens\n", __func__, n_prompt);
 
         const auto * vocab = llama_model_get_vocab(model);
         const auto   n_vocab = llama_vocab_n_tokens(vocab);
@@ -205,7 +207,6 @@ int main(int argc, char ** argv) {
 
     common_params params;
     params.prompt = "";
-    params.n_batch = 100;
     params.out_file = "dump_state.bin";
     params.sampling.seed = 1234;
 

tests/test-speculative-rejection-sampling.cpp

@@ -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;
+}

tools/mtmd/clip-graph.h

@@ -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);

tools/mtmd/clip.cpp

@@ -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);
 }

tools/mtmd/models/ling3vl.cpp

@@ -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
     {

tools/mtmd/models/models.h

@@ -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 {

tools/mtmd/models/qwen2vl.cpp

@@ -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
     {

tools/mtmd/models/qwen3vl.cpp

@@ -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
     {

tools/server/README-dev.md

@@ -286,7 +286,7 @@ Instead of building everything from the ground up (like what most AI agents will
 
 The flow for downloading a new model:
 - POST request comes in --> `post_router_models` --> validation
-- A new `llama-server` subprocess will be spawned with special `SERVER_CHILD_MODE_DOWNLOAD`
+- A new `llama-server` subprocess will be spawned with the `SERVER_ROLE_DOWNLOAD` role (see `server_role_type` in `server-common.h`)
 - Child process runs the download and report status back to router via stdin/out
 - If a stop request comes in, the router asks the child process to stop (same mechanism as running a model in child process)
 - Otherwise, upon completion, we call `load_models()` to refresh the list of models

tools/server/server-common.cpp

@@ -16,7 +16,9 @@
 #include <limits>
 #include <cstring>
 #include <type_traits>
+#include <atomic>
 #include <chrono>
+#include <iostream>
 #include <thread>
 
 #ifdef _WIN32
@@ -2210,3 +2212,64 @@ void server_subproc::waiter::wait(const std::vector<server_subproc *> & procs, s
     }
 #endif
 }
+
+//
+// server_role
+//
+
+server_role::server_role() {
+    if (is_managed_by_router()) {
+        cmd_out = server_reserve_stdout();
+    }
+}
+
+server_role::~server_role() {
+    if (cmd_out) {
+        fclose(cmd_out);
+    }
+}
+
+bool server_role::is_managed_by_router() {
+    return std::getenv(SERVER_ENV_ROUTER_PORT) != nullptr;
+}
+
+server_role_type server_role::get_role() {
+    const char * role = std::getenv(SERVER_ENV_CHILD_ROLE);
+    return server_role_type_from_str(role ? role : "");
+}
+
+std::thread server_role::setup(const std::function<void(int)> & shutdown_handler) {
+    // setup thread for monitoring stdin
+    return std::thread([shutdown_handler]() {
+        // wait for EOF on stdin
+        SRV_INF("%s", "child server monitoring thread started, waiting for EOF on stdin...\n");
+        bool eof = false;
+        while (true) {
+            std::string line;
+            if (!std::getline(std::cin, line)) {
+                // EOF detected, that means the router server is unexpectedly exit or killed
+                eof = true;
+                break;
+            }
+            if (line.find(CMD_ROUTER_TO_CHILD_EXIT) != std::string::npos) {
+                SRV_INF("%s", "exit command received, exiting...\n");
+                shutdown_handler(0);
+                break;
+            }
+        }
+        if (eof) {
+            SRV_INF("%s", "EOF on stdin detected, forcing shutdown...\n");
+            exit(1);
+        }
+    });
+}
+
+void server_role::notify_to_router(const std::string & state, const json & payload) {
+    json data = {
+        {"state", state},
+        {"payload", payload},
+    };
+    std::lock_guard<std::mutex> lk(mtx_stdout);
+    fprintf(cmd_out, "%s%s\n", CMD_CHILD_TO_ROUTER_STATE, safe_json_to_str(data).c_str());
+    fflush(cmd_out);
+}

tools/server/server-common.h

@@ -18,6 +18,7 @@
 #include <filesystem>
 #include <functional>
 #include <mutex>
+#include <thread>
 #include <queue>
 #include <string>
 #include <vector>
@@ -725,3 +726,97 @@ private:
     intptr_t out_handles[SERVER_SUBPROC_STREAMS] = { -1, -1 }; // fd on POSIX, HANDLE on Windows; taken lazily from sproc
     bool     out_closed [SERVER_SUBPROC_STREAMS] = { false, false };
 };
+
+//
+// router <-> child protocol
+//
+
+// what this llama-server process does. a router spawns the other roles as child processes
+// and talks to them over their stdin/stdout
+enum server_role_type {
+    SERVER_ROLE_ROUTER,   // no model: lists the models, spawns instances on demand and proxies requests to them
+    SERVER_ROLE_INSTANCE, // serves one model, standalone or as a child of a router
+    SERVER_ROLE_DOWNLOAD, // child of a router: download a model, report the progress, exit
+};
+
+static inline const char * server_role_type_to_str(server_role_type role) {
+    switch (role) {
+        case SERVER_ROLE_ROUTER:   return "router";
+        case SERVER_ROLE_INSTANCE: return "instance";
+        case SERVER_ROLE_DOWNLOAD: return "download";
+        default: GGML_ASSERT(false && "invalid server_role_type");
+    }
+}
+
+// the role a router asks of a child; anything it does not name is an instance
+static inline server_role_type server_role_type_from_str(const std::string & str) {
+    if (str == "download") {
+        return SERVER_ROLE_DOWNLOAD;
+    }
+    return SERVER_ROLE_INSTANCE;
+}
+
+// set by the router in the environment of every child
+#define SERVER_ENV_ROUTER_PORT "LLAMA_SERVER_ROUTER_PORT" // its presence makes the process a child
+#define SERVER_ENV_CHILD_ROLE  "LLAMA_SERVER_CHILD_ROLE"  // server_role_type_to_str(), an instance when unset
+
+// lines on the child's stdin (router -> child) and stdout (child -> router),
+// see server_role::setup() and server_role::notify_to_router()
+#define CMD_ROUTER_TO_CHILD_EXIT  "cmd_router_to_child:exit"
+#define CMD_CHILD_TO_ROUTER_STATE "cmd_child_to_router:state:" // followed by json string
+
+// address a child binds to: the router may listen on 0.0.0.0, but reaches its children locally
+// ref: https://github.com/ggml-org/llama.cpp/issues/17862
+#define CHILD_ADDR "127.0.0.1"
+
+// the states a child reports to the router
+enum server_state {
+    SERVER_STATE_DOWNLOADING,
+    SERVER_STATE_LOADING,
+    SERVER_STATE_READY,
+    SERVER_STATE_SLEEPING,
+};
+
+static inline std::string server_state_to_str(server_state state) {
+    switch (state) {
+        case SERVER_STATE_DOWNLOADING: return "downloading";
+        case SERVER_STATE_LOADING:     return "loading";
+        case SERVER_STATE_READY:       return "ready";
+        case SERVER_STATE_SLEEPING:    return "sleeping";
+        default: GGML_ASSERT(false && "invalid server_state");
+    }
+}
+
+static inline server_state server_state_from_str(const std::string & str) {
+    if (str == "downloading") return SERVER_STATE_DOWNLOADING;
+    if (str == "loading")     return SERVER_STATE_LOADING;
+    if (str == "ready")       return SERVER_STATE_READY;
+    if (str == "sleeping")    return SERVER_STATE_SLEEPING;
+    GGML_ASSERT(false && "invalid server_state string");
+}
+
+// the role of this process and, when it is a child of a router, its side of the protocol
+struct server_role {
+    // serializes the notify_to_router writes
+    std::mutex mtx_stdout;
+
+    // in a child, keeps stdout for the commands to the router, so it is created before anything is written;
+    // everything else written to stdout goes to stderr with the logs
+    server_role();
+    ~server_role();
+
+    // true if the current process was spawned by a router
+    bool is_managed_by_router();
+    server_role_type get_role();
+
+    // register the shutdown_handler to be called by the router
+    // return the monitoring thread (to be joined by the caller)
+    std::thread setup(const std::function<void(int)> & shutdown_handler);
+
+    // notify router server for status changes (e.g. loading, downloading, sleeping, etc.)
+    // message will be handled by server_models::handle_child_state() on the router side
+    void notify_to_router(const std::string & state_name, const json & payload);
+
+private:
+    FILE * cmd_out = nullptr; // the stdout the router reads the commands from
+};

tools/server/server-context.h

@@ -58,31 +58,6 @@ struct server_context_meta {
     std::string model_ftype;
 };
 
-enum server_state {
-    SERVER_STATE_DOWNLOADING,
-    SERVER_STATE_LOADING,
-    SERVER_STATE_READY,
-    SERVER_STATE_SLEEPING,
-};
-
-static std::string server_state_to_str(server_state state) {
-    switch (state) {
-        case SERVER_STATE_DOWNLOADING: return "downloading";
-        case SERVER_STATE_LOADING:     return "loading";
-        case SERVER_STATE_READY:       return "ready";
-        case SERVER_STATE_SLEEPING:    return "sleeping";
-        default: GGML_ASSERT(false && "invalid server_state");
-    }
-}
-
-static server_state server_state_from_str(const std::string & str) {
-    if (str == "downloading") return SERVER_STATE_DOWNLOADING;
-    if (str == "loading")     return SERVER_STATE_LOADING;
-    if (str == "ready")       return SERVER_STATE_READY;
-    if (str == "sleeping")    return SERVER_STATE_SLEEPING;
-    GGML_ASSERT(false && "invalid server_state string");
-}
-
 using server_state_callback_t = std::function<void(server_state, json /* payload */)>;
 
 struct server_context {

tools/server/server-models.cpp

@@ -37,9 +37,6 @@ extern char **environ;
 
 #define DEFAULT_STOP_TIMEOUT 10 // seconds
 
-#define CMD_ROUTER_TO_CHILD_EXIT  "cmd_router_to_child:exit"
-#define CMD_CHILD_TO_ROUTER_STATE "cmd_child_to_router:state:" // followed by json string
-
 // note: SIGPIPE is ignored by the server
 static void request_child_exit(server_subproc & proc) {
     FILE * stdin_file = proc.sproc.stdin_file();
@@ -49,10 +46,6 @@ static void request_child_exit(server_subproc & proc) {
     }
 }
 
-// address for child process, this is needed because router may run on 0.0.0.0
-// ref: https://github.com/ggml-org/llama.cpp/issues/17862
-#define CHILD_ADDR "127.0.0.1"
-
 // single-threaded, watching all child processes at once
 struct server_monitor {
     server_monitor(server_models & models) : models(models) {
@@ -65,11 +58,11 @@ struct server_monitor {
     }
 
     // thread-safe
-    void watch(const std::string & name, std::shared_ptr<server_subproc> proc, server_child_mode mode, int port) {
+    void watch(const std::string & name, std::shared_ptr<server_subproc> proc, server_role_type role, int port) {
         child_t c;
         c.name = name;
         c.proc = std::move(proc);
-        c.mode = mode;
+        c.role = role;
         c.port = port;
         if (!c.proc->has_output()) {
             SRV_ERR("failed to get stdout/stderr of child process for name=%s\n", name.c_str());
@@ -87,7 +80,7 @@ private:
     struct child_t {
         std::string name;
         std::shared_ptr<server_subproc> proc;
-        server_child_mode mode = SERVER_CHILD_MODE_NORMAL;
+        server_role_type role = SERVER_ROLE_INSTANCE;
         int port = 0;
         std::string buf[SERVER_SUBPROC_STREAMS]; // partial line of each pipe
         bool eof = false;     // output closed, waiting for the process to be reaped
@@ -241,7 +234,7 @@ private:
                 if (it->eof && !it->proc->is_alive()) {
                     int exit_code = it->proc->join();
                     it->proc->stopped.store(true, std::memory_order_release);
-                    models.on_child_exit(it->name, it->proc, it->mode, exit_code);
+                    models.on_child_exit(it->name, it->proc, it->role, exit_code);
                     SRV_INF("instance name=%s exited with status %d\n", it->name.c_str(), exit_code);
                     it = children.erase(it);
                 } else {
@@ -1324,7 +1317,7 @@ void server_models::load(const std::string & name) {
 
 void server_models::load(const std::string & name, const load_options & opts) {
     // the fake load time of a model; a download worker loads none
-    if (debug_fake_timing && opts.mode == SERVER_CHILD_MODE_NORMAL) {
+    if (debug_fake_timing && opts.role == SERVER_ROLE_INSTANCE) {
         // do not hold the mutex here, other requests must keep making progress
         std::this_thread::sleep_for(std::chrono::seconds(2));
     }
@@ -1345,7 +1338,7 @@ void server_models::load(const std::string & name, const load_options & opts) {
     // same model cannot start in between: a paused download or an idle cached model
     // is downloaded again with its files kept, a download on its way and a running
     // model are refused, and a preset or local model is not a download
-    if (opts.mode == SERVER_CHILD_MODE_DOWNLOAD) {
+    if (opts.role == SERVER_ROLE_DOWNLOAD) {
         auto it = mapping.find(name);
         if (it != mapping.end()) {
             const auto & cur = it->second.meta;
@@ -1389,7 +1382,7 @@ void server_models::load(const std::string & name, const load_options & opts) {
     // releasing its lock and this lock_guard acquiring allows multiple
     // threads to each observe capacity and all proceed to load.
     // Download workers do not use models_max slots.
-    if (opts.mode == SERVER_CHILD_MODE_NORMAL && base_params.models_max > 0) {
+    if (opts.role == SERVER_ROLE_INSTANCE && base_params.models_max > 0) {
         size_t count_active = 0;
         for (const auto & m : mapping) {
             if (m.second.meta.is_running()) {
@@ -1421,11 +1414,11 @@ void server_models::load(const std::string & name, const load_options & opts) {
 
         std::vector<std::string> child_args = inst.meta.args; // copy
         std::vector<std::string> child_env  = base_env; // copy
-        child_env.push_back("LLAMA_SERVER_ROUTER_PORT=" + std::to_string(base_params.port));
+        child_env.push_back(std::string(SERVER_ENV_ROUTER_PORT) + "=" + std::to_string(base_params.port));
 
-        if (opts.mode == SERVER_CHILD_MODE_DOWNLOAD) {
+        if (opts.role == SERVER_ROLE_DOWNLOAD) {
             inst.meta.status = SERVER_MODEL_STATUS_DOWNLOADING;
-            child_env.push_back("LLAMA_SERVER_CHILD_MODE=download");
+            child_env.push_back(std::string(SERVER_ENV_CHILD_ROLE) + "=" + server_role_type_to_str(SERVER_ROLE_DOWNLOAD));
             child_env.push_back("LLAMA_ARG_HF_REPO=" + name);
 
             // the sidecars as -hf would pull them: the draft types, and the mmproj switch
@@ -1460,7 +1453,7 @@ void server_models::load(const std::string & name, const load_options & opts) {
     }
 
     // remember the download: the partial files carry a content hash, not the name
-    if (opts.mode == SERVER_CHILD_MODE_DOWNLOAD) {
+    if (opts.role == SERVER_ROLE_DOWNLOAD) {
         unfinished_downloads.add({ name, opts.sidecars, opts.partials });
     }
 
@@ -1480,7 +1473,7 @@ void server_models::load(const std::string & name, const load_options & opts) {
     auto proc = inst.subproc;
     int  port = inst.meta.port;
     mapping[name] = std::move(inst);
-    monitor->watch(name, proc, opts.mode, port);
+    monitor->watch(name, proc, opts.role, port);
     cv.notify_all();
 }
 
@@ -1493,7 +1486,7 @@ void server_models::request_stop(const std::string & name, bool send_exit) {
     monitor->stop(name, it->second.meta.stop_timeout, send_exit);
 }
 
-void server_models::on_child_exit(const std::string & name, const std::shared_ptr<server_subproc> & proc, server_child_mode mode, int exit_code) {
+void server_models::on_child_exit(const std::string & name, const std::shared_ptr<server_subproc> & proc, server_role_type role, int exit_code) {
     {
         std::lock_guard<std::mutex> lk(mutex);
         auto it = mapping.find(name);
@@ -1502,7 +1495,7 @@ void server_models::on_child_exit(const std::string & name, const std::shared_pt
             return; // entry erased, or a newer instance took the name
         }
     }
-    if (mode == SERVER_CHILD_MODE_DOWNLOAD) {
+    if (role == SERVER_ROLE_DOWNLOAD) {
         // instance will be cleaned up on next load_models() call
         std::lock_guard<std::mutex> lk(mutex);
         stopping_models.erase(name);
@@ -2046,149 +2039,6 @@ void server_models::handle_child_state(const std::string & name, const std::stri
     }
 }
 
-//
-// server_child
-//
-
-server_child::server_child() {
-    if (is_child()) {
-        cmd_out = server_reserve_stdout();
-    }
-}
-
-server_child::~server_child() {
-    if (cmd_out) {
-        fclose(cmd_out);
-    }
-}
-
-bool server_child::is_child() {
-    const char * router_port = std::getenv("LLAMA_SERVER_ROUTER_PORT");
-    return router_port != nullptr;
-}
-
-server_child_mode server_child::get_mode() {
-    const char * mode = std::getenv("LLAMA_SERVER_CHILD_MODE");
-    std::string mode_str(mode ? mode : "");
-    if (mode_str == "download") {
-        return SERVER_CHILD_MODE_DOWNLOAD;
-    } else {
-        return SERVER_CHILD_MODE_NORMAL;
-    }
-}
-
-struct server_download_state : public common_download_callback {
-    server_child * self;
-    std::function<bool()> should_stop;
-    std::atomic<int64_t> last_progress_time{0}; // multiple files downloading in different threads
-    bool is_ok = false;
-
-    server_download_state(server_child * s) : self(s) {}
-
-    bool run(common_params & params) {
-        try {
-            common_models_handler handler = common_models_handler_init(params, LLAMA_EXAMPLE_SERVER);
-            common_models_handler_apply(handler, params, this);
-            is_ok = true;
-        } catch (const std::exception & e) {
-            auto model_name = params.model.get_name();
-            SRV_ERR("download failed for model name=%s: %s\n", model_name.c_str(), e.what());
-            is_ok = false;
-        }
-        return is_ok;
-    }
-    void on_progress(const common_download_progress & p) {
-        json data = {
-            {"url", p.url},
-            {"downloaded", p.downloaded},
-            {"total", p.total},
-        };
-        self->notify_to_router(server_state_to_str(SERVER_STATE_DOWNLOADING), data);
-    }
-    void on_start(const common_download_progress & p) override {
-        on_progress(p);
-    }
-    void on_update(const common_download_progress & p) override {
-        int64_t now = ggml_time_ms();
-        // throttle progress updates to avoid flooding logs
-        if (now - last_progress_time.load(std::memory_order_relaxed) >= 100) {
-            on_progress(p);
-            last_progress_time.store(now, std::memory_order_relaxed);
-        }
-    }
-    void on_done(const common_download_progress & p, bool) override {
-        on_progress(p);
-    }
-    bool is_cancelled() const override {
-        return should_stop ? should_stop() : false;
-    }
-};
-
-int server_child::run_download(common_params & params) {
-    auto cancelled = std::make_shared<std::atomic<bool>>(false);
-
-    // monitor stdin for cancellation command from the router
-    std::thread signal_thread = setup([cancelled](int) {
-        cancelled->store(true, std::memory_order_relaxed);
-    });
-
-    server_download_state dl(this);
-    dl.should_stop = [cancelled]() {
-        return cancelled->load(std::memory_order_relaxed);
-    };
-
-    bool ok = dl.run(params);
-
-    notify_to_router(server_state_to_str(SERVER_STATE_DOWNLOADING), {
-        {"result", ok ? "download_finished" : "download_failed"},
-    });
-
-    // router should send CMD_ROUTER_TO_CHILD_EXIT after receiving the result
-    if (signal_thread.joinable()) {
-        signal_thread.join();
-    }
-
-    SRV_INF("download completed %s\n", ok ? "successfully" : "with errors");
-    return 0;
-}
-
-std::thread server_child::setup(const std::function<void(int)> & shutdown_handler) {
-    // setup thread for monitoring stdin
-    return std::thread([shutdown_handler]() {
-        // wait for EOF on stdin
-        SRV_INF("%s", "child server monitoring thread started, waiting for EOF on stdin...\n");
-        bool eof = false;
-        while (true) {
-            std::string line;
-            if (!std::getline(std::cin, line)) {
-                // EOF detected, that means the router server is unexpectedly exit or killed
-                eof = true;
-                break;
-            }
-            if (line.find(CMD_ROUTER_TO_CHILD_EXIT) != std::string::npos) {
-                SRV_INF("%s", "exit command received, exiting...\n");
-                shutdown_handler(0);
-                break;
-            }
-        }
-        if (eof) {
-            SRV_INF("%s", "EOF on stdin detected, forcing shutdown...\n");
-            exit(1);
-        }
-    });
-}
-
-void server_child::notify_to_router(const std::string & state, const json & payload) {
-    json data = {
-        {"state", state},
-        {"payload", payload},
-    };
-    std::lock_guard<std::mutex> lk(mtx_stdout);
-    fprintf(cmd_out, "%s%s\n", CMD_CHILD_TO_ROUTER_STATE, safe_json_to_str(data).c_str());
-    fflush(cmd_out);
-}
-
-
 //
 // server_models_routes
 //
@@ -2617,7 +2467,7 @@ void server_models_routes::init_routes() {
         SRV_INF("starting download for model '%s'\n", name.c_str());
         {
             server_models::load_options load_opts;
-            load_opts.mode = SERVER_CHILD_MODE_DOWNLOAD;
+            load_opts.role = SERVER_ROLE_DOWNLOAD;
             load_opts.custom_meta = server_model_meta{};
             load_opts.custom_meta->source = SERVER_MODEL_SOURCE_CACHE;
             load_opts.custom_meta->name   = name;

tools/server/server-models.h

@@ -46,11 +46,6 @@ enum server_model_source {
     SERVER_MODEL_SOURCE_CACHE,
 };
 
-enum server_child_mode {
-    SERVER_CHILD_MODE_NORMAL,   // load the model and run normally
-    SERVER_CHILD_MODE_DOWNLOAD, // download the model and exit
-};
-
 static std::string server_model_status_to_string(server_model_status status) {
     switch (status) {
         case SERVER_MODEL_STATUS_DOWNLOADING: return "downloading";
@@ -171,7 +166,7 @@ private:
         server_model_meta meta;
         int req_count = 0; // number of active proxy requests
 
-        // ask the child to exit (it handles the command on its stdin, see server_child::setup)
+        // ask the child to exit (it handles the command on its stdin, see server_role::setup)
         void request_exit() const;
     };
 
@@ -276,7 +271,7 @@ private:
     void request_stop(const std::string & name, bool send_exit = true);
 
     // called by the monitor once a child exited and was reaped
-    void on_child_exit(const std::string & name, const std::shared_ptr<server_subproc> & proc, server_child_mode mode, int exit_code);
+    void on_child_exit(const std::string & name, const std::shared_ptr<server_subproc> & proc, server_role_type role, int exit_code);
 
     // notify SSE clients
     void notify_sse(const std::string & event, const std::string & model_id, const json & data = nullptr);
@@ -320,7 +315,7 @@ public:
     std::vector<server_model_meta> get_all_meta();
 
     struct load_options {
-        server_child_mode mode = SERVER_CHILD_MODE_NORMAL;
+        server_role_type role = SERVER_ROLE_INSTANCE;
         // used for spawning a downloading child process
         std::optional<server_model_meta> custom_meta = std::nullopt;
         // download mode: the sidecar kinds pulled with the model, as -hf would, and
@@ -369,7 +364,7 @@ public:
     // forward a websocket connection to the child, returns when either side closes
     void proxy_ws(const server_http_req & req, server_http_context::ws_conn & conn, const std::string & name);
 
-    // handle message sent from server_child::notify_to_router()
+    // handle message sent from server_role::notify_to_router()
     // raw input must starts with CMD_CHILD_TO_ROUTER_STATE, followed by a JSON string
     // called from the monitor thread
     // payload per state:
@@ -383,33 +378,6 @@ private:
     std::unique_ptr<server_monitor> monitor;
 };
 
-struct server_child {
-    // serializes the notify_to_router writes
-    std::mutex mtx_stdout;
-    std::atomic<bool> is_finished_downloading = false; // set by run_download
-
-    // in a child, keeps stdout for the commands to the router, so it is created before anything is written;
-    // everything else written to stdout goes to stderr with the logs
-    server_child();
-    ~server_child();
-
-    // return true if the current process is a child server instance
-    bool is_child();
-    server_child_mode get_mode();
-    int run_download(common_params & params);
-
-    // register the shutdown_handler to be called by the router
-    // return the monitoring thread (to be joined by the caller)
-    std::thread setup(const std::function<void(int)> & shutdown_handler);
-
-    // notify router server for status changes (e.g. loading, downloading, sleeping, etc.)
-    // message will be handled by server_models::handle_child_state() on the router side
-    void notify_to_router(const std::string & state_name, const json & payload);
-
-private:
-    FILE * cmd_out = nullptr; // the stdout the router reads the commands from
-};
-
 struct server_models_routes {
     common_params params;
     json ui_settings = json::object();     // Primary: new name

tools/server/server.cpp

@@ -37,12 +37,17 @@ static inline void signal_handler(int signal) {
     shutdown_handler(signal);
 }
 
+// how the server is run; everything is on when it is the llama-server binary
+struct server_run_opts {
+    bool install_signal_handlers = true; // off when embedded in llama-cli, which handles them itself
+    bool resolve_model           = true; // download / resolve -hf and docker models; off when the caller did it already
+};
+
 // satisfies -Wmissing-declarations (used by llama command)
 int llama_server(int argc, char ** argv);
 
 // to be used via CLI (argc / argv are used by router mode only)
 int llama_server(common_params & params, int argc, char ** argv);
-int llama_server(common_params & params, int argc, char ** argv, server_child & child);
 void llama_server_terminate();
 void llama_server_terminate() {
     if (shutdown_handler) {
@@ -91,72 +96,24 @@ static server_http_context::handler_t ex_wrapper(server_http_context::handler_t
     };
 }
 
-int llama_server(int argc, char ** argv) {
-    server_child child;
+//
+// shared by every role
+//
 
-    std::setlocale(LC_NUMERIC, "C");
-
-#ifndef _WIN32
-    // Ignore SIGPIPE so the server does not crash if a child (MCP server, tools runtime) exits while we are writing to its stdin
-    signal(SIGPIPE, SIG_IGN);
-#endif
-
-    // own arguments required by this example
-    common_params params;
-
-    common_init();
-
-    // start the stream session manager GC right after common init, before any HTTP route can
-    // touch it. lifecycle is symmetric, stop_gc() runs in clean_up() before backend free
-    server_stream_session_manager_start();
-
-    if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_SERVER)) {
-        return 1;
+static server_role_type server_role_detect(const common_params & params, server_role & role) {
+    if (role.is_managed_by_router()) {
+        return role.get_role();
     }
-
-    SRV_INF("%s", "initializing ...\n");
-
-    llama_backend_init();
-    llama_numa_init(params.numa);
-
-    const int result = llama_server(params, argc, argv, child);
-    common_log_flush(common_log_main());
-    return result;
+    // a server given no model is a router
+    const bool has_model = !params.model.path.empty()
+                        || !params.model.hf_repo.empty()
+                        || !params.model.docker_repo.empty();
+    return has_model ? SERVER_ROLE_INSTANCE : SERVER_ROLE_ROUTER;
 }
 
-int llama_server(common_params & params, int argc, char ** argv) {
-    server_child child;
-    return llama_server(params, argc, argv, child);
-}
-
-int llama_server(common_params & params, int argc, char ** argv, server_child & child) {
-    bool is_run_by_cli = (argv == nullptr);
-
-    common_models_handler models_handler;
-
-    // note: router mode also accepts -hf remote-preset, so we need to check that first
-    if (!is_run_by_cli && !params.model.hf_repo.empty()) {
-        try {
-            models_handler = common_models_handler_init(params, LLAMA_EXAMPLE_SERVER);
-            if (common_models_handler_is_preset_repo(models_handler)) {
-                // apply the preset and start the server in router mode
-                common_models_handler_apply(models_handler, params);
-            }
-        } catch (const std::exception & e) {
-            SRV_ERR("failed to fetch model metadata: %s\n", e.what());
-            return 1;
-        }
-    }
-
-    // router server never loads a model and must not touch the GPU
-    const bool is_router_server = params.model.path.empty()
-                               && params.model.hf_repo.empty()
-                               && params.model.docker_repo.empty();
-
-    // skip device enumeration so the CUDA primary context stays uncreated
-    common_params_print_info(params, !is_router_server);
-
-    if (!is_router_server) {
+// adjust the params for the role; returns false on a configuration error
+static bool server_prepare_params(common_params & params, server_role_type role_type) {
+    if (role_type != SERVER_ROLE_ROUTER) {
         // validate batch size for embeddings
         // embeddings require all tokens to be processed in a single ubatch
         // see https://github.com/ggml-org/llama.cpp/issues/12836
@@ -197,7 +154,7 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         const std::string reason = server_connect::unavailable_reason(params);
         if (!reason.empty()) {
             SRV_ERR("--connect is not available: %s\n", reason.c_str());
-            return 1;
+            return false;
         }
         params.server_connect_code = server_connect::resolve_code(params.server_connect_code);
     } else if (!params.server_connect_code.empty()) {
@@ -206,77 +163,11 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         params.server_connect_code.clear();
     }
 
-    // note: this is guaranteed to out-live ctx_http and tools
-    server_mcp mcp_mgr;
-
-    // struct that contains llama context and inference
-    server_context ctx_server;
-
-    //
-    // Router
-    //
-
-    // register API routes
-    server_routes routes(params, ctx_server);
-    server_tools tools;
-
-    std::optional<server_models_routes> models_routes{};
-
-    server_http_context ctx_http;
-    if (!ctx_http.init(params)) {
-        SRV_ERR("%s", "failed to initialize HTTP server\n");
-        return 1;
-    }
-
-    if (is_router_server) {
-        // setup server instances manager
-        try {
-            models_routes.emplace(params, argc, argv);
-        } catch (const std::exception & e) {
-            SRV_ERR("failed to initialize router models: %s\n", e.what());
-            return 1;
-        }
-
-        // proxy handlers
-        // note: routes.get_health stays the same
-        routes.get_metrics                 = models_routes->proxy_get;
-        routes.post_props                  = models_routes->proxy_post;
-        routes.post_completions            = models_routes->proxy_post;
-        routes.post_completions_oai        = models_routes->proxy_post;
-        routes.post_chat_completions       = models_routes->proxy_post;
-        routes.post_control                = models_routes->proxy_post;
-        routes.post_responses_oai          = models_routes->proxy_post;
-        routes.post_transcriptions_oai     = models_routes->proxy_post;
-        routes.post_tts                    = models_routes->proxy_post;
-        routes.ws_realtime                 = models_routes->proxy_ws;
-        routes.post_anthropic_messages     = models_routes->proxy_post;
-        routes.post_anthropic_count_tokens = models_routes->proxy_post;
-        routes.post_infill                 = models_routes->proxy_post;
-        routes.post_embeddings             = models_routes->proxy_post;
-        routes.post_embeddings_oai         = models_routes->proxy_post;
-        routes.post_rerank                 = models_routes->proxy_post;
-        routes.post_systemone              = models_routes->proxy_post;
-        routes.post_tokenize               = models_routes->proxy_post;
-        routes.post_detokenize             = models_routes->proxy_post;
-        routes.post_apply_template         = models_routes->proxy_post;
-        routes.post_chat_completions_tok   = models_routes->proxy_post;
-        routes.post_responses_tok_oai      = models_routes->proxy_post;
-        routes.get_lora_adapters           = models_routes->proxy_get;
-        routes.post_lora_adapters          = models_routes->proxy_post;
-        routes.get_slots                   = models_routes->proxy_get;
-        routes.post_slots                  = models_routes->proxy_post;
-
-        // custom routes for router
-        routes.get_props                   = models_routes->get_router_props;
-        routes.get_models                  = models_routes->get_router_models;
-
-        ctx_http.post("/models",               ex_wrapper(models_routes->post_router_models));
-        ctx_http.post("/models/load",          ex_wrapper(models_routes->post_router_models_load));
-        ctx_http.post("/models/unload",        ex_wrapper(models_routes->post_router_models_unload));
-        ctx_http.get ("/models/sse",           ex_wrapper(models_routes->get_router_models_sse));
-        ctx_http.del ("/models",               ex_wrapper(models_routes->del_router_models));
-    }
+    return true;
+}
 
+// the API of every server; a router points the handlers in routes at its proxies before calling this
+static void server_register_routes(server_http_context & ctx_http, server_routes & routes) {
     ctx_http.get ("/health",                   ex_wrapper(routes.get_health)); // public endpoint (no API key check)
     ctx_http.get ("/v1/health",                ex_wrapper(routes.get_health)); // public endpoint (no API key check)
     ctx_http.get ("/metrics",                  ex_wrapper(routes.get_metrics));
@@ -322,27 +213,30 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
     ctx_http.get ("/slots",                    ex_wrapper(routes.get_slots));
     ctx_http.post("/slots/:id_slot",           ex_wrapper(routes.post_slots));
 
-    // resumable streaming: a child binds the local session factories, the router binds
-    // proxies that resolve the owning child, see server-stream.h
-    server_http_context::handler_t stream_get_h;
-    server_http_context::handler_t streams_lookup_h;
-    server_http_context::handler_t stream_delete_h;
-    if (is_router_server) {
-        stream_get_h     = models_routes->router_stream_get;
-        streams_lookup_h = models_routes->router_streams_lookup;
-        stream_delete_h  = models_routes->router_stream_delete;
-    } else {
-        stream_get_h     = server_stream_make_get_handler();
-        streams_lookup_h = server_stream_make_lookup_handler();
-        stream_delete_h  = server_stream_make_delete_handler();
-    }
-    ctx_http.get ("/v1/stream",                ex_wrapper(stream_get_h));
-    ctx_http.post("/v1/streams/lookup",        ex_wrapper(streams_lookup_h));
-    ctx_http.del ("/v1/stream",                ex_wrapper(stream_delete_h));
-
     // Google Cloud Platform (Vertex AI) compat
     ctx_http.register_gcp_compat();
+}
 
+// resumable streaming, see server-stream.h: an instance binds the local session factories,
+// a router binds proxies that resolve the owning child
+static void server_register_stream_routes(
+        server_http_context & ctx_http,
+        const server_http_context::handler_t & stream_get,
+        const server_http_context::handler_t & streams_lookup,
+        const server_http_context::handler_t & stream_delete) {
+    ctx_http.get ("/v1/stream",                ex_wrapper(stream_get));
+    ctx_http.post("/v1/streams/lookup",        ex_wrapper(streams_lookup));
+    ctx_http.del ("/v1/stream",                ex_wrapper(stream_delete));
+}
+
+// the features behind a flag, and the warning about the ones that must not face untrusted clients
+// returns false when one of them fails to start
+static bool server_setup_features(
+        server_http_context & ctx_http,
+        const common_params & params,
+        server_mcp & mcp_mgr,
+        server_tools & tools,
+        server_role_type role_type) {
     // return 403 for disabled features
     server_http_context::handler_t res_403 = [](const server_http_req &) {
         auto res = std::make_unique<server_http_res>();
@@ -360,9 +254,8 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         SRV_WRN("%s", "security: no API key is set and CORS allows all origins (see https://github.com/ggml-org/llama.cpp/pull/25655)\n");
     }
 
-    // CORS proxy (EXPERIMENTAL, only used by the Web UI for MCP)
     std::vector<std::string> warn_names;
-    if (is_router_server) {
+    if (role_type == SERVER_ROLE_ROUTER) {
         warn_names.push_back("router mode");
     }
 
@@ -370,6 +263,7 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         warn_names.push_back("peer-to-peer tunnel (llama-connect, experimental)");
     }
 
+    // CORS proxy (EXPERIMENTAL, only used by the Web UI for MCP)
     if (params.ui_mcp_proxy) {
         ctx_http.get ("/cors-proxy",      ex_wrapper(proxy_handler_get));
         ctx_http.post("/cors-proxy",      ex_wrapper(proxy_handler_post));
@@ -383,7 +277,7 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         mcp_mgr.start(params);
     } catch (const std::exception & e) {
         SRV_ERR("MCP starting failed: %s\n", e.what());
-        return 1;
+        return false;
     }
 
     if (!params.server_tools.empty() || !mcp_mgr.empty()) {
@@ -391,7 +285,7 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
             tools.setup(params.server_tools, mcp_mgr, params.server_tools_runtime);
         } catch (const std::exception & e) {
             SRV_ERR("tools setup failed: %s\n", e.what());
-            return 1;
+            return false;
         }
         ctx_http.get ("/tools",           ex_wrapper(tools.handle_get));
         ctx_http.post("/tools",           ex_wrapper(tools.handle_post));
@@ -418,15 +312,267 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         SRV_WRN("security: %s enabled - do not expose to untrusted environments\n", features.c_str());
     }
 
-    //
-    // Handle downloading model
-    //
+    return true;
+}
 
-    if (child.is_child() && child.get_mode() == SERVER_CHILD_MODE_DOWNLOAD) {
-        return child.run_download(params);
-    } else if (!is_router_server && !is_run_by_cli) {
-        // single-model mode (NOT spawned by router)
-        // if this is invoked by CLI, model downloading should be already handled
+static void server_install_signal_handlers() {
+#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__))
+    struct sigaction sigint_action;
+    sigint_action.sa_handler = signal_handler;
+    sigemptyset (&sigint_action.sa_mask);
+    sigint_action.sa_flags = 0;
+    sigaction(SIGINT, &sigint_action, NULL);
+    sigaction(SIGTERM, &sigint_action, NULL);
+#elif defined (_WIN32)
+    auto console_ctrl_handler = +[](DWORD ctrl_type) -> BOOL {
+        return (ctrl_type == CTRL_C_EVENT) ? (signal_handler(SIGINT), true) : false;
+    };
+    SetConsoleCtrlHandler(reinterpret_cast<PHANDLER_ROUTINE>(console_ctrl_handler), true);
+#endif
+}
+
+// once the HTTP server listens: log the addresses and spawn llama-connect if asked, so the port is
+// bound and the child startup health check reaches us. the destructor of connect_proc stops it
+static bool server_on_listening(server_http_context & ctx_http, const common_params & params, server_connect & connect_proc) {
+    for (const auto & address : ctx_http.listening_addresses) {
+        SRV_INF("listening on %s\n", address.c_str());
+    }
+
+    if (params.server_connect && !connect_proc.start(params, ctx_http.port)) {
+        SRV_ERR("%s", "exiting due to llama-connect error\n");
+        return false;
+    }
+    return true;
+}
+
+//
+// download worker: a child of the router that pulls a model into the cache and reports the progress
+//
+
+struct server_download_state : public common_download_callback {
+    server_role * role;
+    std::function<bool()> should_stop;
+    std::atomic<int64_t> last_progress_time{0}; // multiple files downloading in different threads
+    bool is_ok = false;
+
+    server_download_state(server_role * r) : role(r) {}
+
+    bool run(common_params & params) {
+        try {
+            common_models_handler handler = common_models_handler_init(params, LLAMA_EXAMPLE_SERVER);
+            common_models_handler_apply(handler, params, this);
+            is_ok = true;
+        } catch (const std::exception & e) {
+            auto model_name = params.model.get_name();
+            SRV_ERR("download failed for model name=%s: %s\n", model_name.c_str(), e.what());
+            is_ok = false;
+        }
+        return is_ok;
+    }
+    void on_progress(const common_download_progress & p) {
+        json data = {
+            {"url", p.url},
+            {"downloaded", p.downloaded},
+            {"total", p.total},
+        };
+        role->notify_to_router(server_state_to_str(SERVER_STATE_DOWNLOADING), data);
+    }
+    void on_start(const common_download_progress & p) override {
+        on_progress(p);
+    }
+    void on_update(const common_download_progress & p) override {
+        int64_t now = ggml_time_ms();
+        // throttle progress updates to avoid flooding logs
+        if (now - last_progress_time.load(std::memory_order_relaxed) >= 100) {
+            on_progress(p);
+            last_progress_time.store(now, std::memory_order_relaxed);
+        }
+    }
+    void on_done(const common_download_progress & p, bool) override {
+        on_progress(p);
+    }
+    bool is_cancelled() const override {
+        return should_stop ? should_stop() : false;
+    }
+};
+
+static int server_run_download(common_params & params, server_role & role) {
+    auto cancelled = std::make_shared<std::atomic<bool>>(false);
+
+    // monitor stdin for cancellation command from the router
+    std::thread signal_thread = role.setup([cancelled](int) {
+        cancelled->store(true, std::memory_order_relaxed);
+    });
+
+    server_download_state dl(&role);
+    dl.should_stop = [cancelled]() {
+        return cancelled->load(std::memory_order_relaxed);
+    };
+
+    bool ok = dl.run(params);
+
+    role.notify_to_router(server_state_to_str(SERVER_STATE_DOWNLOADING), {
+        {"result", ok ? "download_finished" : "download_failed"},
+    });
+
+    // router should send CMD_ROUTER_TO_CHILD_EXIT after receiving the result
+    if (signal_thread.joinable()) {
+        signal_thread.join();
+    }
+
+    SRV_INF("download completed %s\n", ok ? "successfully" : "with errors");
+    return 0;
+}
+
+//
+// router
+//
+
+static int server_run_router(common_params & params, int argc, char ** argv, const server_run_opts & opts) {
+    SRV_INF("%s", "starting server in router mode. models will be automatically loaded on-demand\n");
+
+    // note: this is guaranteed to out-live ctx_http and tools
+    server_mcp mcp_mgr;
+
+    // the router loads no model, its context stays empty; the routes it does not proxy still need one
+    server_context ctx_server;
+    server_routes  routes(params, ctx_server);
+    server_tools   tools;
+
+    server_http_context ctx_http;
+    if (!ctx_http.init(params)) {
+        SRV_ERR("%s", "failed to initialize HTTP server\n");
+        return 1;
+    }
+
+    // setup server instances manager
+    std::optional<server_models_routes> models_routes;
+    try {
+        models_routes.emplace(params, argc, argv);
+    } catch (const std::exception & e) {
+        SRV_ERR("failed to initialize router models: %s\n", e.what());
+        return 1;
+    }
+
+    // proxy handlers
+    // note: routes.get_health stays the same
+    routes.get_metrics                 = models_routes->proxy_get;
+    routes.post_props                  = models_routes->proxy_post;
+    routes.post_completions            = models_routes->proxy_post;
+    routes.post_completions_oai        = models_routes->proxy_post;
+    routes.post_chat_completions       = models_routes->proxy_post;
+    routes.post_control                = models_routes->proxy_post;
+    routes.post_responses_oai          = models_routes->proxy_post;
+    routes.post_transcriptions_oai     = models_routes->proxy_post;
+    routes.post_tts                    = models_routes->proxy_post;
+    routes.ws_realtime                 = models_routes->proxy_ws;
+    routes.post_anthropic_messages     = models_routes->proxy_post;
+    routes.post_anthropic_count_tokens = models_routes->proxy_post;
+    routes.post_infill                 = models_routes->proxy_post;
+    routes.post_embeddings             = models_routes->proxy_post;
+    routes.post_embeddings_oai         = models_routes->proxy_post;
+    routes.post_rerank                 = models_routes->proxy_post;
+    routes.post_systemone              = models_routes->proxy_post;
+    routes.post_tokenize               = models_routes->proxy_post;
+    routes.post_detokenize             = models_routes->proxy_post;
+    routes.post_apply_template         = models_routes->proxy_post;
+    routes.post_chat_completions_tok   = models_routes->proxy_post;
+    routes.post_responses_tok_oai      = models_routes->proxy_post;
+    routes.get_lora_adapters           = models_routes->proxy_get;
+    routes.post_lora_adapters          = models_routes->proxy_post;
+    routes.get_slots                   = models_routes->proxy_get;
+    routes.post_slots                  = models_routes->proxy_post;
+
+    // custom routes for router
+    routes.get_props                   = models_routes->get_router_props;
+    routes.get_models                  = models_routes->get_router_models;
+
+    ctx_http.post("/models",               ex_wrapper(models_routes->post_router_models));
+    ctx_http.post("/models/load",          ex_wrapper(models_routes->post_router_models_load));
+    ctx_http.post("/models/unload",        ex_wrapper(models_routes->post_router_models_unload));
+    ctx_http.get ("/models/sse",           ex_wrapper(models_routes->get_router_models_sse));
+    ctx_http.del ("/models",               ex_wrapper(models_routes->del_router_models));
+
+    server_register_routes(ctx_http, routes);
+    server_register_stream_routes(ctx_http,
+        models_routes->router_stream_get,
+        models_routes->router_streams_lookup,
+        models_routes->router_stream_delete);
+
+    if (!server_setup_features(ctx_http, params, mcp_mgr, tools, SERVER_ROLE_ROUTER)) {
+        return 1;
+    }
+
+    auto clean_up = [&models_routes, &mcp_mgr]() {
+        SRV_INF("%s: cleaning up before exit...\n", __func__);
+        // stop the session GC first, it finalizes live sessions and wakes pending readers
+        server_stream_session_manager_stop();
+        if (models_routes.has_value()) {
+            models_routes->stopping.store(true); // maybe redundant, but just to be safe
+            models_routes->models.unload_all();
+        }
+        mcp_mgr.shutdown();
+        llama_backend_free();
+    };
+
+    if (!ctx_http.start()) {
+        clean_up();
+        SRV_ERR("%s", "exiting due to HTTP server error\n");
+        return 1;
+    }
+    ctx_http.is_ready.store(true);
+
+    shutdown_handler = [&](int) {
+        if (models_routes.has_value()) {
+            // important to disconnect any SSE clients
+            models_routes->stopping.store(true);
+        }
+        mcp_mgr.shutdown();
+        ctx_http.stop();
+    };
+
+    try {
+        models_routes->models.load_startup_models();
+    } catch (const std::exception & e) {
+        SRV_ERR("failed to load models on startup: %s\n", e.what());
+        ctx_http.stop();
+        ctx_http.join();
+        clean_up();
+        return 1;
+    }
+
+    if (opts.install_signal_handlers) {
+        server_install_signal_handlers();
+    }
+
+    server_connect connect_proc;
+    if (!server_on_listening(ctx_http, params, connect_proc)) {
+        ctx_http.stop();
+        ctx_http.join();
+        clean_up();
+        return 1;
+    }
+
+    if (!params.models_preset_hf.empty()) {
+        SRV_WRN(      "NOTE: using preset.ini from HF repo '%s'\n", params.models_preset_hf.c_str());
+        SRV_WRN("%s", "      please only use presets that you can trust! Unknown presets may be unsafe\n");
+    }
+
+    ctx_http.join(); // keep the main thread alive
+
+    // when the HTTP server stops, clean up and exit
+    clean_up();
+
+    return 0;
+}
+
+//
+// instance
+//
+
+static int server_run_instance(common_params & params, server_role & role, common_models_handler & models_handler, const server_run_opts & opts) {
+    // -hf / docker: download if needed and resolve the local path. a child resolves what the router downloaded
+    if (opts.resolve_model) {
         try {
             common_models_handler_apply(models_handler, params);
         } catch (const std::exception & e) {
@@ -435,161 +581,188 @@ int llama_server(common_params & params, int argc, char ** argv, server_child &
         }
     }
 
-    //
-    // Start the server
-    //
+    // note: this is guaranteed to out-live ctx_http and tools
+    server_mcp mcp_mgr;
 
-    std::function<void()> clean_up;
+    // struct that contains llama context and inference
+    server_context ctx_server;
+    server_routes  routes(params, ctx_server);
+    server_tools   tools;
 
-    if (is_router_server) {
-        SRV_INF("%s", "starting server in router mode. models will be automatically loaded on-demand\n");
-
-        clean_up = [&models_routes, &mcp_mgr]() {
-            SRV_INF("%s: cleaning up before exit...\n", __func__);
-            // stop the session GC first, it finalizes live sessions and wakes pending readers
-            server_stream_session_manager_stop();
-            if (models_routes.has_value()) {
-                models_routes->stopping.store(true); // maybe redundant, but just to be safe
-                models_routes->models.unload_all();
-            }
-            mcp_mgr.shutdown();
-            llama_backend_free();
-        };
-
-        if (!ctx_http.start()) {
-            clean_up();
-            SRV_ERR("%s", "exiting due to HTTP server error\n");
-            return 1;
-        }
-        ctx_http.is_ready.store(true);
-
-        shutdown_handler = [&](int) {
-            if (models_routes.has_value()) {
-                // important to disconnect any SSE clients
-                models_routes->stopping.store(true);
-            }
-            mcp_mgr.shutdown();
-            ctx_http.stop();
-        };
-
-        try {
-            models_routes->models.load_startup_models();
-        } catch (const std::exception & e) {
-            SRV_ERR("failed to load models on startup: %s\n", e.what());
-            ctx_http.stop();
-            ctx_http.join();
-            clean_up();
-            return 1;
-        }
-
-    } else {
-        // setup clean up function, to be called before exit
-        clean_up = [&ctx_http, &ctx_server, &mcp_mgr]() {
-            SRV_INF("%s: cleaning up before exit...\n", __func__);
-            // stop the session GC first, it finalizes live sessions and wakes pending readers
-            server_stream_session_manager_stop();
-            ctx_http.stop();
-            ctx_server.terminate();
-            mcp_mgr.shutdown();
-            llama_backend_free();
-        };
-
-        // start the HTTP server before loading the model to be able to serve /health requests
-        if (!ctx_http.start()) {
-            clean_up();
-            SRV_ERR("%s", "exiting due to HTTP server error\n");
-            return 1;
-        }
-
-        // setup communication child --> router if necessary
-        if (child.is_child()) {
-            ctx_server.set_state_callback([&](server_state state, json payload) {
-                child.notify_to_router(server_state_to_str(state), payload);
-            });
-        }
-
-        if (!ctx_server.load_model(params)) {
-            clean_up();
-            ctx_http.join();
-            SRV_ERR("%s", "exiting due to model loading error\n");
-            return 1;
-        }
-
-        routes.update_meta(ctx_server);
-        ctx_http.is_ready.store(true);
-
-        SRV_INF("%s", "model loaded\n");
-
-        shutdown_handler = [&](int) {
-            mcp_mgr.shutdown();
-            // this will unblock start_loop()
-            ctx_server.terminate();
-        };
+    server_http_context ctx_http;
+    if (!ctx_http.init(params)) {
+        SRV_ERR("%s", "failed to initialize HTTP server\n");
+        return 1;
     }
 
-    // register signal handler if not running by CLI
-    if (!is_run_by_cli) {
-#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__))
-        struct sigaction sigint_action;
-        sigint_action.sa_handler = signal_handler;
-        sigemptyset (&sigint_action.sa_mask);
-        sigint_action.sa_flags = 0;
-        sigaction(SIGINT, &sigint_action, NULL);
-        sigaction(SIGTERM, &sigint_action, NULL);
-#elif defined (_WIN32)
-        auto console_ctrl_handler = +[](DWORD ctrl_type) -> BOOL {
-            return (ctrl_type == CTRL_C_EVENT) ? (signal_handler(SIGINT), true) : false;
-        };
-        SetConsoleCtrlHandler(reinterpret_cast<PHANDLER_ROUTINE>(console_ctrl_handler), true);
-#endif
+    server_register_routes(ctx_http, routes);
+    server_register_stream_routes(ctx_http,
+        server_stream_make_get_handler(),
+        server_stream_make_lookup_handler(),
+        server_stream_make_delete_handler());
+
+    if (!server_setup_features(ctx_http, params, mcp_mgr, tools, SERVER_ROLE_INSTANCE)) {
+        return 1;
     }
 
-    for (const auto & address : ctx_http.listening_addresses) {
-        SRV_INF("listening on %s\n", address.c_str());
+    // setup clean up function, to be called before exit
+    auto clean_up = [&ctx_http, &ctx_server, &mcp_mgr]() {
+        SRV_INF("%s: cleaning up before exit...\n", __func__);
+        // stop the session GC first, it finalizes live sessions and wakes pending readers
+        server_stream_session_manager_stop();
+        ctx_http.stop();
+        ctx_server.terminate();
+        mcp_mgr.shutdown();
+        llama_backend_free();
+    };
+
+    // start the HTTP server before loading the model to be able to serve /health requests
+    if (!ctx_http.start()) {
+        clean_up();
+        SRV_ERR("%s", "exiting due to HTTP server error\n");
+        return 1;
+    }
+
+    // setup communication child --> router if necessary
+    if (role.is_managed_by_router()) {
+        ctx_server.set_state_callback([&](server_state state, json payload) {
+            role.notify_to_router(server_state_to_str(state), payload);
+        });
+    }
+
+    if (!ctx_server.load_model(params)) {
+        clean_up();
+        ctx_http.join();
+        SRV_ERR("%s", "exiting due to model loading error\n");
+        return 1;
+    }
+
+    routes.update_meta(ctx_server);
+    ctx_http.is_ready.store(true);
+
+    SRV_INF("%s", "model loaded\n");
+
+    shutdown_handler = [&](int) {
+        mcp_mgr.shutdown();
+        // this will unblock start_loop()
+        ctx_server.terminate();
+    };
+
+    if (opts.install_signal_handlers) {
+        server_install_signal_handlers();
     }
 
-    // spawn only once listening, so the port is bound and the child startup health check reaches us. the destructor stops it
     server_connect connect_proc;
-    if (params.server_connect && !connect_proc.start(params, ctx_http.port)) {
-        SRV_ERR("%s", "exiting due to llama-connect error\n");
+    if (!server_on_listening(ctx_http, params, connect_proc)) {
         ctx_http.stop();
         ctx_http.join();
         clean_up();
         return 1;
     }
 
-    if (is_router_server) {
-        if (!params.models_preset_hf.empty()) {
-            SRV_WRN(      "NOTE: using preset.ini from HF repo '%s'\n", params.models_preset_hf.c_str());
-            SRV_WRN("%s", "      please only use presets that you can trust! Unknown presets may be unsafe\n");
-        }
+    // optionally, notify router server that this instance is ready
+    std::thread monitor_thread;
+    if (role.is_managed_by_router()) {
+        monitor_thread = role.setup(shutdown_handler);
+        role.notify_to_router(server_state_to_str(SERVER_STATE_READY), routes.get_model_info());
+    }
 
-        ctx_http.join(); // keep the main thread alive
+    // this call blocks the main thread until queue_tasks.terminate() is called
+    ctx_server.start_loop();
 
-        // when the HTTP server stops, clean up and exit
-        clean_up();
-    } else {
-        // optionally, notify router server that this instance is ready
-        std::thread monitor_thread;
-        if (child.is_child()) {
-            monitor_thread = child.setup(shutdown_handler);
-            child.notify_to_router(server_state_to_str(SERVER_STATE_READY), routes.get_model_info());
-        }
+    clean_up();
+    ctx_http.join();
+    if (monitor_thread.joinable()) {
+        monitor_thread.join();
+    }
 
-        // this call blocks the main thread until queue_tasks.terminate() is called
-        ctx_server.start_loop();
-
-        clean_up();
-        ctx_http.join();
-        if (monitor_thread.joinable()) {
-            monitor_thread.join();
-        }
-
-        auto * ll_ctx = ctx_server.get_llama_context();
-        if (ll_ctx != nullptr) {
-            common_memory_breakdown_print(ll_ctx);
-        }
+    auto * ll_ctx = ctx_server.get_llama_context();
+    if (ll_ctx != nullptr) {
+        common_memory_breakdown_print(ll_ctx);
     }
 
     return 0;
 }
+
+//
+// entry points
+//
+
+static int server_run(common_params & params, int argc, char ** argv, server_role & role, const server_run_opts & opts) {
+    common_models_handler models_handler;
+
+    // -hf may name a preset repo, which makes this a router: settle that before deciding the role
+    if (opts.resolve_model && !params.model.hf_repo.empty()) {
+        try {
+            models_handler = common_models_handler_init(params, LLAMA_EXAMPLE_SERVER);
+            if (common_models_handler_is_preset_repo(models_handler)) {
+                // apply the preset and start the server in router mode
+                common_models_handler_apply(models_handler, params);
+            }
+        } catch (const std::exception & e) {
+            SRV_ERR("failed to fetch model metadata: %s\n", e.what());
+            return 1;
+        }
+    }
+
+    const server_role_type role_type = server_role_detect(params, role);
+
+    // only an instance loads a model; the others must not touch the GPU, so skip the
+    // device enumeration for them and the CUDA primary context stays uncreated
+    common_params_print_info(params, role_type == SERVER_ROLE_INSTANCE);
+
+    if (!server_prepare_params(params, role_type)) {
+        return 1;
+    }
+
+    switch (role_type) {
+        case SERVER_ROLE_DOWNLOAD: return server_run_download(params, role);
+        case SERVER_ROLE_ROUTER:   return server_run_router(params, argc, argv, opts);
+        case SERVER_ROLE_INSTANCE: return server_run_instance(params, role, models_handler, opts);
+    }
+
+    GGML_ABORT("invalid server role");
+}
+
+int llama_server(int argc, char ** argv) {
+    server_role role;
+
+    std::setlocale(LC_NUMERIC, "C");
+
+#ifndef _WIN32
+    // Ignore SIGPIPE so the server does not crash if a child (MCP server, tools runtime) exits while we are writing to its stdin
+    signal(SIGPIPE, SIG_IGN);
+#endif
+
+    // own arguments required by this example
+    common_params params;
+
+    common_init();
+
+    // start the stream session manager GC right after common init, before any HTTP route can
+    // touch it. lifecycle is symmetric, stop_gc() runs in clean_up() before backend free
+    server_stream_session_manager_start();
+
+    if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_SERVER)) {
+        return 1;
+    }
+
+    SRV_INF("%s", "initializing ...\n");
+
+    llama_backend_init();
+    llama_numa_init(params.numa);
+
+    const int result = server_run(params, argc, argv, role, server_run_opts{});
+    common_log_flush(common_log_main());
+    return result;
+}
+
+int llama_server(common_params & params, int argc, char ** argv) {
+    // embedded in llama-cli: it handles the signals, and the model is resolved already
+    server_run_opts opts;
+    opts.install_signal_handlers = false;
+    opts.resolve_model           = false;
+
+    server_role role;
+    return server_run(params, argc, argv, role, opts);
+}

tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte

@@ -477,12 +477,9 @@
 		queueMicrotask(() => inputRef?.focus());
 	}
 
-	// A flow outside the chat, e.g. the models manager closing onto it, asks for the
-	// composer back. The dialog's focus scope is still tearing down, so focus on a
-	// frame of its own.
+	// A flow outside the chat, e.g. the models page returning to it, asks for the
+	// composer back. Focus on a frame of its own, once the route has settled.
 	$effect(() => {
-		if (uiStore.manageModelsOpen) return;
-
 		if (!uiStore.composerFocusRequested) return;
 
 		uiStore.composerFocusRequested = false;

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormContextGauge/ContextGaugePopup.svelte

@@ -129,13 +129,13 @@
 {:else if gaugePopup.open}
 	<div
 		bind:this={cardEl}
-		class="absolute z-50 w-64 -translate-x-1/2 rounded-lg border border-border/50 bg-popover p-3 text-sm text-popover-foreground shadow-lg ring-1 ring-foreground/10"
+		class="absolute z-50 w-64 -translate-x-1/2 rounded-xl border border-border bg-popover p-1.5 text-sm text-popover-foreground shadow-md dark:border-border/20"
 		onpointerenter={gaugeCardEnter}
 		onpointerleave={gaugeCardLeave}
 		role="status"
 		style="left: {gaugePopup.centerX}px; bottom: {gaugePopup.bottom}px"
 	>
-		<div class="flex flex-col gap-2">
+		<div class="flex flex-col gap-2 px-2 py-1.5">
 			<div class="flex items-center gap-2">
 				<span class="font-medium">Context</span>
 

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte

@@ -111,13 +111,8 @@
 			{#if showHoverCard}
 				{@const publishDate = formatPublishDate(result.published)}
 				{@const host = hostFor(safeUrl)}
-				<HoverCard.Content
-					align="start"
-					class="bg-popover text-popover-foreground z-50 w-80 max-w-[90vw] rounded-lg border p-0 shadow-lg"
-					side="top"
-					sideOffset={6}
-				>
-					<div class="flex flex-col gap-2 p-3">
+				<HoverCard.Content align="start" class="w-80 max-w-[90vw]" side="top" sideOffset={6}>
+					<div class="flex flex-col gap-2 px-2 py-1.5">
 						<a
 							class="line-clamp-3 text-sm font-medium leading-snug hover:underline"
 							href={safeUrl}

tools/ui/src/lib/components/app/dialogs/DialogDiscoverModels.svelte

@@ -0,0 +1,53 @@
+<script lang="ts">
+	import ModelsDiscover from '$lib/components/app/models/discover/ModelsDiscover.svelte';
+	import * as Dialog from '$lib/components/ui/dialog';
+	import { DISCOVER_ICON } from '$lib/constants';
+	import { uiStore } from '$lib/stores';
+	import { untrack } from 'svelte';
+
+	interface Props {
+		open?: boolean;
+		onOpenChange?: (open: boolean) => void;
+	}
+
+	let { onOpenChange, open = $bindable(false) }: Props = $props();
+
+	// the setting can be turned off while the dialog is open: it closes and drops the
+	// repo it was asked to show
+	$effect(() => {
+		if (open && !uiStore.isDiscoverEnabled) {
+			untrack(() => {
+				uiStore.discoverModelFocus = null;
+				handleOpenChange(false);
+			});
+		}
+	});
+
+	function handleOpenChange(value: boolean) {
+		open = value;
+		onOpenChange?.(value);
+	}
+</script>
+
+<Dialog.Root onOpenChange={handleOpenChange} {open}>
+	<Dialog.Content
+		class="max-md:h-[100dvh]! max-md:w-screen! max-md:max-w-none! max-md:rounded-none! max-md:border-0! md:h-[calc(100vh-4rem)]! md:max-h-240! md:w-[calc(100vw-4rem)]! md:max-w-380! flex flex-col p-4 pb-0"
+		onOpenAutoFocus={(event) => event.preventDefault()}
+	>
+		<Dialog.Header class="flex flex-row items-center p-2 pr-8">
+			<Dialog.Title class="flex min-h-7 items-center gap-2">
+				<DISCOVER_ICON class="h-5 w-5" />
+
+				<span>Discover</span>
+			</Dialog.Title>
+
+			<Dialog.Description class="sr-only">
+				Find models on Hugging Face and download them to the server.
+			</Dialog.Description>
+		</Dialog.Header>
+
+		<div class="grid min-h-0 flex-1 overflow-hidden" style="grid-template-columns: auto 1fr;">
+			<ModelsDiscover />
+		</div>
+	</Dialog.Content>
+</Dialog.Root>

tools/ui/src/lib/components/app/dialogs/DialogManageModels.svelte

@@ -1,207 +0,0 @@
-<script lang="ts">
-	import { ArrowLeft } from '@lucide/svelte';
-	import ModelsDiscover from '$lib/components/app/models/discover/ModelsDiscover.svelte';
-	import ModelsManager from '$lib/components/app/models/ModelsManager/ModelsManager.svelte';
-	import { Button } from '$lib/components/ui/button';
-	import * as Dialog from '$lib/components/ui/dialog';
-	import {
-		DISCOVER_ICON,
-		MODEL_ICON,
-		MODELS_DIALOG_VIEW_FADE_MS,
-		SETTINGS_KEYS
-	} from '$lib/constants';
-	import { ModelsDialogView } from '$lib/enums';
-	import { modelsStore, settingsStore, uiStore } from '$lib/stores';
-	import { untrack } from 'svelte';
-
-	interface Props {
-		open?: boolean;
-		onOpenChange?: (open: boolean) => void;
-	}
-
-	let { onOpenChange, open = $bindable(false) }: Props = $props();
-
-	// what the user asked for, and what is actually rendered: a view change fades the
-	// current one out, swaps, then fades the next one in
-	let view = $state<ModelsDialogView>(ModelsDialogView.MANAGE);
-	let shownView = $state<ModelsDialogView>(ModelsDialogView.MANAGE);
-	let isSwapping = $state(false);
-	/** Set while an entry opens the dialog on an explicit view, so the open reset keeps hands off. */
-	let viewRequested = false;
-
-	// the discover view exists only while its setting is on
-	let discoverEnabled = $derived(
-		Boolean(settingsStore.getConfig(SETTINGS_KEYS.ENABLE_DISCOVER_MODELS))
-	);
-
-	$effect(() => {
-		const next = view;
-
-		if (next === shownView) return;
-
-		// nothing to fade while the dialog is closed, so the view is set before it opens
-		if (!open) {
-			untrack(() => (shownView = next));
-
-			return;
-		}
-
-		untrack(() => (isSwapping = true));
-
-		const timer = setTimeout(() => {
-			untrack(() => {
-				shownView = next;
-				isSwapping = false;
-			});
-		}, MODELS_DIALOG_VIEW_FADE_MS);
-
-		return () => clearTimeout(timer);
-	});
-
-	let title = $derived(view === ModelsDialogView.DISCOVER ? 'Discover' : 'Models');
-
-	// the sidebar's Discover entry opens this dialog on its Discover view
-	$effect(() => {
-		if (!uiStore.discoverModelsOpen) return;
-
-		untrack(() => {
-			uiStore.discoverModelsOpen = false;
-			view = discoverEnabled ? ModelsDialogView.DISCOVER : ModelsDialogView.MANAGE;
-
-			// the dialog may already be open, so no open transition will follow
-			if (open) return;
-
-			viewRequested = true;
-			handleOpenChange(true);
-		});
-	});
-
-	// opening always lands on the manager, unless an entry requested the discover view
-	$effect(() => {
-		if (!open) return;
-
-		untrack(() => {
-			if (viewRequested) {
-				viewRequested = false;
-			} else {
-				view = ModelsDialogView.MANAGE;
-			}
-
-			// sync point: the table opens on cached state and corrects itself when
-			// the list fetch lands; a no-op outside router mode
-			void modelsStore.fetchRouterModels();
-		});
-	});
-
-	// the discover view disappears with its setting; fall back to the manager
-	$effect(() => {
-		if (!discoverEnabled && view === ModelsDialogView.DISCOVER) {
-			untrack(() => (view = ModelsDialogView.MANAGE));
-		}
-	});
-
-	function handleOpenChange(value: boolean) {
-		open = value;
-		onOpenChange?.(value);
-	}
-</script>
-
-<Dialog.Root onOpenChange={handleOpenChange} {open}>
-	<Dialog.Content
-		class="max-md:h-[100dvh]! max-md:w-screen! max-md:max-w-none! max-md:rounded-none! max-md:border-0! md:h-[calc(100vh-4rem)]! md:max-h-240! md:w-[calc(100vw-4rem)]! md:max-w-380! flex flex-col p-4 pb-0"
-		onOpenAutoFocus={(event) => event.preventDefault()}
-	>
-		<Dialog.Header class="relative flex flex-row items-center justify-between p-2 pr-8">
-			<!--
-				The arrow is out of the flow and the title slides, so entering a sub-view
-				moves the title with a transform instead of reflowing the header. min-h
-				keeps the header's height fixed either way.
-			-->
-			<Button
-				aria-label="Back to models"
-				class="absolute top-1/2 left-1 h-7 w-7 -translate-y-1/2 transition-[opacity,visibility] duration-200 ease-[cubic-bezier(0.23,1,0.32,1)] {view ===
-				ModelsDialogView.MANAGE
-					? 'invisible opacity-0'
-					: 'visible opacity-100'}"
-				onclick={() => (view = ModelsDialogView.MANAGE)}
-				size="icon"
-				variant="ghost"
-			>
-				<ArrowLeft class="h-4 w-4" />
-			</Button>
-
-			<Dialog.Title
-				class="flex min-h-7 items-center gap-2 transition-transform duration-200 ease-[cubic-bezier(0.23,1,0.32,1)] {view ===
-				ModelsDialogView.MANAGE
-					? 'translate-x-0'
-					: 'translate-x-8'}"
-			>
-				<!-- the marks cross-fade in one grid cell, so the title never shifts -->
-				<span class="grid h-5 w-5 shrink-0 place-items-center">
-					<MODEL_ICON
-						class="col-start-1 row-start-1 h-5 w-5 transition-opacity duration-150 {view ===
-						ModelsDialogView.MANAGE
-							? 'opacity-100'
-							: 'opacity-0'}"
-					/>
-
-					<DISCOVER_ICON
-						class="col-start-1 row-start-1 h-5 w-5 transition-opacity duration-150 {view ===
-						ModelsDialogView.DISCOVER
-							? 'opacity-100'
-							: 'opacity-0'}"
-					/>
-				</span>
-
-				<span>{title}</span>
-			</Dialog.Title>
-
-			<Dialog.Description class="sr-only">
-				Browse, load and delete the models the server can serve.
-			</Dialog.Description>
-		</Dialog.Header>
-
-		<div class="dialog-view min-h-0 flex-1" data-visible={!isSwapping}>
-			{#if shownView === ModelsDialogView.MANAGE}
-				<ModelsManager class="h-full">
-					{#snippet toolbarEnd()}
-						{#if discoverEnabled}
-							<Button
-								class="gap-1.5"
-								onclick={() => (view = ModelsDialogView.DISCOVER)}
-								size="sm"
-								variant="tertiary"
-							>
-								<DISCOVER_ICON class="h-3.5 w-3.5" />
-
-								Discover Models
-							</Button>
-						{/if}
-					{/snippet}
-				</ModelsManager>
-			{:else if shownView === ModelsDialogView.DISCOVER}
-				<div class="grid h-full overflow-hidden" style="grid-template-columns: auto 1fr;">
-					<ModelsDiscover />
-				</div>
-			{/if}
-		</div>
-	</Dialog.Content>
-</Dialog.Root>
-
-<style>
-	/* a view change reads as one surface swapping, not as content teleporting */
-	.dialog-view {
-		opacity: 0;
-		transition: opacity 150ms cubic-bezier(0.23, 1, 0.32, 1);
-	}
-
-	.dialog-view[data-visible='true'] {
-		opacity: 1;
-	}
-
-	@media (prefers-reduced-motion: reduce) {
-		.dialog-view {
-			transition-duration: 100ms;
-		}
-	}
-</style>

tools/ui/src/lib/components/app/dialogs/index.ts

@@ -500,5 +500,5 @@ export { default as DialogMcpResourcePreview } from './DialogMcpResourcePreview.
  */
 export { default as DialogMermaidPreview } from './DialogMermaidPreview.svelte';
 
-/** **DialogManageModels** - shell for the models manager. */
-export { default as DialogManageModels } from './DialogManageModels.svelte';
+/** **DialogDiscoverModels** - shell for the Hugging Face model discovery view. */
+export { default as DialogDiscoverModels } from './DialogDiscoverModels.svelte';

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManager.svelte

@@ -10,9 +10,10 @@
 		type ModelsTableGroup,
 		splitHiddenQuants
 	} from './utils';
+	import * as Drawer from '$lib/components/ui/drawer';
 	import { LOCAL_BACKEND_ID, type ModalityKey, MODELS_TABLE_GROUP_LABELS } from '$lib/constants';
 	import { ModelCapability, ModelsTableGroupKind } from '$lib/enums';
-	import { conversationsStore, modelsStore, serverStore, uiStore } from '$lib/stores';
+	import { conversationsStore, deviceStore, modelsStore, serverStore, uiStore } from '$lib/stores';
 	import type { ModelOption } from '$lib/types/models';
 	import { filterModelOptions } from '$lib/utils';
 	import { type Snippet, untrack } from 'svelte';
@@ -48,8 +49,8 @@
 			downloads.find((group) => group.base.id === selectedId)?.base ??
 			null
 	);
-	// The pane is laid out before it is ever opened, so the first open only slides a
-	// finished panel in. It renders the selection, else the model it last showed.
+	// The pane renders the selection, else the model it last showed, so its copy stays
+	// in place while it closes.
 	let lastPicked = $state<ModelOption | null>(null);
 	let target = $derived(selected ?? lastPicked ?? allModels[0] ?? null);
 	let shownId = $state<string | null>(null);
@@ -254,10 +255,31 @@
 		await modelsStore.status.load(option.model);
 	}
 
-	/** The chat sits behind the dialog, so the dialog closes and the composer takes focus. */
+	/** Leave the models page for the chat it was opened from, with the composer focused. */
 	function returnToChat(): void {
-		uiStore.manageModelsOpen = false;
 		uiStore.requestComposerFocus();
+		void uiStore.closeModelsManager();
+	}
+
+	/** Overlays that handle their own Escape, so the page leaves them alone. */
+	const OPEN_OVERLAY_SELECTOR = [
+		'[role="dialog"][data-state="open"]',
+		'[role="alertdialog"][data-state="open"]',
+		'[role="menu"][data-state="open"]',
+		'[role="listbox"][data-state="open"]'
+	].join(', ');
+
+	// Escape steps back the way the dialog's did: the pane's drawer closes itself, then
+	// the next Escape leaves the page
+	function handleKeydown(event: KeyboardEvent): void {
+		if (event.key !== 'Escape' || event.defaultPrevented) return;
+
+		if (document.querySelector(OPEN_OVERLAY_SELECTOR)) return;
+
+		// a typed filter is not thrown away by a stray Escape
+		if (event.target instanceof HTMLInputElement && event.target.value) return;
+
+		returnToChat();
 	}
 
 	/** Load the model when the server can take load requests, without waiting for it. */
@@ -292,6 +314,8 @@
 	}
 </script>
 
+<svelte:window onkeydown={handleKeydown} />
+
 <div class={['relative flex min-h-0 flex-1', className]}>
 	<div class="min-h-0 min-w-0 flex-1">
 		<ModelsManagerModelsTable
@@ -306,16 +330,29 @@
 			{toolbarEnd}
 		/>
 	</div>
+</div>
 
-	<!-- on a phone the pane covers the whole dialog: the manager header would only
-	     repeat what the pane's own header says -->
-	<div
-		class="pane-drawer max-md:fixed max-md:inset-0 max-md:z-[60] shrink-0"
-		data-open={selected !== null}
+<!-- the pane is a drawer over the table, from the right on desktop and from the bottom
+     on a phone: Escape, the overlay and the pane's own close button clear the selection -->
+<Drawer.Root
+	direction={deviceStore.isMobile ? 'bottom' : 'right'}
+	onOpenChange={(open) => {
+		if (!open) selectedId = null;
+	}}
+	open={selected !== null}
+	shouldScaleBackground={false}
+>
+	<Drawer.Content
+		class="data-[vaul-drawer-direction=bottom]:h-[85dvh] data-[vaul-drawer-direction=bottom]:max-h-[85dvh] data-[vaul-drawer-direction=right]:w-[30rem] data-[vaul-drawer-direction=right]:max-w-[calc(100vw-2rem)] data-[vaul-drawer-direction=right]:sm:max-w-[30rem]"
 	>
-		<!-- the content box keeps the open width, so it never reflows with the drawer -->
+		<Drawer.Title class="sr-only">Model details</Drawer.Title>
+
+		<Drawer.Description class="sr-only">
+			Configuration, actions and metadata of the selected model.
+		</Drawer.Description>
+
 		<div
-			class="pane-content flex h-full min-h-0 w-[30rem] max-w-[30rem] flex-col border-l border-border/40 max-md:w-full max-md:max-w-none max-md:border-l-0 max-md:bg-background"
+			class="pane-content flex min-h-0 flex-1 flex-col"
 			data-fade={fade}
 			data-visible={selected !== null && !isSwapping}
 		>
@@ -331,33 +368,10 @@
 				{/key}
 			{/if}
 		</div>
-	</div>
-</div>
+	</Drawer.Content>
+</Drawer.Root>
 
 <style>
-	/*
-	 * The drawer moves by width because the table behind it gets that space back, so
-	 * the content box inside holds the open width and only the container changes.
-	 * Opening takes the iOS-like drawer curve; closing is the system responding, so
-	 * it snaps back on the stronger ease-out.
-	 */
-	.pane-drawer {
-		width: 0;
-		overflow: clip;
-		visibility: hidden;
-		transition:
-			width 120ms cubic-bezier(0.23, 1, 0.32, 1),
-			visibility 120ms;
-	}
-
-	.pane-drawer[data-open='true'] {
-		width: 30rem;
-		visibility: visible;
-		transition:
-			width 200ms cubic-bezier(0.32, 0.72, 0, 1),
-			visibility 200ms;
-	}
-
 	.pane-content {
 		opacity: 0;
 		transition: opacity 120ms cubic-bezier(0.23, 1, 0.32, 1);
@@ -376,11 +390,4 @@
 	.pane-content[data-visible='true'][data-fade='swap'] {
 		transition: opacity 120ms cubic-bezier(0.23, 1, 0.32, 1);
 	}
-
-	/* a phone has no room beside the table: the pane covers it instead */
-	@media (max-width: 767px) {
-		.pane-drawer[data-open='true'] {
-			width: auto;
-		}
-	}
 </style>

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerFilters.svelte

@@ -5,6 +5,7 @@
 	import {
 		CAPABILITY_ICONS,
 		CAPABILITY_LABELS,
+		FILTER_TOGGLE_GROUP_CLASS,
 		FILTER_TOGGLE_ITEM_CLASS,
 		FILTER_TRIGGER_CLASS,
 		MODALITY_FLAG_KEYS,
@@ -89,10 +90,12 @@
 </script>
 
 <!-- below md the carousel keeps a row of its own: a toolbar that also holds a call to
-     action would otherwise squeeze the filters out of sight -->
+     action would otherwise squeeze the filters out of sight. The scroll box clips on
+     both axes, so it gets room for the controls' shadows and gives it back outside -->
 <ScrollCarousel
 	alwaysShowArrows
 	class="min-w-0 flex-1 max-md:w-full max-md:flex-none"
+	containerClass="-m-1 p-1"
 	gapSize="2"
 	innerClass="items-center"
 >
@@ -115,11 +118,10 @@
 	</Select.Root>
 
 	<ToggleGroup.Root
-		class="border border-border/30 bg-muted/60 shadow-sm dark:border-border/20 dark:bg-muted/75"
+		class={FILTER_TOGGLE_GROUP_CLASS}
 		onValueChange={setToggles}
 		type="multiple"
 		value={[...capabilities, ...modalities]}
-		variant="outline"
 	>
 		{#each toggles as toggle (toggle.value)}
 			<ToggleGroup.Item

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte

@@ -98,7 +98,7 @@
 		{status}
 	/>
 
-	<div class="min-h-0 flex-1 overflow-y-auto py-4 pl-4 max-md:pr-4">
+	<div class="min-h-0 flex-1 overflow-y-auto p-4">
 		{#if isDownload}
 			<!-- nothing is on disk yet, so the pane keeps to the download itself -->
 			<div class="space-y-2 py-1">

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationActions.svelte

@@ -11,8 +11,7 @@
 		Trash2
 	} from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
-	import { SETTINGS_KEYS } from '$lib/constants';
-	import { conversationsStore, deviceStore, settingsStore } from '$lib/stores';
+	import { conversationsStore, deviceStore, uiStore } from '$lib/stores';
 
 	interface Props {
 		/** A provider that cannot load or unload has no button to offer. */
@@ -48,10 +47,7 @@
 		size = 'sm'
 	}: Props = $props();
 
-	// the discover view exists only while its setting is on
-	let discoverEnabled = $derived(
-		Boolean(settingsStore.getConfig(SETTINGS_KEYS.ENABLE_DISCOVER_MODELS))
-	);
+	let discoverEnabled = $derived(uiStore.isDiscoverEnabled);
 
 	// a chat with anything in it takes the model, a new chat is not what its user asked for
 	let hasChat = $derived(

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationHeader.svelte

@@ -74,7 +74,7 @@
 <!-- only the desktop pane leaves the right edge to the panel border: on a phone the
      pane covers the screen, so its two sides carry the same padding, and it takes a
      little more room around the model it names -->
-<header class="space-y-2.5 pt-3 pl-4 max-md:space-y-4 max-md:pt-4 max-md:pr-4 max-md:pb-4">
+<header class="space-y-2.5 px-4 pt-3 max-md:space-y-4 max-md:pt-4 max-md:pb-4">
 	<div class="flex items-start justify-between gap-3">
 		<div class="flex min-w-0 items-center gap-3">
 			<!-- same geometry as the discover details header: base org, quant org badge -->
@@ -108,15 +108,18 @@
 			</div>
 		</div>
 
-		<Button
-			aria-label="Close details"
-			class="h-7 w-7 max-md:-mt-3 max-md:-mr-3 max-md:h-10 max-md:w-10 max-md:rounded-full"
-			onclick={onClose}
-			size="icon"
-			variant="ghost"
-		>
-			<X class="h-4 w-4" />
-		</Button>
+		<!-- a phone's bottom drawer closes by its handle and overlay, it needs no button -->
+		{#if !isMobile}
+			<Button
+				aria-label="Close details"
+				class="h-7 w-7"
+				onclick={onClose}
+				size="icon"
+				variant="ghost"
+			>
+				<X class="h-4 w-4" />
+			</Button>
+		{/if}
 	</div>
 
 	<!-- a phone floats the actions at the bottom of the pane, where the thumb is -->

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelRow.svelte

@@ -7,7 +7,7 @@
 	import ModelId from '../ModelId.svelte';
 	import ModelsManagerStatusCell from './ModelsManagerStatusCell.svelte';
 	import { modelRowActions } from './row-actions';
-	import { configuredContext, downloadProgressFor } from './utils';
+	import { configuredContext, downloadProgressFor, isRowControlClick } from './utils';
 	import { Heart, MoreHorizontal } from '@lucide/svelte';
 	import { DropdownMenuActions } from '$lib/components/app';
 	import { MODEL_ROW_GRID_CLASS, MODEL_ROW_TRAILING_CELL_CLASS } from '$lib/constants';
@@ -55,13 +55,20 @@
 	}
 </script>
 
+<!-- the whole row opens the model for a pointer; keyboard users reach the same action
+     through the model's own button -->
+<!-- svelte-ignore a11y_click_events_have_key_events -->
+<!-- svelte-ignore a11y_no_static_element_interactions -->
 <div
 	class={[
 		MODEL_ROW_GRID_CLASS,
-		'group relative rounded-md px-2 py-3 transition max-md:px-3 max-md:py-4',
+		'group relative cursor-pointer rounded-md px-2 py-3 transition max-md:px-3 max-md:py-4',
 		isHidden && 'opacity-60',
 		selected ? 'bg-accent text-accent-foreground' : 'hover:bg-muted/40'
 	]}
+	onclick={(event) => {
+		if (!isRowControlClick(event)) activate();
+	}}
 >
 	<div class="flex min-w-0 flex-col">
 		<!-- the row holds a load control and an actions menu, so only the model itself is

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelsTable.svelte

@@ -29,9 +29,11 @@
 		ModelsSection
 	} from '$lib/components/app';
 	import { DialogConfirmDownload } from '$lib/components/app/dialogs';
+	import * as Empty from '$lib/components/ui/empty';
 	import {
 		FAMILY_ROW_WINDOW,
 		type ModalityKey,
+		MODEL_ICON,
 		MODEL_ROW_GRID_CLASS,
 		MODEL_ROW_TRAILING_CELL_CLASS,
 		MODEL_ROW_WINDOW,
@@ -207,7 +209,7 @@
 
 {#snippet sortHeader(key: ModelsTableSortKey, label: string)}
 	<button
-		class="inline-flex cursor-pointer items-center gap-1 uppercase transition hover:text-foreground focus:outline-none"
+		class="inline-flex cursor-pointer items-center gap-1 transition-colors hover:text-foreground focus:outline-none"
 		onclick={() => toggleSort(key)}
 		title={sortTitle(key, label)}
 		type="button"
@@ -270,7 +272,7 @@
 	{@const countLabel = `${family.entries.length} model${family.entries.length === 1 ? '' : 's'}`}
 
 	<div
-		class="{MODEL_ROW_GRID_CLASS} relative group cursor-pointer rounded-md px-2 py-1 transition hover:bg-muted/40 max-md:px-3 max-md:py-2.5"
+		class="{MODEL_ROW_GRID_CLASS} relative group cursor-pointer rounded-md px-2 py-1.5 transition hover:bg-muted/40 max-md:px-3 max-md:py-2.5"
 		onclick={toggle}
 		onkeydown={(event) => handleFamilyKeydown(event, toggle)}
 		role="button"
@@ -286,7 +288,7 @@
 
 			<span class="truncate text-sm font-medium">{family.label}</span>
 
-			<span class="text-sm text-muted-foreground">{countLabel}</span>
+			<span class="shrink-0 text-sm whitespace-nowrap text-muted-foreground">{countLabel}</span>
 		</span>
 
 		<!-- the context and status columns stay empty here: the fold control is this
@@ -350,7 +352,7 @@
 	/>
 
 	<div
-		class="{MODEL_ROW_GRID_CLASS} shrink-0 border-y border-border/40 px-2 py-2 text-[11px] font-semibold tracking-wide text-muted-foreground uppercase max-md:px-3"
+		class="{MODEL_ROW_GRID_CLASS} shrink-0 border-b border-border/30 px-2 pb-2 text-xs font-medium text-muted-foreground max-md:px-3 dark:border-border/20"
 	>
 		<span>{@render sortHeader(ModelsTableSortKey.NAME, 'Model')}</span>
 
@@ -363,7 +365,7 @@
 		<span class="text-center max-md:hidden">Actions</span>
 	</div>
 
-	<div class="min-h-0 flex-1 overflow-y-auto">
+	<div class="mt-2 min-h-0 flex-1 overflow-y-auto">
 		{#each sections as group (group.key)}
 			{#if group.items.length > 0}
 				{#snippet groupIcon()}
@@ -388,34 +390,49 @@
 					label={group.label}
 					persistKey={group.key}
 					revealChevronOnHover
-					sectionHeaderClass="m-0 px-2 py-2 text-[13px] font-semibold text-muted-foreground select-none max-md:px-3"
+					sectionHeaderClass="m-0 rounded-lg px-2 py-2 text-[13px] font-semibold text-foreground/85 select-none transition-colors hover:bg-muted max-md:px-3"
 					sticky
-					stickyClass="sticky z-10 bg-muted/90 backdrop-blur-lg"
+					stickyClass="sticky z-10 bg-muted/70 backdrop-blur-lg"
 				>
-					<GroupedList
-						group={familyRow}
-						groupWindow={FAMILY_ROW_WINDOW}
-						groups={group.flat ? null : familyGroups(group)}
-						initialCollapsed={modelsStore.collapsedGroupsUnder(group.key)}
-						item={listItem}
-						items={group.flat ? group.items : []}
-						keyOf={(entry) => entry.key}
-						more={showMore}
-						onCollapsedChange={(key, collapsed) =>
-							modelsStore.setGroupCollapsed(group.key, key, collapsed)}
-						sectionWindow={MODEL_ROW_WINDOW}
-						stickyClass="bg-muted/30 backdrop-blur-lg"
-						stickyStyle={MODELS_TABLE_GROUP_STICKY_OFFSET}
-						weightOf={(entry) => entry.quants.length}
-					/>
+					<!-- a little air before the next section's band -->
+					<div class="pb-3">
+						<GroupedList
+							group={familyRow}
+							groupWindow={FAMILY_ROW_WINDOW}
+							groups={group.flat ? null : familyGroups(group)}
+							initialCollapsed={modelsStore.collapsedGroupsUnder(group.key)}
+							item={listItem}
+							items={group.flat ? group.items : []}
+							keyOf={(entry) => entry.key}
+							more={showMore}
+							onCollapsedChange={(key, collapsed) =>
+								modelsStore.setGroupCollapsed(group.key, key, collapsed)}
+							sectionWindow={MODEL_ROW_WINDOW}
+							stickyClass="rounded-lg bg-muted/25 backdrop-blur-lg"
+							stickyStyle={MODELS_TABLE_GROUP_STICKY_OFFSET}
+							weightOf={(entry) => entry.quants.length}
+						/>
+					</div>
 				</ModelsSection>
 			{/if}
 		{/each}
 
 		{#if isEmpty}
-			<p class="px-4 py-10 text-center text-sm text-muted-foreground">
-				{hasFilters ? 'No models match these filters.' : 'No models found.'}
-			</p>
+			<Empty.Root class="py-16">
+				<Empty.Header>
+					<Empty.Media variant="icon">
+						<MODEL_ICON />
+					</Empty.Media>
+
+					<Empty.Title>{hasFilters ? 'No matching models' : 'No models yet'}</Empty.Title>
+
+					<Empty.Description>
+						{hasFilters
+							? 'Try a different search or clear the filters.'
+							: 'Models the server can run show up here.'}
+					</Empty.Description>
+				</Empty.Header>
+			</Empty.Root>
 		{/if}
 	</div>
 </div>

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerQuantRow.svelte

@@ -3,7 +3,7 @@
 	import ModelDraftSidecars from '../ModelDraftSidecars.svelte';
 	import ModelsManagerStatusCell from './ModelsManagerStatusCell.svelte';
 	import { modelRowActions } from './row-actions';
-	import { configuredContext } from './utils';
+	import { configuredContext, isRowControlClick } from './utils';
 	import { MoreHorizontal } from '@lucide/svelte';
 	import { DropdownMenuActions } from '$lib/components/app';
 	import { Badge } from '$lib/components/ui/badge';
@@ -28,13 +28,20 @@
 	let quant = $derived(option.parsedId?.quantization ?? option.model);
 </script>
 
+<!-- the whole row opens the quant for a pointer; keyboard users reach the same action
+     through the quant's own button -->
+<!-- svelte-ignore a11y_click_events_have_key_events -->
+<!-- svelte-ignore a11y_no_static_element_interactions -->
 <div
 	class={[
 		MODEL_ROW_GRID_CLASS,
-		'group relative rounded-md px-2 py-2 transition max-md:px-3 max-md:py-3',
+		'group relative cursor-pointer rounded-md px-2 py-2 transition max-md:px-3 max-md:py-3',
 		isHidden && 'opacity-60',
 		selected ? 'bg-accent text-accent-foreground' : 'hover:bg-muted/40'
 	]}
+	onclick={(event) => {
+		if (!isRowControlClick(event)) onSelect(option);
+	}}
 >
 	<!-- the row holds a load control and an actions menu, so only the quant itself is
 	     the button: a button nested in a role="button" row is invalid -->

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerTableToolbar.svelte

@@ -6,7 +6,7 @@
 	import { Button } from '$lib/components/ui/button';
 	import { type ModalityKey } from '$lib/constants';
 	import { ModelCapability } from '$lib/enums';
-	import { deviceStore, uiStore } from '$lib/stores';
+	import { deviceStore } from '$lib/stores';
 	import type { Snippet } from 'svelte';
 
 	interface Props {
@@ -32,12 +32,10 @@
 	let hasFilters = $derived(hasActiveFilters(contextLimit, modalities, capabilities));
 	let filterInput = $state<HTMLInputElement | null>(null);
 
-	// The search takes the focus the dialog would give its first control: the input
-	// mounts a frame or two after the dialog opens (the toolbar renders below the
-	// table in the DOM), so the poll retries until the ref lands or gives up.
+	// The search takes the focus when the models page opens: the input mounts a frame
+	// or two after the page (the toolbar renders below the table in the DOM), so the
+	// poll retries until the ref lands or gives up.
 	$effect(() => {
-		if (!uiStore.manageModelsOpen) return;
-
 		let frames = 0;
 		let handle = requestAnimationFrame(function focusFilter() {
 			if (filterInput) {

tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts

@@ -13,13 +13,18 @@ export function modelRowActions(
 	/** Download state, when the row stands for a tracked download. */
 	download?: ModelRowDownloadState | null
 ) {
-	const viewInDiscover = {
-		icon: Compass,
-		label: 'See in Discover',
-		// the details pane is keyed by repo, not by `<repo>:<quant>`
-		onclick: () => uiStore.openModelsDiscover(repoOf(option.model)),
-		separator: true
-	};
+	// the details pane is keyed by repo, not by `<repo>:<quant>`; the entry exists only
+	// while Discover is on
+	const viewInDiscover = uiStore.isDiscoverEnabled
+		? [
+				{
+					icon: Compass,
+					label: 'See in Discover',
+					onclick: () => uiStore.openModelsDiscover(repoOf(option.model)),
+					separator: true
+				}
+			]
+		: [];
 
 	// a tracked download is paused, resumed or dropped: it is not loaded, hidden or
 	// drafted. An in-flight row cancels, which stops the download and removes the
@@ -40,7 +45,7 @@ export function modelRowActions(
 				separator: true,
 				variant: 'destructive' as const
 			},
-			viewInDiscover
+			...viewInDiscover
 		];
 	}
 
@@ -63,6 +68,6 @@ export function modelRowActions(
 			onclick: () => modelsStore.toggleHidden(option.id),
 			separator: true
 		},
-		viewInDiscover
+		...viewInDiscover
 	];
 }

tools/ui/src/lib/components/app/models/ModelsManager/utils.ts

@@ -81,6 +81,14 @@ export interface ModelsTableGroup {
 	label: string;
 }
 
+/** Controls in a model row that act on their own, so a row click passes them by. */
+const ROW_CONTROL_SELECTOR = 'button, a, input, [role="menuitem"]';
+
+/** Whether a click on a model row landed on one of its own controls. */
+export function isRowControlClick(event: MouseEvent): boolean {
+	return event.target instanceof Element && event.target.closest(ROW_CONTROL_SELECTOR) !== null;
+}
+
 /** Context the model runs with: what a loaded model reports. */
 export function configuredContext(option: ModelOption): number | null {
 	return modelsStore.isModelRunning(option.model)

tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDropdown.svelte

@@ -91,7 +91,7 @@
 	function handleManageModels() {
 		open = false;
 
-		// let the menu finish closing before the dialog takes focus
+		// let the menu finish closing before the models page takes focus
 		setTimeout(() => (onManageModels ? onManageModels() : uiStore.openModelsManager()), 0);
 	}
 </script>

tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte

@@ -5,7 +5,7 @@
 	import {
 		ActionIcon,
 		DialogConversationRename,
-		DialogManageModels,
+		DialogDiscoverModels,
 		DialogMcpServers,
 		DialogModelInformation,
 		DialogSettingsChat,
@@ -469,7 +469,7 @@
 
 <DialogMcpServers bind:open={mcpDialogOpen} />
 
-<DialogManageModels bind:open={uiStore.manageModelsOpen} />
+<DialogDiscoverModels bind:open={uiStore.discoverModelsOpen} />
 
 <DialogModelInformation />
 

tools/ui/src/lib/components/ui/hover-card/hover-card-content.svelte

@@ -21,7 +21,7 @@
 		bind:ref
 		{align}
 		class={cn(
-			'data-open:animate-in data-closed:animate-out data-closed:fade-out-0 data-open:fade-in-0 data-closed:zoom-out-95 data-open:zoom-in-95 data-[side=bottom]:slide-in-from-top-2 data-[side=left]:slide-in-from-right-2 data-[side=right]:slide-in-from-left-2 data-[side=top]:slide-in-from-bottom-2 ring-foreground/10 bg-popover text-popover-foreground w-64 rounded-lg p-2.5 text-sm shadow-md ring-1 duration-100 z-50 origin-(--transform-origin) outline-hidden',
+			'z-50 w-64 origin-(--bits-link-preview-content-transform-origin) rounded-xl border border-border bg-popover p-1.5 text-sm text-popover-foreground shadow-md outline-hidden data-[side=bottom]:slide-in-from-top-2 data-[side=left]:slide-in-from-right-2 data-[side=right]:slide-in-from-left-2 data-[side=top]:slide-in-from-bottom-2 data-[state=closed]:animate-out data-[state=closed]:fade-out-0 data-[state=closed]:fill-mode-forwards data-[state=closed]:zoom-out-95 data-[state=open]:animate-in data-[state=open]:fade-in-0 data-[state=open]:zoom-in-95 dark:border-border/20',
 			className
 		)}
 		data-slot="hover-card-content"

tools/ui/src/lib/constants/css-classes.constants.ts

@@ -54,9 +54,17 @@ export const FILTER_TRIGGER_CLASS = `
     ${INPUT_CLASSES}
 `;
 
-/** One filter toggle of the models manager toolbar. */
+/** The filter toggles' shared box: the same fill, border and shadow as a filter trigger. */
+export const FILTER_TOGGLE_GROUP_CLASS = `
+    h-8
+    overflow-hidden
+    divide-x divide-border/30 dark:divide-border/20
+    ${INPUT_CLASSES}
+`;
+
+/** One filter toggle of the models manager toolbar; its group draws the box. */
 export const FILTER_TOGGLE_ITEM_CLASS =
-	'bg-muted! border-border/30! shadow-none! dark:border-border/20! data-[state=on]:bg-muted-foreground/15! data-[state=on]:text-foreground! dark:data-[state=on]:bg-muted-foreground/25!';
+	'h-full rounded-none hover:bg-muted/80 dark:hover:bg-muted data-[state=on]:bg-muted-foreground/15 data-[state=on]:text-foreground dark:data-[state=on]:bg-muted-foreground/25';
 
 /** Column grid shared by every row of the models manager table. Below md the context,
  *  status and actions columns are dropped, so a phone gives the whole width to the

tools/ui/src/lib/constants/routes.constants.ts

@@ -13,8 +13,13 @@ export const ROUTES = {
 	CHAT: '#/chat',
 	/** MCP servers. */
 	MCP_SERVERS: '#/mcp-servers',
+	/** Models manager. */
+	MODELS: '#/models',
 	/** Search — mobile-only full-page conversation search. */
 	SEARCH: '#/search',
 	/** Root — start of the app. */
 	START: '#/'
 } as const;
+
+/** SvelteKit route id of the models manager, as `page.route.id` reports it. */
+export const MODELS_ROUTE_ID = '/models';

tools/ui/src/lib/constants/ui.constants.ts

@@ -1,4 +1,5 @@
 import { MODEL_ICON } from './icons.constants';
+import { MODELS_ROUTE_ID } from './routes.constants';
 import { Search, Settings, SquarePen } from '@lucide/svelte';
 import McpLogo from '$lib/components/app/mcp/McpLogo.svelte';
 import { SidebarAction, ToolSource } from '$lib/enums';
@@ -7,9 +8,6 @@ import type { DesktopIconStripItem } from '$lib/types';
 /** Delay before a copied affordance resets from its check to the copy icon. */
 export const COPY_FEEDBACK_RESET_MS = 1500;
 
-/** How long a models dialog view fades out before the next one fades in. */
-export const MODELS_DIALOG_VIEW_FADE_MS = 150;
-
 export const FORK_TREE_DEPTH_PADDING = 8;
 export const SYSTEM_MESSAGE_PLACEHOLDER = 'System message';
 
@@ -83,6 +81,7 @@ export const SIDEBAR_ACTIONS_ITEMS: DesktopIconStripItem[] = [
 	{ icon: Search, keys: ['cmd', 'k'], tooltip: 'Search' },
 	{
 		action: SidebarAction.MANAGE_MODELS,
+		activeRouteId: MODELS_ROUTE_ID,
 		icon: MODEL_ICON,
 		tooltip: 'Models'
 	},