* granite-switch: add llama.cpp backend (POC, CPU)
New "granite-switch" architecture: a dense, all-attention Granite-4.1
model with N embedded LoRA adapters selected per-token by control tokens.
- gguf-py schema (arch, KV keys, stacked LoRA tensor names) + writer helpers
- conversion/granite.py: GraniteSwitchModel converter (stacks N adapters +
zero base slot into per-projection A/B tensors; emits switch metadata)
- C++ arch registration (llama-arch.{h,cpp}, llama-model.{h,cpp})
- src/models/granite_switch.cpp: load + per-token switched-LoRA graph via
ggml_mul_mat_id over stacked tensors; sticky per-token index + control-token
substitution in llm_graph_input_switch::set_input
- llm_graph_input_switch in src/models/models.h
Runs end-to-end on CPU: convert 3b checkpoint (842 tensors, stacked dim 13)
and generate on both base and control-token paths. Sticky switch state is
single-sequence (POC); full multi-sequence machinery is a follow-up.
* granite-switch: add Mac (Metal) build + mid-sequence switch demo script
Self-contained script to build llama.cpp on Apple Silicon (Metal),
convert the composed 3b checkpoint, and run the crisp mid-sequence
adapter-switch demos verified on Vela:
- answerability: <|answerability|> mid-seq -> "unanswerable"
- query_rewrite: <|query_rewrite|> mid-seq -> {"rewritten_question": ...}
Each demo runs the same prompt twice, differing only by a control token
placed before the assistant turn, so the per-token switch is visible.
* granite-switch mac demo: add -no-cnv so each run is one-shot
The composed model ships a chat template, so llama-completion auto-enables
interactive conversation mode and halts at a `>` prompt after generating,
stalling the script. -no-cnv disables conversation mode: generate once from
the raw prompt and exit (also prints special tokens, making the switch visible).
* granite-switch: replace global sticky index with in-graph router attention
The POC computed the per-token adapter index on the CPU and carried it
across ubatches in ONE global `mutable int32_t poc_sticky_index`, reset
only when a ubatch contained sequence position 0. That global had two
problems:
1. Concurrency: with multiple sequences in a batch it was last-writer-
wins — one sequence's adapter leaked into the others.
2. Multi-turn: an interactive `ollama run` chat continues one KV cache,
so turn 2 never saw position 0 and the index never reset — the
adapter stayed stuck on across turns.
Port the vLLM/HF backend mechanism faithfully: a single-head causal
"router" attention recovers the adapter index in-graph. Per token, only
dim 0 carries signal — Q[0]=1, K[0]=+gain for a control token / -gain
otherwise, V[0]=adapter slot / 0 — and the causal softmax over the single
visible control token recovers that adapter's slot (readback =
clamp(round(V[0]), 0, n_adapters)). gain=15 matches config.py and is
F16-safe (no F32 cache).
The router's K/V live in the model KV cache at an extra layer
R == hparams.router_layer (== n_layer). We bump n_layer_all to n_real+1
so the cache allocator gives the router its own per-sequence slot, and
set n_layer_nextn=1 so n_layer() stays n_real — the decoder loop and
tensor loading are untouched and never reference layer R. The router K is
exempted from the k-shift RoPE loop (its dim-0 value is a literal
magnitude, not a rotation).
Because the selection now lives in the per-sequence KV cache, CONCURRENT
requests are isolated for free (problem 1 fixed; verified by
scratch/concurrent_switch_test.cpp). set_input becomes stateless pure
per-token maps; the global is gone.
Single-switch contract / known limitation, identical to vLLM & HF: the
gain is flat (no recency), so within one sequence there is no mechanism to
revert to base mid-sequence — once an adapter fires it stays on until that
sequence ends (problem 2 is therefore NOT fixed by a faithful copy; vLLM/HF
avoid it only because each served request is a fresh sequence). A client
continuing one KV cache across turns must start a fresh sequence per turn,
or opt into a recency-biased router (a deliberate divergence, not done
here). Documented in granite_switch.cpp and asserted by
scratch/multiturn_leak_test.cpp.
Verified (CPU): both demos unchanged (answerability -> "unanswerable",
query_rewrite -> rewritten query); concurrent two-sequence isolation
passes; multi-turn carry-over matches the vLLM/HF contract.
* granite-switch: drop scratch tests and mac demo for upstream PR
Remove the local-only development artifacts that should not ship in the
upstream PR:
- granite-switch-mac-demo.sh (local Metal build + demo driver)
- scratch/concurrent_switch_test.cpp
- scratch/multiturn_leak_test.cpp
Also drop the now-dangling reference to the scratch tests from the
granite_switch.cpp header comment. Leaves only the core architecture
support (conversion, gguf constants, llama-arch/model/kv-cache, and the
granite_switch graph).
* granite-switch: trim comments to match native llama.cpp style
* granite-switch: trim conversion comments to match native style
* granite-switch: drop unused adapter_ranks metadata
* granite-switch: rename arch to graniteswitch and drop obid alias
* granite-switch: fix non-ASCII comments and document router gain assumption
* granite-switch: drop section comments from constants.py to match native style
* granite-switch: add functional tensor block comments matching Granite4 Vision style
* granite-switch: clarify n_expert_used comment
State the actual constraint: mul_mat_id needs n_expert_used == 1, and
since the GGUF carries expert_count = 0 the generic loader's
n_expert == 0 => n_expert_used == 0 assertion has already passed by the
time load_arch_hparams runs, so it is forced to 1 here.
* granite-switch: note n_layer_nextn reuse has no MTP
The router carving reuses n_layer_nextn, normally the MTP/next-token
count. Clarify in the comment that it is borrowed here purely as the
trailing-layers lever and that there is no MTP head, to spare readers
the double-take.
* granite-switch: rename source file and apply review nits
* granite-switch: don't force LoRA tensors to F16, follow --outtype instead
* granite-switch: drop redundant _permute_qk wrapper, call LlamaModel.permute directly
* granite-switch: read router gain from GGUF (control_token_gain) instead of hardcoding 15.0
* granite-switch: derive n_slots()
* granite-switch: move llm_graph_input_switch into granite-switch.cpp
* granite-switch: cut AI-style narration comments
* granite-switch: collapse multi-line comments
* granite-switch: rename control_token_* maps to adapter_token_*
* granite-switch: cut noise comments
* granite-switch: rename embedded LoRA tensors to <base>.lora_a/lora_b
* granite-switch: GGML_ASSERT token input to avoid UB on embeddings
* granite-switch: TODO for raw embedding input support
* granite-switch: collapse LoRA tensor constants to .lora_a/.lora_b suffix
* granite-switch: drop n_expert_used hack, guard mul_mat_id buft probe
* granite-switch: stop forcing dense expert counts, read from config
* granite-switch: renamed control_token_gain metadata key to router_gain
* granite-switch: trim header comments to match native style
* granite-switch: collapse LoRA tensors to base name + suffix
* granite-switch: inline suffix checks in tensor op resolution
* granite-switch: drop switch-lora struct comment
* granite-switch: guard router layer index and inline n_slots
* granite-switch: group adapter metadata under {arch}.adapters.* namespace
* granite-switch: add hparams.has_rope(il) for KV-shift rope skipping
* granite-switch: skip arch in test-llama-archs (adapter fixture missing, TODO)
* granite-switch: Keys.Adapters namespace + simplify n_slots
* granite-switch: validate substitute token ids against n_vocab
* granite-switch: bound adapter count and lora rank from GGUF
* granite-switch: reject MTP context type when router_layer is set
* granite-switch: throw on bad adapter metadata instead of GGML_ASSERT
* granite-switch: use ASCII +/- in router K signal comment
* granite-switch: document n_layer_nextn repurpose and its leak points
* granite-switch: gate lora_a/lora_b op mapping on router_layer
* granite-switch: label all three preview model sizes
docker info only proves the daemon answers, so the Windows CI passes
the check and then dies trying to run a linux image. The hosted
Windows runners cannot run one: GitHub states the VMs are not enabled
for nested virtualization and will not be, since they already sit one
level deep and the hypervisor does not support more levels
(https://github.com/orgs/community/discussions/25491). Probing the
image itself skips those tests there, and pulls it before the server
waits for the container id.
The saver called add_kv with LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH twice, the
second time passing n_ff_chexp. gguf_set_val_u32 removes-then-appends, so the second
call clobbers the first: the saved shared_feed_forward_length ends up as n_ff_chexp
(0 for every arch except GroveMoE), and expert_chunk_feed_forward_length is never
written at all.
So a save->load roundtrip of any MoE model with a shared expert loses n_ff_shexp. On
reload the arch falls back to n_ff for the shexp tensor shape, that no longer matches
the saved tensor, and the model FAILS to load. Hits qwen2moe, qwen3-next, granite-moe,
hunyuan-moe, ernie4.5, bailingmoe2, nemotron-h, and the other shared-expert MoEs.
Fix: the second call writes LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH.
test-llama-archs: set expert_shared_feed_forward_length to a value distinct from n_ff
in the MoE setup so the roundtrip exercises it. Without the fix the reload fails on a
shexp tensor-shape mismatch; with it, every arch roundtrips clean.
* Update build-sanitize.yml
* make it run on pr
* fix thread
* Update build-sanitize.yml
* Update build-sanitize.yml
* just run thread on github machine
The picker mounts whenever a cwd-aware builtin tool is enabled, so
it can open while file_glob_search is not served or was disabled by
the user. Every typed query then fired a search that could only
fail with a raw error.
Gate the debounced search on the tool state, the same way the
mention picker does, and show a message in place of the results
list that explains why search is unavailable. Manual entry with
Enter still commits a directory. The Browse button and the search
scope footer are hidden as well: Browse resolves the picked folder
name through file_glob_search, and the client-side toggle would not
stop that call.
* server: report the isolate working directory from get_info
Without an explicit cwd, get_info fell back to the server process
working directory even when a tools runtime was configured. That named a
host path no tool would ever run in, since an isolate starts in a
directory of its own.
It now asks the isolate for its working directory in that case, and
keeps the process one only when the tools run on the host.
* remove redundant comment
---------
Co-authored-by: Xuan-Son Nguyen <thichthat@gmail.com>
The working directory chip showed up as soon as the server exposed any
builtin tool, so a server started with just get_datetime, or a user who
turned every filesystem tool off in the settings, still got a control
that nothing would read.
Tools now declare whether they resolve their paths and run against the
working directory, next to the write permission they already publish in
the /tools listing. The WebUI shows the chip and enables the /cwd
command only when at least one such tool is both served and left
enabled.
get_output runs the waveform work the pipeline defers to it, from a
single trailing window to a full pass depending on the model. Measuring
it keeps the reported total and the audio to process ratio honest.
* feat: Add contenteditable tokenizer for badge/code-chip chat input
* feat: Add source-space undo/redo history for the rich input
* feat: Split text glued to a closing code fence onto its own line
* feat: Add ChatFormContenteditable rich input renderer
* feat : wire the contenteditable into ChatForm with auto-switch gating
* base : slash-command/misc foundation - model icon and focus-selector constants
* feat : slash-command picker and command parsing helpers
* refactor : wire command and @-mention pickers into the chat form
* ui : improve model selector keyboard navigation and load/dismiss
* feat: Unify markdown/raw-text rendering under one setting with migration
* fix: Misc fixes - tool-call subtitle, assistant wrap, progress guards
* feat: Clamp and style numeric settings inputs from registry bounds
ggml_metal_op_norm sized the threadgroup with
`nth = std::min(nth, args.ne00_t)`, which can leave nth not a multiple of
the simdgroup size. The kernels finish their row reduction with a
cross-simdgroup step where each lane of the last simdgroup reads one
per-simdgroup partial sum out of shmem_f32:
if (tiisg == 0) { shmem_f32[sgitg] = sumf; }
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
When the last simdgroup is partial it has fewer lanes than the
threadgroup has simdgroups, so the tail of the partial sums is never
read and the row sum is too small. For ne00_t = 33 nth becomes 33: two
simdgroups, but only one lane in the second, so one of the two partial
sums is dropped. The mean and variance are then wrong for the whole row.
Round ne00_t up to a whole number of simdgroups instead. Rounding up
rather than dropping the clamp keeps the threadgroup as small as
possible: deleting the line would raise nth to the next power of two
(ne00_t = 544 -> 1024 instead of 544), which costs idle lanes on 26 row
lengths below 8192 that were already correct, including 1536 and 3584.
GGML_OP_NORM is affected as well as GGML_OP_RMS_NORM - both dispatch
through ggml_metal_op_norm.
No mainstream LLM hidden size hits this: ne00_t is ne00/4 on the
vectorized path, so 4096, 8192, 2048 and friends all give a multiple of
32. It is reachable from other norm shapes, e.g. 320-channel norms.
Add NORM and RMS_NORM cases for ne0 = 33, 132 and 260 across the
existing eps values. 33 exercises the scalar path and 132/260 the
vectorized one, since only those divide by 4.
Before, on M3 Pro:
test-backend-ops test -b MTL0 -o NORM 25/50
test-backend-ops test -b MTL0 -o RMS_NORM 26/51
After:
test-backend-ops test -b MTL0 -o NORM 50/50
test-backend-ops test -b MTL0 -o RMS_NORM 51/51
test-backend-ops test -b MTL0 13943/13943
* base : @-mention picker foundation - glob search, picker nav, highlight
* feat : @-mention file/folder picker and mention badges in message bubbles
* fix: Imports
* feat : wire the @-mention picker into the chat form
* fix: Bound the glob-search result cache key and prune stale entries
* webui: load the model selected via ?model=
Opening the WebUI with ?model= selects the model but doesn't load it. The load only starts when you send your first message, so you wait for it then.
This loads it as soon as the page opens, while you're still typing your prompt. It's what the model dropdown already does, and it isn't awaited, so the UI still works while the model loads.
This is the path the Llama macOS app uses to open the WebUI, so it's a common way in.
* webui: gate the load behind ?load=true
Loading on landing is opt-in, so a plain ?model= link behaves as before and doesn't allocate memory on its own.
* webui: name the chat URL params
Collects the query params the chat routes read into a URL_PARAMS constant, instead of repeating the literals across three files. NEW_CHAT_PARAM folds into it.
The reference implementation has two mutually exclusive prompt layouts.
In non streaming mode the prefill carries the whole utterance text plus
tts_eos summed with codec_pad, and the trailing text hidden collapses to
a single tts_pad row. In streaming mode the prefill carries only the
first text token and the trailing rows stream the rest of the text
followed by tts_eos.
The pipeline built the non streaming prefill but the streaming overlay,
so the talker saw the utterance a second time during generation and read
it twice before emitting codec_eos.
The overlay is now the single tts_pad row that matches the prefill.
* ggml : add fallback definitions for missing aarch64 HWCAP bits
* ggml : require HWCAP_ASIMDHP for the aarch64 fp16 cpu variants
Also rename has_fp16_va to has_fp16, the field gates the whole FEAT_FP16
extension, scalar and vector half-precision arithmetic together.
* ui: read model modalities from the router model list
The router advertises input modalities for every model, loaded or not.
Reading them at list build time lets the UI accept image and audio
uploads for a model selected through ?model=, which has no /props yet.
* enum
The NVFP4 quantization format stores a scaling factor for every group of
16 weights, packed into a single UE4M3 byte.
The SYCL GPU code was converting these scale values using the E4M3 path,
but that's *signed*, and these are unsigned values.
* tests: add SWIGLU perf cases
perf mode had no GLU coverage. Adds SWIGLU at 17408 columns, 512 and
2048 tokens, f16 and f32, with the operands both fused and split.
* sycl: consolidate fused-GLU kernels
They differed only in which op_* they called, so take the op as an argument and share a common launcher.
Their block sizes were all 256, so launch geometry is unchanged;
SYCL_GELU_BLOCK_SIZE and SYCL_SILU_BLOCK_SIZE lose their last users so are dropped.
* sycl: contiguous fast path for the fused GLU ops
o0 == n and o1 == n collapse the de-interleave index math to the
identity, so dispatch a flat kernel in that case. It fires for
ggml_glu_split with packed operands; a fused [gate|up] tensor keeps the
strided path. test-backend-ops perf -o SWIGLU on an Arc Pro B70: split
+14% f16 and +4% f32, fused unchanged.
- In MSL, declaring an array of matrix types like `threadgroup half4x4` causes
a 'no matching constructor' compilation error because MSL matrix types do not
have zero-argument default constructors and threadgroup variables cannot have
initializers.
- Fix this by declaring a POD `threadgroup half` array instead and casting
to `threadgroup half4x4 *` for matrix indexing.
Signed-off-by: JamePeng <jame_peng@sina.com>
