921 Commits

Include renames

Author	SHA1	Message	Date
Gaurav GargandGeorgi Gerganov	dd1ea52433	llama : support multi-output backend sampling (#25532) ... * Enable backend sampling with token speculation * Clamp the mask sum before converting it into the sampled index * Add a numeric context parameter declaring the maximum outputs one sequence * More fixes * Don't reuse memory for output views. * Match dist between CPU and GPU * Fix CPU and backend sampling mismatches * Simpify some of the changes * Fix tests on Vulkan * More test fixes * Rebase changes * Rebase and address review comments * Address review comments * Address review comments * Update src/llama-sampler.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-08-10 16:58:56 +03:00
Yash Raj Pandey	f8def7fe16	ggml : require contiguous src for ROLL on CUDA and Metal (#25928) ... ggml_roll only asserts nb[0] == ggml_type_size, so a permuted src is a valid input, but the CUDA and Metal roll kernels index by ne alone and never read the nb strides. A non-contiguous src therefore produced silently wrong results. Neither backend declared a contiguity requirement in supports_op, so the scheduler did not fall back to the CPU implementation, which does handle strides correctly. Add the requirement to both backends, matching the existing GGML_OP_ROPE guard, and add a permuted test_roll case.	2026-08-10 15:01:44 +03:00
7 people	62bf73d25c	model: Muse Glimmer Support (#26841) ... * Get started with Onyx * Add architecture * Skip keys handled in super() * Loading tensors * Shorten * Graph * Apply suggestion from @pcuenca * Remove norm now embedding in transformers weights * Add eot * Explicit output_multiplier * Handle post_norm_eps * No super call; unhardcode eot. The pattern `self._set_vocab_gpt2()` seems preferred throughout the codebase, and it allows `set_vocab()` to be called from a different part of the Python class hierarchy: the drafter model converter that we may need eventually. * Register for drafting * DFlash: inherit rope type from the linked target. Another option would be to store it in the gguf file itself. * mmproj conversion Note: some fields to be renamed after the implementation works. We are keeping compatibility with the reference Meta gguf for testing purposes. * "clip" header declarations * Load mmproj * Pre-processing * Graph * Go back to using delimiters. Otherwise our generations are worse. Transformers does not use them. We need to trace inputs to verify whether they are equivalent. * downsample_factor -> merge_size * Add vision graph lol, forgot from a previous commit * Additional renames, align with llama.cpp / transformers * Prefer _size instead of independent _h and _w * Fix token layout Co-authored-by: Young Han <younghan@fb.com> * onyx: bring the chat parser onto the onyx branch common/chat.cpp on this branch has no Onyx handling, so a converted model serves malformed chat: the assistant preamble leaks into content ("to=self<|message|>...") and tool calls fail with HTTP 500 "The model produced output that does not match the expected peg-native format" common_chat_params_init_onyx exists on onyx-fair-patch, added there by 8bb73dd3d. It was never on this branch, so this is not a regression -- the two lines developed independently. The code here is taken verbatim from that commit. It is the clean side of `git merge origin/onyx-fair-patch`: chat.cpp is one of the files that merges without conflict. The full merge is not viable -- it produces 13 conflicts, including add/add on conversion/onyx.py and src/models/onyx.cpp where the q_norm-folding and metadata-scale approaches contradict each other, and #4/#7 are stacked on this branch's side of that. Verified on this branch: builds with 0 errors, converts an Onyx checkpoint, and serving it gives "4" for "What is 2+2?" plus a correct get_weather {"city":"Paris"} tool call, where the unported branch gives the two failures above. No converter or runtime changes are included, so this should not interact with the q_norm work. Co-authored-by: Beto de Paola <betodepaola@meta.com> * Less params, bilinear pos-emb interpolation as a graph op instead of CPU * Map to symbolic V_MMPROJ instead of strings * Make a couple params explicit * Patchify via build_inp() * No param for rope_theta * Small cleanup * Restore blank line * Unpermute, to adapt to the latest transformers checkpoint * Apply norm after token embeddings This follows the latest transformers approach. * Remove duplicated function * build_vit * onyx: use the model rope theta on sliding-window layers * DFlash: conversion from transformers drafter * Revert rope_type derivation from target NOTE: this breaks compatibility with Meta's distributed DFlash GGUFs, as the Q/K are stored in "NEOX" (rotated half) format, like in transformers. * Apply suggestion from @pcuenca * Set model type * Remove comment that will become obsolete * Hardcode post_norm_rms_eps instead of new param * Derive SWA+RoPE pattern from gguf array or scalar * Fix model type <-> number of layers * Reorder * Rename * Fix typo * DFlash: seed the draft KV cache from multimodal embedding batches `common_speculative_impl_draft_dflash::process()` returned early on any batch carrying embeddings, so an image prefill never had its target-layer features fused through the DFlash encoder and injected into the draft's KV cache. That left a hole spanning the image's positions, and the next injection at a post-image position failed to initialize its batch: ``` decoding image batch 1/1, n_tokens_batch = 256 decode: failed to initialize batch llama_decode: failed to decode, ret = -1 process: llama_decode(ctx_dft) failed rc=-1 (n_tokens=17, offset=0) srv decode: failed to process speculative batch ``` Every image request with `--spec-type draft-dflash` failed with HTTP 500. Text-only was unaffected, since those batches carry token ids and were let through. Restore the earlier condition, which admits a batch that is either tokens or embeddings and skips only the degenerate neither/both cases. The rest of `process()` is already layout-agnostic -- it gathers features via `llama_get_embeddings_layer_inp()` and indexes `batch_in.pos[]` / `batch_in.seq_id[]`, none of which assume token ids -- so this is the whole fix. Validated against `muse-glimmer-30B-bf16.gguf` + `mmproj-muse-glimmer-30B-bf16.gguf` + a DFlash draft head, on an image describe-the-shapes request: - before: HTTP 500, `failed to process speculative batch` - after: HTTP 200, draft acceptance 0.34012 (167 accepted / 491 generated), mean len 3.04 Output equivalence holds, which is the property that matters: at temperature 0 the drafted response is byte-identical to the same request served with no draft attached (1213/1213 chars), so the draft is drafting correctly through the image context rather than merely not crashing. * Conversion: prefer rewrite to mapping * Revert "Conversion: prefer rewrite to mapping" This reverts commit a92d0ac584d315e876741e85b6dad3dbc8b23bf7. * fix lint * sliding_window metadata is not optional * disable state save/load * Apply suggestion from @pcuenca --------- Co-authored-by: Young Han <younghan@fb.com> Co-authored-by: Beto de Paola <betodepaola@meta.com> Co-authored-by: Daniel Han <michaelhan2050@gmail.com> Co-authored-by: ruanrms <ruanslv@gmail.com> Co-authored-by: Xuan Son Nguyen <son@huggingface.co> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>	2026-08-10 13:07:27 +02:00
Guido Imperiale	a52077c4ca	chat : Align Laguna-S-2.1 chat template to huggingface (#26232)	2026-08-10 05:20:59 -05:00
Bar Haim	157b81fe6d	model : Granite-Switch Architecture (#25107) ... * 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	2026-08-10 09:53:46 +02:00
Caleb DeLeeuw	0377426cef	model-saver : fix expert shared/chunk FFN length key clobber (#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.	2026-08-10 09:32:01 +03:00
Rafail Giavrimis	687e778927	CUDA: fuse rms_norm + mul + rope (+ view + set_rows) (#26767) ... * 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	2026-08-09 00:32:37 +08:00
Rafail Giavrimis	69bf643791	CUDA: fix thread/block count in quantized cpy kernel launches (#26731) ... * CUDA: fix thread/block count in quantized cpy kernel launches * tests: add uneven block count cpy case	2026-08-08 07:40:04 +03:00
robertomeroni	a194a75b7e	metal : fix NORM/RMS_NORM for row lengths that leave a partial simdgroup (#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	2026-08-07 21:09:07 +03:00
Masato Nakasaka	42e98813e4	Mitigate crashing issue on Windows MSYS2 UCRT64 environment (GCC 16.1.0) (#26555)	2026-08-07 11:17:16 +02:00
Titaniumtown	6b5c2efb4e	sycl: *glu flat path (#26354) ... * 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.	2026-08-07 08:24:40 +03:00
Xuan-Son Nguyen	15586e2d71	mtmd: add chunk save/load function (#26645) ... * mtmd: add chunk save/load function * nits * add tests * rn _MAX --> _COUNT	2026-08-06 19:46:40 +02:00
5 people	a1f96d4fc2	ci : onboard AMD ROCm CI with gfx1151 fixes (#26544) ... * ci: prepare for amd rocm ci Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ci: fix editorconfig-checker Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ci: fix device not recognised Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ci: rename gpu-amd to gpu-hip Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ci: gpu-hip to gpu-rocm haha Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * CUDA: allow integrated-GPU host output buffer in debug assert On integrated GPUs (APUs), the scheduler can legitimately place a graph node's output on the host-visible buffer, which ggml_cuda_compute_forward already handles. The debug assert in ggml_cuda_graph_evaluate_and_capture required every node output to be on the device buffer, so a debug build aborts on such a node (e.g. attn_residual ADD -> ROCm_Host on RDNA3.5). The source-tensor assert directly below already permits this via the integrated + cuda_host exception; apply the same exception to the node's own output buffer. Debug-only; no effect on release/compute. Fixes test-recurrent-state-rollback on gfx1151 (Strix Halo). * ci: enable unified memory for ROCm gfx1151 job Work around a coherence issue on integrated RDNA3.5 (gfx1151) where GPU kernels reading mmap-loaded weights can return incorrect output, which makes test-llama-archs (and real inference) intermittently wrong. GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 uses managed memory, which restores coherence. Remove once the underlying ROCm/HIP issue is fixed. * test-llama-archs: skip jamba on HIP backend jamba produces incorrect output (~0.55 NMSE vs CPU) on the HIP backend on RDNA3.5 (gfx1151); the SSM kernels need separate investigation. Skip it for now, matching the existing per-backend carve-outs (WebGPU), so the ROCm CI can run the test for the remaining architectures. * ci: use HIP_LAUNCH_BLOCKING for ROCm gfx1151 job The gfx1151 ROCm CI job produced incorrect inference output (qwen3 perplexity ~88 vs ~9.4) due to an async-execution correctness issue in the HIP path. Serializing kernel launches with HIP_LAUNCH_BLOCKING=1 restores correctness. This replaces the earlier GGML_CUDA_ENABLE_UNIFIED_MEMORY workaround, which did not fix batched inference. * test-backend-sampler: skip top-k subtests on HIP backend The ROCm backend does not support the TOP_K/ARGSORT op at vocab scale (no CUB; bitonic argsort is capped at ncols <= 1024), so top-k/top-p backend samplers cannot be offloaded. The penalties, set_sampler, mixed, and top_p subtests assert that offload happened, so they fail on HIP. Skip them until TOP_K is supported on the ROCm backend. * Update tests/test-backend-sampler.cpp Co-authored-by: Aaron Teo <taronaeo@gmail.com> * Update tests/test-backend-sampler.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> Co-authored-by: Aaron Teo <aaron.teo1@ibm.com> Co-authored-by: Jim Wu <ywu@xilinx.com> Co-authored-by: Aaron Teo <taronaeo@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-08-06 10:43:26 +02:00
Niklas Wenzel	360e1349f0	tests: re-enable MiniMax M3 in test-llama-archs (#26633)	2026-08-05 17:58:34 +02:00
Aldehir Rojas	cd0fa6051a	grammar : degrade max repetition >= 2000 to unbounded (#26613)	2026-08-05 07:39:10 -05:00
Oliver Simons	a6aa6f5450	sampler : remove "full-context windows" from history-based samplers (#26524) ... * Resolve -1 to 1024 instead of ctx-len for samplers Because of backend-sampling we initialize samplers before the complete llama_context is there. Therefore, we cannot infer the resolved context length yet at the time we construct the samplers. * Shared default of 64 for history-based samplers, remove context_size	2026-08-04 21:28:55 +03:00
Xuan-Son NguyenandPascal	0713275082	mtmd: support Qwen3-TTS (note: breaking change to llama-tts binary) (#26254) ... * convert text model * main model load ok * convert encoder ok * speaker encoder loading ok * speaker enc graph * adapt vocab for backbone (with some tricks) * add suppress_tokens * poc new mtmd gen api * convert code_predictor to gguf * load gen_code model ok * add clip_encode * wire up * code gen cgraph init version Co-authored-by: Pascal <admin@serveurperso.com> * code2wav convert to gguf * code2wav graph ok * wire up in/out * (wip) subgraph * wire up * wip, correct code2wav * demo (to be removed) * code2wav preserve kv between calls * demo voice clone * llama: add llama_model_get_tok_embd * mtmd_helper_gen_audio API * fix clamp cold prefix Co-authored-by: Pascal <admin@serveurperso.com> * fuse snake op Co-authored-by: Pascal <admin@serveurperso.com> * demo: use proper sampling * update dev docs * polymorphism helper * revamp llama-tts binary * update docs * fix compile * fix lint * nits * add guide + docs * more timings info * clean up code comments * security fixes * update docs * use ggml_build_forward_select, clean up comments * fix ci * use ISO 639-1 language code * rename CODE2WAV --> GEN_WAV, update docs * clean up * clean up tts.cpp * add seq_id * add step_prompt() * mtmd_helper_model_can_chat * clean up comments --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-08-04 17:26:15 +02:00
Oliver Simons	935cad6497	llama : move n_vocab from llama_sampler_data to penalty_sampler (#26520) ... This matches how it is done for logit_bias and mirostat samplers, see https://github.com/ggml-org/llama.cpp/pull/25262#discussion_r3703951151	2026-08-04 09:02:49 +03:00
Thiago Padilha	0ef6e55edb	chat : add new template for DeepSeek V4 Flash 0731 (#26398) ... * common/chat: update DeepSeek V4 templates Align the DeepSeek V4 templates with the official encoders while keeping parser behavior out of this change. - Default drop_thinking for DeepSeek V4 history so prior thinking is omitted unless preserve_reasoning is requested or tools are present. - Add structured output response-format instructions to the V4 templates and pass the schema into template rendering. - Add a separate Flash 0731 template for the updated high and max reasoning effort mapping. - Cover reasoning effort, drop_thinking, structured output prompts, preserved reasoning, continuations, and empty tool arguments in template rendering tests. Official references: https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash/blob/main/encoding/encoding_dsv4.py https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash-0731/blob/main/encoding/encoding_dsv4.py Assisted-by: Codex * Fix deepseek v4 0731 template selection * remove unneeded lower normalization * Fix DSML parser to consume the tool call separator * address aldehir requests * address aldehir comment	2026-08-03 17:59:11 -05:00
PascalandXuan Son Nguyen	ee0445c99c	tests: add model resolution test on synthetic repo listings (#26172) ... * tests: add model resolution test on synthetic repo listings Include download.cpp and arg.cpp inside a namespace with hf_cache monkey patched to serve hardcoded listings, so the resolution and the model handler assembly are tested end-to-end through the real CLI parsing, without modifying the tested code and without network access. Covers the primary, shard, mmproj, sidecar and preset resolution on layouts mimicking real vendor conventions, replays every case on permutations of the listing to assert determinism, and asserts the final wired paths for the spec type auto-selection, the -md precedence and the fallback suppression. * tests: keep model resolution checks active and let the handler resolve Replace assert with a REQUIRE macro alive in Release builds, key the fake hf_cache by repo id so the real handler init resolves every plan itself, check the exact shard sets, restrict the permutation exception to the order dependent picks, and cover dflash and eagle3. * tests: fix model resolution build on fatal warnings CI and Windows The namespaced copy of the sources leaves many static functions unused in this TU, exempt it from the unused warnings. Pre-include the windows headers so arg.cpp does not pull them inside the namespace. Declare the renamed copies of the download.h functions, verbatim from the header and renamed in sync by the macros, so missing declaration and missing prototype warnings are satisfied on every toolchain. * tests: fix winsock inclusion order for the model resolution test WIN32_LEAN_AND_MEAN and winsock2.h before windows.h, so http.h does not redefine the socket types afterwards. * tests: link cpp-httplib to the model resolution test The test compiles its own copy of download.cpp, which calls httplib directly, and the private link of llama-common does not propagate the symbols under lld-link. * common_http_client * common: finish the http client wrapper Add the virtual Head, Get and Post methods and the passthrough setters to the common_http_client skeleton, move follow_location into the constructor, expose the underlying client for the ranged pull path, and rename the missed common_http_client_init call sites. * tests: rewrite model resolution on the http client stub Replace the namespace inclusion of the sources by a plain TU: the common_http_client factory returns a stub serving hardcoded HF API responses, so the real hf_cache parsing, resolution and CLI handler run against synthetic listings in an isolated cache directory. Failures print the named case, the reordering and the actual versus expected values, the assembly cases use the full command line as context, and the empty result cases are checked once to keep the logs short. * tests: fix the model resolution on Windows and the builds without TLS Assert the exact expected paths composed like the cache does instead of suffix matching on forward slashes, set the environment portably, and serve the stub through an http endpoint so the builds rejecting the https scheme still reach it. Pause the log so the negative cases can be replayed on every reordering. * tests: make the model resolution failures self explanatory Resume the paused log before the failure report so the CI shows why the tested code bailed, and format the stub oids portably. * common: hold the http client factory behind exported functions The factory was an inline variable, and the Windows shared builds export functions but not data, so the executable and the DLL each had their own instance: the stub installed by the test was invisible to the library, which reached for the real endpoint and resolved nothing. Route the creation through functions compiled into the library and format the stub oids portably. * common: add the http client factory source missed in the previous commit * common: typedef the http client factory callback Address review from @ngxson * tests: serve the model resolution repos over the loopback Replace the client stub by a real httplib server bound to the loopback, so no C++ object crosses the module boundary anymore and the library exercises its own client and transport end to end. The debug shared build on Windows crashed inside the stubbed path. * common: add portable common_get_env and common_set_env helpers Address review from @ngxson * common: drop the http client factory left without a caller The loopback server made the stub substitution unnecessary, the client init builds the real client directly again. * common: read the model endpoint through the env helpers * nit: drop the stub leftovers from the model resolution test * common: align common_set_env and isolate the test cache per run The POSIX branch now behaves like _putenv_s, so the helper has a single contract on every platform, and common_get_env already reads an unset and an empty variable alike. The model resolution test keys its cache directory on the loopback port, where two concurrent runs on the same machine used to share one directory and the initial cleanup of either wiped the other. * tests: move the model resolution server into main * tests: support the DSpark sidecar resolution * common: revert the http client to the plain httplib client address review from @ngxson --------- Co-authored-by: Xuan Son Nguyen <son@huggingface.co>	2026-08-03 18:58:15 +02:00
Konrad Moren	96278e39fc	CUDA: Add backend sampler for penalties sampler (#25262) ... * sampling: enhance penalty handling in common_sampler_init - Set default value for penalty_last_n based on model context if not specified. - Ensure penalty_last_n and n_prev are non-negative. - Update llama_sampler_penalties structure to inherit from llama_sampler_backend and add backend input handling for penalties. - Implement backend initialization and application logic for penalties, including frequency and presence adjustments. * tests: add backend penalties sampling tests and utility functions - Introduced `accept_prompt` and `unique_prompt_tokens` functions to handle prompt acceptance and token uniqueness. - Implemented `compare_penalties_logits` to compare logits from backend and CPU samplers with penalties. - Added `test_backend_penalties_sampling` to validate backend penalties with various configurations. - Enhanced the test suite for better coverage of penalty handling in sampling. * sampling: add support for top-k penalties in backend sampling * sampling: add fix to ensure stable numerical results. Preserve masked logits as -Inf and no longer generate NaN. * sampling: enhance penalty comparison tests with masking penalties logic * add comments on padding * sampling: add comments on modifications * add the unit test to cover masked-out token as -INF * validate repeat penalty to ensure it is finite and greater than 0; add tests for invalid values * refactor: test functions to share logic and be less verbose * add test to cover case where previously penalized token is not part of candidates * remove comments * remove redundant penalty_last_n initialization and validation in common_sampler_init * add support for penalties in sampler chain with configurable positions * add validation for penalty parameters and enhance tests for non-finite values * add context parameter to common_sampler_init and set default for penalty_last_n * add llama_n_ctx parameter to common_sampler_init for improved sampler initialization * replace penalty_last_n x n_candidates comparison matrix with a vocabulary-sized count tensor * add tests for backend penalties sampling without filler entries , token_count.size() == n_active == n_max == 64 * add test for backend penalties sampling after top-p with large history window * remove as unused * add is_disabled method, tensor logits reshape, add rest review suggestions * clarify comment	2026-08-03 14:26:09 +02:00
3 people	1464c62d88	metal: implement DSv4 Lightning Indexer (#25893) ... * metal: implement F16 Lightning Indexer - Implement GGML_OP_LIGHTNING_INDEXER for 128-dimensional, 64-head inputs with F32 queries and weights plus F16 keys and masks. - Add tiled and tail kernels and test KV lengths around 8- and 64-element boundaries. llama-bench (--mmap 1, -fa 1, -p 512, -n 128; d=0/10k/20k/30k): Before: - pp512: 153.73 ± 0.87 t/s - tg128: 8.91 ± 0.04 t/s - pp512 @ d10000: 73.90 ± 0.39 t/s - tg128 @ d10000: 8.66 ± 0.03 t/s - pp512 @ d20000: 45.83 ± 0.18 t/s - tg128 @ d20000: 8.26 ± 0.03 t/s - pp512 @ d30000: 33.40 ± 0.21 t/s - tg128 @ d30000: 7.94 ± 0.01 t/s After: - pp512: 155.19 ± 0.91 t/s - tg128: 8.95 ± 0.04 t/s - pp512 @ d10000: 86.95 ± 0.69 t/s - tg128 @ d10000: 9.00 ± 0.05 t/s - pp512 @ d20000: 62.01 ± 0.45 t/s - tg128 @ d20000: 8.68 ± 0.04 t/s - pp512 @ d30000: 49.18 ± 0.33 t/s - tg128 @ d30000: 8.60 ± 0.02 t/s Assisted-by: Codex * metal: stage Lightning Indexer K tiles - Stage and dequantize K in F16 threadgroup memory before simdgroup matrix loads. - Zero-fill partial tiles and guard stores so all KV segments use the same numerical path. - Support F32, F16, BF16, Q4_0, Q4_1, Q5_0, Q5_1, and Q8_0 K caches. llama-bench (--mmap 1, -fa on, -p 512, -n 128; d=0/10k/20k): - pp512: 160.38 +/- 1.01 t/s - tg128: 9.08 +/- 0.03 t/s - pp512 @ d10000: 88.37 +/- 0.46 t/s - tg128 @ d10000: 9.07 +/- 0.04 t/s - pp512 @ d20000: 62.53 +/- 0.46 t/s - tg128 @ d20000: 8.84 +/- 0.03 t/s Assisted-by: Codex * dedup Lightning Indexer constants, fix flaky test * cont : fix whitespace --------- Co-authored-by: forforever73 <690105611@qq.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-08-03 07:33:37 +03:00
Georgi GerganovandThiago Padilha	fffbcbdb9d	metal: implement DeepSeek V4 hyper-connections (#26459) ... - Implement GGML_OP_DSV4_HC_COMB, GGML_OP_DSV4_HC_PRE, and GGML_OP_DSV4_HC_POST with SIMDgroup register and shuffle optimized kernels. - Add Metal dispatch and support plumbing and test the production Sinkhorn iteration count and embedding width. Assisted-by: Codex Co-authored-by: Thiago Padilha <thiago@padilha.cc>	2026-08-02 21:06:02 +03:00
Aman Gupta	596a5795bd	DeepseekV4 MTP + DSpark (#25784)	2026-08-02 20:55:34 +08:00
Aldehir RojasandBart de Boer	f5919bf458	chat : add qwen3 specialized parser (#26252) ... * Add tagged thinking tool parser * chat : refactor and add permute helper * cont : add support for <tool_call> omission * cont : update tool delimiters * cont : add comment for qwen3-coder * cont : fix trigger pattern for <function --------- Co-authored-by: Bart de Boer <bart.deboer@gmail.com>	2026-08-02 04:13:20 -05:00
Masashi Yoshimura	7a2db1a0cf	ggml-webgpu: add support for f16 repeat (#26307)	2026-08-02 08:28:31 +02:00
Xuan-Son Nguyen	11924d4c17	test: fix some CI errors (#26415)	2026-08-02 00:16:29 +02:00
Piotr Wilkin (ilintar)	ddd4ec1428	chat : enable tool call in thinking for DS4 (#26269)	2026-08-01 00:13:07 -05:00
Anand Patil	876a432116	vulkan: add POOL_1D op (#25431) ... * vulkan : add pool1d push constants and pipeline field Declared data structures needed for POOL1D OP, which are the vk_op_pool1d_push_constants struct and pipeline_pool1d_f32 field. * vulkan : add pool1d compute shader Added pool1d.comp for Vulkan backend mirroring the existing pool2d shader. * vulkan : add full GGML_OP_POOL_1D support Added pipeline creation and op dispatch for 1D pooling in the Vulkan backend. * vulkan : fix pool1d shader logic Registered pool1d_f32 in vulkan-shaders-gen.cpp and fixed tensor dimension indices and avg pool scale. * vulkan : fix pool1d end boundary crash and expand test coverage Fixed an issue where the shader crashed when the end boundary was negative when k0 < p0. Also, added more test cases related to this fix.	2026-07-31 16:48:58 +02:00
3 people	9d9a6d29f6	SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc… (#25025) ... * SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing * fattn-mkl: fix interleaved dst layout in normalize kernel - Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include <cfloat> - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * fattn-mkl: enable MKL FA for all KV cache types Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types. * fattn-mkl: rename mkl_disable -> mkl_enable for clarity * fattn-mkl: refine MKL FA dispatch gates Three changes: 1. Remove quantized-only restriction - MKL FA activates for all KV cache types (F16 default, BF16, F32, quantized). The MKL kernel converts non-F16 K/V via to_fp16_sycl before GEMM. 2. Rename mkl_disable -> mkl_enable to match env var (GGML_SYCL_ENABLE_MKL_FA). 3. Replace batch-size threshold with Q->ne[1] >= 32 gate. Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where fused kernel beats MKL launch overhead. Routes all multi-token prefill through XMX-accelerated GEMM. Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse, 128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s. * ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (#25222) which surfaced the same insight: always normalize inputs to contiguous F16 before GEMM. Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> * fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests Adding K>=1024 flash-attn test cases surfaced several MKL bugs: - Quant K/V with a padded seq-view (real KV cache) used the wrong strides in the dequant path... only the true Gemma interleave layout should reconstruct strides. nb[2] vs ne[1]*nb[1] - Gate was firing on shapes the kernel doesn't handle: head_dim < 64 or not a multiple of 64, MHA, attention sinks, and bf16 decode... fell through to vec which no bf16 case. Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512 (has to be a multiple of 64) with matching K/V head size, mask, no sinks/alibi/softcap... everything else falls back to tile. Covers Qwen Dense/MoE and Gemma4 Dense/MoE Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass. Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using Qwen 27b q5_k_xl * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review. * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners --------- Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>	2026-07-31 10:43:16 +03:00
pmaybank	958d9c0b61	Test support for alternative conv layout (#25617) ... * add bool cwhn = true to conv_2d test cases * add layout check at graph building time * extend layout checks for conv2d.cu kernel * in CPU back-end kernel needs to be stored contiguously to prevent test failures with cwhn=1 * trim white space * do op support check in vulkan backend * fix CI failure and vulkan run-time assert failure by introducing new graph build-time check in ggml_backend_vk_device_supports_op * add additional check in support_op function for Vulkan to fix run-time assert failure	2026-07-31 01:14:16 +08:00
Georgi Gerganov	47f686f53f	tests : avoid building get-model.cpp many times (#26317) ... * tests : remove get-model.cpp * tests : fix quant type selection	2026-07-30 19:34:04 +03:00
Robert Esclapez	e1a1abb787	ggml-cuda: Allow transpose-free gemmv computation (#26171) ... When matrix's weights are shaped 1xK is leverage a transpose-free computation to use mat_mul_vec_f.	2026-07-30 21:39:46 +08:00
Pasha Khosravi	9b2a088819	CUDA: add Q2_0 support (#25707)	2026-07-30 12:33:25 +03:00
Niklas WenzelandYiChen Lv	d0bfb19812	metal: fix memory unwire if model is freed without any GPU operations (#26082) ... * metal: fix memory leak if model is freed without any GPU operations * metal: run dummy work only if residency sets are used * metal: wrap function in #if defined * metal: measure system-wide wired memory in test * metal: always build regression test Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>	2026-07-30 11:11:27 +03:00
Neo ZhangandGeorgi Gerganov	a6a77bc48d	[UT] enhance UT to show all real unsupported backends (#25234) ... * enhance UT to show real unsupported backends * cont : simplify --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-07-30 14:04:58 +08:00
Reese Levine	bc71c24c9d	ggml-webgpu: Fix some binding alias issues to support all archs, fix recurrent-state-rollback test (#25931) ... * Add overlap glu variant to support all archs, fix recurrent-state-rollback test * format * Fix all arch overlapped ranges * format * diagnose bus error on apple ci * More testing * more testing * more targeted testing * Fix bug in alignment for > 4gb buffer offsets * Fix bug in view offsets * Try avoiding multi_buffers * not fixed yet, more logging :( * Handle edge case in set_rows * Try looking at view source * Skip deepseek32 for now and clean up trace infrastructure * simplify skipping * last cleanup * actually final cleanup * update handling of overlap * format * try skipping other failing model	2026-07-28 21:13:06 +03:00
Bhavik Sharda	b62b350981	ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration (#22675) ... * ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration * cuda: added SSD CICD fixes for CUDA / HIP / MUSA / MSVC. * ggml-cuda: review comments fixed. * ggml-cuda: Fuse M matrix materialization into pre_matmul kernel and enabled test. * ggml-cuda: test updates and fixes * ggml-cuda: test updates to remove hardcoding of tensor initialise data limits. * ggml-cuda: ssd minor review comment fixed. * ggml-cuda: ssd minor CICD fixed. * CUDA SSD: Fixes correctness by promoting s0_stride_seq to int64_t, improves memory coalescing in ssm_ssd_prepare_dt_kernel, and boosts efficiency by merging B_weighted and C_scaled; also addresses prior review comments. * cuda: fix sdata read-write race in prepare_dt fallback scan loop	2026-07-28 17:33:42 +05:30
Aldehir Rojas	6ba5ef2470	common/chat: add specialized minimax m3 parser (#26210)	2026-07-28 04:27:20 -05:00
meatposesandNeo Zhang	d6b61ac0d3	sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path (#25880) ... * sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path The scale was uploaded with an async memcpy sourced from a stack local. On the in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale. Output then collapses to a single repeated token and the KV cache is poisoned for the rest of the session. Short contexts win the race by accident, and test-backend-ops caps FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the symptom only by keeping the frame alive across the copy at the cost of a host sync on every FA call. Fix: cache one device scalar per (device, value) -- the scale is constant per model -- and upload it synchronously once. The single-device fast path (no per-call host sync) is then safe: every device-side hazard already serializes on the in-order queue. The multi-GPU conservative wait is kept unchanged. Also: - GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that routes very long sequences to the native FA kernel. - test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind spot. Note the race itself needs a live multi-op pipeline to reproduce; single-op runs pass even on broken builds. Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output byte-identical at temp 0 to the native FA path through 32k-deep prefill, with prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth). Assisted-by: Claude Fable 5 * sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it Review feedback on #25880: - read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it) - document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md runtime table Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536 shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels (long-sequence fp16 accumulation drift, present before this PR) and would fail CI for an unrelated reason. Assisted-by: Claude Fable 5 * sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled Assisted-by: Claude Fable 5 * sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly Assisted-by: Claude Fable 5 * Update ggml/src/ggml-sycl/fattn-onednn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * sycl: write the SDPA scale from a kernel instead of caching it The per-(device, value) scale cache was a function-local static unordered_map with no synchronization, so concurrent backend instances could access and rehash it at the same time. Write the scalar with a single_task instead. The value is captured into the command, so no host memory has to outlive the call -- which is what the use-after-return fix needed in the first place. That removes the shared container, the leaked device allocation and the string key, and it also closes the remaining async-memcpy-from-a-stack-local on the first flash-attention call. Ordering does not rely on timing: the queue is created with sycl::property::queue::in_order and the dnnl stream wraps that same queue, so the write completes before the SDPA reads the scalar. The multi-GPU wait_and_throw() branch is unchanged. Also drop the <cstdlib> include, which is unused. Assisted-by: Claude Opus 5 --------- Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>	2026-07-28 11:37:25 +03:00
3 people	f95de9776b	ggml-metal: FWHT kernel for metal backend (#25924) ... * metal fwht wip * shape guard and formatting * formatting * Formatting and typos Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * fix narrowing issue Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * cont : minor style --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-07-28 10:44:06 +03:00
Georgi Gerganov	7ef790f90a	tests : remove unnecessary sync in test-save-load-state (#26166)	2026-07-27 13:11:20 +03:00
Aaron Teo	ad256ded30	args: add -lm mlock where it mlocks but doesnt mmap (#26135) ... * arg: add `-lm mlock` where it mlocks but doesnt mmap Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: rm unwanted docs changes Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: revert auto-formatting Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * bench: fix automated review point 3 Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: revert the meaning of --mlock to non-mmap'ed mlock Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: update docs Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: remove extra changes from `llama-gen-docs` Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>	2026-07-27 16:44:08 +08:00
Xuan-Son Nguyen	d2a818231e	common: add subproc.h wrapper, disabled on android/ios (#26102) ... * add common/subproc.h|cpp * add compile flag LLAMA_SUBPROCESS * disabled by default on android and ios * test-jinja: use common subproc * mtmd: disable video if subproc is not set * disable subproc on wasm * make is_created atomic * migrate server-mcp	2026-07-26 20:54:25 +02:00
4 people	b1d4c65524	model: Add MiniMax-M3 (MSA: MiniMax Sparse Attention) support (#24908) ... * Add preliminary MiniMax-M3 support Text-only port that re-uses existing components: MiniMax-M2 style GQA with per-head QK-norm and partial rotary, DeepSeek-V3 style leading-dense and routed/shared experts, and swigluoai activation. Sparse attention is not yet supported (dense fallback); vision tower and MTP heads are dropped. * MiniMax-M3 vision tower (mmproj + clip graph) * Delete m3_vision_ref.py * Update clip.cpp * MSA * Update constants.py * Update minimax.py * Cache creation. Working withotu flash attention * Added flash attention for sparse layers * Decomposed slow cpu OP into GPU + CPU ops. Massive speedup over long ctx * Rewrote indexer op to be cuda native. Modified flash attention to match per group block picking * Implement sparse attention calc out of stock ops. * Fix a cache allocation and cont issue * Fixed -fa auto crash, flagged debug spots * Delete vocab.json * Delete model.safetensors.index.json * Delete generation_config.json * Delete Minimax directory * Handled multi stream case to fall back on Dense Attention * Development scaffolding cleanup. No functional change to the decode or 4-way paths. Full debug harness remains at <8136a9c68ed7a5eb009aa67bba3fda8062f4648f> for reproducing the selection-parity validation. * Remove redundant comment from minimax-m3.cpp * Changed 3 Gelu Ops for vision into Gelu_erf ops * Assert that n_kv is multiple of 128 * Rename MSA index tensors to indexer convention Note: All GGUFs generated before this change will need to be regenerated. * Fix incorrect Assert * Review driven changes (#3) * Remove comment from conversion minimax.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Remove whitespaces from constants.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Tighten comment in minimax.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * inherit MiniMax-M3 from MiniMax-M2 * drop dead text_config fallbacks * Add indexer writer methods * Reuse LLM_FFN_SWIGLU_OAI_MOE * Remove duplicate indexer setters, add only block_size/local_blocks, follow value naming convention * Fix conversion error /gguf_writer.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Update gguf-py/gguf/gguf_writer.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Update gguf-py/gguf/tensor_mapping.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Update conversion/minimax.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Update conversion/minimax.py Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Remove whitespace in src/llama-kv-cache.cpp Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Remove Whitespace in Update src/llama-model.h Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Remove whitespace in src/llama-hparams.h Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * remove multimodal code upon maintainer request. Will be made as a separate PR * Whitespace clean in tensor_mapping.py * Log cache size on launch, block ctx shift, support prompt caching Log indexer cache size on launch Disallow ctx shift Support prompt caching * Update minimax-m3.cpp * Optimize implementation, add multi stream support. Fully rewrote minimax-m3.cpp for speed and buffer size gains: Unified the 4-way + decode, 1 FA call per layer instead of 4, with the groups mapped onto ne[3] Custom CPU op now emits block-level mask, expanded on GPU, which causes CPU to GPU transfer to shrinks at prefill Decode: ~25 nodes/layer vs ~50, no per-group concats/conts Unified selection semantics, so both regimes rank bs + local bias (position-anchored local force), which means prefill/decode can no longer disagree on selection can_reuse on the MSA bias input. Graph reuse at decode restored (was rebuilding the full graph every token) In-place mask adds, shrinking compute buffer ~6.8 to ~4.2 GiB at ub2048/62k Multi-stream: MSA now runs with -np N when kv_unified=false. Decode stays batched across streams (still 1 FA call), prefill loops per stream. dense fallback only for --kv-unified + multi-seq Measured effect on expert offload bound setup: decode 6.2(4WAY)–7.15(MSA_decode) -> 7.7~7.8 t/s, flat from 5k to 60k+. prefill around 10% faster. buffer about 20% smaller, multi-user support. * set default cache type to F32 * Fix potential DSA double indexer cache allocation bug, only allocate in-cache k_idx for archs that opt in * remove F16 downcasts in MSA attention, force F32 indexer score accum * Add Minimax eos to llama vocab * Guard edge case where idx cache can become stale after a tail trim * Update llama-kv-cache.h * Update llama-kv-cache.cpp * Update llama-kv-cache.cpp * Update llama-kv-cache.h * Update llama-kv-cache.cpp * Review driven changes * style fix * indexer hparams are required * fix tests * fix lint --------- Co-authored-by: Daniel Han <danielhanchen@gmail.com> Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> Co-authored-by: Xuan Son Nguyen <son@huggingface.co>	2026-07-26 19:43:45 +02:00
Aldehir Rojas	910196f6b3	common : add support for multiple end sequences in the reasoning budget sampler (#25544) ... * common : extract trie/ac to a separate file * common : support multiple token sequences in the reasoning budget sampler * common/trie : return matched word index * common/trie : rename "word" to "pattern" * common/reasoning-budget : expose matched end sequence * common/sampling : replay end sequence when reasoning budget is done * cont : update to use multiple end sequences * cont : clean up	2026-07-25 11:58:09 +02:00
helanfxz	d67c0b4107	tests: synchronize save-load-state generation (#26056)	2026-07-25 10:23:31 +02:00
Matt Thompson	8f5ab832ca	cohere2 moe template parser: enforce JSON schema for text responses if a response schema is provided (#26018)	2026-07-24 12:54:47 +02:00
Aaron Teo	e6dd0e29a6	args: refactor mlock/mmap/directio into load-mode (#20834) ... * args: overhaul mmap/mlock/dio into single arg Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: update docs with llama-gen-docs Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * chore: satisfy code quality Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * args: make the `+` sign an actual modifier now Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * chore: general code clean up + comments Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: fix deprecated flags support Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: quick sanity check Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * bench: sync llama-bench argument parsing Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * fix: bugfix variable behaviour + llama-bench lm column size Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: inverse commands should do the opposite instead of doing nothing Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * bench: fix incorrect dash Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * bench: fix missing modifiers for deprecated flags Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * llama: switch back to thread_local Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: switch back to single enum Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: update arg docs Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * chore: fix missing `mlock` from llama_load_mode_from_str + cleanup llama-bench Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * llama: fix mlock not activating Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: add deprecation warning when old and new flags are combined Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * arg: cont add comment for todo in the future Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: sync with upstream Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: re-sync with upstream again Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>	2026-07-23 20:32:56 +08:00
Piotr Wilkin (ilintar)	f534da26e4	Fix DeepSeek4 crafted template (#25414) ... * chat: fix DS4 template to explicitly follow reference behavior * Support DeepSeekv4 flag (`drop_reasoning`). * fix: hook DS3.2 parser for DS4 as well * fix: add tool result reordering * fix: post-merge	2026-07-22 12:54:40 +02:00