Sync fork master with upstream (2026-08-13) - #85
Open
jimw567 wants to merge 108 commits into
Open
Conversation
* mtmd: fix longest_edge ignoring min/max pixels * nits
* 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
…oup (ggml-org#26708) 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
test-backend-ops perf -o SSM_CONV on an Arc Pro B70, interleaved A/B against master, 6 reps, us/run: ne_a=[515,3328,1,1] ne_b=[4,3328,1,1] n_t=512 97.68 -> 52.95 1.85x ne_a=[937,8192,1,1] ne_b=[4,8192,1,1] n_t=934 516.16 -> 276.13 1.87x ne_a=[4,3328,1,1] ne_b=[4,3328,1,1] n_t=1 2.73 -> 2.71 flat llama-bench on qwen35 27B Q4_K - Medium (48 of its 64 blocks run ssm_conv), -ngl 99 -fa 1 -ctk f16 -ctv f16, interleaved passes of r=3: -b 2048 -ub 2048 pp2048 1045.1 / 1043.5 / 1043.7 -> 1069.5 / 1066.3 / 1065.9 +2.2% -b 2048 -ub 512 pp2048 771.8 / 772.7 -> 785.5 / 786.6 +1.8% -b 2048 -ub 512 tg128 23.81 / 23.88 -> 23.87 / 23.86 flat
* 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
* tests : speed-up test suite 3x * cont : print 30 slowest tests
* 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
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.
…rg#26731) * CUDA: fix thread/block count in quantized cpy kernel launches * tests: add uneven block count cpy case
* server: add initial tool isolation support (via docker) * add docs * adapt get_info * py: fix type check * cont * separate tools_io_sandbox / tools_io_docker * rename sandbox --> isolate * x-tool-docker --> x-tool-runtime --------- Co-authored-by: Pascal <admin@serveurperso.com>
…-org#26762) 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.
* CUDA: fuse rms_norm + mul + rope (+ view + set_rows) * tests: add broadcast weight case to rms_norm_mul_rope * CUDA: check memory ranges before rms_norm rope fusion * CUDA: check memory ranges in rope set_rows fusion
…26773) * 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>
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
…l-org#26811) 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.
* 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
…g#26693) 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.
…gml-org#26826) 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.
* 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
* model: add MTP support for Nemotron Nano model * model: add mtp_flags for nemotron model * address review comments
* ci: Add support for CUDA 13.4 ARM64 builds for Windows Added an architecture-specific CUDA 13.4 Windows build entry targeting ARM64. Added a CMake configuration to enable ARM64 CUDA cross-compilation from an x64 Windows environment using the x64-hosted CUDA and MSVC toolchain while linking against the ARM64 CUDA import libraries to produce ggml-cuda.dll. Validated the self-hosted Windows x64 workflow, including toolkit acquisition, CMake configuration, ARM64 CUDA cross-compilation, and packaging. Runtime validation was performed separately on a native ARM64 RTX Spark system using TinyLlama 1.1B Q4_K_M to verify the generated binaries. The ARM64 CUDA job builds only the ggml-cuda.dll backend (LLAMA_BUILD_SERVER=OFF). The release consists of two packages: the main ARM64 release package, which combines the existing ARM64 CPU outputs with ggml-cuda.dll, and a separate runtime package containing the required CUDA runtime libraries (cudart64_13.dll, cublas64_13.dll, and cublasLt64_13.dll). The CUDA 13.4 setup uses NVIDIA Developer Preview component archives instead of the GA component downloads used by the existing CUDA setups and will require updates once CUDA 13.4 reaches GA. * ci: cleans up to align with x64 CUDA setup - Moves CUDA-specific CMake options into matrix defines. - Keeps the CUB 3DOT2 option only for CUDA 12.4. - Removes runtime argument construction and the unnecessary server option. - Aligns ARM64 CUDA runtime packaging with the existing robocopy approach. - Generalizes the ARM64 release label from CUDA 13.4 to CUDA 13. * ci: Set CUDA job name as version-architecture pair * mark as preview Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Restore quantization of mmprojs This was lost in the refactor undertaken in ggml-org#22004. * add noreturn --------- Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* disable ubuntu-rocm * link PR
* test address on Intel-LNL-U7-258V * retry * run address on github * use native build for cpu * this should be runnable everywhere multicore * disable ccache --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* common: Add CLI > ENV > models-presets > INI precedence
1. CLI flags have the highest precedence
2. ENV vars have the second-highest precedence
3. System and User configs have the lowest precedence
- Linux/BSD/Mac
- /etc/llama.cpp/config.ini < ${XDG_CONFIG_HOME:-~/.config}/llama.cpp/config.ini
- Windows
- %PROGRAMDATA%\llama.cpp\config.ini < %APPDATA%\llama.cpp\config.ini
* fix UB
* use common_get_env
* ignore_unknown_keys
* nits
* add docs
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* refactor: Constants * refactor: Constants/Enums cleanup * refactor: Constant objects instead of multiple single value constants * refactor: Cleanup constants
* refactor: Stores barrel imports + SSR gates * refactor: Drop agenticStore wrapper exports * refactor: Drop chatStore wrapper exports * refactor: Drop modelsStore wrapper exports * refactor: Drop serverStore wrapper exports * refactor: Drop unused mcpStore wrapper exports * refactor: Drop mcpResourceStore wrapper exports * refactor: Drop conversationsStore wrapper exports + move buildConversationTree to utils * refactor: Drop settingsStore wrapper exports * refactor: Drop unused toolsStore wrapper exports * refactor: Fix lint errors from store wrapper removal * fix: Missing change * refactor: Cleanup * refactor: Context Stats store
…gml-org#26951) * refactor: Remove dead context for Chat Settings and create a new one for Chat Messages Actions * refactor: Contexts & types
It enables -fassociative-math, which reassociates FP reductions and can flip greedy argmax on RDNA3.5 (e.g. MTP speculative decode diverging from the non-speculative baseline). Drop it so HIP builds are IEEE-conformant. Co-authored-by: Jim Wu <ywu@xilinx.com>
* server: refactor metrics * move most fields to server_slot_stats * cont * rm result_timings * tie stats to batch * cont * nits: move place in code * exclude first generated token * more accurate batch metrics tracking * n_predict --> n_gen * metrics_on_prediction * metrics_flush_idle * metrics: seperate cache/processed prompt tokens * refactor server_task_result_metrics * add test * nits * fix flush before reset() * cont * rm dead code * nits
* Add DMMV Q4_K and Q6_K ESIMD kernels Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Refactor ESIMD kernels to share common code Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Move control of ESIMD from compile to runtime Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Use ESIMD by default when available Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Fix possible error when using ESIMD by default While not an issue in the current version, this will become an issue when additional QK ESIMD kernels are added (such as Q2_K). Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Add explicit unroll to ESIMD kernels Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Tidy up ESIMD kernels a bit Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Add DMMV Q3_K ESIMD kernel Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
Measured on Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, -fa on, f16 KV, -b 2048 -ub 2048, llama-bench -r 3, three interleaved A/B rounds: pp2048 1014.70 -> 1018.56 t/s (+0.38%, within run-to-run spread) tg128 23.73 -> 23.86 t/s (+0.57%) tg128 @ d4096 22.71 -> 22.86 t/s (+0.62%)
…ml-org#26372) * sycl: use automatic fp16 promotion in gemm * sycl: remove redundant comment
…ggml-org#26947) Co-authored-by: jinzihao <jinzihao.jzh@alibaba-inc.com>
* dflash: enable backend sampling for both dflash & dspark * enable p_min > 0 in backend sampling and add guard * cont : add TODO --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* common : auto-detect spec type from draft GGUF metadata When -md loads a local draft model without --spec-type, the sidecar inference in common_models_handler_apply only checks HF repo sidecars and misses local files. The draft model loads into VRAM but speculative decoding never activates (types stays NONE). Read general.architecture from the draft GGUF header and map: dflash + markov_w1.weight tensor -> draft-dspark dflash without markov head -> draft-dflash Assisted-by: opencode * common : address review feedback on spec-type auto-detect PR - Fix comment spacing to match surrounding style (/* .x = */ not /*.x =*/) - Add LOG_INF when auto-detection fires so users can see why spec decoding enabled - Document single-file assumption for split-GGUF edge case Addresses bot review feedback on ggml-org#26814. * common : move spec-type GGUF auto-detect into speculative module - add common_speculative_types_from_gguf() in speculative.cpp/.h - use gguf_context_ptr (RAII) from ggml-cpp.h - reduce comments to a single line per AGENTS.md style Addresses review feedback on ggml-org#26814 * common : add doc note and join SPC_INF line in spec-type auto-detect Assisted-by: opencode
* metal: add TQ2_0 support Add support for the GGML_TYPE_TQ2_0 (ternary, 2 bits per element) type in the Metal backend. Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731 * cont : optimize mul_mv kernel - float ops over integer ops - precalculate sums - hoist coef out of the inner loop - contiguous y loads llama.cpp:DeepSeek-v4-Flash-0731
index.html was served with `max-age=31536000, immutable` like the hashed assets, but its name is stable while its contents change every build, so a cached copy pins the UI to an old build. It now revalidates via its existing ETag, which keeps the 304 for unchanged builds.
Assisted-by: Claude Opus 5
…Access (ggml-org#26789) * support host pinned mem, ggml_backend_sycl_host_buffer_type_get_max_size, * fix the thread-safe issue
…back (ggml-org#26952) * OpenVINO backend: 1) enable gpt-oss moe on OV bk; 2) enable mxfp4 support * OpenVINO backend: disable TOPK_MOE op test * OpenVINO Backend: Add op FILL support * OpenVINO backend: enable set rows with multi dims * fix the name missmatch in setrow + view * OpenVINO backend: enable op GGML_UNARY_OP_SIGMOID * OpenVINO Backend: enable SQR & SQRT * OpenVINO backend: 1) ensure unique node names for OpenVINO; 2) add org_src to recorde the src ggml tensor for OpenVINO dynamic shape infer * OpenVINO backend: enable fallback for openVINO to CPU backend * OpenVINO backend: fix accurace issue in gemma3n arch test * fix mpt failed case * OpenVINO backend: clean nodeinfo * OpenVINO Backend: enable zero-size copy for view * add concat ssm_conv in compute_dynamic_dim enable qwen35 Fix after rebase remove logging * OpenVINO backend: disable EXP with FP32, which failed in op test. Root reason: the backend test initializes unary op inputs over a wide range, [-150, 150]. For FP32, exp(x) overflows around x ~= 88.7, so this test can randomly generate values right in or beyond the overflow region * OpenVINO backend: fix CPY op test failed issue * OpenVINO backend: fix GATED_DELTA_NET op test failed issue * handle in-place op, handle qwen35 dynamic clearing of cache in cgraph * handle qwen35 dynamic clearing of cache correctly * Enable qwen35 dense multi seq * Fix qwen35 9b gqa * Fix after rebase * Disable SOLVE_TRI * openvino: fix NEOX RoPE accuracy on GPU stateful (mixed-rank Multiply) In stateful mode the NEOX RoPE branch fed rank-3 data ([S, n_heads, head_size]) into the Multiply against the rank-4 cos/sin tables ([1, S, 1, n_dims/2]). That mixed-rank broadcast is miscomputed by the OpenVINO GPU plugin, corrupting the rotated Q/K and producing garbage output (e.g. Phi-3-mini). Lift the data to rank-4 before the split/ Multiply so the operands are equal-rank, matching what the TYPE_NORMAL branch already does. CPU and stateless paths are unaffected. Phi-3-mini-Q4_K_M, wiki.test perplexity, GPU stateful: before: PPL = 27120.43 after: PPL = 6.2263 (CPU reference: 6.2251) * OpenVINO backend: 1) remove the unique name in llama.cpp; 2) add new ov name in ov bk; 3) fix issue in arch test & op test with latest code update * OpenVINO Backenb: remove changes in llama.cpp * Doc change (use x64 Native Tools Command Prompt for VS) * Cleaner sentence Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * OpenVINO Backend: cache key upgrade includes all src name * OpenVINO Backend: enable llama arch test on ci * OpenVINO Backend: move parameter node creating from decoder into translate * OpenVINO Backend: create extra input ov node move from decoder to translate * fix for op regression due to is_model_splitted * openvino: fix CPY writeback for recurrent state rollback Detect the rollback conv/gdn state writeback CPY nodes structurally instead of by tensor name, since the rollback path in build_conv_state does not call cb() and left the nodes unnamed. Add per-node runtime offsets (rs_slot_begin_*, rs_src_begin_*) so the cached IR handles any kv head, sequence count and snapshot slot for both the conv state and the GDN state writeback. Assisted-by: GitHub Copilot * qwen35 moe * optimize MoE expert aggregation with ReduceSum * Skip GET_ROWS inaccurate test * openvino: fallback dynamic MUL_MAT_ID shapes * OpenVINO Backend: fix error in arch test model mpt * fix error caused by cpy in arch test model kimi-linear * OpenVINO Backend: fix error in arch test model minimax-m3 * openvino: fix GPU mul_mat_id op tests * ggml-openvino: add GGML_OPENVINO_RELEASE_WEIGHTS to reclaim host weight RSS on GPU The OpenVINO weight Constants are zero-copy views into host buffers allocated by the backend (ggml_aligned_malloc, anonymous memory). On GPU the plugin holds its own device copy after compile_model, so these host pages are dead weight for inference. For a 1B Q4_K_M model this leaves ~850 MB of host RSS resident that the GPU path never reads again. Add an opt-in GGML_OPENVINO_RELEASE_WEIGHTS mode that madvise(MADV_DONTNEED)s the registered host weight buffers once the model is compiled, dropping their resident pages while keeping the mappings valid (ggml still owns the lifetime; tensors still point in). Measured steady-state RSS drops from ~1555 MB to ~710 MB on Llama-3.2-1B-Q4_K_M (Arc iGPU) with unchanged throughput and correct output. The GPU backend uses a single dynamic-shape model for both prefill and decode, so a graph is compiled once and reused; the only event that forces a recompile is clear_caches() on backend teardown. The change therefore: - releases on the first cache-hit (model compiled, plugin has its copy); - pins the compiled-model cache across backend teardown so a later context reuses it instead of recompiling against the dropped pages; - fails loud (GGML_ABORT) on a cache-miss recompile or on a second model load, both of which would otherwise read zeroed weights or silently reuse the wrong compiled graph. Scope/limitations (all fail loud, never silently wrong): GPU only (the CPU plugin reads the host Constants at inference time), one model per process, and stable graph shapes. This reduces steady-state RSS, not the transient compile-time peak. All changes are confined to the OpenVINO backend. * ggml-openvino: stream weight requantization to cut the compile-time RSS peak requantize_to_buffers() dequantized the entire tensor to a temporary std::vector<float> of n_elements before requantizing. For token_embd.weight (128256 x 2048) that transient is ~1 GB (1B model) / ~2 GB (8B), and it is the single largest contributor to the OpenVINO compile-time memory peak -- it also fires twice for token_embd (once at load, once at graph build, because token_embd is loaded via a CPU/mmap buffer and not cached as an OV weight extra). Stream the dequant instead: process a fixed window of complete rows (CHUNK_ROWS=256) into a small scratch buffer and quantize/convert each chunk straight into the output buffers. The transient F32 footprint is now CHUNK_ROWS*ne0 floats regardless of tensor size. quantize_q8_0/q8_1 gain an optional block_offset arg (default 0) so a chunk writes its weights/scales/zp at the correct block. Streaming is applied to the Q8_0_C / Q8_1_C / F16 targets (the large requant cases); the u4 (Q4_0) path keeps the whole-array call because it packs two weights per byte with running zp ORs, and a fallback handles any future target whose block size does not divide a row. Measured peak RSS (cold compile, GPU): 1B 2868 -> 1809 MB (-1.06 GB); 8B 11618 -> 9608 MB (-2.0 GB). Output verified unchanged ("capital of France is Paris"); throughput unchanged. Unlike GGML_OPENVINO_RELEASE_WEIGHTS this reduces the transient peak, not just steady-state, and needs no env flag. All changes confined to the OpenVINO backend. * ggml-openvino: avoid redundant token_embd requantization at compile token_embd.weight is referenced twice in the graph path: as the GET_ROWS embedding (a CPU/mmap-buffer tensor) it was re-extracted/re-requantized on every weight-node build, and is_model_splitted() built a full (naive) set of weight nodes just to test name membership — each requant is a ~1-2 GB F32 dequant of the 262M-element embedding. Two changes: - Add collect_weight_names(): a name-only collector for topology checks. is_model_splitted() now uses it instead of create_weight_nodes(cgraph, true), so the splitted-check no longer triggers any weight extraction. - Memoize weight nodes built from non-OpenVINO buffers in a process-lifetime cache keyed by tensor->data. These tensors have no OV buffer context to own a cached extra, so without this they were rebuilt on every (re)compile; prefill and decode graphs now share one build (verified: 2nd graph hits the cache instead of re-requantizing). Peak RSS is unchanged (the streaming-requant commit already removed the F32 transient); this removes redundant compile-time work. Output verified unchanged ("capital of France is Paris"). Confined to the OpenVINO backend. * ggml-openvino: gate compile-memory optimizations behind GGML_OPENVINO_REDUCE_COMPILE_MEM The streaming requantization and the non-OpenVINO-buffer weight-node cache (plus the name-only is_model_splitted path that pairs with it) are now opt-in via GGML_OPENVINO_REDUCE_COMPILE_MEM. When unset, requantize_to_buffers() fully materializes the F32 buffer and weights are rebuilt per compile exactly as before; when set, the streaming path and the cross-compile weight cache are used. Default off keeps behavior identical to upstream unless explicitly enabled. Verified: flag off -> peak RSS 2800 MB (original), flag on -> 1810 MB; output "capital of France is Paris" in both modes. (GGML_OPENVINO_RELEASE_WEIGHTS, added earlier, remains a separate opt-in for the steady-state release.) * ggml-openvino: add frontend model cache (GGML_OPENVINO_MODEL_CACHE_DIR) The plugin-level ov::cache_dir caches the compiled blob keyed by the OV model, but producing that model still runs the full frontend every time: weight requantization (incl. the large token_embd F32 transient) and the ggml->OV graph conversion. This adds an opt-in frontend cache keyed off a fingerprint computed directly from the ggml cgraph, so a hit imports a previously exported CompiledModel and skips requant + convert + compile entirely. Key (model-cache.{h,cpp}) = 64-bit FNV-1a of: graph topology (n_nodes + per node op/name), a sampled per-weight fingerprint (name/shape/type + bounded head+tail byte sample), and blob-affecting config (device, flash-attn, rope params, REDUCE_COMPILE_MEM/stateful flags, OpenVINO version). A sidecar manifest stores every weight's fingerprint and is re-verified on load, so a sampled-hash collision cannot cause a wrong-model hit (verified: two different quantizations of the same model produce distinct cache entries). Flow (dynamic single-model path only; split models defer to ov::cache_dir): on a verified hit, core.import_model() restores the CompiledModel and a lightweight decoder is built with a names-only weight map (membership is all the decoder needs for I/O mapping; weights live in the imported model). On a miss, compile as usual then export the blob (atomic temp+rename, manifest written first). The frontend cache supersedes ov::cache_dir, so CACHE_DIR/ CACHE_MODE are stripped from the config used for the cached compile and the import — a blob compiled with cache_dir set cannot be re-imported. Measured 8B Q4_K_M (GPU): full requant+convert+compile 15.3s -> import 6.3s (~2.4x faster compile phase). Output verified unchanged on cold and warm, standalone and combined with REDUCE_COMPILE_MEM + RELEASE_WEIGHTS. Default off; confined to the OpenVINO backend. * ggml-openvino: harden frontend model cache correctness The frontend model cache imports a previously exported CompiledModel keyed by a fingerprint of the ggml graph, weights, and blob-affecting config. The original key covered device, stateful execution, REDUCE_COMPILE_MEM, RoPE params, OpenVINO version, topology, and sampled weights, but missed runtime/frontend toggles that can change the lowered graph or the I/O binding contract. That made it possible to reuse a blob produced under a different OpenVINO backend configuration. Add a small extra-config helper for the dynamic model-cache path and fold in the effective values of GGML_OPENVINO_DISABLE_KV_SLICE and GGML_OPENVINO_MANUAL_GQA_ATTN. MANUAL_GQA_ATTN is keyed by the behavior that actually takes effect: an explicit env value wins, otherwise GPU defaults to enabled and other devices default to disabled. This matches flash_attn_ext lowering and avoids unnecessary cache splits for equivalent configurations while separating genuinely different attention graphs. DISABLE_KV_SLICE is also included because it changes the KV-cache tensor shape/output binding strategy used around imported models. Even when weights and graph topology are identical, switching this flag should not inherit a CompiledModel cache entry created for a different binding mode. Also make cache artifact publication cleaner: write manifest.tmp and blob.tmp, publish the blob first, and publish the manifest last. Cache hits already require both blob and a verified manifest, so making the manifest the final visible artifact avoids leaving an apparently complete manifest for a failed or interrupted blob export. Temporary files are removed on the handled failure paths. While touching this path, fix the indentation of the non-imported compile branch so the cache miss flow is easier to review. Behavior is otherwise unchanged: verified hits still import, misses still create weights, convert, compile, export, and create the infer request normally. * ggml-openvino: add memory optimization umbrella switch Add GGML_OPENVINO_MEMORY_OPTIMIZE as a single opt-in switch for the OpenVINO backend memory-saving paths. The existing fine-grained GGML_OPENVINO_REDUCE_COMPILE_MEM and GGML_OPENVINO_RELEASE_WEIGHTS variables remain supported and explicitly override the umbrella switch when set, so users can still bisect or disable one side of the optimization independently. Centralize the policy in ggml_openvino_reduce_compile_mem_enabled() and ggml_openvino_release_weights_enabled(device). The umbrella switch enables compile-memory reductions everywhere REDUCE_COMPILE_MEM is used today: streaming requantization, non-OV weight-node caching, split-model weight-name collection, and the frontend model-cache fingerprint. On GPU it also enables host weight-buffer release unless GGML_OPENVINO_RELEASE_WEIGHTS is explicitly set. Keep host weight release GPU-only because it relies on the plugin holding its own device copy after compile_model. Update the fail-fast diagnostic and comments to mention GGML_OPENVINO_MEMORY_OPTIMIZE, so users who enable the umbrella switch get accurate guidance if a later cache-miss recompile would read released host weight pages. * ggml-openvino: rename compiled model cache env Rename the frontend export/import cache environment variable from GGML_OPENVINO_MODEL_CACHE_DIR to GGML_OPENVINO_COMPILED_MODEL_CACHE_DIR. The cache stores blobs produced by ov::CompiledModel::export_model() and restores them with core.import_model(), so the new name distinguishes it from GGML_OPENVINO_CACHE_DIR, which configures OpenVINO plugin-level ov::cache_dir. Update the registered env var, the cache-directory lookup, and comments around the frontend compiled-model cache. The old GGML_OPENVINO_MODEL_CACHE_DIR name is removed rather than kept as a fallback so there is a single spelling for the new option. * docs: document OpenVINO memory optimization env vars Add runtime configuration entries for the newly recognized OpenVINO environment variables. Document GGML_OPENVINO_COMPILED_MODEL_CACHE_DIR as the frontend compiled-model cache used to export and import compiled blobs for matching single-graph models. Document GGML_OPENVINO_MEMORY_OPTIMIZE as the umbrella switch, including how GGML_OPENVINO_REDUCE_COMPILE_MEM and the GPU-only GGML_OPENVINO_RELEASE_WEIGHTS override or inherit from it. * ggml-openvino: fix Qwen3VL crash and deepstack correctness bug 1. GGML_OP_PAD was missing from compute_node_dynamic_dims(), causing a crash on decode for models that pad the token embedding (n_embd -> n_embd_inp). PAD never reorders/merges dims, so it keeps the same dynamic dim index as its source. 2. process_view_input_new() chained VIEW inputs through src[0] (the immediate op-graph parent) using offsets treated as relative to that parent. But ggml_tensor::view_offs is always absolute from the true root allocation (ggml collapses VIEW-of-VIEW chains internally). For the per-layer deepstack view ("embd (view)", whose src[0] is "embd" - itself an already-narrowed, zero-offset VIEW of the padded root, with the SAME ggml shape as the deepstack view but a different absolute offset), this caused an out-of-bounds re-slice that silently fell back to returning the wrong (already-resolved sibling) tensor. In practice every deepstack ADD ended up adding the real base token embedding into the residual stream instead of zero, corrupting generation ("Hello my name is 1000000..." instead of coherent text). Fixed by detecting this pattern (same shape as the immediate src, different absolute offset) and re-slicing directly from the untouched root tensor using the innermost view's absolute offset. Also adds a GGML_OPENVINO_DEBUG_NODE=<name1>,<name2>,... env var that attaches extra debug Result nodes for arbitrary intermediate tensors, without binding them to any ggml buffer (avoiding the risk of reading a ggml buffer that has since been overwritten by a later in-place op). This was instrumental in diagnosing bug #2 above and is left in as a general-purpose debugging aid. * ggml-openvino: fix IMROPE inp_pos padding for NPU static shapes IMROPE's inp_pos tensor packs 4 stacked t/h/w/e position planes into ne[0] = 4*n_tokens instead of one value per token. On NPU's static-shape path, inp_pos was padded/shaped as if it held a single plane, which interleaved padding across the 4 planes and desynced later reshapes from the rest of the (chunk_size-wide) graph. - add GgmlOvDecoder::get_inp_pos_n_planes() to detect IMROPE's 4-plane layout - get_graph_input_shape(): size inp_pos as n_planes * chunk_size (prefill) or n_planes (decode) instead of assuming 1 value per token - get_ov_input_tensor_static_prefill(): pad each plane to chunk_size independently instead of one flat block - get_ov_input_tensor_static_decode(): copy n_planes contiguous values instead of asserting/copying a single scalar * disable test-llama-archs tests. * openvino: gate fallback with env var * Revert changes in test-llama-archs * Apply editor config * reject CPY with quantized destination as unsupported --------- Co-authored-by: Xuejun <Xuejun.Zhai@intel.com> Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com> Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: suryasidd <surya.siddharth.pemmaraju@intel.com> Co-authored-by: Mustafa Cavus <mustafacavus@intel.com> Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Merge ggml-org/llama.cpp master into AMD-Ecosystem fork master.
No other fork-local changes; all remaining files taken from upstream.