Compare commits

..
18 Commits
Author SHA1 Message Date
Xuan-Son Nguyen f3269ef96a server: refactor server entrypoint (#31113)
* server_role

* cont
2026-10-24 20:04:28 +02:00
Georgi Gerganov 3f03b17724 metal : remove mul_mm_id amax rescaling (#30924)
The src1 power-of-two rescale in kernel_mul_mm_id is superseded by the
per expert-token amax scaling done in the compute graph
(ref: https://github.com/ggml-org/llama.cpp/pull/30861), which keeps the
MoE down-projection activations within f16 range before the op runs.

Drop the two-stage reduction kernels, their scratch buffer and the
GGML_PREC_F32 dispatch added in https://github.com/ggml-org/llama.cpp/pull/26223,
as well as the test cases that drive MUL_MAT_ID src1 out of f16 range.

Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
2026-10-24 18:40:19 +03:00
Katara 9a0fde9967 rpc : add bounds check in rdma_recv (#30399) 2026-10-24 18:34:46 +03:00
Aleksander Grygier e168e041ec ui: Models Manager as Route + UI improvements (#30988)
* ui : move models manager to its own route, keep Discover as a dialog

Models now lives at #/models instead of a dialog. Discover opens as a
standalone dialog over it, and leaving the page returns to the chat it
was opened from.

Assisted by: Claude (Opus 5.5 Medium)

* ui : polish the models page layout and filters

Give the models page a page header with the Discover action, calmer
column labels, tinted section and family headers, and an empty state.
Load the model list on a direct visit, slide the detail pane over the
table below xl, and draw the filter toggles as one box matching the
context select, with room for their shadow.

Assisted by: Claude (Opus 5.5 Medium)

* ui : round the models table family headers

Assisted by: Claude (Opus 5.5 Medium)

* ui : fix models page navigation and Discover gating

Record the models page's way back after every navigation, so Back,
a typed URL or a reload no longer return to a stale chat. Escape
closes the detail pane, then leaves the page. Gate every Discover
entry point on one shared check, recheck the API key on /models,
skip the duplicate router fetch on a direct load, and share the
route load between pages.

Assisted by: Claude (Opus 5.5 Medium)

* ui : show model details in a drawer

Open the models manager's detail pane in a drawer, from the right on
desktop and from the bottom on a phone, instead of an inline side
panel and a full-screen cover. Pad the pane on both sides, drop its
close button on a phone, and assert that filtered rows are gone from
the page rather than merely hidden.

Assisted by: Claude (Opus 5.5 Medium)

* ui : make whole models table rows clickable

A click anywhere on a model or quant row opens it, unless it lands on
one of the row's own controls (load toggle, actions menu). The name
button stays the keyboard path to the same action.

Assisted by: Claude (Opus 5.5 Medium)
2026-10-24 18:34:11 +03:00
Aleksander Grygier ea1d7d66f3 fix(ui): align hover card and context popup styling with dropdown menu (#30978)
Use the dropdown menu surface (border, radius, background, shadow,
padding) for the hover card and the context gauge popup, and fix the
hover card open/close animation to use bits-ui data-state attributes.

Assisted-by: Claude (Opus 5.5 Medium)
2026-10-24 18:33:27 +03:00
Georgi Gerganov 71255fefc0 ci : pass per-runner results dir to self-hosted jobs on T4 runners (#31100)
The two T4 runners in the ggml-ci group (ggml-6-x86-vulkan-t4,
ggml-8-x86-nvidia-a10) share the same labels and can both pick up the
same jobs, so extend the per-runner results dir from #31046 to them.

Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
2026-10-24 17:17:04 +03:00
Satyam Thakurandorigin-labs-ai 73adf89e38 minimax-m3 : forward quant scales to the graph (fixes #31063) (#31097)
The generic scale pass in llama-model.cpp loads the per-expert
NVFP4 scale2 tensors (ffn_{up,gate,down}_exps_s) and the per-tensor
MXFP8 scales for the dense FFN and the shared expert, but
minimax-m3.cpp never forwarded any of them to the graph.

build_moe_ffn() defaults these to nullptr, so build_lora_mm_id()
skipped its w_s branch entirely and the routed-expert weights came
out global_scale times too large, swamping the residual (reported
G ~= 15123 for M3.1 Flash) and producing token soup.

Forward all scale tensors at the three call sites, matching the
pattern used by flint and the other NVFP4-capable MoE models.

Co-authored-by: origin-labs-ai <origin-labs-ai@users.noreply.github.com>
2026-10-24 16:09:02 +02:00
Alexey Kopytko df9beac715 tests : fix test-llama-context for large n_batch (#31083) 2026-10-24 14:14:58 +03:00
Georgi Gerganov 4d32b29aaa ci : exclude test-speculative-rejection-sampling from macos-latest-arm64 (#31088)
Add test-speculative-rejection-sampling to the ctest -E exclusion list of the macos-latest-arm64 CI job.

Assisted-by: dsh:llama.cpp/ggml-org/Flint-GGUF (RTX 5090 32GB, DDR5 128GB)
2026-10-24 14:14:31 +03:00
Georgi Gerganov 3032209758 llama : deny -sm tensor for arches with moe amax rescaling (#31081)
The amax POOL_1D reduction spans the whole matmul row, so a device cannot
compute it from an axis-0 split. The Meta split-state case reported
MIRRORED, which gives each device a partial max of its own slice. Drop
that case and deny -sm tensor for the arches which build the amax graph,
see [TAG_MOE_AMAX_SCALE].

Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
2026-10-24 13:48:32 +03:00
Xuan-Son Nguyenandyomaytk 701fd25b3a ggml : allow MUL_MAT_ID -1 index to skip the computation (#26631)
* ggml : allow MUL_MAT_ID -1 index to skip the computation

* nits

* update vulkan code

* ggml-webgpu: ensure all threads reach workgroup barriers in mul_mat_id_vec

(cherry picked from commit d8a4b9efb6e2b9abf3a866e847756b7e6bbc6b69)

* fix cuda

* openvino

* exclude CUDA from this PR

* add TODO tag

---------

Co-authored-by: yomaytk <yoshimura.masashi.frbs@gmail.com>
2026-10-24 12:15:36 +02:00
Daniel Bevenius 3642238cc1 tests : add speculative rejection draft sampling test (#30213)
* tests : add speculative rejection draft sampling test

This commit adds a test for speculative rejection sampling test that
excercises the speculative rejection sampling mechanism in
common/speculative.cpp.

The motivation for this is have a test that verifies that the
speculative draft models emitts tokens that  match the target models
distribution. So when rejection sampling occurs it should still
produce the same probability distribution as if we directly sampled
from the target model.

Running the test:
```console
$ ctest --test-dir build -R "^test-speculative-rejection-sampling$" \
    --n_trials 5000 -V
```
This will generate two models, one target and one draft model.

For subsequent runs use the "fixture exclude any" option to avoid
recreating the models:
```console
$ ctest --test-dir build-cpu-debug -R \
"^test-speculative-rejection-sampling$" \
    -FA "^speculative-models$" --n_trails 5000 -V
```

* remove draft model
2026-10-24 12:47:49 +03:00
Georgi Gerganov 6b9ba4f2e8 ci : fix data race between runners sharing the model cache (#31046)
* ci : fix data race between runners sharing the model cache

Two runners per mac mini share HOME, so concurrent jobs corrupt each
other downloads, converted models and venv under the shared MNT cache.

- gg_wget: download in a temp dir, publish with mv, so readers only
  ever see complete files (-N cannot be combined with -O)
- convert, quantize and wikitext unzip now produce under a per-runner
  scratch dir in OUT, published to MNT via hardlink when missing
- gguf-split and quantize tests get a per-runner TMP_DIR
- python venv moved from MNT to OUT
- metal, vulkan and webgpu mac jobs pass a per-runner OUT path

Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)

* ci : clean up scratch models after tests

Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)

* ci : fix q4_0 publish when a cached model exists

cp -n skips an existing file with exit code 1, so the ln fallback
aborted the qwen3 job under set -e. Publish through a unique temp
file with an atomic mv: readers never see a partial file and a
stale cached q4_0 gets replaced.

Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)

* ci : cache-first gg_wget to avoid re-downloads

HF sends no Last-Modified, so wget -N re-downloads the full files on
every run. Skip the download when a non-empty cache file exists: the
HF CI artifacts are static. Downloads still go through a temp dir and
an atomic mv, so concurrent runners never read partial files.

Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)
2026-10-24 09:21:24 +03:00
Ruben Ortlam 8e5a58df87 vulkan: ggml src1 precision support (#30656)
* vulkan: update accumulator precision selection

* handle fp32 overflows through scaling

* handle overflows in copy-to-fp16 activations

* add missing cases

* simplify code

* add tests

* add ggml prec get helpers, use in Vulkan backend

* clean up comments

* revert comment edit

* trim tests

* also handle fp32 src0, fix test scaling src0 values too

* don't scale if disabled

* remove fp16 scaling

* add check to mul_mat

* fix inconsistency

* rename to acc_prec
2026-10-24 06:05:05 +02:00
Arseniy Obolenskiy 1ddffe9f70 HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable (#30915)
* HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable

amdgcnspirv rejects target-specific builtins unless they are guarded by `__builtin_amdgcn_is_invocable`, so 66 of 145 HIP files failed to build

Wrap the perm calls in a helper that keeps v_perm_b32 when the target has it and falls back to a portable byte select otherwise

* move perm helper to vendors/hip.h
2026-10-24 12:04:05 +08:00
Xuan-Son Nguyen 7894e8b924 mtmd: move build_inp_temporal_merge to base class (#31065) 2026-10-24 03:38:25 +02:00
Yuri KhrustalevandXuan Son Nguyen edc353842c mtmd : lfm2 resize matches torchvision (#30943)
* mtmd : lfm2 bilinear resize matches torchvision

Lfm2VlImageProcessor and d1-omni vision.py resize uint8 images with
torchvision (antialias=True). That is Pillow's algorithm, except that the
weights are int16 with as many fractional bits as the largest weight allows,
where Pillow uses 22 bits. The Pillow port was 1-2 levels off on about 1% of
the values of the d1-omni and LFM2.5-VL crops.

RESIZE_ALGO_BILINEAR_TORCH rounds the weights of the Pillow port to that
precision and scales them back to 22 bits, which gives the output of
torchvision with the same passes, and "bilinear" maps to it for LFM2 and
d1-omni. Other models keep the Pillow algorithms.

* mtmd : lfm2 bicubic resize matches torchvision

Lfm2VlImageProcessor resizes with torchvision for the bicubic checkpoints too
(d1-3B, LFM2-VL), so "bicubic" maps to RESIZE_ALGO_BICUBIC_TORCH for LFM2: the
Pillow bicubic filter (a = -0.5, as torchvision with antialias) with the int16
weights of torchvision. The Pillow port was 1-2 levels off on about 0.25% of
the values of the d1-3B crops.

* use bitfield, add hparams setter

---------

Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
2026-10-24 02:21:06 +02:00
Hans 9ae73e68f5 mtmd : add Soprano TTS support (#28607)
Add Soprano-80M and Soprano-1.1-80M (https://huggingface.co/ekwek/Soprano-1.1-80M)
as an audio generation model: a Qwen3 backbone plus a Vocos decoder that runs
as a gen-only mmproj.

- convert: detect Soprano checkpoints and export the backbone as qwen3 and the
  decoder.pth Vocos weights as the mmproj
- vocab: add the soprano pre-tokenizer (lowercase, whitespace collapse,
  per-digit split, [UNK] fallback)
- mtmd: add the Vocos graph and ISTFT, allow loading a gen-only mmproj for
  Soprano and include the gen-audio context in mtmd_get_memory_usage
- mtmd-helper: add the Soprano pipeline; the decoder runs over the whole
  utterance, so stream mode sends the audio in one chunk after flush()
- tts: document Soprano usage

Assisted-by: Codex
Assisted-by: Claude Opus 5.5
2026-10-24 01:09:50 +02:00
127 changed files with 2801 additions and 1380 deletions
+1 -1
View File
@@ -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
+1 -1
View File
@@ -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
+1 -1
View File
@@ -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
+3 -3
View File
@@ -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]
+1 -1
View File
@@ -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
+88 -91
View File
@@ -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
+2
View File
@@ -286,6 +286,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
"SeedOssForCausalLM": "olmo",
"SmallThinkerForCausalLM": "smallthinker",
"SmolLM3ForCausalLM": "llama",
"SopranoTextModel": "soprano",
"Spark2_5ForCausalLM": "spark2_5",
"SolarOpenForCausalLM": "glm",
"StableLMEpochForCausalLM": "stablelm",
@@ -380,6 +381,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
"Qwen3ASRForConditionalGeneration": "qwen3vl",
"Qwen3OmniMoeForConditionalGeneration": "qwen3vl",
"PocketTTSModel": "pockettts",
"SopranoModel": "soprano",
"Qwen3TTSForConditionalGeneration": "qwen3tts",
"Qwen3VLForConditionalGeneration": "qwen3vl",
"Qwen3VLMoeForConditionalGeneration": "qwen3vl",
+116
View File
@@ -0,0 +1,116 @@
from __future__ import annotations
import json
from pathlib import Path
from typing import Any
import torch
from .base import ModelBase, MmprojModel, gguf, logger
from .qwen import Qwen3Model
def _is_soprano_checkpoint(dir_model: Path) -> bool:
if not (dir_model / "decoder.pth").is_file() or not (dir_model / "config.json").is_file():
return False
with open(dir_model / "config.json", encoding="utf-8") as f:
config = json.load(f)
if config.get("model_type") != "qwen3" or config.get("bos_token_id") != 3 or config.get("eos_token_id") != 3:
return False
if not (dir_model / "tokenizer.json").is_file():
return False
with open(dir_model / "tokenizer.json", encoding="utf-8") as f:
added = {t["id"]: t["content"] for t in json.load(f).get("added_tokens", [])}
return [added.get(i) for i in range(4)] == ["[UNK]", "[TEXT]", "[START]", "[STOP]"]
@ModelBase.register_hparams_loader(_is_soprano_checkpoint)
def _load_soprano_hparams(dir_model: Path) -> dict[str, Any]:
logger.info("gguf: detected Soprano checkpoint")
hparams = ModelBase.load_hparams(dir_model, False, guess=False)
# the text model stays Qwen3, the mmproj holds the decoder
hparams["text_config"] = {"architectures": ["SopranoTextModel"]}
hparams["architectures"] = ["SopranoModel"]
return hparams
@ModelBase.register("SopranoTextModel")
class SopranoTextModel(Qwen3Model):
model_arch = gguf.MODEL_ARCH.QWEN3
def get_vocab_base_pre(self, tokenizer) -> str:
return "soprano"
def set_vocab(self):
super().set_vocab()
self.gguf_writer.add_unk_token_id(0)
@ModelBase.register("SopranoModel")
@ModelBase.example("ekwek/Soprano-80M", "ekwek/Soprano-1.1-80M")
class SopranoModel(MmprojModel):
has_vision_encoder = False
has_audio_encoder = False
def _load_decoder(self) -> dict[str, torch.Tensor]:
state = torch.load(self.dir_model / "decoder.pth", map_location="cpu", mmap=True, weights_only=True)
return state.get("state_dict", state)
def get_audio_config(self):
state = self._load_decoder()
n_embd, n_inp, _ = state["decoder.embed.weight"].shape
n_ff = state["decoder.convnext.0.pwconv1.weight"].shape[0]
n_layer = sum(1 for name in state if name.endswith(".pwconv1.weight"))
if n_inp != self.n_embd_text:
raise ValueError(f"Soprano decoder input dim {n_inp} does not match the backbone dim {self.n_embd_text}")
return {"num_hidden_layers": n_layer, "hidden_size": n_embd, "intermediate_size": n_ff}
def set_gguf_parameters(self):
self.gguf_writer.add_file_type(self.ftype)
self.gguf_writer.add_clip_has_gen_audio_encoder(True)
self.gguf_writer.add_clip_gen_audio_projector_type(gguf.VisionProjectorType.SOPRANO)
self.gguf_writer.add_gen_audio_projection_dim(self.n_embd_text)
self.gguf_writer.add_gen_audio_embedding_length(self.hparams["hidden_size"])
self.gguf_writer.add_gen_audio_feed_forward_length(self.hparams["intermediate_size"])
self.gguf_writer.add_gen_audio_block_count(self.hparams["num_hidden_layers"])
self.gguf_writer.add_gen_audio_head_count(1)
self.gguf_writer.add_gen_audio_attention_layernorm_eps(1e-6)
def get_tensors(self):
yield from self._load_decoder().items()
def tensor_force_quant(self, name, new_name, bid, n_dims):
if "dwconv.weight" in new_name or (name == "decoder.embed.weight" and n_dims == 3):
return gguf.GGMLQuantizationType.F16
return super().tensor_force_quant(name, new_name, bid, n_dims)
def modify_tensors(self, data_torch, name, bid):
if name == "head.istft.window":
# the ISTFT in mtmd builds its own window
if not torch.equal(data_torch.float(), torch.hann_window(data_torch.shape[0])):
raise ValueError("Soprano requires a periodic Hann window")
return
if name == "decoder.embed.weight" and data_torch.shape[-1] == 1:
data_torch = data_torch.squeeze(-1)
T = gguf.MODEL_TENSOR
if name.startswith("decoder.convnext."):
_, _, blk, key, *suffix = name.split(".")
block_map = {
"dwconv": T.A_GEN_WAV_UP_DWCONV,
"norm": T.A_GEN_WAV_UP_NORM,
"pwconv1": T.A_GEN_WAV_UP_PW1,
"pwconv2": T.A_GEN_WAV_UP_PW2,
"gamma": T.A_GEN_WAV_UP_GAMMA,
}
yield self.format_tensor_name(block_map[key], int(blk), suffix="".join("." + s for s in suffix)), data_torch
return
tensor_map = {
"decoder.embed": T.A_GEN_WAV_INPUT,
"decoder.norm": T.A_GEN_WAV_NORM,
"decoder.final_layer_norm": T.A_GEN_WAV_OUTPUT_NORM,
"head.out": T.A_GEN_WAV_OUTPUT,
}
key, suffix = name.rsplit(".", 1)
yield self.format_tensor_name(tensor_map[key], suffix="." + suffix), data_torch
+12
View File
@@ -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`
+1 -10
View File
@@ -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;
+6 -1
View File
@@ -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;
+10 -5
View File
@@ -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));
}
}
+6 -1
View File
@@ -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;
+6 -1
View File
@@ -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;
+6 -6
View File
@@ -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);
+12 -12
View File
@@ -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);
+19
View File
@@ -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)
}
+30 -6
View File
@@ -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;
+10 -39
View File
@@ -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);
+1 -2
View File
@@ -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);
+5 -8
View File
@@ -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;
+16 -51
View File
@@ -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;
-1
View File
@@ -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);
-1
View File
@@ -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:
{
@@ -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) {
+25 -84
View File
@@ -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);
}
}
}
+13 -3
View File
@@ -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,
+35
View File
@@ -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
+8
View File
@@ -34,6 +34,14 @@ kernel void kernel_add_id(
global float * dst_row = (global float *)((global char *)dst + i1*nb1 + i2*nb2);
global float * src0_row = (global float *)((global char *)src0 + i1*nb01 + i2*nb02);
if (i11 == -1) { // skipped slot, add nothing
for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
dst_row[i0] = src0_row[i0];
}
return;
}
global float * src1_row = (global float *)((global char *)src1 + i11*nb11);
for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
@@ -89,6 +89,9 @@ __kernel void kernel_gemv_moe_mxfp4_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -188,6 +191,9 @@ __kernel void kernel_gemv_moe_mxfp4_f32_ns_wimg(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f;
@@ -44,6 +44,9 @@ __kernel void kernel_gemv_moe_q4_0_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q4_1_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -61,6 +61,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int num_superblocks = ne00 / QK_K;
int num_subblocks = ne00 / 32;
@@ -180,6 +183,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns_wimg(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int num_superblocks = ne00 / QK_K;
int num_subblocks = ne00 / 32;
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q5_0_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -46,6 +46,9 @@ __kernel void kernel_gemv_moe_q5_1_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
@@ -62,6 +62,9 @@ __kernel void kernel_gemv_moe_q5_k_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int num_superblocks = ne00 / QK_K;
int num_subblocks = ne00 / 32;
@@ -45,6 +45,9 @@ __kernel void kernel_gemv_moe_q6_k_f32_ns(
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int num_superblocks = ne00 / QK_K;
int num_subblocks = ne00 / 32; // 8 sub-blocks of 32 per super-block
@@ -18,6 +18,9 @@ __kernel void kernel_moe_histogram(
}
int expert_id = input[n * n_experts + k];
if (expert_id == -1) { // skipped slot, owned by no expert
return;
}
atomic_inc(&hist[expert_id]);
}
@@ -63,6 +66,9 @@ __kernel void kernel_moe_scatter(
}
int val = input[n * n_experts + k];
if (val == -1) { // skipped slot, owned by no expert
return;
}
int local_slot = atomic_inc(&slot_counter[val]);
@@ -74,6 +80,32 @@ __kernel void kernel_moe_scatter(
emap[tile_idx] = val;
}
// the gemm only writes the rows an expert owns, so the skipped rows (negative ids) are zeroed here
__kernel void kernel_moe_zero_dst(
__global const int * input,
__global float * dst,
uint N,
uint topK,
uint n_experts,
uint ne0
) {
uint n = get_global_id(0);
uint k = get_global_id(1);
if (n >= N || k >= topK) {
return;
}
if (input[n * n_experts + k] != -1) {
return;
}
__global float * dst_row = dst + ((ulong)n * topK + k) * ne0;
for (uint i = get_local_id(2); i < ne0; i += get_local_size(2)) {
dst_row[i] = 0.0f;
}
}
// Deterministic replacement for kernel_moe_scatter.
//
// kernel_moe_scatter takes each token's slot from atomic_inc(slot_counter[expert]),
@@ -172,6 +172,9 @@ kernel void kernel_mul_mv_id_mxfp4_f32(
const int idx = get_group_id(2)%ne20;
int i02 = ((global int *) (src2 + iid1*nb21))[idx];
if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int i11 = idx % ne11;
int i12 = iid1;
@@ -101,6 +101,9 @@ kernel void kernel_mul_mv_id_mxfp4_f32_flat(
const int idx = get_group_id(2) % ne20;
uint i02 = ((global uint *) (src2 + offset2 + iid1 * nb21))[idx];
if ((int)i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int i11 = idx % ne11;
@@ -267,6 +267,9 @@ kernel void kernel_mul_mv_id_q4_0_f32_8x_flat(
const int idx = get_group_id(2)%ne20;
const int i02 = ((global int *)(src2 + iid1*nb21))[idx];
if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
const int i11 = idx%ne11;
const int i12 = iid1;
@@ -73,6 +73,9 @@ kernel void kernel_mul_mv_id_q8_0_f32(
int idx = get_group_id(2)%ne20;
int i02 = ((global int *) (src2 + iid1*nb21))[idx];
if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int i11_ = idx % ne11;
int i12_ = iid1;
@@ -72,6 +72,9 @@ kernel void kernel_mul_mv_id_q8_0_f32_flat(
int idx = (int)get_group_id(2)%ne20;
int i02 = ((global int *) (src2 + iid1*nb21))[idx];
if (i02 == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
int i11_ = idx % ne11;
int i12_ = iid1;
@@ -9,8 +9,12 @@
#include <openvino/op/constant.hpp>
#include <openvino/op/convert.hpp>
#include <openvino/op/gather.hpp>
#include <openvino/op/greater_eq.hpp>
#include <openvino/op/maximum.hpp>
#include <openvino/op/multiply.hpp>
#include <openvino/op/reshape.hpp>
#include <openvino/op/shape_of.hpp>
#include <openvino/op/unsqueeze.hpp>
namespace ov {
namespace frontend {
@@ -52,8 +56,16 @@ OutputVector translate_add_id(const NodeContext & context) {
ids = std::make_shared<ov::op::v0::Convert>(ids, ov::element::i32);
}
// an id of -1 marks a skipped slot: gather row 0 for it and zero the bias with a mask
auto zero = ov::op::v0::Constant::create(ids.get_element_type(), ov::Shape{}, {0});
ov::Output<ov::Node> keep = std::make_shared<ov::op::v1::GreaterEqual>(ids, zero);
keep = std::make_shared<ov::op::v0::Convert>(keep, bias.get_element_type());
keep = std::make_shared<ov::op::v0::Unsqueeze>(keep, ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, {-1}));
ids = std::make_shared<ov::op::v1::Maximum>(ids, zero);
auto gather_axis = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{}, {0});
ov::Output<ov::Node> selected_bias = std::make_shared<ov::op::v8::Gather>(bias, ids, gather_axis);
selected_bias = std::make_shared<ov::op::v1::Multiply>(selected_bias, keep);
selected_bias = std::make_shared<ov::op::v1::Reshape>(
selected_bias, std::make_shared<ov::op::v3::ShapeOf>(input, ov::element::i64), false);
+5
View File
@@ -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;
+9
View File
@@ -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;
+68 -10
View File
@@ -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();
+18 -2
View File
@@ -3149,7 +3149,6 @@ static void mul_mat_vec_q_moe(
const int expert_idx = item_ct1.get_group(1);
const int i02 = ids_dev[expert_idx];
const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
const char * vy = (const char *) vy_base + (size_t) expert_idx * src1_row_stride;
float * dst = (float *) ((char *) dst_base + (size_t) expert_idx * dst_row_stride);
@@ -3159,6 +3158,15 @@ static void mul_mat_vec_q_moe(
return;
}
if (i02 == -1) { // skipped slot
if (item_ct1.get_local_id(2) == 0) {
dst[row] = 0.0f;
}
return;
}
const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
const int blocks_per_row = ncols / qk;
constexpr int blocks_per_warp = (vdr * WARP_SIZE + qi - 1) / qi;
@@ -3355,7 +3363,6 @@ static void mul_mat_vec_q_moe_reorder(
const int expert_idx = item_ct1.get_group(1);
const int i02 = ids_dev[expert_idx];
const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
const char * vy = (const char *) vy_base + (size_t) expert_idx * src1_row_stride;
float * dst = (float *) ((char *) dst_base + (size_t) expert_idx * dst_row_stride);
@@ -3366,6 +3373,15 @@ static void mul_mat_vec_q_moe_reorder(
const auto sg = item_ct1.get_sub_group();
if (i02 == -1) { // skipped slot
if (sg.leader()) {
dst[row] = 0.0f;
}
return;
}
const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
constexpr int vdr_mmvq = reorder_vec_dot_q_sycl::gtype == GGML_TYPE_Q8_0 ?
VDR_Q8_0_Q8_1_MMVQ : block_traits::vdr_mmvq;
const int blocks_per_row = ncols / block_traits::qk;
@@ -222,6 +222,9 @@ struct vk_op_count_experts_push_constants {
uint32_t ne00L;
uint32_t row_tile_size;
uint32_t row_ids_offset;
uint32_t ne0_dst;
uint32_t stride_dst;
uint32_t batch_stride_dst;
};
struct vk_op_glu_push_constants {
+56 -23
View File
@@ -3585,9 +3585,9 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_count_equal_i32, "count_equal_i32", count_equal_i32_len, count_equal_i32_data, "main", 3, sizeof(vk_op_push_constants), {512, 1, 1}, { device->subgroup_size }, 1);
if (device->subgroup_arithmetic && device->subgroup_vote && device->subgroup_require_full_support) {
ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_subgroup_len, count_experts_subgroup_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true, true);
ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_subgroup_len, count_experts_subgroup_data, "main", 3, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true, true);
} else {
ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 3, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
}
// comb holds a token's 4x4 matrix in one 16-lane slice of a subgroup, so it
@@ -6005,20 +6005,49 @@ void ggml_vk_matmul(
ctx->prealloc_split_k_need_sync = true;
}
static bool ggml_vk_get_mul_mat_mat_f16acc(ggml_backend_vk_context * ctx, ggml_type src0_type, ggml_type src1_type, ggml_prec prec) {
static bool ggml_vk_src1_prec_allows_q8(const ggml_tensor * dst) {
switch (ggml_prec_get_src(dst, 1)) {
case GGML_PREC_UNDEFINED:
case GGML_PREC_Q8:
case GGML_PREC_Q4:
return true;
case GGML_PREC_F32:
case GGML_PREC_BF16:
case GGML_PREC_F16:
return false;
}
return false;
}
static bool ggml_vk_prec_allows_f16acc(ggml_prec prec) {
switch (prec) {
case GGML_PREC_F32:
case GGML_PREC_BF16:
return false;
case GGML_PREC_UNDEFINED:
case GGML_PREC_F16:
case GGML_PREC_Q8:
case GGML_PREC_Q4:
return true;
}
return false;
}
static bool ggml_vk_get_mul_mat_mat_f16acc(ggml_backend_vk_context * ctx, ggml_type src0_type, ggml_type src1_type, ggml_prec acc_prec) {
if (src0_type == GGML_TYPE_F32 || src0_type == GGML_TYPE_BF16) return false;
if (src1_type == GGML_TYPE_Q8_1) return false;
const bool prec_allows_f16 = ggml_vk_prec_allows_f16acc(acc_prec);
if (src0_type == GGML_TYPE_F16) {
return prec == GGML_PREC_DEFAULT && ctx->device->fp16 && !(ctx->device->coopmat_support && !ctx->device->coopmat_acc_f16_support);
return prec_allows_f16 && ctx->device->fp16 && !(ctx->device->coopmat_support && !ctx->device->coopmat_acc_f16_support);
}
// quant types
if (ctx->device->coopmat2) {
return prec == GGML_PREC_DEFAULT;
return prec_allows_f16;
}
if (ctx->device->coopmat_support) {
return ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec == GGML_PREC_DEFAULT;
return ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec_allows_f16;
}
return ctx->device->fp16 && prec == GGML_PREC_DEFAULT;
return ctx->device->fp16 && prec_allows_f16;
}
static const std::vector<vk_matmul_pipeline_pair>* ggml_vk_get_mul_mat_mat_pipeline_map(
@@ -6386,26 +6415,26 @@ static void ggml_vk_mul_mat_q_f16(ggml_backend_vk_context * ctx, vk_context& sub
const bool is_coopmat = ctx->device->coopmat_support || ctx->device->coopmat2;
const bool can_quantize_y = (ctx->device->integer_dot_product || ctx->device->coopmat_int_support) &&
src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0;
src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst);
// coopmat shaders are fast enough that memory bandwidth reduction from f16 activations helps
const bool prefer_f16_b = is_coopmat && ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32;
const ggml_type y_kernel_type = ggml_vk_mul_mat_b_type(ctx, src0->type, src1, f16_type,
(ggml_prec)dst->op_params[0], false, can_quantize_y, prefer_f16_b);
ggml_prec_get_acc(dst), false, can_quantize_y, prefer_f16_b);
const bool quantize_y = y_kernel_type == GGML_TYPE_Q8_1;
const bool y_f32_kernel = y_kernel_type == GGML_TYPE_F32;
const bool y_needs_reformat = !quantize_y && (!ggml_vk_dim01_contiguous(src1) || y_kernel_type != src1->type);
const std::vector<vk_matmul_pipeline_pair>* mmp_map =
ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, (ggml_prec)dst->op_params[0]);
ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, ggml_prec_get_acc(dst));
const bool qx_needs_dequant = mmp_map == nullptr || x_non_contig;
const bool qy_needs_dequant = y_needs_reformat;
if (qx_needs_dequant) {
// dequant src0 to f16
mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, (ggml_prec)dst->op_params[0]);
mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, ggml_prec_get_acc(dst));
}
GGML_ASSERT(mmp_map != nullptr);
@@ -6736,7 +6765,7 @@ static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context&
const ggml_type y_kernel_type = widen_bf16 ? GGML_TYPE_F32 : src1->type;
const bool f16_f32_kernel = y_kernel_type == GGML_TYPE_F32;
bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_should_use_mmvq(ctx->device, ne01, ne11, ne10, src0->type);
bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && (swap_inputs || ggml_vk_src1_prec_allows_q8(dst)) && ggml_vk_should_use_mmvq(ctx->device, ne01, ne11, ne10, src0->type);
vk_pipeline to_fp16_vk_0 = nullptr;
vk_pipeline to_fp16_vk_1 = nullptr;
@@ -7435,26 +7464,26 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
const bool y_decode_vector_staging = false;
#endif
const bool can_quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 &&
ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0;
ggml_is_contiguous(src1) && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst);
// Intel coopmat1 matrix cores prefer f16 activations for quantized weights
const bool prefer_f16_b = ctx->device->coopmat_support && !ctx->device->coopmat2 &&
ctx->device->vendor_id == VK_VENDOR_ID_INTEL &&
ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32;
const ggml_type y_kernel_type = ggml_vk_mul_mat_b_type(ctx, src0->type, src1, f16_type,
(ggml_prec)dst->op_params[0], true, can_quantize_y, prefer_f16_b);
ggml_prec_get_acc(dst), true, can_quantize_y, prefer_f16_b);
const bool quantize_y = y_kernel_type == GGML_TYPE_Q8_1;
const bool y_f32_kernel = y_kernel_type == GGML_TYPE_F32;
const std::vector<vk_matmul_pipeline_pair>* mmp_map =
ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, (ggml_prec)dst->op_params[0], true);
ggml_vk_get_mul_mat_mat_pipeline_map(ctx, src0->type, y_kernel_type, ggml_prec_get_acc(dst), true);
const bool qx_needs_dequant = mmp_map == nullptr || x_non_contig;
if (qx_needs_dequant) {
// dequant src0 to f16
mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, (ggml_prec)dst->op_params[0], true);
mmp_map = ggml_vk_get_mul_mat_mat_pipeline_map(ctx, f16_type, y_kernel_type, ggml_prec_get_acc(dst), true);
}
GGML_ASSERT(mmp_map != nullptr);
@@ -7614,16 +7643,20 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
ggml_vk_sync_buffers(ctx, subctx);
}
{
// this also zeroes the dst rows of the skipped slots (negative ids)
vk_op_count_experts_push_constants pc = { (uint32_t)nei0,
(uint32_t)nei1,
(uint32_t)(nbi0 / ggml_type_size(ids->type)),
(uint32_t)(nbi1 / ggml_type_size(ids->type)),
(uint32_t)(get_misalign_bytes(ctx, ids) / ggml_type_size(ids->type)),
(uint32_t)n_as,
0, 0, bn, (uint32_t)row_ids_offset };
0, 0, bn, (uint32_t)row_ids_offset,
(uint32_t)ne01,
(uint32_t)(dst->nb[1] / sizeof(float)),
(uint32_t)(dst->nb[2] / sizeof(float)) };
init_pushconst_fastdiv(pc);
ggml_vk_dispatch_pipeline(ctx, subctx, count_experts,
{ d_ids, expert_count_buf }, pc,
{ d_ids, expert_count_buf, vk_subbuffer{ d_D, d_buf_offset, d_sz } }, pc,
{ 1, 1, 1 });
}
@@ -7749,7 +7782,7 @@ static void ggml_vk_mul_mat_vec_id_q_f16(ggml_backend_vk_context * ctx, vk_conte
const bool y_non_contig = !ggml_vk_dim01_contiguous(src1);
const bool f16_f32_kernel = src1->type == GGML_TYPE_F32;
bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_should_use_mmvq(ctx->device, ne01, ne12, ne10, src0->type);
bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0 && ggml_vk_src1_prec_allows_q8(dst) && ggml_vk_should_use_mmvq(ctx->device, ne01, ne12, ne10, src0->type);
vk_pipeline to_fp16_vk_0 = nullptr;
vk_pipeline to_fp16_vk_1 = nullptr;
@@ -8128,7 +8161,7 @@ void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx, const
uint32_t workgroups_y = (uint32_t)neq2;
uint32_t workgroups_z = (uint32_t)neq3;
const bool f32acc = !ctx->device->fp16 || dst->op_params[3] == GGML_PREC_F32 || k->type == GGML_TYPE_BF16;
const bool f32acc = !ctx->device->fp16 || !ggml_vk_prec_allows_f16acc(ggml_prec_get_acc(dst)) || k->type == GGML_TYPE_BF16;
// dequant K/V once into an f16 scratch, reordered KV layout so FA can read without a stride
auto is_dense_kv_cache = [](const ggml_tensor * t) {
@@ -15516,9 +15549,9 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
// If there's not enough shared memory for row_ids and the result tile, fallback to CPU
return false;
}
if (ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) {
return false;
}
}
if (ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
return false;
}
switch (src0_type) {
case GGML_TYPE_F32:
@@ -27,14 +27,22 @@ void main() {
const uint i1 = gl_WorkGroupID.x;
const uint i2 = gl_WorkGroupID.y;
const uint i11 = data_c[i1 + i2 * p.s21];
const int i11 = data_c[i1 + i2 * p.s21];
const uint s1 = p.ne0;
const uint s2 = p.ne0 * p.ne1;
const uint d0 = i1 * s1 + i2 * s2;
const uint a0 = i1 * p.s01 + i2 * p.s02;
const uint b0 = i11 * p.s11;
if (i11 == -1) { // skipped slot, add nothing
for (uint i0 = gl_LocalInvocationID.x; i0 < p.ne0; i0 += BLOCK_SIZE) {
data_d[d0 + i0] = data_a[a0 + i0];
}
return;
}
const uint b0 = uint(i11) * p.s11;
for (uint i0 = gl_LocalInvocationID.x; i0 < p.ne0; i0 += BLOCK_SIZE) {
data_d[d0 + i0] = data_a[a0 + i0] + data_b[b0 + i0];
@@ -23,6 +23,9 @@ layout (push_constant) uniform parameter
uint32_t ne00L;
uint32_t row_tile_size;
uint32_t row_ids_offset;
uint32_t ne0_dst;
uint32_t stride_dst;
uint32_t batch_stride_dst;
} p;
#define BLOCK_SIZE 256
@@ -32,6 +35,7 @@ layout(local_size_x = BLOCK_SIZE, local_size_y = 1, local_size_z = 1) in;
layout (binding = 0) readonly buffer A {uint data_a[];};
layout (binding = 1) writeonly buffer D {uint data_d[];};
layout (binding = 2) writeonly buffer DST {float data_dst[];};
shared uint vals[EXPERTS_PER_CHUNK];
shared uint offsets[EXPERTS_PER_CHUNK];
@@ -47,6 +51,26 @@ void write_tiles(uint expert, uint count, uint tile_begin, uint n_tiles) {
}
}
// the matmul writes only the rows that an expert owns, so zero the rows of the skipped slots
void zero_skipped_rows() {
const uint num_elements = p.ne00 * p.ne01;
const uint tid = gl_LocalInvocationID.x;
for (uint idx = 0; idx < num_elements; idx++) {
const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
const uint i00 = idx - i01 * p.ne00;
if (int(data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00]) != -1) {
continue;
}
const uint d_off = i00 * p.stride_dst + i01 * p.batch_stride_dst;
for (uint i = tid; i < p.ne0_dst; i += BLOCK_SIZE) {
data_dst[d_off + i] = 0.0f;
}
}
}
void main() {
const uint num_elements = p.ne00 * p.ne01;
const uint tid = gl_LocalInvocationID.x;
@@ -123,4 +147,6 @@ void main() {
if (tid == 0) {
data_d[0] = total_tiles;
}
zero_skipped_rows();
}
@@ -42,6 +42,7 @@ layout (push_constant) uniform parameter
#ifdef MUL_MAT_ID
uint expert_id;
bool expert_skip; // an id of -1 skips the slot, the result row is zeroed
#endif
void get_offsets(out uint a_offset, out uint b_offset, out uint d_offset) {
@@ -63,7 +64,9 @@ void get_offsets(out uint a_offset, out uint b_offset, out uint d_offset) {
batch_idx_a = i03 * p.ne02 + i02;
}
#else
expert_id = data_ids[expert_i0 + p.expert_i1 * p.nbi1];
const int expert = data_ids[expert_i0 + p.expert_i1 * p.nbi1];
expert_skip = expert == -1;
expert_id = expert_skip ? 0 : uint(expert); // clamp, so that the offsets below stay in range
#endif
a_offset =
@@ -113,6 +116,9 @@ void reduce_result(inout FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t
const uint expert_i0 = gl_GlobalInvocationID.y;
temp[j][n] *= FLOAT_TYPE(data_fuse1[expert_i0]);
}
if (expert_skip) {
temp[j][n] = FLOAT_TYPE(0);
}
#else
if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
temp[j][n] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
@@ -168,6 +174,9 @@ void reduce_result(FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t d_offs
const uint expert_i0 = gl_GlobalInvocationID.y;
temp[j][n] *= FLOAT_TYPE(data_fuse1[expert_i0]);
}
if (expert_skip) {
temp[j][n] = FLOAT_TYPE(0);
}
#else
if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
temp[j][n] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
@@ -213,6 +222,9 @@ void reduce_result(FLOAT_TYPE temp[NUM_COLS][NUM_ROWS], const in uint32_t d_offs
const uint expert_i0 = gl_GlobalInvocationID.y;
tmpsh[j][n][0] *= FLOAT_TYPE(data_fuse1[expert_i0]);
}
if (expert_skip) {
tmpsh[j][n][0] = FLOAT_TYPE(0);
}
#else
if ((p.fusion_flags & MAT_VEC_FUSION_FLAGS_BIAS0) != 0) {
tmpsh[j][n][0] += FLOAT_TYPE(data_fuse0[j*p.batch_stride_d + d_offset + first_row + n]);
+4
View File
@@ -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({
+14 -2
View File
@@ -46,12 +46,24 @@ fn main(@builtin(workgroup_id) wg_id: vec3<u32>,
let i2 = wg_linear / params.ne1;
let i1 = wg_linear % params.ne1;
let i11 = u32(ids[params.offset_ids + i1 * params.nb20 + i2 * params.nb21]);
let i11 = i32(ids[params.offset_ids + i1 * params.nb20 + i2 * params.nb21]);
let src0_row = params.offset_src0 + i1 * params.nb01 + i2 * params.nb02;
let src1_row = params.offset_src1 + i11 * params.nb11;
let dst_row = params.offset_dst + i1 * params.ne0 + i2 * (params.ne0 * params.ne1);
if (i11 == -1) { // skipped slot, add nothing
#ifdef INPLACE
// src0 already holds the result
#else
for (var i = thread_id;i < params.ne0; i += WG_SIZE) {
dst[dst_row + i] = src0[src0_row + i];
}
#endif
return;
}
let src1_row = params.offset_src1 + u32(i11) * params.nb11;
for (var i = thread_id;i < params.ne0; i += WG_SIZE) {
#ifdef INPLACE
src0[src0_row + i] = src0[src0_row + i] + src1[src1_row + i];
@@ -2,19 +2,24 @@ enable f16;
struct MulMatIdGatherParams {
offset_ids: u32,
offset_dst: u32,
n_expert: u32,
n_expert_used: u32,
n_tokens: u32,
stride_ids_1: u32,
m: u32,
};
@group(0) @binding(0) var<storage, read_write> ids: array<i32>; // [n_expert_used, n_tokens]
@group(0) @binding(1) var<storage, read_write> ids_slot_pos: array<u32>; // [n_expert_used * n_tokens]
@group(0) @binding(2) var<storage, read_write> ids_tiles: array<u32>; // [n_tiles, (expert, first slot, n slots) * n_tiles]
@group(0) @binding(3) var<uniform> params: MulMatIdGatherParams;
@group(0) @binding(3) var<storage, read_write> dst: array<f32>; // [m, n_expert_used, n_tokens]
@group(0) @binding(4) var<uniform> params: MulMatIdGatherParams;
// per expert: the slot count, replaced by the first slot of the expert in ids_slot_pos by the prefix sum, then advanced as its write cursor
var<workgroup> ids_count: array<atomic<u32>, N_EXPERT>;
@@ -36,8 +41,10 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
for (var i = thread_id; i < ids_slots_num; i += WG_SIZE) {
let row = i / params.n_expert_used;
let col = i % params.n_expert_used;
let expert = u32(ids[params.offset_ids + row * params.stride_ids_1 + col]);
atomicAdd(&ids_count[expert], 1u);
let expert = ids[params.offset_ids + row * params.stride_ids_1 + col];
if (expert != -1) { // skipped slot, owned by no expert
atomicAdd(&ids_count[u32(expert)], 1u);
}
}
workgroupBarrier();
@@ -71,8 +78,23 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
for (var i = thread_id; i < ids_slots_num; i += WG_SIZE) {
let row = i / params.n_expert_used;
let col = i % params.n_expert_used;
let expert = u32(ids[params.offset_ids + row * params.stride_ids_1 + col]);
let pos = atomicAdd(&ids_count[expert], 1u);
ids_slot_pos[pos] = i;
let expert = ids[params.offset_ids + row * params.stride_ids_1 + col];
if (expert != -1) {
let pos = atomicAdd(&ids_count[u32(expert)], 1u);
ids_slot_pos[pos] = i;
}
}
// the matmul writes only the rows that an expert owns, so zero the rows of the skipped slots
for (var i: u32 = 0; i < ids_slots_num; i += 1) {
let row = i / params.n_expert_used;
let col = i % params.n_expert_used;
if (ids[params.offset_ids + row * params.stride_ids_1 + col] != -1) {
continue;
}
let dst_row = params.offset_dst + i * params.m;
for (var r = thread_id; r < params.m; r += WG_SIZE) {
dst[dst_row + r] = 0.0f;
}
}
}
@@ -70,11 +70,18 @@ fn main(
let own_expert = ids[params.offset_ids + slot];
let row_base = (wg_linear % output_groups) * OUTPUTS_PER_WG;
let src0_batch_offset = params.offset_src0 + own_expert * params.stride_02;
// a skipped slot (id -1) keeps the zero accumulator, so the reduction below writes zeros
// no early return: the barriers below must stay in uniform control flow
let skipped = i32(own_expert) == -1;
let src0_batch_offset = params.offset_src0 + select(own_expert, 0u, skipped) * params.stride_02;
let src1_idx_base = params.offset_src1 + (slot % params.b_ne1) * params.stride_11;
let dst_idx_base = params.offset_dst + slot * params.m + row_base;
let acc = accumulate_vec_dot(thread_id, row_base, src0_batch_offset, src1_idx_base);
var acc: array<array<f32, OUTPUTS_PER_WG>, NUM_COLS>;
if (!skipped) {
acc = accumulate_vec_dot(thread_id, row_base, src0_batch_offset, src1_idx_base);
}
#ifdef USE_SUBGROUP_REDUCTION
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
+6 -1
View File
@@ -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++;
+32
View File
@@ -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,
+14
View File
@@ -1272,6 +1272,11 @@ class MODEL_TENSOR(IntEnum):
A_GEN_WAV_SEANET_RES_CONV1 = auto()
A_GEN_WAV_SEANET_RES_CONV2 = auto()
A_GEN_WAV_SEANET_SCALE_CONV = auto() # strided upsample convtr
# soprano: Vocos decoder, the ConvNeXt blocks reuse A_GEN_WAV_UP_*
A_GEN_WAV_INPUT = auto()
A_GEN_WAV_NORM = auto()
A_GEN_WAV_OUTPUT_NORM = auto()
A_GEN_WAV_OUTPUT = auto()
A_MMPROJ = auto()
A_MMPROJ_FC = auto()
A_MM_NORM_PRE = auto()
@@ -2084,6 +2089,10 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1: "a.gen.wav.seanet.blk.{bid}.res_conv1",
MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2: "a.gen.wav.seanet.blk.{bid}.res_conv2",
MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV: "a.gen.wav.seanet.blk.{bid}.scale_conv",
MODEL_TENSOR.A_GEN_WAV_INPUT: "a.gen.wav.input",
MODEL_TENSOR.A_GEN_WAV_NORM: "a.gen.wav.norm",
MODEL_TENSOR.A_GEN_WAV_OUTPUT_NORM: "a.gen.wav.output_norm",
MODEL_TENSOR.A_GEN_WAV_OUTPUT: "a.gen.wav.output",
MODEL_TENSOR.A_MMPROJ: "mm.a.mlp.{bid}",
MODEL_TENSOR.A_MMPROJ_FC: "mm.a.fc",
MODEL_TENSOR.A_MM_NORM_PRE: "mm.a.norm_pre",
@@ -2462,6 +2471,10 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV1,
MODEL_TENSOR.A_GEN_WAV_SEANET_RES_CONV2,
MODEL_TENSOR.A_GEN_WAV_SEANET_SCALE_CONV,
MODEL_TENSOR.A_GEN_WAV_INPUT,
MODEL_TENSOR.A_GEN_WAV_NORM,
MODEL_TENSOR.A_GEN_WAV_OUTPUT_NORM,
MODEL_TENSOR.A_GEN_WAV_OUTPUT,
MODEL_TENSOR.A_ENC_CONV_NORM_MEAN,
MODEL_TENSOR.A_ENC_CONV_NORM_VAR,
MODEL_TENSOR.A_ENC_MEL_FILTERS,
@@ -6447,6 +6460,7 @@ class VisionProjectorType:
QWEN3TTS_GEN = "qwen3tts_gen" # audio generation: code_predictor
POCKETTTS_SPKENC = "pockettts_spkenc" # audio: mimi encoder as voice-prompt encoder
POCKETTTS_GEN = "pockettts_gen" # audio generation: flow-matching decoder + mimi decoder
SOPRANO = "soprano" # audio generation: Vocos decoder
HUNYUANVL = "hunyuanvl"
PARAKEET = "parakeet" # audio
MINIMAXM3 = "minimax_m3"
+2
View File
@@ -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:
+1
View File
@@ -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 };
+33 -1
View File
@@ -572,6 +572,10 @@ struct llm_tokenizer_bpe : llm_tokenizer {
"(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])|[^\\r\\n\\p{L}\\p{N}]?(?:\\p{L}|\\p{M}|\\u200C|\\u200D)+|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+",
};
break;
case LLAMA_VOCAB_PRE_TYPE_SOPRANO:
regex_exprs = { "\\p{N}", "\\s+|[\\p{L}\\p{N}_]+|[^\\p{L}\\p{N}_\\s]+" };
byte_encode = false;
break;
case LLAMA_VOCAB_PRE_TYPE_WHITESPACE:
// whitespace pre-tokenizer (jinaai/jina-embeddings-v2-base-zh)
regex_exprs = {
@@ -647,6 +651,22 @@ struct llm_tokenizer_bpe_session {
replace_all(normalized_text, "[" + zwnj, "[");
replace_all(normalized_text, "<" + zwnj, "<");
word_collection = unicode_regex_split(normalized_text, tokenizer.regex_exprs, tokenizer.byte_encode);
} else if (tok_pre == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
std::string normalized_text;
bool space = false;
for (uint32_t cpt : unicode_cpts_from_utf8(text)) {
if (unicode_cpt_flags_from_cpt(cpt).is_whitespace) {
if (!space) {
normalized_text += ' ';
}
space = true;
} else {
// Unicode full lowercase expands LATIN CAPITAL LETTER I WITH DOT ABOVE.
normalized_text += cpt == 0x0130 ? "i\xcc\x87" : unicode_cpt_to_utf8(unicode_tolower(cpt));
space = false;
}
}
word_collection = unicode_regex_split(normalized_text, tokenizer.regex_exprs, tokenizer.byte_encode);
} else {
word_collection = unicode_regex_split(text, tokenizer.regex_exprs, tokenizer.byte_encode);
}
@@ -769,7 +789,9 @@ struct llm_tokenizer_bpe_session {
const std::string str = std::string(symbol.text, symbol.n);
const auto token = vocab.text_to_token(str);
if (token == LLAMA_TOKEN_NULL) {
if (token == LLAMA_TOKEN_NULL && tok_pre == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
output.push_back(vocab.token_unk());
} else if (token == LLAMA_TOKEN_NULL) {
for (auto j = str.begin(); j != str.end(); ++j) {
llama_token token_multibyte = LLAMA_TOKEN_NULL;
if (tokenizer.byte_encode) {
@@ -2730,6 +2752,9 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
pre_type = LLAMA_VOCAB_PRE_TYPE_NEMOTRON_3_5_EMBD;
ignore_merges = true;
add_bos = true;
} else if (
tokenizer_pre == "soprano") {
pre_type = LLAMA_VOCAB_PRE_TYPE_SOPRANO;
} else {
throw std::runtime_error(format("unknown pre-tokenizer type: '%s'", tokenizer_pre.c_str()));
}
@@ -3283,6 +3308,10 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
}
}
if (pre_type == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
special_eog_ids.insert(special_eos_id);
}
// sanity checks
if (special_eos_id != LLAMA_TOKEN_NULL && special_eog_ids.count(special_eos_id) == 0) {
special_eog_ids.insert(special_eos_id);
@@ -4096,6 +4125,9 @@ int32_t llama_vocab::impl::token_to_piece(llama_token token, char * buf, int32_t
return _try_copy(token_text.data(), token_text.size());
}
if (attr & LLAMA_TOKEN_ATTR_NORMAL) {
if (pre_type == LLAMA_VOCAB_PRE_TYPE_SOPRANO) {
return _try_copy(token_text.data(), token_text.size());
}
if (escape_whitespaces) {
// SPM-style BPE: tokens contain ▁ for spaces
std::string result = token_text;
+1
View File
@@ -73,6 +73,7 @@ enum llama_vocab_pre_type {
LLAMA_VOCAB_PRE_TYPE_FLINT = 62,
LLAMA_VOCAB_PRE_TYPE_DOMYN_SMALL = 63,
LLAMA_VOCAB_PRE_TYPE_NEMOTRON_3_5_EMBD = 64,
LLAMA_VOCAB_PRE_TYPE_SOPRANO = 65,
};
// HF Metaspace pre-tokenizer prepend_scheme
+10 -7
View File
@@ -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);
+20
View File
@@ -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
+96 -40
View File
@@ -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));
}
}
}
}
+4 -3
View File
@@ -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;
+76
View File
@@ -5,6 +5,7 @@
#include "mtmd-internal.h"
#include <cmath>
#include <cstdint>
#include <iostream>
#include <stdexcept>
#include <string>
@@ -95,6 +96,45 @@ static std::string image_size_to_str(const clip_image_size & size) {
return std::to_string(size.width) + "x" + std::to_string(size.height);
}
static std::string resize_algo_to_str(resize_algo algo) {
std::string str;
switch (algo & ~RESIZE_ALGO_FLAG_TORCH) {
case RESIZE_ALGO_BILINEAR: str = "bilinear"; break;
case RESIZE_ALGO_BICUBIC: str = "bicubic"; break;
case RESIZE_ALGO_LANCZOS: str = "lanczos"; break;
default: str = "unknown"; break;
}
if (algo & RESIZE_ALGO_FLAG_TORCH) {
str += "_torch";
}
return str;
}
// pseudo-random pixels: s = s * 1103515245 + 12345 from s = 1, one byte (s >> 16) & 255 per channel
static clip_image_u8 make_test_image(const clip_image_size & size) {
clip_image_u8 img;
img.set_size(size, false);
uint32_t s = 1;
auto next = [&s]() {
s = s * 1103515245u + 12345u;
return static_cast<uint8_t>(s >> 16);
};
for (int y = 0; y < size.height; y++) {
for (int x = 0; x < size.width; x++) {
img.set_pixel(x, y, { next(), next(), next() });
}
}
return img;
}
// FNV-1a of the pixels
static uint32_t hash_pixels(const clip_image_u8 & img, uint32_t hash = 2166136261u) {
for (uint8_t v : img.get_ro_buf()) {
hash = (hash ^ v) * 16777619u;
}
return hash;
}
MAKE_TEST(test_image_preprocessor_lfm2) {
clip_hparams hparams;
hparams.patch_size = 16;
@@ -158,6 +198,42 @@ MAKE_TEST(test_image_preprocessor_lfm2) {
std::string(mtmd_image_preprocessor_lfm2::should_tile(hparams, { 1024, 977 }) ? "tiled" : "single"));
}
MAKE_TEST(test_image_preprocessor_lfm2_torch_resize) {
clip_hparams hparams;
hparams.patch_size = 16;
hparams.n_merge = 2;
hparams.set_limit_image_tokens(64, 256);
hparams.preproc_min_tiles = 2;
hparams.preproc_max_tiles = 10;
hparams.preproc_tile_size = 512;
hparams.preproc_max_pixels_tolerance = 2.0f;
// Lfm2VlImageProcessor stretches the image to the tile grid
hparams.image_pad_rf = PAD_NONE;
// { resize algo, image size, expected hash of the tiles and the thumbnail from Lfm2VlImageProcessor (torchvision 0.29, native uint8 resize: arm64, or x86 with AVX2/AVX512) }
const std::vector<std::tuple<resize_algo, clip_image_size, uint32_t>> cases = {
{ RESIZE_ALGO_BILINEAR, { 620, 410 }, 0xf5e9c941 }, // 608x416
{ RESIZE_ALGO_BILINEAR, { 2000, 1500 }, 0x3331f56c }, // 3x2 tiles, 576x416 thumbnail
{ RESIZE_ALGO_BICUBIC, { 620, 410 }, 0x2b34eac9 },
{ RESIZE_ALGO_BICUBIC, { 2000, 1500 }, 0xb035d942 },
};
for (const auto & [base_algo, size, expected] : cases) {
const resize_algo algo = static_cast<resize_algo>(base_algo | RESIZE_ALGO_FLAG_TORCH);
hparams.image_resize_algo_rf = algo;
hparams.image_resize_algo_ov = algo;
const mtmd_image_preprocessor_lfm2 preproc(hparams);
const auto sliced = preproc.slice_image(make_test_image(size), preproc.get_slice_instructions(size));
uint32_t hash = 2166136261u;
for (const auto & slice : sliced.slices) {
hash = hash_pixels(slice, hash);
}
hash = hash_pixels(sliced.overview, hash);
t.assert_equal("pixels of " + std::to_string(size.width) + "x" + std::to_string(size.height) + " with resize algo " + resize_algo_to_str(algo), expected, hash);
}
}
//
// mtmd temporal merge
//
@@ -0,0 +1,229 @@
#include "sampling.h"
#include "llama.h"
#include "common.h"
#include "arg.h"
#include "log.h"
#include <cmath>
#include <cstdlib>
#include <cstring>
#include <optional>
#include <random>
#include <unordered_map>
// Tests that common_sampler_sample_and_accept_n_rejection preserves the target
// distribution regardless of what the draft distribution Q looks like.
static int test_speculative_rejection_sampling(common_params & params, int n_trials) {
params.sampling.backend_sampling = false;
auto llama_init = common_init_from_params(params);
auto * model_tgt = llama_init->model();
auto * ctx_tgt = llama_init->context();
if (!model_tgt || !ctx_tgt) {
LOG_ERR("failed to initialize target model and context\n");
return 1;
}
const llama_tokens prompt = { 1, 2, 3, 1, 2 };
// Process the prompt with the target model to get its true distribution for
// the position right after the prompt.
common_batch batch = common_batch_get_one(ctx_tgt, prompt);
if (llama_process(ctx_tgt, LLAMA_PROCESS_TYPE_DECODE, batch.get()) != 0) {
LOG_ERR("failed to decode the test prompt\n");
return 1;
}
common_params_sampling sampling;
sampling.seed = 18;
sampling.samplers = { COMMON_SAMPLER_TYPE_TEMPERATURE };
sampling.temp = 1.0f;
// Initialize the sampler for the target model.
common_sampler_ptr smpl(common_sampler_init(model_tgt, sampling));
if (!smpl) {
LOG_ERR("failed to initialize target sampler\n");
return 1;
}
// Sample once so the sampler chain runs and populates cur_p. We don't need
// the returned token id itself so the return value is ignored.
common_sampler_sample(smpl.get(), ctx_tgt, -1);
// Get the target candidates distribution in sorted order.
const llama_token_data_array * dist_tgt = common_sampler_get_candidates(smpl.get(), true);
// Clone smpl before the trial loop starts. We need the dist_tgt pointer to stay valid.
common_sampler_ptr smpl_verify(common_sampler_clone(smpl.get()));
if (!smpl_verify) {
LOG_ERR("failed to clone target sampler\n");
return 1;
}
const llama_vocab * vocab = llama_model_get_vocab(model_tgt);
const int n_vocab = llama_vocab_n_tokens(vocab);
std::mt19937 rng(18);
std::uniform_real_distribution<float> uni(0.0f, 1.0f);
std::unordered_map<llama_token, int> token_counts;
std::vector<llama_token_data> q(n_vocab);
for (int i = 0; i < n_vocab; ++i) {
q[i].id = i;
q[i].logit = 0.0f;
}
for (int trial = 0; trial < n_trials; ++trial) {
float q_sum = 0.0f;
for (auto & td : q) {
td.p = uni(rng);
q_sum += td.p;
}
const float inv_sum = 1.0f / q_sum;
for (auto & td : q) {
td.p *= inv_sum;
}
float u = uni(rng);
// Fallback to the last bucket to absorb residual probability from FP precision loss
llama_token draft_tok = q.back().id;
for (const auto & e : q) {
u -= e.p;
if (u <= 0.0f) {
draft_tok = e.id;
break;
}
}
// Both indices reuse the prompt's last output logits. The second entry
// is the bonus draw slot; it is always stale and its result is ignored.
const llama_tokens draft = { draft_tok };
const std::vector<std::vector<llama_token_data>> draft_q = { q };
const std::vector<int> idxs = { -1, -1 };
const std::vector<llama_token> result =
common_sampler_sample_and_accept_n_rejection(smpl_verify.get(), ctx_tgt, idxs, draft, draft_q, false);
token_counts[result[0]]++;
}
LOG("\n");
LOG("Empirical vs target distribution after %d trials:\n", n_trials);
bool ok = true;
size_t covered = 0;
auto get_count_for = [&](llama_token id) -> int {
const auto it = token_counts.find(id);
return it != token_counts.end() ? it->second : 0;
};
// Iterate over the target distribution
for (size_t i = 0; i < dist_tgt->size; ++i) {
const llama_token id = dist_tgt->data[i].id;
const float p_t = dist_tgt->data[i].p;
const int count = get_count_for(id);
const float p_e = (float) count / n_trials;
const float diff = std::fabs(p_t - p_e);
// We only deal with ids in dist_tgt, and if each trial result was one
// of dist_tgt's candidates then covered ends up equal to the number of
// trials. If not this inconsistency tells us how many trials produced
// a token outside of the target's distribution.
covered += count;
const float standard_error = std::sqrt(p_t * (1.0f - p_t) / n_trials);
const float tolerance = 6.0f * standard_error;
LOG("token %3d : target = %f, empirical = %f, diff = %f, tol = %f\n", id, p_t, p_e, diff, tolerance);
if (diff > tolerance) {
LOG_ERR("token %d: empirical probability %f too far from target %f (diff %f > tol %f)\n",
id, p_e, p_t, diff, tolerance);
ok = false;
}
}
if (covered != (size_t) n_trials) {
LOG_ERR("%zu of %d trials produced a token outside dist_tgt's candidates\n",
(size_t) n_trials - covered, n_trials);
ok = false;
}
if (!ok) {
return 1;
}
return 0;
}
struct parsed_args {
int n_trials;
std::vector<char *> argv;
};
static std::optional<parsed_args> parse_n_trials_arg(int argc, char ** argv, int default_n_trials) {
parsed_args result;
result.n_trials = default_n_trials;
result.argv.push_back(argv[0]);
for (int i = 1; i < argc; ++i) {
if (strcmp(argv[i], "--n-trials") == 0) {
if (i + 1 >= argc) {
LOG_ERR("%s: --n-trials requires a number argument\n", __func__);
return std::nullopt;
}
int value = std::atoi(argv[i + 1]);
if (value <= 0) {
LOG_ERR("%s: --n-trials requires a positive number argument\n", __func__);
return std::nullopt;
}
result.n_trials = value;
i++;
} else {
result.argv.push_back(argv[i]);
}
}
result.argv.push_back(nullptr);
return result;
}
int main(int argc, char ** argv) {
std::setlocale(LC_NUMERIC, "C");
common_params params;
params.n_ctx = 256;
common_init();
auto parsed = parse_n_trials_arg(argc, argv, 5000);
if (!parsed) {
return 1;
}
const int fargc = (int) parsed->argv.size() - 1;
if (!common_params_parse(fargc, parsed->argv.data(), params, LLAMA_EXAMPLE_COMMON)) {
return 1;
}
LOG("Using target model '%s'\n", params.model.path.c_str());
llama_backend_init();
const int result = test_speculative_rejection_sampling(params, parsed->n_trials);
if (result == 0) {
LOG("\nAll tests passed.\n");
} else {
LOG("\nSome tests failed.\n");
}
llama_backend_free();
return result;
}
+1
View File
@@ -67,6 +67,7 @@ add_library(mtmd
models/pockettts-seanet.cpp
models/pockettts-spkenc.cpp
models/pockettts-gen.cpp
models/soprano.cpp
models/step3vl.cpp
models/siglip.cpp
models/whisper-enc.cpp
+4
View File
@@ -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);
+8
View File
@@ -300,6 +300,12 @@
#define TN_A_GEN_WAV_SEANET_RES_CONV2 "a.gen.wav.seanet.blk.%d.res_conv2.%s"
#define TN_A_GEN_WAV_SEANET_SCALE_CONV "a.gen.wav.seanet.blk.%d.scale_conv.%s"
// soprano: Vocos decoder, the ConvNeXt blocks use TN_A_GEN_WAV_UP_*
#define TN_A_GEN_WAV_INPUT "a.gen.wav.input.%s"
#define TN_A_GEN_WAV_NORM "a.gen.wav.norm.%s"
#define TN_A_GEN_WAV_OUTPUT_NORM "a.gen.wav.output_norm.%s"
#define TN_A_GEN_WAV_OUTPUT "a.gen.wav.output.%s"
// cogvlm
#define TN_MM_POST_FC_NORM "mm.post_fc_norm.%s"
#define TN_MM_H_TO_4H "mm.up.%s"
@@ -547,6 +553,7 @@ enum projector_type {
PROJECTOR_TYPE_QWEN3TTS_GEN,
PROJECTOR_TYPE_POCKETTTS_SPKENC,
PROJECTOR_TYPE_POCKETTTS_GEN,
PROJECTOR_TYPE_SOPRANO,
PROJECTOR_TYPE_MUSE_GLIMMER,
PROJECTOR_TYPE_SENSENOVA_U1,
PROJECTOR_TYPE_COHERE2V,
@@ -623,6 +630,7 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
{ PROJECTOR_TYPE_QWEN3TTS_GEN, "qwen3tts_gen"},
{ PROJECTOR_TYPE_POCKETTTS_SPKENC, "pockettts_spkenc"},
{ PROJECTOR_TYPE_POCKETTTS_GEN, "pockettts_gen"},
{ PROJECTOR_TYPE_SOPRANO, "soprano"},
{ PROJECTOR_TYPE_MUSE_GLIMMER, "muse-glimmer"},
{ PROJECTOR_TYPE_SENSENOVA_U1, "sensenova_u1"},
{ PROJECTOR_TYPE_COHERE2V, "cohere2v"},
+24 -1
View File
@@ -29,11 +29,14 @@ enum patch_merge_type {
PATCH_MERGE_SPATIAL_UNPAD,
};
// all algos are Pillow-compatible (matching PIL.Image.resize output)
// the low bits select the algo, Pillow-compatible (matching PIL.Image.resize output)
// RESIZE_ALGO_FLAG_TORCH matches torchvision's antialiased resize of uint8 images on the CPU (int16 weights), bilinear and bicubic only
enum resize_algo {
RESIZE_ALGO_BILINEAR,
RESIZE_ALGO_BICUBIC,
RESIZE_ALGO_LANCZOS,
RESIZE_ALGO_FLAG_TORCH = 1 << 8,
};
// Padding style for img_tool::resize
@@ -207,6 +210,11 @@ struct clip_hparams {
int32_t custom_image_min_tokens = -1;
int32_t custom_image_max_tokens = -1;
void set_image_resize_algo(resize_algo algo, bool torch = false) {
GGML_ASSERT(!(torch && algo == RESIZE_ALGO_LANCZOS) && "torchvision has no lanczos");
image_resize_algo = static_cast<resize_algo>(torch ? algo | RESIZE_ALGO_FLAG_TORCH : algo);
}
void set_limit_image_tokens(int n_tokens_min, int n_tokens_max) {
const int patch_area = patch_size * patch_size * n_merge * n_merge;
image_min_pixels = (custom_image_min_tokens > 0 ? custom_image_min_tokens : n_tokens_min) * patch_area;
@@ -579,6 +587,18 @@ struct clip_code2wav {
ggml_tensor * dac_post_conv_b = nullptr;
};
struct clip_vocos {
ggml_tensor * input_w = nullptr;
ggml_tensor * input_b = nullptr;
ggml_tensor * norm_w = nullptr;
ggml_tensor * norm_b = nullptr;
ggml_tensor * output_norm_w = nullptr;
ggml_tensor * output_norm_b = nullptr;
ggml_tensor * output_w = nullptr;
ggml_tensor * output_b = nullptr;
std::vector<clip_code2wav::upsample_block> blocks;
};
struct clip_model {
clip_modality modality = CLIP_MODALITY_VISION;
projector_type proj_type = PROJECTOR_TYPE_MLP;
@@ -845,6 +865,9 @@ struct clip_model {
ggml_tensor * gen_upsample_w = nullptr; // depthwise convtr, frame rate -> encoder frame rate
std::vector<clip_layer> gen_tfm_layers; // mimi decoder_transformer
// soprano: Vocos decoder, hidden states -> STFT coefficients
clip_vocos vocos;
// cogvlm
ggml_tensor * mm_post_fc_norm_w = nullptr;
ggml_tensor * mm_post_fc_norm_b = nullptr;
+82 -3
View File
@@ -584,6 +584,36 @@ ggml_tensor * clip_graph::build_inp_raw(int channels) {
return inp_raw;
}
// qwen-vl style temporal merge: Conv3D(kt=2) split into 2 Conv2D
ggml_tensor * clip_graph::build_inp_temporal_merge() {
ggml_tensor * inp_raw = build_inp_raw();
GGML_ASSERT(img.nx() % (patch_size * 2) == 0);
GGML_ASSERT(img.ny() % (patch_size * 2) == 0);
const size_t nb1 = ggml_row_size(inp_raw->type, img.nx());
const size_t nb2 = ggml_row_size(inp_raw->type, img.nx() * img.ny());
if (n_batch == 1) {
// still image input
return ggml_add(ctx0,
ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_raw, patch_size, patch_size, 0, 0, 1, 1),
ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_raw, patch_size, patch_size, 0, 0, 1, 1));
} else if (n_batch == 2) {
// 2 frames input (video input)
ggml_tensor * inp_0 = ggml_view_3d(ctx0, inp_raw,
img.nx(), img.ny(), 3, nb1, nb2, 0);
ggml_tensor * inp_1 = ggml_view_3d(ctx0, inp_raw,
img.nx(), img.ny(), 3, nb1, nb2,
nb2 * 3); // move to the second frame
return ggml_add(ctx0,
ggml_conv_2d(ctx0, model.patch_embeddings_0, inp_0, patch_size, patch_size, 0, 0, 1, 1),
ggml_conv_2d(ctx0, model.patch_embeddings_1, inp_1, patch_size, patch_size, 0, 0, 1, 1));
} else {
GGML_ASSERT(false && "n_batch > 2 is not supported");
}
}
static std::string get_attn_mask_name(int idx) {
return idx == 0 ? "attn_mask" : "attn_mask_" + std::to_string(idx);
}
@@ -1168,6 +1198,11 @@ static std::unique_ptr<clip_graph> clip_get_graph_builder(clip_ctx * ctx, const
const float top_p = params ? params->top_p : 1.0f;
builder = std::make_unique<clip_graph_qwen3tts_gen>(ctx, img, gen_process, top_k, top_p);
} break;
case PROJECTOR_TYPE_SOPRANO:
{
const int n_frames = params && params->feats ? (int) (params->feats->size() / ctx->model.hparams.projection_dim) : 2;
builder = std::make_unique<clip_graph_soprano>(ctx, img, n_frames);
} break;
case PROJECTOR_TYPE_YOUTUVL:
{
builder = std::make_unique<clip_graph_youtuvl>(ctx, img);
@@ -1584,12 +1619,13 @@ struct clip_model_loader {
// default for older GGUFs
std::string resize_algo = "bilinear";
get_string(KEY_IMAGE_RESIZE_ALGO, resize_algo, false);
// Lfm2VlImageProcessor and d1-omni vision.py resize bilinear and bicubic with torchvision
if (resize_algo == "bilinear") {
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.set_image_resize_algo(RESIZE_ALGO_BILINEAR, true);
} else if (resize_algo == "bicubic") {
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.set_image_resize_algo(RESIZE_ALGO_BICUBIC, true);
} else if (resize_algo == "lanczos") {
hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
hparams.set_image_resize_algo(RESIZE_ALGO_LANCZOS);
} else {
throw std::runtime_error(string_format("%s: unsupported image resize algo: %s\n", __func__, resize_algo.c_str()));
}
@@ -2073,6 +2109,12 @@ struct clip_model_loader {
hparams.flow_n_step = 1;
hparams.gen_eos_threshold = -4.0f;
} break;
case PROJECTOR_TYPE_SOPRANO:
{
// SopranoDecoder defaults, n_fft comes from the head tensor
hparams.audio_sample_rate = 32000;
hparams.audio_hop_len = 512;
} break;
case PROJECTOR_TYPE_PADDLEOCR:
{
hparams.n_merge = 2;
@@ -2452,6 +2494,7 @@ struct clip_model_loader {
model.proj_type != PROJECTOR_TYPE_GEMMA3NV &&
model.proj_type != PROJECTOR_TYPE_QWEN3TTS_SPKENC &&
model.proj_type != PROJECTOR_TYPE_POCKETTTS_GEN &&
model.proj_type != PROJECTOR_TYPE_SOPRANO &&
model.proj_type != PROJECTOR_TYPE_MUSE_SPARK_AUDIO);
// layers
@@ -3391,6 +3434,31 @@ struct clip_model_loader {
c2w.dac_post_conv_b = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV, "bias"));
}
} break;
case PROJECTOR_TYPE_SOPRANO:
{
auto & v = model.vocos;
v.input_w = get_tensor(string_format(TN_A_GEN_WAV_INPUT, "weight"));
v.input_b = get_tensor(string_format(TN_A_GEN_WAV_INPUT, "bias"));
v.norm_w = get_tensor(string_format(TN_A_GEN_WAV_NORM, "weight"));
v.norm_b = get_tensor(string_format(TN_A_GEN_WAV_NORM, "bias"));
v.output_norm_w = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT_NORM, "weight"));
v.output_norm_b = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT_NORM, "bias"));
v.output_w = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT, "weight"));
v.output_b = get_tensor(string_format(TN_A_GEN_WAV_OUTPUT, "bias"));
v.blocks.resize(hparams.n_layer);
for (int il = 0; il < hparams.n_layer; ++il) {
auto & b = v.blocks[il];
b.dwconv_w = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "weight"));
b.dwconv_b = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "bias"));
b.norm_w = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM, il, "weight"));
b.norm_b = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM, il, "bias"));
b.pw1_w = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1, il, "weight"));
b.pw1_b = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1, il, "bias"));
b.pw2_w = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2, il, "weight"));
b.pw2_b = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2, il, "bias"));
b.gamma = get_tensor(string_format(TN_A_GEN_WAV_UP_GAMMA, il));
}
} break;
case PROJECTOR_TYPE_VOXTRAL:
case PROJECTOR_TYPE_VOXTRAL_RT:
{
@@ -5679,6 +5747,10 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
set_input_f32("inp_noise", noise);
}
} break;
case PROJECTOR_TYPE_SOPRANO:
{
set_input_f32("inp_feats", *params->feats);
} break;
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
{
@@ -6329,6 +6401,11 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
out_audio.resize(ggml_nelements(audio));
ggml_backend_tensor_get(audio, out_audio.data(), 0, ggml_nbytes(audio));
if (model.proj_type == PROJECTOR_TYPE_SOPRANO) {
std::vector<float> spectrum = std::move(out_audio);
clip_graph_soprano::decode_spectrum(model, spectrum, out_audio);
}
// drop the tail audio that comes from the code-0 rear padding
const int64_t n_codes = params->codes ? model.gen_code_head_w->ne[2] + 1 : 0;
const int64_t n_frames_w = hparams.wav_tfm_swa;
@@ -6536,6 +6613,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
return ctx->model.spk_proj_w->ne[1];
case PROJECTOR_TYPE_POCKETTTS_GEN:
return ctx->model.gen_input_lin_w->ne[1];
case PROJECTOR_TYPE_SOPRANO:
return ctx->model.hparams.projection_dim;
case PROJECTOR_TYPE_PARAKEET:
return ctx->model.mm_1_w->ne[1];
default:
+1 -1
View File
@@ -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
{
+9 -1
View File
@@ -42,7 +42,6 @@ struct clip_graph_deepseek4v : clip_graph {
struct clip_graph_qwen2vl : clip_graph {
clip_graph_qwen2vl(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
ggml_tensor * build_inp_with_temporal_merge();
};
struct clip_graph_qwen3vl : clip_graph_qwen2vl {
@@ -420,6 +419,15 @@ std::vector<c2w_state_slot> list_c2w_state_slots(const clip_hparams & hparams, c
// same, for the streaming mimi decoder (pocket-tts GEN_WAV)
std::vector<c2w_state_slot> list_pockettts_state_slots(const clip_hparams & hparams, const clip_model & model);
struct clip_graph_soprano : clip_graph {
clip_graph_soprano(clip_ctx * ctx, const clip_image_f32 & img, int n_frames) : clip_graph(ctx, img), n_frames(n_frames) {}
ggml_cgraph * build() override;
static void decode_spectrum(const clip_model & model, const std::vector<float> & spectrum, std::vector<float> & pcm);
// SopranoDecoder default
static constexpr int upscale = 4;
int n_frames;
};
struct clip_graph_kimik25 : clip_graph {
clip_graph_kimik25(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
+1 -30
View File
@@ -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
{
+1 -1
View File
@@ -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
{
+78
View File
@@ -0,0 +1,78 @@
#include "models.h"
#include "../mtmd-audio.h"
// Soprano Vocos: normalized Qwen3 hidden states -> STFT coefficients.
ggml_cgraph * clip_graph_soprano::build() {
const auto & v = model.vocos;
const int n_inp = hparams.projection_dim;
const int n_up = upscale * (n_frames - 1) + 1;
ggml_tensor * cur = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_inp, n_frames);
ggml_set_name(cur, "inp_feats");
ggml_set_input(cur);
cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
cur = ggml_interpolate(ctx0, cur, n_up, n_inp, 1, 1, GGML_SCALE_MODE_BILINEAR | GGML_SCALE_FLAG_ALIGN_CORNERS);
if (ggml_n_dims(v.input_w) == 3) {
const int k = v.input_w->ne[0];
cur = ggml_conv_1d(ctx0, v.input_w, cur, 1, k / 2, 1);
cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
} else {
cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
cur = build_mm(v.input_w, cur);
}
cur = ggml_add(ctx0, cur, v.input_b);
cur = build_norm(cur, v.norm_w, v.norm_b, NORM_TYPE_NORMAL, eps, -1);
for (size_t il = 0; il < v.blocks.size(); ++il) {
const auto & b = v.blocks[il];
const int k = b.dwconv_w->ne[0];
ggml_tensor * residual = cur;
cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
cur = ggml_conv_1d_dw(ctx0, b.dwconv_w, cur, 1, k / 2, 1);
cur = ggml_cont(ctx0, ggml_transpose(ctx0, cur));
cur = ggml_add(ctx0, cur, b.dwconv_b);
cur = build_norm(cur, b.norm_w, b.norm_b, NORM_TYPE_NORMAL, eps, il);
cur = ggml_add(ctx0, build_mm(b.pw1_w, cur), b.pw1_b);
cur = ggml_gelu_erf(ctx0, cur);
cur = ggml_add(ctx0, build_mm(b.pw2_w, cur), b.pw2_b);
cur = ggml_add(ctx0, residual, ggml_mul(ctx0, cur, b.gamma));
cb(cur, "vocos_block", il);
}
cur = build_norm(cur, v.output_norm_w, v.output_norm_b, NORM_TYPE_NORMAL, eps, -1);
cur = ggml_add(ctx0, build_mm(v.output_w, cur), v.output_b);
cb(cur, "vocos_head", -1);
// the head holds n_fft/2+1 log-magnitudes followed by n_fft/2+1 phases
const int n_bins = v.output_w->ne[1] / 2;
ggml_tensor * mag = ggml_view_2d(ctx0, cur, n_bins, n_up, cur->nb[1], 0);
ggml_tensor * phase = ggml_view_2d(ctx0, cur, n_bins, n_up, cur->nb[1], n_bins * sizeof(float));
mag = ggml_clamp(ctx0, ggml_exp(ctx0, ggml_cont(ctx0, mag)), 0.0f, 100.0f);
phase = ggml_cont(ctx0, phase);
ggml_tensor * re = ggml_mul(ctx0, mag, ggml_cos(ctx0, phase));
ggml_tensor * im = ggml_mul(ctx0, mag, ggml_sin(ctx0, phase));
re = ggml_reshape_3d(ctx0, re, 1, n_bins, n_up);
im = ggml_reshape_3d(ctx0, im, 1, n_bins, n_up);
cur = ggml_concat(ctx0, re, im, 0);
ggml_set_name(cur, "out_audio");
ggml_set_output(cur);
ggml_build_forward_expand(gf, cur);
return gf;
}
void clip_graph_soprano::decode_spectrum(const clip_model & model, const std::vector<float> & spectrum, std::vector<float> & pcm) {
const int n_frame_vals = model.vocos.output_w->ne[1];
const int n_fft = n_frame_vals - 2;
const int hop_length = model.hparams.audio_hop_len;
const size_t n_frames = spectrum.size() / n_frame_vals;
mtmd_audio_streaming_istft istft(n_fft, hop_length);
pcm.clear();
for (size_t i = 0; i < n_frames; ++i) {
auto frame = istft.process_frame(spectrum.data() + i * n_frame_vals);
pcm.insert(pcm.end(), frame.begin(), frame.end());
}
auto tail = istft.flush();
pcm.insert(pcm.end(), tail.begin(), tail.end());
// the shared ISTFT removes (n_fft - hop) / 2 samples, torch.istft(center=True) removes n_fft / 2
pcm.erase(pcm.begin(), pcm.begin() + hop_length / 2);
pcm.resize((n_frames - 1) * hop_length);
}
+187
View File
@@ -1076,12 +1076,199 @@ private:
bool wav_header_sent = false;
};
// Soprano uses the hidden state that predicts each non-EOS token, including the prompt's last state.
class soprano_gen_audio_pipeline : public mtmd_gen_audio_pipeline {
public:
using mtmd_gen_audio_pipeline::mtmd_gen_audio_pipeline;
void reset() override {
prompt.clear();
features.clear();
h_next.clear();
audio.clear();
wav.clear();
pos = 0;
prompt_pos = 0;
decoded = false;
sent = false;
}
int32_t set_input(const mtmd_helper_gen_audio_inp * inp) override {
reset();
if (inp->speaker_ref || !inp->prompt || inp->prompt_len == 0 || inp->prompt_len > INT32_MAX ||
inp->seq_id < 0 || (uint32_t) inp->seq_id >= llama_n_seq_max(lctx)) {
LOG_ERR("mtmd_helper_gen_audio: soprano requires text, without a speaker reference\n");
return 1;
}
seq_id = inp->seq_id;
out_type = inp->out_type;
stream = inp->stream;
const llama_token text = find_special_token(vocab, "[TEXT]");
const llama_token start = find_special_token(vocab, "[START]");
const llama_token stop = find_special_token(vocab, "[STOP]");
if (text == LLAMA_TOKEN_NULL || start == LLAMA_TOKEN_NULL || stop == LLAMA_TOKEN_NULL) {
LOG_ERR("mtmd_helper_gen_audio: incompatible soprano vocabulary\n");
return 1;
}
const int required = llama_tokenize(vocab, inp->prompt, inp->prompt_len, nullptr, 0, false, false);
if (required >= 0 || (uint32_t) (-required + 3) >= llama_n_ctx_seq(lctx)) {
LOG_ERR("mtmd_helper_gen_audio: soprano prompt does not fit in the context\n");
return 1;
}
const int n = -required;
prompt.resize(n + 3);
prompt[0] = stop;
prompt[1] = text;
if (llama_tokenize(vocab, inp->prompt, inp->prompt_len, prompt.data() + 2, n, false, false) != n) {
return 1;
}
prompt.back() = start;
if (!llama_memory_seq_rm(llama_get_memory(lctx), seq_id, 0, -1)) {
LOG_ERR("mtmd_helper_gen_audio: cannot reset soprano sequence\n");
return 1;
}
if (!init_gen_state(inp->seed)) {
return 1;
}
return 0;
}
int32_t step_prompt(int32_t n_batch) override {
if (n_batch <= 0 || prompt.empty()) {
return -1;
}
const int n = std::min(n_batch, (int) prompt.size() - prompt_pos);
if (n == 0) {
return 0;
}
if (!decode(prompt.data() + prompt_pos, n)) {
return -1;
}
prompt_pos += n;
return (int) prompt.size() - prompt_pos;
}
int32_t step_gen(llama_token sampled, const float * h_state_in, const float ** h_state_out, bool * out_stop) override {
*h_state_out = nullptr;
*out_stop = false;
if (sampled < 0 || sampled >= llama_vocab_n_tokens(vocab) || !h_state_in || prompt_pos != (int) prompt.size() || prompt.empty()) {
return 1;
}
if (llama_vocab_is_eog(vocab, sampled) || (uint32_t) pos >= llama_n_ctx_seq(lctx)) {
*out_stop = true;
return 0;
}
features.insert(features.end(), h_state_in, h_state_in + n_embd);
if (!decode(&sampled, 1)) {
features.resize(features.size() - n_embd);
return 1;
}
const float * h = llama_get_embeddings_ith(lctx, -1);
if (!h) {
return 1;
}
h_next.assign(h, h + n_embd);
*h_state_out = h_next.data();
return 0;
}
int32_t get_output(int32_t * sample_rate, const char ** data, size_t * data_len, int64_t * n_samples) override {
*sample_rate = info.sample_rate;
*data_len = 0;
if (n_samples) { *n_samples = 0; }
if (stream) {
// the decoder needs the whole utterance, so the audio comes out in one chunk after flush()
if (!decoded || sent) {
return 0;
}
sent = true;
} else if (!decode_audio()) {
return 1;
}
if (n_samples) { *n_samples = audio.size(); }
if (out_type == MTMD_HELPER_GEN_AUDIO_OUTTYPE_PCM) {
*data = (const char *) audio.data();
*data_len = audio.size() * sizeof(float);
} else {
wav.clear();
if (!write_wav16(wav, audio, info.sample_rate)) {
return 1;
}
*data = wav.data();
*data_len = wav.size();
}
return 0;
}
// the whole utterance is decoded at once
int32_t flush() override {
return decode_audio() ? 0 : 1;
}
private:
bool decode_audio() {
if (decoded) {
return true;
}
if (features.size() < 2 * (size_t) n_embd) {
LOG_ERR("mtmd_helper_gen_audio: soprano needs at least two generated frames\n");
return false;
}
mtmd_gen_inp inp = mtmd_gen_inp_default(mctx);
inp.type = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
inp.feats = features.data();
inp.n_feats = features.size();
mtmd_gen_out out{};
if (mtmd_gen_audio_process(mctx, gen_state.get(), &inp, &out) != 0) {
return false;
}
audio.assign(out.audio, out.audio + out.n_samples);
decoded = true;
return true;
}
bool decode(llama_token * tokens, int n) {
llama_batch batch = llama_batch_init(n, 0, 1);
batch.n_tokens = n;
for (int i = 0; i < n; ++i) {
batch.token[i] = tokens[i];
batch.pos[i] = pos + i;
batch.n_seq_id[i] = 1;
batch.seq_id[i][0] = seq_id;
batch.logits[i] = i == n - 1;
}
const int ret = llama_decode(lctx, batch);
llama_batch_free(batch);
if (ret != 0) {
LOG_ERR("mtmd_helper_gen_audio: soprano backbone decode failed\n");
return false;
}
pos += n;
return true;
}
llama_seq_id seq_id = 0;
llama_pos pos = 0;
int prompt_pos = 0;
std::vector<llama_token> prompt;
std::vector<float> features;
std::vector<float> h_next;
std::vector<float> audio;
std::vector<char> wav;
mtmd_helper_gen_audio_outtype out_type = MTMD_HELPER_GEN_AUDIO_OUTTYPE_WAV;
bool stream = false;
bool decoded = false;
bool sent = false;
};
static std::unique_ptr<mtmd_gen_audio_pipeline> make_pipeline(llama_context * lctx, mtmd_context * mctx) {
switch (mtmd_gen_audio_get_info(mctx).type) {
case MTMD_GEN_AUDIO_TYPE_QWEN3TTS:
return std::unique_ptr<mtmd_gen_audio_pipeline>(new qwen3tts_gen_audio_pipeline(lctx, mctx));
case MTMD_GEN_AUDIO_TYPE_POCKETTTS:
return std::unique_ptr<mtmd_gen_audio_pipeline>(new pockettts_gen_audio_pipeline(lctx, mctx));
case MTMD_GEN_AUDIO_TYPE_SOPRANO:
return std::unique_ptr<mtmd_gen_audio_pipeline>(new soprano_gen_audio_pipeline(lctx, mctx));
default:
return nullptr;
}
+25 -4
View File
@@ -205,6 +205,8 @@ private:
// 1. Separable filtering: horizontal pass followed by vertical pass
// 2. Pre-computes normalized filter coefficients for each output pixel
// 3. Fixed-point integer arithmetic (22 fractional bits) for speed and determinism
//
// RESIZE_ALGO_FLAG_TORCH is the same, except for the int16 weights of torchvision (aten/src/ATen/native/cpu/UpSampleKernel.cpp)
static bool resize_pillow(
const clip_image_u8 & img,
clip_image_u8 & dst,
@@ -215,6 +217,13 @@ private:
// This allows encoding fractional weights as integers: weight * 2^22
const int PRECISION_BITS = 32 - 8 - 2;
const bool int16_weights = algo & RESIZE_ALGO_FLAG_TORCH;
algo = static_cast<resize_algo>(algo & ~RESIZE_ALGO_FLAG_TORCH);
if (int16_weights && algo == RESIZE_ALGO_LANCZOS) {
GGML_ABORT("unsupported resize algo: torchvision has no lanczos");
}
// Filter support radius
double filter_support;
switch (algo) {
@@ -222,11 +231,11 @@ private:
case RESIZE_ALGO_BICUBIC: filter_support = 2.0; break;
case RESIZE_ALGO_LANCZOS: filter_support = 3.0; break;
default:
throw std::runtime_error("Unsupported resize algorithm");
GGML_ABORT("unsupported resize algo: %d", (int) algo);
}
// Returns filter weight for distance x from pixel center
// Note: for bicubic, Pillow uses a = -0.5 while GGML/PyTorch use a = -0.75
// Note: for bicubic, Pillow (and PyTorch with antialias) uses a = -0.5 while GGML (and PyTorch without antialias) uses a = -0.75
auto resample_filter = [algo](double x) -> double {
if (algo == RESIZE_ALGO_LANCZOS) {
if (-3.0 <= x && x < 3.0) {
@@ -351,12 +360,24 @@ private:
// Formula: int32 = round(float * 2^PRECISION_BITS)
weights.resize((size_t) outSize * ksize);
const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS
// torchvision rounds the weights to int16, with as many fractional bits as the largest weight allows;
// scaled back to PRECISION_BITS, they give the same output in the passes below
int precision = PRECISION_BITS;
if (int16_weights) {
const double weight_max = *std::max_element(pre_weights.begin(), pre_weights.end());
for (precision = 0; precision < PRECISION_BITS; precision++) {
if (static_cast<int>(0.5 + weight_max * (1 << (precision + 1))) >= (1 << 15)) {
break;
}
}
}
const double fxp_scale = std::ldexp(1.0, precision); // 1.0 * 2^precision
for (size_t i = 0; i < (size_t) outSize * ksize; i++) {
// Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
weights[i] = static_cast<int32_t>(rounded);
weights[i] = static_cast<int32_t>(rounded) * (1 << (PRECISION_BITS - precision));
}
return ksize;
+2
View File
@@ -31,6 +31,7 @@ struct mtmd_image_preprocessor {
const clip_hparams & hparams;
mtmd_image_preprocessor(const clip_ctx * ctx): hparams(*clip_get_hparams(ctx)) {}
mtmd_image_preprocessor(const clip_hparams & hparams): hparams(hparams) {}
virtual ~mtmd_image_preprocessor() = default;
virtual mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const = 0;
@@ -59,6 +60,7 @@ struct mtmd_image_preprocessor {
*/
struct mtmd_image_preprocessor_llava_uhd : mtmd_image_preprocessor {
mtmd_image_preprocessor_llava_uhd(const clip_ctx * ctx) : mtmd_image_preprocessor(ctx) {}
mtmd_image_preprocessor_llava_uhd(const clip_hparams & hparams) : mtmd_image_preprocessor(hparams) {}
mtmd_image_preproc_out preprocess(const clip_image_u8 & img) const override;
struct slice_coordinates {
+23 -2
View File
@@ -624,7 +624,9 @@ struct mtmd_context {
ctx_v = res.ctx_v;
ctx_a = res.ctx_a;
ctx_gen_a = res.ctx_gen_a;
if (!ctx_v && !ctx_a) {
// soprano has no speaker encoder, its mmproj only holds the decoder
const bool gen_only = ctx_gen_a && clip_get_projector_type(ctx_gen_a) == PROJECTOR_TYPE_SOPRANO;
if (!ctx_v && !ctx_a && !gen_only) {
throw std::runtime_error(string_format("Failed to load CLIP model from %s\n", mmproj_fname));
}
@@ -641,7 +643,8 @@ struct mtmd_context {
// since we already validate n_embd of vision and audio mmproj,
// we can safely assume that they are the same
int n_embd_clip = clip_n_mmproj_embd(ctx_v ? ctx_v : ctx_a);
// a gen-only mmproj is validated below
int n_embd_clip = ctx_v || ctx_a ? clip_n_mmproj_embd(ctx_v ? ctx_v : ctx_a) : n_embd_text;
if (n_embd_text > 0 && n_embd_text != n_embd_clip) {
throw std::runtime_error(string_format(
"mismatch between text model (n_embd = %d) and mmproj (n_embd = %d)\n"
@@ -2082,6 +2085,10 @@ mtmd_gen_audio_info mtmd_gen_audio_get_info(const mtmd_context * ctx) {
info.type = MTMD_GEN_AUDIO_TYPE_POCKETTTS;
info.sample_rate = 24000;
break;
case PROJECTOR_TYPE_SOPRANO:
info.type = MTMD_GEN_AUDIO_TYPE_SOPRANO;
info.sample_rate = 32000;
break;
default:
info.type = MTMD_GEN_AUDIO_TYPE_NONE;
break;
@@ -2109,6 +2116,9 @@ mtmd_gen_inp mtmd_gen_inp_default(const mtmd_context * ctx) {
inp.top_p = 1.0f;
inp.temp = 0.7f;
break;
case PROJECTOR_TYPE_SOPRANO:
inp.type = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
break;
default:
break;
}
@@ -2170,6 +2180,14 @@ static int32_t mtmd_gen_audio_process_impl(mtmd_context * ctx, mtmd_gen_state *
*out = {};
if (clip_get_projector_type(ctx_clip) == PROJECTOR_TYPE_SOPRANO) {
const size_t n_embd = (size_t) clip_n_mmproj_embd(ctx_clip);
if (inp->type != MTMD_GEN_PROCESS_TYPE_GEN_WAV || !inp->feats || inp->n_feats < 2 * n_embd || inp->n_feats % n_embd != 0 || inp->codes) {
LOG_ERR("%s: soprano requires at least 2 frames of features, without codes\n", __func__);
return 1;
}
}
if (inp->type == MTMD_GEN_PROCESS_TYPE_GEN_CODE) {
const size_t n_embd = (size_t) clip_n_mmproj_embd(ctx_clip);
@@ -3046,6 +3064,9 @@ mtmd_memory_usage mtmd_get_memory_usage(const char * mmproj_fname,
if (ctx->ctx_a) {
merge(ctx->ctx_a);
}
if (ctx->ctx_gen_a) {
merge(ctx->ctx_gen_a);
}
mtmd_memory_usage res;
res.backend_mem_usage = std::move(total_mem);
res.image_max_tokens = ctx->ctx_v ? clip_get_image_max_tokens(ctx->ctx_v) : -1;
+1
View File
@@ -427,6 +427,7 @@ enum mtmd_gen_audio_type {
MTMD_GEN_AUDIO_TYPE_NONE, // not supported
MTMD_GEN_AUDIO_TYPE_QWEN3TTS,
MTMD_GEN_AUDIO_TYPE_POCKETTTS,
MTMD_GEN_AUDIO_TYPE_SOPRANO,
};
struct mtmd_gen_audio_info {
+1 -1
View File
@@ -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
+63
View File
@@ -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);
}
+95
View File
@@ -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
};
-25
View File
@@ -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 {
+15 -165
View File
@@ -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;
+4 -36
View File
@@ -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
+473 -300
View File
@@ -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);
}
+25
View File
@@ -59,3 +59,28 @@ The [upstream repository](https://huggingface.co/kyutai/pocket-tts) holds one co
python convert_hf_to_gguf.py path/to/pocket-tts/languages/english --outfile pocket-tts.gguf
python convert_hf_to_gguf.py path/to/pocket-tts/languages/english --mmproj --outfile mmproj-pocket-tts.gguf
```
## Soprano
Available params:
- `--tts-speaker-file` is not supported; the model uses a fixed voice
- Note: `lang` is not used. Supply normalized English text
Example usage:
```sh
llama-tts -m soprano.gguf \
-mm mmproj-soprano.gguf \
-p "Hello world" \
--temp 0 \
--output out.wav
```
**Note for GGUF conversion:**
Both [Soprano-80M](https://huggingface.co/ekwek/Soprano-80M) and [Soprano-1.1-80M](https://huggingface.co/ekwek/Soprano-1.1-80M) are supported. Convert both files from the same model directory, including `decoder.pth`:
```sh
python convert_hf_to_gguf.py path/to/Soprano-1.1-80M --outtype f32 --outfile soprano.gguf
python convert_hf_to_gguf.py path/to/Soprano-1.1-80M --mmproj --outtype f16 --outfile mmproj-soprano.gguf
```
@@ -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;
@@ -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>
@@ -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}
@@ -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>
@@ -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>
@@ -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';
@@ -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>

Some files were not shown because too many files have changed in this diff Show More