The hidden desktop strip relied on inherited `visibility: hidden` to keep
its content from catching clicks. Since the sidebar's sections fold
(#31936), an open section's CollapsibleRegion sets `visibility: visible`
again, and its rows' context-menu triggers carry an inline
`pointer-events: auto`, so they caught clicks meant for the page.
Make the hidden strip's content inert, which no descendant can undo, and
add a client test that hit-tests the hidden sidebar's area.
Fixes#31971
Assisted-by: Claude (Opus 5.5 Medium)
* media : remove the legacy llama0 and llama1 brand assets
The brand assets live in ggml-org/llama.brand; none of these is referenced
in the repository.
Assisted by: Claude (Opus 5.5 Medium)
* readme : use the new llama.cpp cover from llama.brand
Assisted by: Claude (Opus 5.5 Medium)
* ui: show audio in Discover for models whose mmproj has an audio encoder
Discover took audio only from the pipeline tag, so omni models tagged
any-to-any, such as Gemma 4 and LFM 3 Omni, showed vision but never audio,
though their mmproj carries an audio encoder.
Read the encoders from the mmproj's GGUF metadata instead: a ranged read of
its first 64 KB, once per repo, for clip.has_vision_encoder and
clip.has_audio_encoder. The Discover list and details now show audio when
either the pipeline tag or the mmproj says so.
Assisted by: Claude (Opus 5.5 Medium)
* ui: take Discover modalities from the mmproj's encoders in one place
Review fixes for the mmproj header read:
- drop vision for audio-only mmprojs: the converter writes
clip.has_vision_encoder only when true, so a header read in full without
it has no vision encoder; a header cut short leaves vision as it was
- accept only a 206 to the ranged read, so a server that ignores Range
never streams the whole mmproj
- encode the mmproj path in the resolve URL
- derive the list and details modalities with one HuggingFaceService
helper, so both views agree
- point the key comment at tools/mtmd/mtmd-impl.h, where the keys live now
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep a gated repo's mmproj from being read again on every mount
A gated, private or missing mmproj answers 401, 403 or 404 to every read,
so caching nothing made each mount of its Discover row request it again and
log an error. Cache it as unknown instead: no audio, and vision from the
mmproj being there, as before the header was read.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : move Models, MCP and Settings to a sidebar footer
The sidebar's actions split in two: New chat, Search and Projects stay at
its top, and Models, MCP and Settings move to a footer at its bottom, under
a thin rule, in the expanded sidebar and its collapsed strip alike. On a
phone the footer keeps clear of the home indicator, and Models stays
desktop-only as before. The search dialog still lists every action.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : group the sidebar footer in an inset surface
The design system gains inset-group-surface, for rows grouped inside a glass
panel: white on the light glass and a faint tint of the text colour on the
dark one, with the sidebar's own ambient edge, no inner highlight, a faint
wide shade under it and the card radius. It is matte-glass pointed at its
own component tokens, like the sidebar and chat form surfaces, and reads as
part of the panel: no tight contact shadow or solid border like a button,
no inset shade like a field.
The sidebar footer uses it in place of the rule above Models, MCP and
Settings. DESIGN.md and the Surfaces page document it.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : give the sidebar's logo button the add button's key look
Whenever the sidebar shows, open or peeking, its logo button takes the same
key as the chat form's add button (matte-key, white in light mode, a light
tint in dark), in place of its own muted fill. While the sidebar is hidden
the logo has no fill and no shadow, and hover and focus never change the
fill. It keeps its round shape and its own size.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : put search in a round key at the top of the phone sidebar
On a phone the sidebar's Search row moves to a round key at the right end
of its top row, across from the logo, with the same key look as the logo
and the chat form's add button. It opens search as the row did. Desktop
keeps the Search row.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : size the phone sidebar for a thumb and list its footer with the actions
On a phone the sidebar's actions and rows are 48px tall at 17px, flush one under
the other with 16px corners, its section headings 15px and grey, and its groups
28px apart. MCP and Settings join the top actions there: the footer is a desktop
pattern.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : mark the selected sidebar row with a tinted inset surface
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the phone sidebar's round keys white on the grey page
Assisted-by: Claude (Opus 5.5 Medium)
* ui : open Settings and MCP as pages on a phone, dialogs elsewhere
On a phone the sidebar opens them at #/settings and #/mcp-servers, with the menu key and a
title like any phone page. Pinned under the status bar, Settings' section tabs fill the
screen above them in the page's colour.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : open a phone sidebar's pages behind it, settled before they slide back
A tap in the phone sidebar now opens its page while the current one is still aside and slides
back once it has rendered, instead of swapping pages mid-slide. The slid page's fixed bars are
pinned again as the page under them changes, the lock ends with the slide itself, and a
phone's new chat shows its form and greeting at once, without the move and entry they have
on a desktop.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: style the sidebar footer group and the selected row with master's classes until the design system lands
Assisted-by: Claude (Opus 5.5 Medium)
* ui: add GlassKey with master's styles for the phone sidebar's round keys
Assisted-by: Claude (Opus 5.5 Medium)
* ui: center the top edge of a new chat's form on the screen
The form rested 42px below the middle up to 2xl and 6px above it past
that; one drop, from the resting card's height, now puts its top edge on
the middle at every width from md up.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: round the sidebar's action rows concentric with the footer group
The sidebar stacks four surfaces: the page, the sidebar, the footer group
inset in it and a row in the group. app.css now names the group's radius
and padding, and derives the row radius from them (16px less 4px), so the
footer's rows run parallel to its corners; the actions at the top of the
sidebar share that radius. The footer moves to its own component, the
README documents the pattern, and a client test holds the rule on the
rendered footer.
A selected action row was bg-accent, the same grey as a hovered row in
light and darker than the footer's fill in dark, so it read as hovered or
not at all; it is now raised and edged, while a hover only tints.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: import demo chats through the Settings page on a phone
A phone opens Settings as a page now, not a dialog, so the demo harness
waited for a dialog that never came and every phone demo with traces
failed; it now imports through whichever of the two opens.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: count the sidebar's rows from their items in the surface test
The test expected three footer rows and two top actions, which the
stacked PRs change as they add Projects, Library and Servers; it now
expects as many rows as the sidebar has items.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: add a line on Meta+Enter off macOS too
The chat form takes Meta+Enter as Ctrl+Enter and leaves the newline to
the editor, but the editor bound only Mod-Enter, which is Ctrl off
macOS, so on Linux Meta+Enter neither sent nor added a line. This failed
the client test on CI's Linux runners, on master too.
Assisted-by: Claude (Opus 5.5 Medium)
This commit removes the unused ids and type_src0 parameters from
quantize_row_q8_1_cuda.
The motivation for this is that these parameters were required in the
past because this function was used in combination with the function
pointer quantize_cuda_t, which has been now been removed, and this
function no longer needs to adher to the function pointer signature.
Refs: https://github.com/ggml-org/llama.cpp/pull/31884.
* Add FP8 KV-cache support
* Add tests for scaled, FP8 KV-cache to test-backend-ops
* Fix get_*_scale mapping and dequnatize FP8 also before rope shift
* Normalizes cache-scale names independently of projection layout.
While fused KQV is currently impractival for non-shared scales, this is
an improvement for future use regardless
* Enable Gemma 4 Assistant with FP8 by storing and fetching scale ptrs
* Update comment wording
* Move E2E FP8 KV cache tests to test-llama-archs
They belong there as they are not just a test for ggml's backend ops
* Increase support surface, remove seed
* make scaled kv generic
* Fix CI failings for test-kv-cache
https://github.com/ggml-org/llama.cpp/actions/runs/38718763447/job/116367842578?pr=28392
* CI: execute kv-cache-scale tests
* exclude fp8-kv tests from OpenVINO
* Fix linting errors of test_kv_cache_scales
* Disable kv-ffload=false + multi-GPU tests for FP8
See https://github.com/ggml-org/llama.cpp/actions/runs/38952722084/job/117148353699?pr=28413
* Adjust test label for test-kv-cache.cpp
* Fix conflict-resolution mistakes in test-kv-cache.cpp
https://github.com/ggml-org/llama.cpp/actions/runs/40554674463/job/122367426461?pr=28392
* Move KV tests to test-fp8.cpp
https://github.com/ggml-org/llama.cpp/pull/27512#discussion_r4410518158
* Remove tmp file for test-kv-cache-scales.py on Windows as well
* Fix fp8 kv cache test by population kv_source_impl
* Exclude models that consume KV cache directly
* Remove obsolete comment
* Allow BF16 scales as well for KV cache
* Rename *_cache_scale to *_cache_s
https://github.com/ggml-org/llama.cpp/pull/28392#discussion_r4517194999
* Move `is_type_scaled()` to `llama-kv-cache.cpp`
* ui: let a menu action open a submenu
Group and project pickers open as flyouts from a chat's menus. Master's menu
list has no submenus, so bring them back next to its tooltips, with a close
callback for picks made inside them.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : store groups and let conversations join them
Groups live in their own IndexedDB table (schema version 2) with a name.
groupsStore creates, renames and deletes them; deleting a group keeps its
chats, out of any group. A conversation records the group it is sorted
into, and a group picked before the first message is kept pending and
applied when the conversation is created.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : show each group as a sidebar section
Every group is a sidebar section next to Pinned and Recent, folding,
showing its first six chats and dragged into place like them. A chat
dropped on a group joins it and leaves Pinned; dropped on Recent, it
leaves its group. Recent is called Ungrouped once a group exists, and its
menus create a group in a small naming dialog; so does a right-click
anywhere in the list that no row or heading takes, down to the empty space
under the last section. A group's heading starts a new chat in it, and its
menus rename or delete it; deleting keeps its chats, back in Ungrouped.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : move chats into groups from their menus
A chat's menus, in the sidebar and on right-click, move it into a group
or take it out of its group. Desktop menus show a searchable group flyout
whose New group entry names a group in the group dialog and moves the chat
into it; a phone gets a picker dialog instead, which creates a group from
what is typed.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: open the group picker from a phone's hold menu
The hold menu is anchored to the row, so it showed the desktop flyout,
which ran off the left edge of the screen and whose New group closed
without naming one. On a phone, Move to group now opens the picker
dialog, as intended.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: open a new chat in a tab of its own
With conversation tabs on, New chat from the sidebar, a group's heading
or the shortcut replaced the open chat's tab. It now adds a tab, as the
tab bar's "+" does, and the open chat keeps its own.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: list a phone's groups under the picker's search
On a phone the picker fills the screen, and its grid spread the rows
over the height, so the list sat in the middle. It is now a column: the
list starts under the search and scrolls in the rest of the screen.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fix gaps in chat groups found in review
- A phone shows a group heading's New chat and menu, and the Ungrouped
heading's menu, which were reachable only by hover or right-click.
- A pinned chat moved into a group from its menu leaves Pinned, as a
drag onto the group already did, so the move shows.
- Only the sidebar's New chat, a group's New chat and the new-chat
shortcut open a new tab; stepping past the last chat, search and the
models drawer show the new-chat screen in the current tab again.
- An export carries each chat's group, and an import sorts the chat
into the group of that id or name, or creates it; a group the file
does not carry is dropped.
- The group picked with a group's New chat survives leaving the
new-chat tab and coming back.
- A phone's menus skip any submenu in the shared menu list, not only
the group one.
- Deleting a group forgets its section's place and fold.
- Groups load newest first by creation, so a rename no longer moves one
after a reload.
- groupsStore looks groups up in a map and answers which live group a
chat is in, for the three places that checked it by hand.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep a forked chat in its source chat's group
A fork copied the working directory but not the group, so it landed in
Ungrouped.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: move a chat with its whole fork family between groups
Moving a chat into a group, or out of one, takes its parents and its
children along, so a fork tree never splits across sidebar sections. A
pinned member that joins a group leaves Pinned with the rest. Group
sections render the tree with its depth, as Ungrouped does.
* ui: show the group name prompt as a small centered card on phones
The new and rename group prompt opened full screen on a phone, like a
large dialog. It is now built on the alert dialog, as the chat rename
prompt is, so a phone shows it as a card in the middle of the screen.
Its Cancel is a plain button, so Enter in the name field saves.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: space a phone's group sections like Ungrouped
On a phone a group or Pinned section put its first chat right under the
heading, where Ungrouped leaves a row gap. They now leave the same gap,
and the sections sit a little further apart than the rows inside them.
Assisted-by: Claude (Opus 5.5 Medium)
---------
Co-authored-by: Pascal <admin@serveurperso.com>
This commit removes the unused quantize_cuda_t from quantize.cuh.
The motivation for this is that the usage of this typedef as a function
parameter was removed in Commit 74976e1aef
("CUDA: remove -sm row, refactor cuBLAS (#24216)") and there is no longer
any usages of it.
The sampling chains are read on the host after the graph compute, and the
graph addresses are reused across builds with different selected rows: the
token-generation reserve graph has one output row per sequence, while the
actual block draft has n_max rows per sequence. The reserve allocates the
unselected chains without keeping them, so their buffers alias, and the
actual graph then reuses those addresses by node index.
set_outputs only marked the tensors stored for the selected rows, so the
unselected per-(sequence, row) chains were freed and reused. Collect every
chain tensor in t_sampling_all and mark all of them as outputs there.
Reproduces with -np 2 (non-unified KV) and 2 concurrent requests for
draft-dspark, draft-dflash and draft-eagle3; with different prompts one
sequence's acceptance collapses (eagle3 0.15 vs 0.40). After the fix:
dspark 0.71/0.78, dflash 0.56/0.44, eagle3 0.35/0.42.
Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4.1-Flash-1017-GGUF (Mac Studio, M5 Ultra, 256GB)
* ui: expand the chat form into a large editor for longer messages
The chat form expands into a large editor once a message has markdown
beyond plain paragraphs or reaches four lines, and a Maximize button in its
top-right corner toggles it by hand. The expanded form starts around half the
screen height and grows with the text up to the space below the tabs bar,
animating its height by mirroring the editor's content height.
On a new chat the expanded form keeps the collapsed form's top edge and the
greeting only fades out, waiting for a collapse to settle before it fades
back in. In a conversation the messages follow the form as it grows, so the
latest one stays in view above it.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: make the expanded chat form a full-screen sheet on phones
In the expanded editor Enter adds a line and Cmd/Ctrl+Enter sends.
On a phone the expanded form is a full-screen sheet with no corners or
border. The page around it and Safari's bars take its colour, and the
sidebar's logo button and the working-directory chip step aside while it is
open. It moves by relative offsets and a clip on the main thread, so its
edges, text, actions and chip stay in step in every engine, and its content
keeps the collapsed card's places and widths, so nothing reflows when it
settles. Expanding does not focus the editor on a phone.
On a new chat on desktop the form is lifted by its sticky bottom offset
instead of a translate, and the expanded size is built on the collapsed
card's real height, so its top edge holds still through any sequence of
quick toggles.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: hide the phone menu toggle and the sheet's focus edge while expanded
On a phone the menu toggle rides on the page, above the chat form: it now
fades out while the expanded form covers the screen, and comes back when
it collapses. The expanded sheet also drops the glass rim and inner ring,
so focusing the editor no longer outlines the whole screen.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the expanded phone sheet clear of the status bar
The expanded chat form fills the screen on a phone, under the status bar
an installed app draws over the page. Its text now starts below the bar
(pt-safe) and its collapse button sits below it too (top-safe-2.5).
Assisted-by: Claude (Opus 5.5 Medium)
* ui: hide the phone's context gauge while the chat form is expanded
The gauge pinned to the phone's top right sat over the expanded sheet's
collapse key. Like the menu toggle across from it, it now fades out while
the expanded form covers the screen and comes back when it collapses.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: lift the expanded phone sheet's actions above the iOS keyboard
Assisted-by: Claude (Opus 5.5 Medium)
* ui: move the chat form with a short, quick drift, without sliding after it lands
Its position transition came back on before the browser had laid out the new position,
so it slid the form there after the fade; the drift is now 10px each way and quicker.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: style the expanded chat form with master's surfaces until the design system lands
Assisted-by: Claude (Opus 5.5 Medium)
* ui: tighten the expanding chat form after review
- Expand on markdown only in finished lines: a lone `-`, `#` or `1.` on the
line being typed no longer opens the large editor.
- Show "Cmd/Ctrl+Enter to send" by the send button while expanded.
- Move the expansion state, the phone sheet and the held collapsed height
into useChatFormExpansion, and the sheet's geometry into pure functions.
- Drop the working-directory chip's phone animation, which never ran since
the phone chip moved to the + menu.
- Clearing or sending mid-move cancels the sheet's animation.
- A click only continues at the end of the text when it lands below it.
- Reuse ActionIcon for the toggle and INPUT_CLASSES for the card, derive the
centered form's offset from its drop, name the shared easing and fade.
- The greeting owns its fade behind the expanded form.
- MarkdownEditor's view is reactive, so following its content starts once it
exists.
Assisted-by: Claude (Fable 5.2)
* ui: keep the expanded chat form to the message being written
- The form expands on its own only for what is written or pasted, and
another conversation starts collapsed: a restored draft no longer opens
the large editor, or the phone sheet, on arriving at a chat.
- In the expanded editor the pickers open by the line being written, above
it or, short of room, below, instead of above the form's top edge.
- The sidebar's swipe is off while the phone sheet covers the screen.
- The send hint shows only beside the send button, where the row has room.
- The inner area keeps master's focus shadow, except on the phone sheet.
- The card and the form's position honour reduced motion.
- The phone greeting's place accounts for the wrapper it is counted from.
- The sheet's animation tolerates the form going away mid-move.
Assisted-by: Claude (Fable 5.2)
* ui: turn the typed pickers off on phones, and scroll long messages
- On a phone, typing `/` or `@` opens no picker for now: each mobile
browser moves the page its own way as the keyboard opens, so a list by
the form can not stay clear of the text being written. A note marks it
for later.
- The editor scrolls inside the form again: MarkdownEditor's own box sat
between the input and the editor without taking its height, so a message
longer than the form ran under the actions instead of scrolling.
Assisted-by: Claude (Fable 5.2)
* ui: settle the expanding chat form's review findings
- A phone expands the form only by hand: the sheet switching in mid-typing
moves the text under the keyboard.
- The typed pickers are off wherever a finger is the main pointer, not
below a viewport width, so a narrow desktop window keeps them and an
iPad loses them. The prompt picker opened from the + menu floats by the
sheet's actions on a phone.
- A message put back after a failed send expands the form again.
- Only the left button moves the caret from a click below the text.
- The sheet's fill is one token, used by the sheet, the page and Safari's bars.
- The collapsed-height hold is its own module and no longer finds the
editor's elements by class name.
- Shorter comments.
Assisted-by: Claude (Sonnet 5.5)
* ui: keep a one-line message on the chat form's actions row
On a desktop the main chat form is a single pill while its message is one
line: the add button, the text, then the model selector and send. A second
line, or a line that wraps beside the actions, brings back the two rows,
and they stay until the text is shorter than where it wrapped, so the
wider row can not flip it straight back. The editor stays mounted: the
actions join the text's row by CSS alone.
The single row's corners are `--chat-form-single-row-radius` in app.css, a
pill by default; the "Scrapbook / ChatForm single row" story shows the real
form with other radii.
Assisted-by: Claude (Opus 5.5)
* ui: give the one-line chat form 1.25rem corners
Assisted-by: Claude (Opus 5.5)
* ui: pick the working directory from the + menu on desktop too
The chip under the chat form is gone: the + menu offers the working
directory everywhere but a phone's chat, whose header shows it. The menu
item names the directory once one is set, and the picker gains a row to
clear it, which the chip's X used to do.
Assisted-by: Claude (Opus 5.5)
* ui: put the single-row chat form behind a Display setting
"Single-row chat form for short messages", on by default.
Assisted-by: Claude (Opus 5.5)
* ui: send the chat form on Enter, and add a line with Shift or Ctrl/Cmd
Enter sends wherever the caret is, the expanded editor and code blocks
included; Shift+Enter and Ctrl/Cmd+Enter add a line, continuing lists and
quotes. With "Send message on Enter" off the two swap as before. The
expanded editor's send hint and the code-block rule are gone with it.
Assisted-by: Claude (Opus 5.5)
* ui: start the chat form on the row layout it will settle on
Assisted-by: Claude (Opus 5.5)
* ui: show a set working directory under the chat form again
On desktop the chip under the form is back once a directory is set, and
the + menu offers the directory only until then.
Assisted-by: Claude (Opus 5.5)
* ui: add a line on Enter with an on-screen keyboard
A phone's Return has no Shift, so on a touch screen Enter adds a line and
the send button sends, as "Send message on Enter" turned off does; on a
desktop Enter still sends.
Assisted-by: Claude (Opus 5.5)
* cuda : Lowering VMM pool reservation size to 8 GB on GPUs with combined memory due to CUDA out of memory errors
Assisted-by: Gemini Argon
* cuda : Halve VMM pool reservation size and retry on failure
Assisted-by: Gemini Argon
* cuda : only shrink VMM pool on CUDA_ERROR_OUT_OF_MEMORY
Assisted-by: Gemini Argon
* cuda : log warning when VMM pool reserves less than max size
Assisted-by: Gemini Argon
ggml_cuda_graph, ggml_cuda_concurrent_event and ggml_cuda_stream_context
are only used by ggml-cuda.cu, but keeping them in common.cuh meant any
change to them rebuilt every CUDA translation unit. Move them to graph.cuh
and leave forward declarations in common.cuh. The context constructor,
cuda_graph() and any_cuda_graph_*() are now defined in ggml-cuda.cu, and
the stream context is held through a unique_ptr.
Also define ggml_cuda_kernel_can_use_pdl once in ggml-cuda.cu instead of
as a static function in the header, so there is one PDL cache instead of
one per translation unit, and fold the GGML_CUDA_PDL env check into it.
* ui: use a live-preview markdown editor for the chat form input
Replace the contenteditable rich input with a CodeMirror editor that shows
the markdown source of the block being edited and renders every other block
through the MarkdownContent pipeline and stylesheet, so they match chat
messages. Plain text types exactly as before, the value stays the exact
markdown source, and mention links show as atomic badges.
The editor lives in a shared MarkdownEditor component, and the highlight.js
theme loader is shared between MarkdownContent and SyntaxHighlightedCode.
The old rich input, its tokenizer, source history and their tests go away.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: give the production build more memory, note what mdsvex is for
The single-bundle build now includes CodeMirror and runs past Node's default
heap limit on CI runners. Raise it in the build script, so local and CI builds
behave the same on every platform.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the live preview's blocks while only the cursor moves
Moving the cursor re-walked the syntax tree and rebuilt every block's
decoration. Keep the blocks with the decorations and rebuild only when the
text changes or the cursor enters another block.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: point the chat form story at the editor and drop the engraved placeholder until the design system lands
Assisted-by: Claude (Opus 5.5 Medium)
* ui: name the markdown editor for screen readers
CodeMirror's content is a textbox with no name of its own, so the chat form
input was announced as an unnamed text box, and could not be found by its name
as the old input could. It takes an aria-label, its placeholder by default,
which makes the chat form's "Type a message..." again.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: format the chat form input's classes
Assisted-by: Claude (Opus 5.5 Medium)
* ui: share the markdown styles and diagram drawing between messages and the editor
The editor and MarkdownContent each @imported markdown-content.css into their own
<style>, so the bundle carried two scoped copies of it. Both now render into a
MarkdownSurface, the one element that carries the stylesheet; the few rules for
MarkdownContent's own template elements are global under .markdown-content.
Drawing mermaid and svg blocks and the actions of their header buttons move to
markdown-diagrams.ts, so the editor's rendered blocks can draw diagrams the same
way messages do.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: harden the live preview editor after review
- onChange skips only the editor's own syncs from `value`, marked by an
annotation, instead of reporting only edits that carry a user event
- Enter and the mention badges ask the syntax tree whether they are in code, so
~~~ fences and indented code count, and a closing fence ends a block
- a picked mention is one undoable edit, through replaceText
- rendered blocks draw mermaid and svg diagrams with their header buttons, a $$
math block with blank lines renders as one block, and a link reference
definition stays raw instead of rendering to nothing
- rendered blocks render again when the mention path settings change
- a drag in a rendered block selects its text and copy takes it, across blocks;
a click still opens the block
- the block margins are the ones a chat message renders (each heading is a first
child there), checked by a test; the blank lines after a raw block are its gap
- the tree comes from syntaxTree, without forcing a full parse on every edit
- placeholder, label, extensions, mention widgets and render key follow prop
changes after mount, through compartments
- the editor uses the markdown language alone, which drops the html, css and
javascript parsers markdown() pulls in for html tag completion (-70 KB gzip)
Assisted-by: Claude (Opus 5.5 Medium)
* ui: pin the CodeMirror packages and set the build heap in .npmrc
The new dependencies are pinned like the rest of package.json. The build heap
moves from a node call on vite's bin path to npm's node-options, so the build
script is plain `vite build` again and the setting reaches it on every platform.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: update the live preview block by block and estimate block heights
An edit used to build every block's decorations and widgets again, which the
editor then compared one by one. The previous decorations, and the mention badges
of raw blocks, now move through the edit, and only the blocks whose range, text,
raw or rendered state, or gap changed get new ones; the texts of the blocks an
edit left alone are reused instead of sliced again. A test drives random edits,
cursor moves and focus changes and checks each result against a full build.
Rendered blocks report the height they were last measured at, or a guess from
their text, so CodeMirror's height map for blocks off screen in a long document
stays close and scrolling stays steady.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: skip the editor's redundant syncs, reconfigures and theme reloads
- syncing `value` compares it with the text the editor last reported or synced, so
a keystroke coming back through onChange no longer serializes the document again
- the prop effects reconfigure the editor only when a prop changed, not once more
right after it is created (which rebuilt the whole preview)
- loadHighlightTheme leaves the theme's style element alone while the mode stays,
so a message or an editor mounting no longer restyles the whole page
Tests check that typing in one block leaves the other rendered blocks without a
single DOM mutation, and that mounting reconfigures nothing.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: simplify the live preview to a plain rebuild in three small modules
Rebuilding the decorations of every block on an edit measures the same as
patching only the changed ones (0.09 ms against 0.11 ms per keystroke at
200 blocks, 0.9 ms against 1.1 ms at 2000), and equal widgets keep their
DOM either way, so the incremental update and its test go.
- markdown-blocks.ts collects the document's blocks, markdown-block-widget.ts
renders one, markdown-live-preview.ts decides which ones show raw text
- one bounded cache holds a block's HTML and its measured height
- the render key is a plain option: a new one reconfigures the preview,
which keeps the editor's focus
- MarkdownEditor follows its props with one helper instead of five effects
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep setext headings raw in the live preview editor
`test` over `---` is a setext heading, so the preview rendered it as an h2
and the underline vanished once the cursor left the block. The editor now
shows a setext heading's text as typed at all times, like a link
reference definition.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: add a setting to turn the live markdown preview off
"Live markdown preview in the message input" (General) is on by default.
Off, the input shows the plain markdown text, with file mention badges
still in place; switching it takes effect without a reload.
Assisted-by: Claude (Opus 5.5 Medium)
normalize_quotes_to_json copies raw control characters verbatim while
converting pythonic tool arguments to JSON, so calls like
python(code='a<LF>b') produce invalid JSON in the arguments returned to
clients; replaying that history then fails hard. Escape C0 control
characters inside quoted strings via escape_json_string_inner.
* sycl: Add device-side MoE scheduler
Assisted-by: Codex
* Add width-heuristic override bit and Q2_K and Q3_K types
Assisted-by: Codex
* tests: cover skipped MUL_MAT_ID slots on the device MoE schedule
Assisted-by: Claude Opus 5.5
* sycl: keep graphs on for lazily reordered MUL_MAT and reject BN=64 combos on the device MoE schedule
The graph check disabled SYCL graphs until every MUL_MAT weight with a
reorder extra was reordered, which never happens for some (dense MXFP4,
the fused GLU plain path), so graphs stayed off and the check reran on
every compute. Upstream already reorders those during recording with
the async memory extension, so only MUL_MAT_ID weights are reordered
ahead of recording now.
ggml_sycl_grouped_dequant_gemm_device_supported() accepted any XMX
combo, but the device schedule only runs with BN == GGML_SYCL_FG_BN, so
a 32x64 combo (GGML_SYCL_XMX_GATHER_SHAPES=4 or 8) made the graph check
pass and the compute fall back to the host sync path while recording.
Assisted-by: Claude Opus 5.5
* sycl: keep the ggml_sycl_mul_mat_id_needs_sync name for the graph check
Assisted-by: Claude Opus 5.5
* fix: stop erroneously reordering host-offloaded tensors
With -ncmoe, prompt processing offloads the CPU-resident expert matmuls
to the GPU: the scheduler uploads the experts into a copy in the compute
buffer before every run. init_tensor gave that copy a reorder record, so
the MUL_MAT_ID gather path reordered every offloaded expert tensor on
every ubatch, and the next upload overwrote the result.
Qwen3-30B-A3B Q4_K_M, 1x B60, llama-perplexity -c 4096 -ub 2048
-ncmoe 48, 8 chunks, on master ac8e1694:
master 116 s, PPL 7.739
master XMX_GATHER_TYPES=0 67 s, PPL 7.738
with this change 67 s, PPL 7.729
Refs ggml-org/llama.cpp#31262
Assisted-by: Claude Opus 5.5
* Remove SYCL-specific test comments
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* sycl : name the MUL_MAT_ID path predicates after what they decide
ggml_sycl_mul_mat_id_mmvq_ok -> ggml_sycl_mul_mat_id_should_use_mmvq
ggml_sycl_mul_mat_id_mmvq_shape_ok -> ggml_sycl_mul_mat_id_mmvq_eligible
ggml_sycl_mul_mat_id_grouped_shape_ok -> ggml_sycl_mul_mat_id_device_sched_eligible
node_is_device_local -> graph_node_is_in_local_device
ggml_sycl_gg_rows -> ggml_sycl_grouped_gemm_rows
ggml_sycl_gg_tile -> ggml_sycl_grouped_gemm_tile (from #29245, renamed to match)
"eligible" is a necessary condition; the type and width checks follow
elsewhere. "should_use" means the path runs. No functional change.
Assisted-by: Claude Opus 5.5
* sycl : reorder Q3_K MoE experts and load them as vectors in the grouped GEMM
Assisted-by: Claude Opus 5.5
---------
Co-authored-by: cwriter <cwriter@localhost>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* vulkan: tiled transpose for dim-0 concat with a transposed source
The delta-net conv-state path does ggml_transpose() straight into a dim-0
ggml_concat(), so the generic concat kernel reads src1 with the transposed
stride, one element per thread, and neighbouring lanes land on different cache
lines. Stage a 32x32 tile through shared memory instead, which keeps both the
load and the store coalesced. 33 columns pads away the bank conflicts.
Gated narrowly: dim 0, 2D, unquantized 4-byte elements, contiguous src0/dst
rows, and src1 actually transposed. Everything else falls through to the
existing kernel.
Qwen3.6-35B-A3B on gfx1151/RADV, 30 gated-delta-net layers: the CONCAT op goes
11877 us -> 957 us, and pp2048 at ub2048 goes 1224 -> 1628 t/s. Generated text
is unchanged.
The approach follows the tiled transposed concat Gaetan Puleo wrote for the HIP
backend in his Strix Halo fork.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
* enable by default, add smaller shader sizes
* clean up comments
* improve shader
* Add non-contiguous concat tests
Co-authored-by: ravel7524 <58877666+ravel7524@users.noreply.github.com>
* handle x or y axis overflow
---------
Co-authored-by: Nathan Wilson <nath.flagman151@passmail.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: ravel7524 <58877666+ravel7524@users.noreply.github.com>
* ggml-rpc: fail gracefully on unreachable endpoints and failed handshakes
rpc_dispatcher::start() aborted the whole process (GGML_ABORT -> SIGABRT) when
the endpoint could not be parsed, the transport failed to initialise, the socket
could not connect, or the HELLO handshake failed. A typo in --rpc or a
temporarily down node killed the host application, and the reason went to a log
stream the host may never see.
Make the connect-time failure reportable instead of fatal: start() returns bool
(LOG_ERROR + false; the dispatcher thread only starts after a successful
handshake), negotiate_hello() returns false on a half-open socket,
get_dispatcher() returns nullptr and does not cache failed attempts, and every
consumer treats nullptr as "endpoint did not answer" with the conservative
behaviour of its domain (allocation -> nullptr, alloc-size -> lower bound,
device memory -> 0/0, events -> no-op, device count -> 0, and add_server()
already maps 0 devices to nullptr).
Mid-stream failures keep the existing RPC_STATUS_ASSERT by design (fail fast
against silent corruption) and are intentionally out of scope.
Field motivation: on Android (llama.rn + phone RPC cluster) an unreachable
worker crashed the app; the crash buffer showed
abort <- ggml_abort <- rpc_dispatcher::start <- ggml_backend_rpc_add_server.
* tests : add a regression test for unreachable rpc endpoints
Calls ggml_backend_rpc_add_server with a socket that nothing listens on and with a
malformed endpoint, and expects nullptr from both. Without the fix in this branch
the first call aborts the process, so the test cannot even fail cleanly.
* ggml-rpc remove low-value comments
* tests : drop the internal device-count check after the rpc registry refactor
---------
Co-authored-by: figkcoci-debug <figkcoci-debug@users.noreply.github.com>
With head sizes above 128 and KV >= 1024, scalar flash attention uses
16 rows per workgroup. On the Qualcomm Adreno driver this variant
spills registers to scratch memory (264 bytes per invocation on the
X1-85, none at 8 rows) and runs about 40x slower. The largest
test-backend-ops case then becomes a single dispatch of several
seconds, past the Windows GPU timeout, and the test never completes on
the X1 Elite.
Use at most 8 rows per workgroup on Qualcomm when hsk or hsv > 128.
Smaller heads keep 16 rows, which is faster for them on the X2 Elite.
Adreno X1-85, FLASH_ATTN_EXT hsk=256 kv=4096: nb=64 711 ms -> 17 ms,
nb=512 no result -> 125 ms; test-backend-ops completes. Adreno X2-90:
Gemma 4 E4B pp512 at depth 4096 with flash attention 81-90 -> 117 t/s.
On a device without mmap support, --lazy-mode auto resolved to off even
with --load-mode mmap, so llama_prefetch_rows never ran and the rows of
the large tensors faulted in one page at a time. Resolve it to off only
when mmap is off; the default iGPU load path is unchanged.
Fixes#31943
* server: add oai-compat /v1/audio/speech endpoint
* add e2e pytest
* server: share the s16 rounding, type speech form fields
The WAV writer of the audio helper truncates while pcm_f32_to_s16
rounds, so wav and pcm differ by one LSB. Both now round.
pcm_f32_to_s16 takes a pointer and a count and converts the TTS
bytes without an intermediate copy.
/tts and /v1/audio/speech parse their body with parse_tts_body, so
multipart fields such as speed get their type before the speech
conversion checks them.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* mtmd: support gemma 4 input audio stream
* use 480ms step
* server: enable realtime embeddings only for audio generation
The realtime slot forced embeddings on, which marks every row of a
batch as an output. Only a model that generates audio reserves
n_batch outputs, so a chat model with an audio stream (Gemma 4)
aborted on its first batch of audio rows. The hidden state is only
read by the audio generation pipeline.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
A cold dev server reloads the page once it has optimized its
dependencies, and that reload can land right after the harness opened
the phone sidebar, closing it again before Settings is clicked. Repeat
the pass over the menu keys up to three times while Settings is not
reachable.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: reveal the mobile sidebar by sliding the page aside
On a phone the sidebar now stays put under the page, and the page slides
right to uncover it, as a rounded card with the card border. The menu
button rides on the page; a swipe from the left edge opens the sidebar,
and a swipe back or a tap on the page closes it.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: polish the mobile sidebar reveal and add a hold menu to rows
Holding a conversation row on a phone opens its action menu anchored to
the row, in place of the hidden three-dot trigger. The page no longer
widens the document when it slides aside, so iOS stops panning it on the
first open.
The page and sidebar backgrounds start in the page color and blend into
the sidebar tint as the sidebar opens, following a drag frame by frame,
and a theme-color tag follows them for the browser bars. The sidebar's
contents fade in on a log curve and scale up slightly, and sit at the
same inset as the menu toggle, which now uses a two-line icon in a
filled circle.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: pin the context gauge on phones and lock the slid-aside page
On a phone the context gauge leaves the chat form for the top right of
the chat screen, in a round button across from the menu toggle. Both
take the scroll-down button's frosted look.
While the sidebar is out, and until the page is fully back in place,
the page neither scrolls nor takes taps; a tap or drag on it still
closes the sidebar.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: open the mobile sidebar with a swipe from anywhere
A rightward swipe from anywhere on the page now opens the phone sidebar
too, beside the edge swipe that competes with the browser's own back and
forward gesture. It starts after a short sideways move, skips text
fields and content that scrolls sideways, and keeps the page from
scrolling under it. Letting go opens or closes the sidebar past 30% of
the way or on a flick.
The top buttons and the chat form now hold their place on screen while
the page is slid aside, instead of jumping to the far ends of a
scrolled conversation.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: hide the model org name on mobile unless a setting keeps it
A phone's model selector is narrow, so it now drops the organization
from model names by default. A new Display setting, "Show organization
name on mobile", brings it back; it applies only while the organization
name is shown at all.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: close the mobile sidebar before following a tap in it
On a phone, picking a conversation, New chat, Search, MCP or Settings
in the open sidebar now slides it shut first and acts 180ms later, so
the page or dialog appears once the sidebar is out of the way. With the
sidebar already closed, or on a desktop, the action runs at once.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: make the app feel native on phones and run it under the status bar
Applies the mobile-native checklist: no tap highlight or landscape text
inflation, instant taps on buttons and links, no label selection or
callout on a held control, hover styles only where a pointer can hover,
and 16px textareas on phones so focusing them does not zoom the page.
The page now runs under the notch (viewport-fit=cover), so the installed
iOS app tints its status bar from the page and the sidebar. The menu
toggle, the context gauge, the sidebar and the top of the chat and
models pages pad themselves clear of it, and the locked, slid page
reaches down by the status bar's height, which iOS leaves out of it.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: add a viewport height readout for debugging on phones
With VITE_DEBUG set, a dev build shows the window, viewport unit and
safe-area heights, and the elements that reach past the window's bottom
and make the page scroll. It is loaded on demand, so production builds
leave it out.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fill the screen from launch in the installed iOS app
iOS 26 gives an installed app with viewport-fit=cover and a translucent
status bar a viewport short of that bar (WebKit bug 301108), clipping
the bottom of the page and lifting the chat form, but only while the
page is shorter than the screen. In the installed app the page now
reaches the screen plus the status bar, capped at the screen, so it
gets the full screen from the start without ever scrolling.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: raise the phone's corner controls in the installed app
The installed app's status bar already sets the menu toggle, the context
gauge and the open sidebar's logo apart from the screen's edge, so they
sit just below it rather than a full 1rem lower. They share one value,
which stays at 1rem below the inset in a browser tab.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: add safe-area utilities for what is pinned to a screen edge
The page runs under the status bar and over the home indicator, which an
installed iPhone app always shows. New pt-safe, pb-safe, top-safe and
bottom-safe utilities, each with a spacing step (pt-safe-12 is 3rem plus
the top inset), keep pinned content clear of them instead of hand-written
calc()s. A --safe-area-bottom joins --safe-area-top, and the chat screen and
Models page use pt-safe-12.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: open search in a full-height drawer on phones
A phone used to leave for the /search page. It now opens the same search
as the desktop dialog, in a bottom drawer that fills the screen below the
status bar. Keyboard shortcuts, the Enter hint and the keyboard highlight
stay on desktop, and so does the models manager, as in the phone sidebar.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: show every dialog full screen on phones, clear of the status bar
A phone now shows every dialog full screen, whatever size it asks for on a
desktop, instead of each full-screen dialog repeating its own classes. Its
top and bottom borders take the status bar's and the home indicator's
height, transparent, so an installed app's dialogs keep their content clear
of both with no change to their padding. The attachment preview, which has
its own full-screen content, pads itself and lowers its close button the
same way.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fade the page out under the status bar of the installed app on phones
Assisted-by: Claude (Opus 5.5 Medium)
* ui: show search in the full-screen dialog on phones instead of a drawer
Assisted-by: Claude (Opus 5.5 Medium)
* ui: start a phone page's title a row below the corner controls
Assisted-by: Claude (Opus 5.5 Medium)
* ui: pin the slid page's bottom bars against the screen in the installed iOS app
Assisted-by: Claude (Opus 5.5 Medium)
* ui: cast a soft shadow from the slid page onto the phone sidebar
Assisted-by: Claude (Opus 5.5 Medium)
* ui: hide the scroll-down key at the real bottom on iOS and show it only well above
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fade the status-bar fade out as the page slides aside for the sidebar
Assisted-by: Claude (Opus 5.5 Medium)
* ui: draw the slid page's shadow above its rounded frame, so it follows the corners
Assisted-by: Claude (Opus 5.5 Medium)
* ui: show a phone chat's title and working directory at the top, and pick it in a drawer
Assisted-by: Claude (Opus 5.5 Medium)
* ui: give iOS Safari a solid top and fade the bottom of the screen on phones
Assisted-by: Claude (Opus 5.5 Medium)
* ui: let iOS Safari find the top strip, which it skips when it takes no taps
Assisted-by: Claude (Opus 5.5 Medium)
* ui: lengthen the top fade and lift the chat form a little in phone browsers
Assisted-by: Claude (Opus 5.5 Medium)
* ui: give iOS Safari's bottom toolbar the page's colour, not the chat form's
Assisted-by: Claude (Opus 5.5 Medium)
* ui: draw the phone's context gauge key like the menu key across from it
Assisted-by: Claude (Opus 5.5 Medium)
* ui: confirm opening the phone sidebar, its hold menu and its rows with a haptic tap
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the screen-edge fades on the page as it slides aside for the sidebar
Fading them out with the slide made the page flicker; the rounded frame over the
slid page already clips them at its corners.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the working directory chip optional for hosts with their own trigger
The phone layout opens the picker from the chat's header, so ChatForm hides the
chip there. Master's version of the chip has no such switch; bring it back.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: style layer 5's additions with master's tokens until the design system lands
Assisted-by: Claude (Opus 5.5 Medium)
* ui: drop the design system's stylesheets and fonts from app.css until it lands
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep master's font stacks until the design system brings its fonts
Assisted-by: Claude (Opus 5.5 Medium)
* ui: restore master's colour palette in app.css until the design system lands
Resolving the safe-area commit took the design system's app.css, which drops
the palette for its token files and aliases the shadcn names to semantic
tokens. With the tokens gone those aliases point at themselves, so every
colour went invalid: transparent surfaces and currentColor borders.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: center alert dialogs on phones instead of sliding them up as drawers
An alert is a short question, not a drawer: on a phone it is now a centered
card, 90% of the screen wide and as tall as its text, scrolling only past the
screen's height. It zooms in as on desktop, and its footer drops the margin
the sheet needed.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: space the working directory drawer for a phone
The drawer reused the popover's compact picker: Browse sat against the search
field, the divider ran past the rows and each row had its own inset. In the
drawer the picker now takes the header's inset, gives Browse room, lines its
rows up with the search results and sizes its text for a phone.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: tap the haptic engine when a message sends, and once per hold
A phone now confirms a sent message by feel, at the moment the send goes
through (not on every press of the button), and before any await, so it still
runs inside the tap or keypress that iOS needs to play its tick. The hold that
opens a row's menu now taps the engine from the long-press hook itself, once:
the hook vibrated and the row tapped as well, which buzzed twice on Android.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: name the gauge dial sizes, the chat form's bottom offset and the sidebar fade
Review follow-ups: the context gauge dial reads its grid, radius, stroke and
track from one table per size instead of ternaries; the chat form's distance
from the screen's bottom per platform and the sidebar content's fade duration
are named constants. The viewport debug readout moves out of src, next to the
tests, as it is a debug tool the app only loads in a dev build.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: say why the microphone can't record instead of failing silently
A llama-server opened over plain HTTP on the local network or over Tailscale
is not a secure context, where browsers hide the microphone, so a tap on the
mic did nothing but log to the console. The button now shows a crossed-out mic
in the warning tone with the reason in its tooltip, and a tap says it in a
toast, as a phone shows no tooltip. A start that fails names the browser's
reason (access denied, no microphone, in use by another app) the same way,
and the next tap tries again. The recorder passes the browser's error on
instead of replacing it with a generic one.
Assisted-by: Claude (Opus 5.5 Medium)
* model : allow loading external PLE files
Assisted-by: Codex
* convert : add separate PLE exports
Assisted-by: Codex
* tests : fix PLE export test lint
Assisted-by: Codex
* model : make --ple an alias of --shard, allow multiple shard locations
Reuse the multi-shard loading path for external PLE files instead of a
separate option.
* --ple is an alias of --shard: both take a comma separated list and may
be given more than once, and the shards may live anywhere
* metadata comes from the main file; numbered shards are still
auto-discovered and strictly validated
* explicit files outside the discovered set are appended as extra shards,
exempt from split.no and split.count checks
* -isnm/--ignore-shard-num-mismatch loads only the given files, in order,
with no auto-discovery and no split metadata validation
* ple_path, path_ple and llama_model_loader::add_ple_file are removed;
common_model_load() uses llama_model_load_from_splits for two or more
paths
* common/fit.{h,cpp} take the shard path vector, so --fit and the router
memory measurement load every shard
Assisted-by: GPT-6-Astra
* rm -isnm
* make sure to dedup paths
* rm test_convert_ple
* address comments
* allow ple shard to override tensor
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* Windows unbuffered io
* Fix read len calculation
* Add test and model load timing
* chunked fallback read + removing filepointer semantic divergence between linux and windows
* tests : remove unused mmap test arguments
Assisted-by: Codex
* remove timing, set file pointer after chunk read, test refinement
* Update tests/test-llama-mmap.cpp
Co-authored-by: Oliver Simons <osimons@nvidia.com>
* Update src/llama-mmap.cpp
Co-authored-by: Oliver Simons <osimons@nvidia.com>
* test refinement
Assisted-by: Codex
* remove duplicate constexpr
Assisted-by: Codex
* adding label behavior to tests
Assisted-by: Codex
---------
Co-authored-by: Oliver Simons <osimons@nvidia.com>
* server : functionality test to make sure every time when we add new spec dec method the existing ones are not broken.
* server: use compact models for speculative decoding CI
* ci: avoid GHA cache for speculative models
* ci: run speculative tests on DGX Spark
* server: address speculative test review
* server: use consolidated speculative models
* server: refine speculative test calibration
* docs: record speculative Spark validation
* ci : run the server-speculative steps with bash
* server: recalibrate speculative test floors on GB10
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* server : accept top-level schema in response_format with type json_schema
The OpenAI-compatible chat completions endpoint only applied the schema
when it was nested inside response_format.json_schema.schema. When the
schema was provided directly as response_format.schema (as documented in
the server README, and used by some clients), it was silently ignored
and generation ran unconstrained.
Also mention json_schema in the error message for invalid response_format
types, document the OpenAI wrapper form in the README, and add a test
case for the top-level schema form.
Fixes#10732
* server : let the json_schema wrapper take precedence by presence
Pick the schema source based on whether response_format has a
json_schema key, instead of falling back when the parsed wrapper schema
is empty. An empty schema {} is valid (any JSON value), so the old check
let a top-level schema override an explicit {} in the wrapper. Use
json_value() for the top-level schema too, like the rest of the
function, and add tests for the precedence.
---------
Co-authored-by: linsen <251731047+linsen458-spec@users.noreply.github.com>
* ggml: implement the backward pass for TRI
GGML_OP_TRI had no case in the backward switch, so a trainable graph containing it aborted with "unsupported ggml op for backward pass: TRI". The op is linear, so its gradient is masked the same way. The incoming gradient is not necessarily contiguous (the backward of a transpose hands over a view) and ggml_tri requires contiguity, so it is made contiguous first.
test_tri gains a post_transpose variant (out = transpose(tri(a))) that aborts as above without this change.
* ggml : simplify the TRI backward pass
Pass the gradient directly to ggml_tri and drop the transpose test. The existing TRI cases cover the backward pass for contiguous F32 gradients.
The dummy-model generator fills every tensor from a plain normal pool.
`blk.N.ffn.act.scales` is a divisor in the mpt graph, so with seed
1224489258 it draws |v| ~ 1e-6, and `ffn_act = gelu(...) / scales` reaches
2.7e5. Metal stages both operands of `kernel_mul_mm_f32_f32` as half, so
those activations overflow to inf, the norm downstream turns the inf into
nan and the logits come out nan, while the CPU reference stays in f32 and
stays finite. Both NMSE values then become nan and the check fails.
Give `ffn.act` the same treatment as `rope_factors` and keep its magnitude
away from zero.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 32GB, DDR5 128GB)
* ui : fold the sidebar's conversation sections
Pinned and Recent conversations fold under their heading. The whole heading
row is the toggle; its chevron follows the title, shows on hover and stays
while the section is folded. Folded sections are remembered across reloads.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : show a sidebar section's first six chats
Pinned and Recent list their first six chats and end with a Show more
row that opens the whole section; Show less, sticking to the bottom of
the sidebar while the section scrolls under it, closes it back. The
expanded sections last for the session.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : rearrange the sidebar's sections and chats by dragging
Pinned and Recent can be dragged by their title. The real section follows
the pointer, tinted like the sidebar's resize handle and edged with a
faint border, the other sliding aside as it passes; let go, it glides from
there into its new slot. The order is remembered across reloads, and a
section never moved keeps its default place. All sections now scroll
together.
A chat dragged onto Pinned is pinned, onto Recent unpinned. A lifted copy
of the row follows the pointer while the row slides out of its section.
The section under the pointer takes the same tint and border; a section
ignores its own chats. Let go, the row slides in where its section sorts
it, the rows below making room, and the copy glides into it; anywhere
else, it slides back into its own section.
While dragging, the rows and headings passed over take no hover. Dragging
needs a mouse or a pen, so a finger still scrolls the sidebar; the sidebar
scrolls near its edges, and Escape cancels.
Assisted-by: Claude (Opus 5.5 Medium)
* ui : set the sidebar's sections apart from their rows
A section heading now reads above its chats: larger, semibold and in the
full foreground colour. Chat titles drop to a regular weight in a slightly
softened colour, full again when the chat is open or hovered, and "Show
more" drops to a regular weight in the muted colour. Chats inside a
section sit closer together, sections further apart, and the sidebar's
actions get more room above and below.
Assisted-by: Claude (Opus 5.5 Medium)
* fix(rocm,ci): Update ci workflow to use devlab ci runners for gfx1151
* ci : move gpu-rocm job to Strix Halo devlab runners
Assisted-by: Claude
* ci : ROCm 10 wheels with per-GPU device packages for gpu-rocm; add gfx1201
The runners' system ROCm 7.2.1 cannot compile the current HIP code for
gfx1151 ("Invalid dpp_ctrl value" in mmf-instance-ncols_*.cu). Go back
to the TheRock ROCm 10 wheels with the device-<gfx> extra, so rocBLAS
has the target's kernels, and keep ci/run.sh's results and models under
$RUNNER_TEMP instead of the runners' RAM-backed home. Add a gfx1201
(Radeon AI PRO R9700) row with HIP graphs disabled.
Assisted-by: Claude
---------
Co-authored-by: David Hekimian <dhekimian@users.noreply.github.com>
* metal: make ggml_tri select instead of scaling by the comparison
The kernel multiplied by the 0/1 comparison, which turns a +-inf outside the triangle into NaN, while the CPU kernel writes a plain zero there. Select, like the CPU does.
Adds test_tri_special, which puts +-inf into the cells that have to be discarded, for all four triangle types. It fails against the previous kernel and passes with this change.
* Just trim down the repetitive comments.
* mtmd, server : add temperature to /tts, default Pocket TTS to 0.3
Assisted-by: Claude Opus 5.5
* feat: add temp parameter in llama-tts
* temp: fix default hardcoded values
* fix: remove too narrow test
* fix: addressing PR comments
The rotation angle of a position p is rounded in f32 to about
p*FLT_EPSILON radians, and the CPU reference and the GPU kernels reach
it through different rounding sequences (cumulative product versus
pow). With positions up to 8192 this puts the NMSE floor between
backends near 8e-8, so the default 1e-7 threshold fails at random on
Metal. Derive the tolerance from n_ctx instead: it stays at 1e-7 for
n_ctx=512 and becomes 9.5e-7 for the YaRN cases at n_ctx=8192.
* ggml : fix for the radix TOP_K thread timing bug
* make the key an ordered 64-bit pair
* a rewrite without shift-skip logic
* fix casts
* warp-contiguous layout
* reduce the scope removing vulkan+metal changes
* turn it off for Metal/ROCm, note about cub::DeviceTopK
* fix the hip build
* skip Vulkan as well
* popc lambda from fattn
* do block_cols math in int64_t
* make it allowlist bc SYCL looks non deterministic too
* ggml : make Metal radix Top-K tied selections deterministic
The radix top-k kernel gathered the selected indices with atomic
counters, so the order of the tied candidates at the threshold (and,
with more ties than remaining slots, the subset itself) depended on
the thread scheduling.
Replace the two atomic compaction loops with a deterministic scheme:
count the elements above and tied with the threshold in contiguous
per-thread chunks, rank them with a simdgroup scan and a serial scan
over the (now unused) radix histogram, and write each kept element
exactly once in column order. The output is now byte-identical
across runs.
This is also faster than the atomic version: the strided atomic
loops become contiguous passes, reducing top-k times by ~20-45% on
M5 Ultra, with tie-heavy shapes within a few percent.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* tests : enable deterministic top-k check on Metal, add radix-path tie cases
Metal now produces deterministic tied selections for the radix top-k
path, so the determinism self-check runs for the MTL backend as well.
Add eval cases covering the radix path (ncols > 2048) with more ties
at the threshold than remaining slots, and tie-heavy perf cases to
monitor the cost of the deterministic tie handling.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* ggml : use coalesced accesses in the Metal radix Top-K compaction
The deterministic compaction walked contiguous per-thread chunks with
scalar sequential loads, which the memory system of the M2 Ultra can
not overlap well: the radix Top-K times regressed by ~1.5-2.4x on the
large-ncols shapes.
Give each warp a contiguous chunk of the row instead and let the lanes
interleave within it with stride 32, so every access is fully
coalesced. The lanes write their kept elements at lane-major ranks,
which is deterministic, and only a tie overflow falls back to a
per-step cross-lane ranking to keep the lowest tied indices in column
order.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* tests : remove the extra top-k tie and perf cases
The radix-path eval and tie-heavy perf cases were added to validate and
benchmark the Metal deterministic tie handling. Keep only the enabling
of the determinism self-check for the MTL backend.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* metal: scan the radix bins and warp totals with a simdgroup
The radix top-k selected the threshold bucket with a serial scan of
the 256 bins by thread 0 in each of the 4 passes, and ranked the warps
of the compaction with two serial scans by thread 0. The first
simdgroup now does both: each lane sums 8 bins, an exclusive prefix
sum and a ballot find the lane that reaches the k-th value, and the
warp totals get their exclusive bases from simd_prefix_exclusive_sum.
* tests: enable the deterministic top-k check on ROCm
HIP has no cub::DeviceTopK, so every row goes through the radix or
bitonic path, which break ties in the same way in each run. The check
covers all tie cases on ROCm, including the 1 and 2 row ones that CUDA
skips.
* cuda: remove the redundant radix top-k counter reset
Address review from ORippler. top_k_radix_init already zeroes
greater_count for each chunk of rows and top_k_radix_select keeps it,
so the reset kernel before the gather is one launch with no effect.
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Pascal <admin@serveurperso.com>
Zamba2 and similar architectures use a state size of 64 in their
recurrent-scan step, which fell outside the GPU kernel's existing
fast paths (96, 128, 256). Every scan step was falling back to CPU,
bottlenecking generation to 6-7 tok/s against the 50-70 tok/s band
other architectures achieve on the same hardware.
Add a 4th template instantiation for d_state=64, following the
existing pattern; no change to the underlying math. Verified against
the existing CPU-reference correctness harness with 3 new test cases
built from Zamba2's real hyperparameters (28/28 passing).
Measured ~35% throughput improvement (11 -> 15.1 tok/s) on Zamba2-7B-
Instruct. Full parity with other architectures not yet reached - a
separate, unrelated fraction of scan calls (~28%) still falls back to
CPU for a cause not yet diagnosed.
Assisted-by: Claude Sonnet 5
Co-authored-by: Sargentmajord <331229907+sargentmajord@users.noreply.github.com>
* ggml : switch to F32 src1 in the CPU backend
Change the default computation for GGML_TYPE_F16 to use F32 for src1
data in matrix multiplications: vec_dot_type(F16) is now F32 and the dot
product goes through the new ggml_vec_dot_f16_f32 (x = F16 src0, y = F32
src1), so src1 is no longer quantized to F16 and worst-case inputs that
only fit in the F32 FP range are handled correctly.
- conv_2d/conv_3d: store im2col patches (the inner mul_mat src1) as F32
and pass separate src0/src1 types to the inner mul_mat; the previous
F16 patches made the inner mul_mat widen src1 into the conv work
buffer, corrupting the patches and crashing on large inputs
- conv_transpose_1d/2d (F16 kernel): keep src1 in F32 in the work buffer
and use ggml_vec_dot_f16_f32; update work buffer planning accordingly
- supports_op: restore MUL_MAT support for F16 src1, which is now
widened to F32
- ggml_vec_dot_f16_f32: add GGML_SIMD support by converting chunks of x
to F32 and reusing the F32 SIMD dot
- mul_mat/mul_mat_id: add TODOs for a more efficient F16 x F16 path that
can be used if ggml_prec_get_src(dst, 1) allows converting src1 to F16
MUL_MAT f16 x f32 perf on M5 Ultra (GFLOPS):
- m=2048, n=512, k=16 : 120.9 -> 101.7
- m=2048, n=512, k=64 : 150.5 -> 344.2
- m=4096, k=14336 GEMV: 63 -> 343-415
- m=4096, n=512, k=14336: 65.6 -> 770.2
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* cont : add some F16 x F32 vectorized paths
* cont : enable the fused F16 x F32 dot on AVX without FMA
The 256-bit fused path of ggml_vec_dot_f16_f32 required AVX2 and FMA,
so CPUs with AVX and F16C only fell back to the chunked conversion and
F16 decode was 4x slower than before. Enable it on __AVX__ && __F16C__
and accumulate with GGML_F32x8_FMA, which does mul + add without FMA.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* llama, mtmd: follow the reference precision of ggml_prec_set_acc calls
Requests that the reference implementation does not keep in F32 move
to GGML_PREC_BF16: the GLM4 family down and output projections, the
MiMoVL vision matmuls and the Qwen3-TTS speaker encoder conv1d. The
DeepSeek V4/V4.1 routers and the MiniMax-M3 indexer stay in F32 with
a link to the reference line that computes them in F32.
* ggml: handle GGML_PREC_BF16 the same way as GGML_PREC_F32 in backends
Backends that reject GGML_PREC_F32 on the src1 of MUL_MAT_ID now also
reject GGML_PREC_BF16 (CUDA, Hexagon, OpenCL, WebGPU, Vulkan), and
backends that accept GGML_PREC_F32 for flash attention now also accept
GGML_PREC_BF16 with the same F32 accumulation (CPU, SpaceMIT, ET).
Vulkan keeps the F32 transfer between devices for both. The CUDA
cuBLAS path already handles GGML_PREC_BF16 and is unchanged.
* tests: cover GGML_PREC_BF16 in the flash attention tests
The flash attention cases run with GGML_PREC_BF16 next to F32 and
DEFAULT, and var_to_str prints bf16 so these cases keep their own name.
* llama: request GGML_PREC_BF16 for the KQ multiplication
* cont : add TODO in clip.cpp
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* opencl: add kernel_gemv_noshuffle_q8_0_f32_mc
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: small-n MoE GEMV for q5_1, q4_k
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: adreno kernel k floor 256 for x2
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: small-n MoE GEMV for more quants
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: rename function
* opencl: disable bin FA for small verify batch
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* hexagon: use absolute skel path
* hexagon: support non-ascii dir and long path in skel-abs path
* hex-libdir: add support for FASTRPC_MANAGE_DSP_LIBRARY_PATH and GGML_HEXAGON_LIBHTP_DIR on all platforms
* hex-libdir: keep extra fwd slash in abs path on Android and Linux
* hex-libpath: always start with absolute path then fallback to session ctl then env
* hex-libdir: add CDSP1,2,3 to the fallback, needed for the IQ10
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* hex-nhvx: apply nhvx updates to the latest master
* hex-nhvx: set N_HVX and N_THREADS based on the arch to avoid overallocation
* hex-build: add support for enabling HTP_DEBUG via build.py
* hex-trace: move tracing state/helpers into new htp-trace.h header
* hex-dcache: cleanup dcache #defines
* hex-hnvx: minimize the diff by reflowing the macros
* hex-ci: fix python type/linter warnings
* hex-nhvx: consitent logging
* mtmd: support MageVL-Cowork model
* mtmd: keep magevl rope_freqs at model level and check its shape
The rope factors are shared by every layer, so they live in clip_model
instead of a pointer copied into each clip_layer. The loader now checks
that v.rope_freqs is F32 with d_head elements and that d_head splits
into the 4:6:6 pairs the graph expects.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* moe-cache : compute part of the misses on the CPU
--moe-cache-cpu-share F computes a share F of the MoE cache misses on the CPU
from the host experts instead of uploading them. The graph builds the expert
FFN twice: on the device from the cache slots, where the misses left to the CPU
read zero slots, and on the CPU from the host experts, where the experts
computed on the device have id -1. The two results are added before the down
bias.
- ggml-cpu: MUL_MAT_ID skips negative ids and writes zero rows
- sched: the copy callback is called for all the inputs of a split, and before
the host sync when there are no events
- the slot lookup callback reads the inputs of the CPU part together with the
selected experts, so the CPU part does not wait for the device, and its
result is copied back asynchronously
- each token uses distinct zero slots, mul_mat_id does not support an expert
selected twice by a token
- LLAMA_MOE_CACHE_TRACE writes the selected experts of each lookup, for a
simulator of the cache
* scripts : add moe-cache-tune.py to pick the MoE cache arguments
The script wraps a llama-server command: it starts the server at two cache
sizes and three CPU shares, fits the upload and CPU speeds of the machine, and
replays the routing of the runs through a copy of the cache logic to predict
other sizes and shares. It prints the recommended --moe-cache-mib and
--moe-cache-cpu-share.
docs/moe-cache.md describes the MoE cache arguments and the script.
* add am17an as codeowner for MoE cache
* moe-cache : upload all the misses until the cache is full
With a high --moe-cache-cpu-share, few misses are uploaded, so an empty cache
fills slowly and more misses than needed go to the CPU. The CPU share now
starts once all the slots of a cache group hold an expert. moe-cache-tune.py
replays the same rule.
* add backend sync for async transfer queue
* moe-cache : let the scheduler copy the results of the CPU part
* Fix MSVC compilation
* remove sync change
* llama-graph : extract the CPU share condition into can_cpu_share
* llama-graph : make llm_graph_params::moe_cache const
The graph construction no longer updates the MoE cache directly: the CPU
parts of the layers are accumulated in the graph result and applied to
the cache in llama_context::process_ubatch, after the graph inputs are
set and before the graph is computed.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* scripts : replay the ClassLRU policy in moe-cache-tune.py
* moe-cache : pick CPU misses by CRF instead of random (#31853)
* moe-cache : remove the backend logic of the CPU part from the graph
The CPU part reads the host experts, so the scheduler places it on the
CPU without set_tensor_backend. The cache gives no CPU map for batches
that a device would offload. Graphs with a CPU part look up their own
slot map, so the graph does not register the part with the cache.
ggml-backend : the inputs of the next split that are already computed
are copied to host memory before the current split runs, so a CPU split
does not wait for the device split before it. This replaces the stash
of the cache, and the copy callback is again called for host weights
only.
* llama-graph : return the MoE cache slots and CPU ids in a struct
The scheduler prefetch no longer checks that the next split runs on
another backend: the inputs of a split are never on its own backend,
so the input loop already skips that case.
---------
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Alexey Kopytko <alexey@kopytko.com>
The kernel_group_norm_f32 kernel and its host dispatch ignored the
batch dimension (ne03) of the source tensor: only ngrp threadgroups
were dispatched and each normalized the flat range
[tgpig*gs, (tgpig+1)*gs) with gs = ne00*ne01*ceil(ne02/ngrp). For
tensors with ne03 > 1, all but the first ne00*ne01*ne02 elements were
left unwritten (exposed by the new GROUP_NORM test with
ne=[3,1,1,65537], num_groups=1).
Dispatch one threadgroup per (group, i03) pair (ngrp in x, ne03 in z,
same work distribution as the Vulkan backend) and compute per-batch
statistics, matching the CPU reference. Also clamp the group element
count for the last group when ne02 is not divisible by ngrp and use
int64_t indexing.
The kernel now only requires ggml_is_contiguous_rows(op->src[0])
instead of full contiguity: it addresses src0 and dst with their own
strides (nb03 and the dst strides are added to the kernel args - dst
can be a strided view of src0 with ggml_group_norm_inplace) and falls
back from the flat fast path to a strided path that decodes the linear
element index when rows or channels are padded. The flat/strided
selection is done on the host in
ggml_metal_library_get_pipeline_group_norm and compiled into one of
two specialized pipeline variants via a FC_group_norm_flat function
constant, so the dead path is eliminated at pipeline creation.
GROUP_NORM support is also gated on F32 src and dst in supports_op -
the kernel only implements f32 and an f16 input would silently be read
as f32.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* server: visit parent slots before child slots
A parent slot wakes its waiting child slots once it finishes the shared
prompt prefix, while slots are visited in id order. When a child task
lands on a lower slot id than its parent, the child is visited before
the parent wakes it and waits an extra batch to process its suffix.
Non-child slots are now visited before active child slots, so every
child woken by its parent joins the same batch.
* server: check systemone test events through test tags
The systemone tests matched free-form debug messages, including the
formatted f_keep value of the prompt cache. The server now emits
__TEST_TAG_WAIT_PARENT_SLOT__ when a child slot waits for its parent,
__TEST_TAG_PROMPT_CACHE_RESTORE__ on every prompt cache restore and
__TEST_TAG_PROMPT_CACHE_FULL_RESTORE__ when the whole cached prompt is
a prefix of the new one, compared on token counts.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* Only run webgpu CI on my fork
* Add webgpu only workflow
* Add WebGPU group_norm kernel
* clean up
* resolve merge conflicts
* parameterize data type; clean up cmake changes
* add nullptr check
* clean up webkit experiments
* more cleanup
* more cleanup
* clang format
* Remove unused feature check for WebGPU
* clean up supports_op condition and update documentation
* add perf tests for group_norm
* clean up doc with script and undo change in test-backend-ops
* Do 2d dispatch of workgroups
* Do 2d dispatch of workgroups
* Unstage unrelated changes
* undo docs churn
* add new test case
* undo small change
* undo small change
* ggml-backend-meta : minimum number of compute views per buffer
Assisted-by: Claude Opus 5.5
* tests : revert the deepseek41 test model to 6 layers
Assisted-by: Claude Opus 5.5
* hexagon: fuse post-norm residual add into RMS_NORM_MUL
* hex-fusion: refactor RMS_NORM fusion and add dedicated RMS_NORM_MUL_ADD opcode
Improves performance by keeping things in HVX registers.
Also simplifies overall fusion logic and logging.
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* model: add Kimi-K3.1-Pro text model
K3 with the full-attention MLA layers replaced by DSA: on this checkpoint 24
of the 93 layers carry a lightning indexer (32 heads x 128, top_k 2048, no
rope). The KDA (linear-attention) layers and the rest of the K3 skeleton -
cross-layer residuals, latent MoE, situ, MLA output gate, full-rank KDA gate -
are unchanged, so this reuses the K3 graph and converter and adds the indexer.
Indexer:
- the selection is an additive mask over the whole cache, not a gather of
the top-k rows, so it reuses the existing sparse flash attention path;
off that path the mask stays correct but still costs a kq over the cache
- the indexer key LayerNorm takes its eps from the model's rms_norm_eps,
because dsa_config declares no eps of its own
Graph and cache:
- new arch with four indexer tensors per DSA layer, and a hybrid memory
that keeps a K-only MLA cache for those layers, recurrent state for the
KDA layers, and the indexer key cache
- the graph node budget shares the K3 entry
Conversion:
- KimiK31Model(KimiK3Model), picked by KimiK31ForConditionalGeneration
- dsa_config becomes attention.indexer.{head_count,key_length,top_k,types}
plus the indexer LayerNorm eps; the layer list is 1-indexed like
linear_attn_config's, and the converter asserts that it covers exactly
the full-attention layers
- the routed experts stay MXFP4 and are repacked as in K3, not dequantized
- the vision tower and mm_projector are skipped: text only for now
- the arch key follows the checkpoint's own model_type ("kimi_k31") and the
loader registers KimiK31ForConditionalGeneration; the release name carries
a -Pro suffix that nothing in the config or the tensor names uses
Quantisation: the indexer's q_b, k and proj are kept out of the quant
overrides, since the top-k selection is sensitive to their precision.
Verified on an 8-layer slice against the HF checkpoint: 38 tensors checked
exactly, the indexer metadata against the config, and the selection against
full attention - identical PPL at top_k beyond the context length, divergent
once top_k is forced below it. Arch test fixture and the full arch suite pass.
* converter: avoid clutter in the k3 base
* narrow comment to match model
* update metal fusion expectations
* restore
* sync with PR 31733 refactoring
* hexagon: Generalization of CONV_2D implementation for CONV_3D
- New op HTP_OP_CONV_3D is wired through htp-ops.h, main.c (shares op_conv_2d) and the op remap in htp-kparams.cpp.
- conv_3d runs on the existing HMX tiled (implicit GEMM) conv2d path. KD and KH are folded into one virtual kernel row (KD*KH), and output rows are counted as N*OD*OH.
- htp_kparams_solve_conv3d checks the conv_3d layout and calls the tiled solver htp_kparams_conv2d_tiled, which takes an is3d flag.
- The tiled kernel takes the sizes and strides that differ between conv_2d and conv_3d from a struct htp_conv_geom. Only the input stage unfolds depth.
- 256-entry DMA queue can no longer overflow
Assisted-by: Claude Code
* hexagon: modify checking to eliminate divide-by-zero hazards
* hexagon: conv2d tiled: compute the conv_3d unfold divisors on the host
The KH and OH fastdiv values used to unfold conv_3d kh into (kd, kh) and
oh into (od, oh) were built per op on the device, each with a 64-bit
divide. Pass them in the kernel params instead; they fit in the space the
window variant already reserves, so the struct size does not change.
Assisted-by: Claude Code
* hexagon : preserve fold scheduling in conv2d for small size
hexagon : support arbitrary positive conv2d strides
Support positive horizontal strides for WHCN and CWHN inputs. Preserve
load/shuffle paths for strides 1/2 and use gather for larger strides.
hexagon : improve conv2d tile size selection
Prefer the largest feasible output-channel block and derive spatial
chunk capacity from VTCM usage. Compare capacity and thread-aligned
candidates using per-thread work and chunk overhead.
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
* hex-conv: fix compiler warning
---------
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* tests : qualify isnan and signbit with std::
The unqualified calls rely on the C math.h names also being present in the
global namespace, which is not the case for the NetBSD 11.99.9/x86_64 gcc
14.4.0 build reported in the issue. Prefix the calls with std:: to select
the standard C++ overloads declared by <cmath>.
ref: https://github.com/ggml-org/ggml/issues/1681
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* tests : qualify C math calls with std::
Follow-up to the isnan / signbit fix: make the remaining math calls in
test-quantize-fns.cpp consistent with the std:: isfinite / max / numeric_limits
already used in the same file.
The float variants are overloads in C++, not distinct names, so the mapping is
fabsf -> std::fabs, copysignf -> std::copysign, nextafterf -> std::nextafter -
not std::fabsf & friends, which do not exist in libstdc++ and would break the
gcc build. NAN and INFINITY are macros and are left unqualified.
Only behavioural difference: array_rmse() fed a double accumulator to sqrtf(),
now std::sqrt() keeps the computation in double. The result is more accurate
and differs from the previous value by at most 1.4e-7 relative, three orders of
magnitude below the tightest tolerance in this test.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
The self-hosted CPU runners are containers with 8 CPUs, but
std::thread::hardware_concurrency reports the cores of the host. The
table mode of test-llama-context sizes its pool from that count, so it
runs 4 models x 8 threads on an 8-CPU job and hangs.
Add TEST_ARG_HARDWARE_CONCURRENCY and TEST_ARG_PARALLEL_WORKERS as the
defaults for the auto-selected counts of test-llama-context and of
test-backend-ops, and set them to 8 and 4 for cpu-x64-high-perf.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 32GB, DDR5 128GB)
* moe-cache : LRFU/MQ-inspired hybrid cache
* make the clock advance by N tokens: use the real token count (ne[1] is always 1, ids have n_expert_used entries per token)
* sel is a RESHAPE node, the orig input is in src[0], so sel->src[0]->ne[1] is the count
* c for the class index, last_use for last use time
* explain the roles of these
* compute the key once
* use the key directly instead of re-reading
* show how the score is computed from the age
* remove the branch and clarify exp2 and log2 relationship
* reword the header comment and elaborate on unused slots and sorting
* not a copy, not just skip
* std::array maybe
* Revert "std::array maybe"
This reverts commit 7649b1f172.
* half life it is
* add key and iteration helpers
* extract decayed_crf
* bump classes to 8, document the formula, display the cache size in token worth count
* assert the computed number of classes matches the half life
* opencl: use the multi-column q6_K vocab head GEMV for 2..4 tokens on X2E
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: rename GGML_OPENCL_Q6K_MC3 to GGML_OPENCL_Q6_K_MC3
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: rename GGML_OPENCL_Q40/Q41/Q4K_MC3 to GGML_OPENCL_Q4_0/Q4_1/Q4_K_MC3
GGML_OPENCL_Q4K_MC3_DQ and GGML_OPENCL_Q4K_MC3_LDS become GGML_OPENCL_Q4_K_MC3_DQ
and GGML_OPENCL_Q4_K_MC3_LDS.
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: remove unnecessary debug probe
Co-authored-by: Shangqing Gu <shawngu@qti.qualcomm.com>
Assisted-by: Claude Opus 5.5
* opencl: clean up
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* hex-repack: move repacking/tiling into graph-optimize
* hex-tiling: simplify tiling rules and remove obsolete helpers
* hex-ops: use a hash for de-duplicating tensor descriptors
* hex-bufs: improve tensor hash
* hex-ops: accept empty (zero-elements) ops
* hex-bufs: enable hostbuf by default and reject tensors that do not fit into max_vmem
* hex-bufs: improve error message and update readme
* hex-bufs: make sure to sync flush opbatch in all sync set_tensor paths
* feat : added ricursiveai ripple support
* fix : mandatiry SWA, scale through helper, comments back
* fix : metal
* model : make ricursive RoPE/NoPE explicit and validate the layer pattern
Write attention.rope_pattern and read it as optional, falling back to the
sliding window pattern so existing GGUFs keep loading. Refuse a
sliding_window_pattern that does not cover every layer, a model with no
sliding layer, and a gate or value residual present in only some layers.
The converter derives layer_types when the config omits it, ignores a stale
hybrid_layer_pattern and drops tensors of features the config disables.
Assisted-by: Claude Opus 5.5
* model : make ricursive swiglu_clamp per-layer
Store swiglu_clamp as a per-layer array, read it with get_key_or_arr so the
scalar in existing GGUFs still applies to every layer, and refuse a negative
or non-finite value. Add a test-llama-archs fixture with mixed per-layer
clamps and no rope_pattern, so the saver roundtrip checks both.
Assisted-by: Claude Opus 5.5
* model : use the ricursive per-tensor scales and RoPE parameters
Pass the attn_gate and output weight scales to build_lora_mm, take
ext_factor and attn_factor from the graph instead of hardcoding them, and
copy rope_freq_scale_train to the SWA layers instead of pinning it to 1.
Cite the reference lines for the value residual and the NoPE layers.
Assisted-by: Claude Opus 5.5
* convert : drop the ricursive chat_template.jinja override
The precedence between tokenizer_config.json and chat_template.jinja is to
be changed in SpecialVocab for all models instead. Until then a new
conversion stores the short tokenizer_config.json template; GGUFs that
already carry the tool-aware template are not affected.
Assisted-by: Claude Opus 5.5
* fix : remove NVFP4 checkpoint quantization block
* match #31733 refactor
---------
Co-authored-by: Piotr Wilkin <ilintar@gmail.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* RDNA3/4 Vulkan Large-Tile Split-K Optimization
* Add single-operator performance test cases
* Update the policy to use driver-reported LDS usage
* Add comments
* Use the simpler split_k rule from review
Assisted-by: GitHub Copilot
* tests : drop the split_k test cases
* vulkan : simplify the RDNA3/4 split_k condition
* metal : support scale-bias routing norm in the TOPK_MOE fusion
The DeepSeek-V4.1 graph normalizes the routing weights with a scale+bias
(sum + 1e-20) instead of the lower clamp used by the other models. Add
norm_bias/norm_scale_bias fusion chains that accept a SCALE node with a
unit scale factor at the clamp position, pass an additive bias in the
kernel args, and apply it as max(wt_sum, clamp) + bias so both variants
share one kernel. Cover the new chains in test-backend-ops.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* metal : fuse sqrt(softplus) MoE routing with a selection bias
Add TOPK_MOE chains anchored at the softplus node (UNARY+SQRT+ADD+
ARGSORT+GET_ROWS plus optional norm/scale) so the DeepSeek-V4/V4.1
routing fuses on Metal. The kernel scores with sqrt(softplus) and ranks
selection-biased keys while the gathered weights and the norm stay
unbiased, matching the reference graph. Dispatch hooks at the softplus
node, cover the norm-bias combination in test-backend-ops and record the
new deepseek4/deepseek41 fusion counters.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* metal : fuse sigmoid MoE routing with a selection bias
The sigmoid-gated MoE archs (GLM, DeepSeek-V2/V3.2, ...) rank
sigmoid(logits) + exp_probs_b with an optional weights-norm, but only
the softmax and sqrt(softplus) routing chains were fusable.
- add 6 sigmoid-anchored TOPK_MOE chains: sigmoid + ADD + ARGSORT +
GET_ROWS with optional SCALE and SUM_ROWS + CLAMP/SCALE + DIV norm
tail; the selection-bias ADD is required, same as the sqrt(softplus)
chains
- generalize the gating function constant from a sqrt(softplus) bool
to an int (0 = softmax, 1 = sqrt(softplus), 2 = sigmoid); sigmoid
scores are 1 / (1 + exp(-logits)), matching the CPU, Metal unary and
CUDA kernels
- anchor the unary dispatch hook at sigmoid nodes too
- test-backend-ops: 4 sigmoid routing cases with and without a
weights scale
- re-record tests/fusion/MTL.csv: the sigmoid archs (glm5-next,
deepseek2/32, kimi, minimax, ...) now fuse their routing, and
muse-spark prefill drops one MOE_REDUCE fire because the reordered
graph allocates the routing weights inside the reduce output range,
which the dst/input overlap check correctly declines
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* cont : remove low-signal comments
* ocd : valign ggml_metal_fusions
* CUDA/HIP: matrix-core (MFMA) lightning indexer for CDNA2 (gfx90a)
* Addressing reviewers comments, mma.cuh and physical warp size
* CUDA/HIP: single mma.cuh matrix-core lightning indexer ::CDNA::
* Support RDNA3
* CUDA/HIP: stream heads through the mma lightning indexer kernel
Process 16 heads at a time with double-buffered Q staging instead of keeping
all heads resident in shared memory, and drop the GGML_INDEXER_NO_MFMA env.
* enable for CUDA
* CUDA: Remove wmma lightning-indexer, also run head 32 through mma kernel
* CUDA: Nv mma dosent transpose accumulator
---------
Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* vulkan: add F16 input to the FWHT
The Vulkan FWHT accepts F32 input only. This adds an F16 source variant of the
shader, for both the subgroup and the shared-memory paths, so the kernel reads
an F16 source directly instead of requiring a converted copy.
The F16 shader needs shader-float16, so the pipelines are only created when the
device reports it. A null pipeline makes ggml_vk_can_use_fwht fall back, which
is the same path an unsupported width already takes.
supports_op and the dispatch both already route through ggml_vk_can_use_fwht,
so there is a single place asking whether an op qualifies; nothing to
consolidate here the way the CUDA/Metal ports needed.
test-backend-ops on an M5 Pro (MoltenVK, fp16: 1): MUL_MAT 1128/1128,
including all 24 Hadamard cases (18 existing F32, 6 new F16).
Assisted-by: Codex
* vulkan: FWHT kernels for block widths above 512
The Vulkan FWHT covers block widths 64 to 512; wider blocks fall back to the
dense matmul. Extend the width list to 8192 and pick the shader per width: the
subgroup shader keeps n/subgroup_size values per invocation, so widths above
512 (or a subgroup narrow enough to exceed 64 values per invocation) use the
shared-memory shader with one row per workgroup and a 256-wide block. The rows
per workgroup become a specialization constant so the dispatch matches.
Pipelines for a width are only created when the device's shared-memory and
workgroup-invocation limits allow it; otherwise that width keeps the fallback.
This is the Vulkan counterpart to #29100 (CUDA). Stacked on #29101; the last
commit is the one to review.
Assisted-by: Codex
* checkpoints: reduce overhead
Assisted-by: Codex
* rebase
Adapt the slot save/restore checkpoint appendix from #26004 to the
shared_ptr checkpoint list and the common_state_data buffers.
Assisted-by: Claude Opus 5.5
---------
Co-authored-by: cwriter <cwriter@localhost>
* ggml : support `truncate=false` for YaRN
YaRN calculates correction intervals to switch frequency interpolation
methods, but the setting to not round the boundaries of these intervals
to integers (`truncate=false`) was not supported.
Models that use `truncate=false` include gpt-oss and some PLaMo-3
models; to support them correctly, the GGML operations and CPU backend
are modified to handle calculations without rounding. The default
behavior is `truncate=true`, and existing calls not explicitly
configured will maintain their previous behavior.
* support CUDA backend
* add tests
* move RoPE helpers to internal header
* format
* conversion : resolve quantized_layers from source tensor names
This commit changes how the per layer NVFP4 quant_algo are resolved from
hf_quant_config.json quantized_layers.
The keys in quantized_layers are module paths in the source checkpoint,
and these can be either be modules like 'lm_head' or parent modules
like 'layers.0.ffn.experts'. Currently these keys are passed through
filter_tensors, but they are not tensor names. So a filter_tensors that
keeps an explicit list of tensor names, like Nemotron's with --mtp
would drop them. That means there was no `_nvfp4_layer_algo entry`, and
`_tag_prec_a4` fell back to `_nvfp4_global_algo`.
This commit records the mapping from source tensor name to filtered
tensor name in index_tensors. Then in prepare_tensors each source
tensor is resolved to its closest quantized_layers parent, and
the algo is stored under the filtered tensor name, which is what
_tag_prec_a4 later looks up.
Converting Nemotron 3.5 Lightning NVFP4 with --mtp now produces:
```console
general.tensor_extra.name = ['output.weight']
general.tensor_extra.prec_a4 = [False]
```
whereas on master no general.tensor_extra fields are written. For the
non-MTP conversion, the general.tensor_extra fields are identical to
master.
Also manually verified (using pdb) that this works for
DeepSeek-V4.1-Flash-1017-NVFP4.
Refs: https://github.com/ggml-org/llama.cpp/pull/31602#issuecomment-6603337403
* conversion : shorten comment
* conversion : also remove '.bias' suffixes
Prevent malformed realtime session updates from throwing an uncaught
exception and terminating the server
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
* fit : account for skipped NextN layers when setting n_gpu_layers
Without MTP, fit counts n_layer + 1 layers while the model loader counts
n_layer_all + 1 for n_gpu_layers, including the NextN layers which hold no
weights. A fit result that offloads all layers then leaves layer 0 on the CPU
("offloaded 65/66 layers" for Qwen3.8-27B on two GPUs).
Add the skipped NextN layers to n_gpu_layers and to the tensor split of the
device that holds the output layer so the loader assigns the layers that
fit measured.
Fixes#31531
* fit : pass hp_ngl_skip by reference like the other hparams
* fit : add hp_ngl_skip to common_get_device_memory_data
* graph: store the conv state of every rollback arch in one set_rows
The rollback slots of the causal conv history were written with one
copy per slot in qwen4exp, glm5-next, kimi-k3, bailingmoe3, lfm and
mamba2, while delta_net already wrote them with a single set_rows.
build_conv_state_store now does this for all of them: one copy without
rollback, otherwise one set_rows of the overlapping windows into the
min(n, n_rs_seq + 1) slots a rollback can read.
The slot order {n_rs_seq, ..., 0} is computed once per graph and each
conv layer takes a view of its tail. Its size does not depend on the
ubatch, so the reserved graph covers every ubatch shape, and the graph
input that held it is gone.
A model that is not causal never rolls back, so the context clamps
n_rs_seq to 0 for it, like for an arch without rollback support.
* nits: compute the conv slot order in the graph and keep the n_head name of bailingmoe3
* graph: build the conv slot order in each conv layer
The slot order {K-1, ..., 0} of the conv state store was one tensor
shared by every conv layer of the graph, built on the device of the
first one, so each layer on another device copied the view of its
tail. Each conv layer now builds its own order next to its set_rows,
and llm_graph_input_rs no longer holds it.
* tests: fix CI flakes of TOPK_MOE caused by near ties
With uniform random logits, two experts sometimes land within one f32
ulp of each other at a routing boundary, mostly with sqrt_softplus
gating whose output range is narrow. A backend whose gating differs by
a few ulps from the CPU then picks the other expert and the test fails.
Draw each row of logits again until the group selection and the final
top-k both have a relative margin of 1e-4, computed on the host in
double precision.
* tests: build the TOPK_MOE logits by inverting the gating function
Draw the selection scores (probs plus bias) of each row as a random
permutation of a grid, move the experts of the best groups up by half a
step, and set the logits to the inverse of the gating function at these
scores. Every expert and group decision is then one grid step apart on
all backends, and the host no longer recomputes the routing.
The softmax probs are built to sum to one so that the softmax keeps the
scores, and the views of the logits are not initialized again after the
logits are set.
The first command buffer takes the first 10% of the graph nodes and the
rest is split evenly between the extra command buffers. A fusion cannot
span two command buffers, so a boundary that falls inside a fusable
group runs that group unfused, and the output then depends on the node
count of the graph.
Move each boundary to the end of the fusable group it falls into, so
the encoders fuse the same nodes as with a single command buffer.
This commit adds nvfp4 scales to qwen2's build_ffn call which are
currently passed as NULL.
The motivation for this is that a model converted to NVFP4 will
currently crash because the scales are loaded but never applied in
the FFN:
```console
llama-completion: llama.cpp/src/llama-sampler.cpp:1218:
void llama_sampler_dist_apply(llama_sampler*, llama_token_data_array*):
Assertion `found' failed.
```
The attention already applies its scales but ffn is missing them.
* opencl: use the small-N f32 GEMV by default on X2E
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
* opencl: use the q5_K multi-column verify GEMV by default on X2E
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* cuda: fuse NVFP4 gate/up and GLU in MMQ
* cuda: honor the requested src1 precision in fused gate/up MMQ
* cuda: share NVFP4 scale checks, skip unused fused FP4 kernels
* cuda: test NVFP4 fused MMQ at W4A8 without scales and with batches, clean up naming
* cuda: check gate_scale before loading the fused NVFP4 gate scale
* llama : fix on-device state copy sizes for quantized KV caches
The on-device sequence state copies (LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) computed
their element count as size / ggml_element_size(). For a quantized type such as
q8_0 or q4_0 that is the number of blocks, not elements, so only part of each
range was copied: the restored state was wrong, or the copy crashed.
Compute the count as size / ggml_type_size() * ggml_blck_size() for the save,
the restore and the copy between differently chunked ranges, and assert that
each range covers whole blocks.
* llama : drop the one-line comments
* llama : drop the redundant cast in the state copy size
* server : realtime: send an error event instead of crashing on exceptions
* server : route websocket handlers through ex_wrapper, remove the realtime test
* server : realtime: keep the previous tools when the chat template rejects new ones
* server : realtime: reject tool schemas the grammar cannot convert, drop ws_ex_wrapper
* server : realtime: clarify the tool schema check comment
* opencl: use more workgroups for q4_K gate up + GLU
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
* opencl: handle FA for GQA 2 and 6 with head size 256
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
* opencl: handle FA for GQA 16 with head size of 128, 512
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
* opencl: handle FA for GQA 2 with head size 512
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* ui : share directory resolution and the cwd row
Move the lookup of a browser-picked folder to ToolsService.resolveDirectoryPath, let the working directory picker run without its chip, and split the 'Set working directory to' row into its own component, so other views can pick and show a folder the same way.
Assisted by: Claude (Opus 5.5 Medium)
* ui : open an item's actions on right-click
Sidebar conversations and models manager rows open the same actions as
their "more actions" menu on right-click. Both menus render one shared
list of actions, so they cannot drift apart. Phones keep the browser's
long-press; the dropdown's drawer already holds the actions there.
An action can render a flyout submenu in place of its row, through the
kit of the menu it is in, and close the menu after a pick.
Adds the shadcn-svelte context-menu component, mirroring the
dropdown-menu one.
Assisted by: Claude (Opus 5.5 Medium)
* ui : close conversation tabs from a right-click menu
A tab's right-click menu closes it, the other tabs or the tabs to its
right, and stops a response the tab is generating. Closing tabs that
include the active one keeps the tab the menu was opened on.
Assisted by: Claude (Opus 5.5 Medium)
* ui : open a message's actions on right-click
User and assistant messages open the same actions as their action icons
on right-click. Selected text, links, media, code and fields keep the
browser's own menu, so copying them works as before. The fork dialog
opens from the actions context, like the delete one, so both the icon
and the menu can open it.
Assisted by: Claude (Opus 5.5 Medium)
* ui : scope right-click menu shortcuts and selection to their item
Show the conversation shortcuts only on the active row, as they act on
the active conversation, and keep the browser menu only for a selection
inside the message. Name the close-tab shortcut, and drop the unused
chip toggle, menu submenus and shared cwd row.
Assisted-by: Claude (Opus 5.5 Medium)
* fix(ui): show extracted text in PDF attachment preview
Assisted-by: Claude (Opus 5.5 Medium)
* feat(ui): show page thumbnails for PDF attachments sent as images
Assisted-by: Claude (Opus 5.5 Medium)
* ui: stop retrying PDF preview conversion after it fails
A failed PDF-to-image conversion or text extraction reset its loading
flag, which re-ran the effect and retried forever. Remember the failed
item for each view and only retry when the item changes, so the error
state stays on screen.
Add a PdfConversionError story that feeds a broken PDF.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep the text of PDFs sent as images and tidy the PDF preview
Store the extracted text of a PDF sent as images in its content, so the
Text view reads it instead of parsing the PDF again; only messages stored
with the old placeholder extract it. The model still receives the images
only. Guard retries on the error state, as the item remounts per id, skip
an empty base64 payload, and share the status markup between views.
Assisted-by: Claude (Opus 5.5 Medium)
* gguf-py : validate every tokenizer merge pair, not only the first
The transformers 4.45 pair format was detected by checking merges[0],
then every element was converted unchecked: a non-pair element raised
a bare TypeError, and pairs with the wrong arity or non-string parts
were silently mangled into invalid merges.
Validate each element (2-string list) up front and raise the existing
"Unknown tokenizer merges format" error for anything else.
Assisted-by: Devin
* gguf-py : remove test_vocab.py per review
The RISC-V vector intrinsic specification requires __riscv_v_intrinsic
to be defined when the compiler supports the API, even when the vector
extension is not selected. Clang therefore defines it for rv64gc,
causing ggml to enter RVV branches where vector types are rejected
because the required ISA extensions are absent.
Use __riscv_v for optional RVV implementations; this macro represents
the selected target ISA. Retain __riscv_v_intrinsic checks in SpaceMIT
sources, where they validate mandatory compiler API support rather than
select optional vector code.
Closes: https://github.com/ggml-org/ggml/issues/1535
* ggml: bound views by the bytes their strides reach
A view was checked against its source with the size of a contiguous
tensor of the same shape, computed before its strides were set. Views
whose rows overlap failed that check although they stay inside their
source, and views with wider strides passed it while reaching past
the end. The view functions now pass their strides to the tensor
constructor, which checks the bytes they actually reach, like
ggml_nbytes. The conv windows of delta_net are created with all their
rows at once instead of being widened after the check.
* ggml: assert that a strided view starts from contiguous elements
The view functions set nb[0] to the type size, so the source must have
contiguous elements along dim 0 for the view to read them.
The gguf writer staged each tensor into a temporary vector before
writing it. For host-buffer tensors the data pointer is directly
writable, so skip the staging copy and write it straight to the file.
This avoids the per-tensor allocation, zero-init and memcpy.
Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* common : fix -hf tag matching inside longer quant names
":F16" picked a BF16 file when the repo had both, since "BF16"
contains "F16". The same applied to ":Q4_XS" vs "IQ4_XS". The tag
now has to start at a token boundary in the path.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* common : accept underscore as -hf tag boundary
Some repos separate the quant with an underscore, e.g.
stepfun-ai/Step-3.5-Flash-GGUF-Q4_K_S ships
step3p5_flash_Q4_K_S-00001-of-00012.gguf, which :Q4_K_S no longer
matched. Add _ to the token boundary and pin the tag and default
lookup cases.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* hexagon: initial support for LIGHTNING_INDEXER op
* hex-lght-idxr: hmx for prefill
* hex-lght-idxr: q8_0 decode gather transpose, two-block decode compute
* hex-lght-idxr: pipelined hmx path for decode and prefill
* hex-lid: add support for K (f32)
* hex-lid: add support for K (f32)
* hex-lid: add support for q4_0
* hex-lid: remove scalar VTCM reads in the HMX q4_0 K prep and remove vector spills
* hex-lid: host/kernel kparams alighnment and remove sw devices in kernel
* hex-lit: update dma pipeline to use canonical dbl buffering in all paths
* hex-lit: add GGML_HEXAGON_LI_SELECT to force hmx/hvx paths
* hex-lit: add support for dims collapse
* hex-lit: make kparams constistent with other ops
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* Optimizing Gemma4 graph for prompt processing
TODO:
- Currently complete microbatch needs to contain tokens which does not require the logits; need to optimize this further
- There are a few more operations that are performed after the kv cache is calculated, need to safely optimize them.
* Not filtering the rows when they are unmasked
* Cropping out the attention mask too.
* gemma-4: improved pruning, mask cropping and PLE scaling
Prune before the first KV-shared layer with the single-stream, causal, non-embedding safety gates.
Use cropped masks only during shared-layer attention.
Preserve requested layer-input embeddings in the KV prefix.
Remove the PLE contiguous materialization and fold equivalent PLE scaling into RMSNorm parameters/weights.
* spec-mtp: Enable masked h_nextn extraction only when the draft context shares the target KV cache
This lets Gemma 4 prune unused prompt rows before its shared layers so that we keep the full perf P.P gains with MTP.
Avoid extracting target hidden states for output-free prompt chunks in that shared-MTP path.
When a batch has a selected final token, extract only that row unless the whole verification batch is selected.
Keep ordinary MTP and classic speculative decoding behavior unchanged: they still use unmasked h_nextn because the draft
context must catch up on all target rows.
speculative-simple needs an update to select the final prompt row only for shared-KV MTP,
so the optimization works there without changing other speculative modes.
And it looks like the server already selects its final prompt row for first-token sampling.
* spec-mtp: make sure to not drop first draft round
* gemma4: enable pruning based on static cparams.kv_unified/n_seq_max=1
* gemma4-mtp: simplify prunning logic and avoid graph branches/re-allocs
* gemma4: update pruning gates to use only n_seq_max==1, kv_unified can now be used for multi-seq
* spec-simple: revert spec-simple changes for now
* gemma4: remove causal_attn from pruning gate (last source of reallocs)
* cont : adapt to llm_attn_mask changes
* cont : add TODO
---------
Co-authored-by: Meet Amin <meetamin378@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* hexagon: faster TOP_K for the backend sampler (skip tiles, then block scan)
Skip the bitonic sort for tiles whose max cannot reach the running k-th
best value, then replace the tile sort with top_k_scan(): reject 256-float
blocks with one HVX max-reduce and insert only surviving elements into a
sorted k-entry list. Used for k in 2..64 and rows wider than 1024, in both
the single-row and multi-row paths.
Ties keep the first-seen index. NaN is treated as -INF.
* hex-top-k: eliminate scalar reads from vtcm
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Playwright's bundled Chromium lags behind Chrome security releases
(1.56.1 ships 141, and even 1.64.1 ships 156.0.8078.4, before the
156.0.8078.73 fixes), and Playwright launches it with --no-sandbox by
default. Outside CI, use the installed, auto-updated Chrome with the
sandbox enabled. CI keeps the bundled Chromium, so the self-hosted
runner's preinstalled 1.56.1 browsers still work.
Assisted-by: Claude (Opus 5.5 Medium)
* hexagon: add support for DSV4_HC_PRE
* hexagon: hsv4-hc-pre plain/gated split kernel and l2 prefetch optimization
* hex-dsv4-hc-pre: double-buffered DMA for prefill
* hex-dsv4-hc-pre: refactor DMA block setup
* hex-dsv4-hc-pre: precompute kernel params
* hex-dsv4-hc-pre: support for n_hc != 4
* hexagon: add support for DSV4_HC_POST
* hex-dsv4-hc-pre: refactor gated PRE prefill DMA to dst writeback after each block
* hex-dsv4-hc-pre: remove in-loop vector spills and redundant device-side checks
* hex-run: add support for test-llama-archs tool
* hex-dsv4: refactoring and optimizations
Remove non-dma paths.
Clean up DMA pipelines and dma_addr_t casts.
Add dims collapsing support.
Tested with test-llama-archs -a qwen4ext and test-backend-ops
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* speculative : save and restore MTP hidden state
Assisted-by: Codex
* tests : remove MTP state test
Assisted-by: Codex
* speculative : simplify MTP checkpoint state
Save only the boundary hidden row; verification rows are rebuilt on the supported restore paths.
Assisted-by: GPT-6
---------
Co-authored-by: zouyongxing <zouyongxing@bytedance.com>
* mtmd : add Nemotron 3.5 Transcribe support
Assisted-by: Claude
* convert : fix vocab of draft models for Nemotron-H targets
Assisted-by: Claude
* convert : reject non-BF16 Transcribe checkpoints
Assisted-by: Claude
* mtmd : make speaker FIFO predictions local
Assisted-by: Claude
* mtmd : normalize the mel in encode_windows
Assisted-by: Claude
* convert : map the generated speaker kernel
Assisted-by: Claude
* mtmd : validate Nemotron Transcribe mmproj shapes
Assisted-by: Claude
build_conv_state_at makes a contiguous copy of the conv input tail
before writing it into the recurrent state, so every GDN layer issues
two device copies per generated token where one is enough. On
Qwen3.8-Flash-Next (36 GDN layers) that is 72 copies per token, now 36.
Speculative decoding adds one more per rollback slot. ggml_cpy reads
the strided view directly, as the shared build_conv_state already does,
so the extra copy is dropped.
* ui: add Playwright harness for recording demo videos
Scripted walkthroughs under tools/ui/demo are recorded with Playwright and
exported as H.264 mp4 for pull requests (demo/pr) and docs (demo/docs).
Fixture MCP servers and seeded settings keep recordings deterministic, a
drawn cursor and captions make the flow readable, and DEMO_UI_DIR points
the recording at another checkout such as a PR worktree.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: record demo scripts from PR descriptions and a configurable storage
Demo scripts no longer live in the repo. `npm run demo:pr -- <number>`
records the `ts demo` block from a PR description against the PR head,
uploads the mp4, a GIF preview and the script to DEMO_STORAGE, and links
the preview in a `## Demo` section of the PR description.
`npm run demo:docs` records the docs scripts kept in DEMO_STORAGE.
DEMO_STORAGE (hf:// buckets or file://) and the public URL are read
from tools/ui/.env.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: import generated conversations into demo recordings
Demos can now start with existing chats: `test.use({ traces })` imports
llama.app JSONL traces through Settings > Import/Export before the video
starts. Built-in traces are generated deterministically, other names are
downloaded from DEMO_TRACES_DATASET, and `demo:traces` pushes the built-in
ones there. Adds `demo.scroll` for wheel scrolling.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: keep demo captions clear of the chat form and isolate concurrent runs
Captions move to the top of the frame, since at the bottom they covered
the chat form a demo is often about. Each recording run gets its own
Playwright output folder, so a run starting elsewhere no longer clears
another run's video before it is saved.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: record demos at 2x and 60 fps with human-like timing
Recording goes through Chromium's screencast instead of Playwright's
video, which encoded VP8 at about 1 Mbit/s at 1x. Frames are captured at
the device pixel ratio (DEMO_SCALE, default 2) and encoded into a 60 fps
H.264 mp4, starting once the page has settled.
The cursor now glides faster along a slight arc with Fitts-like
durations, settles briefly before each click, and typing and pauses vary
a little. The randomness is seeded per demo, so re-recording gives the
same motion.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: document the demo harness prerequisites
The demo README now lists the tools the harness needs and what each is
for, and walks through creating the Hugging Face bucket and the optional
traces dataset, with the matching .env settings.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: version linked demo URLs so re-recorded previews refresh
GitHub's image proxy caches a preview by URL, so a demo re-recorded at
the same PR head kept showing the old GIF. The Demo section now links
the files with a content hash in the query string.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: label PRs that feature a video demo
Adds a `video demo` label and a workflow that applies it when a PR
description has a demo written by the UI demo pipeline (the managed
demo:start/demo:end markers, or a preview linking to a demo video),
and removes it when that section is gone.
Assisted-by: Claude (Opus 5.5 Medium)
* ci: label video demos with gh instead of github-script
The label workflow now runs the gh CLI in a shell step, like
pr-draft-label.yml, so it needs no action outside the repository's
allowed list. The video link check only counts videos uploaded for the
PR itself, so a description linking other PRs' demos isn't labeled.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: zoom demos in on details with a two-finger pinch
demo.zoom pinches in on a control or a point, holds, and pinches back out
to 1x, so a detail too small to read in the frame can be shown up close.
It drives Chromium's own pinch gesture, so the page's fixed bars stay put,
and works at desktop and phone sizes alike.
The pinch is evened out where Chromium misses: one anchored near the top
does not zoom at all, so it starts lower, and one near the bottom
overshoots its ratio, so a small pinch from the middle corrects it. The
way out reads the real scale until the page is back at 1x, which the next
gesture needs.
The overlay also draws a fingertip per touch, so a pinch on a touch
device shows as two fingers.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: wait out a cold dev server before a demo starts
A fresh checkout's dev server optimizes its dependencies on the first
load and reloads the page, which can take well past an action's timeout,
so a demo failed before it started. The demo now waits up to 90 seconds
for the app; the recording only starts once it is ready.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: record a video per test in a PR demo, shown in a table
A large PR can split its demo into several shorter videos: each test()
in the ts demo block is recorded as its own video (<sha>-1.mp4, ...),
and the Demo section lays their previews out in a table, each column
headed by its test's title.
GIF previews now keep their longer side at DEMO_PREVIEW_WIDTH and shrink
until they are under 5 MB, since GitHub's image proxy refuses anything
larger and shows a broken image instead.
Importing traces works on a phone viewport too, where Settings sits in
the sidebar under the page until the menu key slides it aside.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fix demo trace imports on desktop and lay landscape demos out two to a row
Importing traces failed on a desktop whose collapsed sidebar sits under
the page: it now expands the sidebar from the page's menu key on a phone
or from the sidebar's logo on a desktop, whichever makes Settings
reachable. Afterwards it forgets the expanded sidebar, so a demo starts
with it collapsed, as on a first visit.
A PR demo with several videos shows up to four portrait (phone) videos
to a row, or two landscape ones, so each preview stays readable.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: only record other people's PR demos at a reviewed head commit
demo:pr runs the PR's code and demo script on this machine, with the
gh and hf logins, so a PR by someone else now needs --trust <sha>
naming the head commit that was read, and is refused again after new
pushes. Recordings also launch Chromium with its sandbox on, which
Playwright turns off by default.
Assisted-by: Claude (Opus 5.5 Medium)
ggml_roll validated each shift with abs(shift), but abs(INT_MIN) is signed
integer overflow UB evaluated before the bounds check runs. Compare with a
signed-safe test instead.
This commit updates the _flush_nvfp4_experts to take a new parameter
which is a dictionary mapping the key to the original model tensor name.
The motivation for this is that I noticed that the tensor names being
looked up were not the original model tensor names, for example:
```console
model.layers.1.mlp.experts.up_proj.weight
```
But _nvfp4_layer_algo uses the original per-expert names:
```console
(Pdb) p self._nvfp4_layer_algo
{
'backbone.layers.1.mixer.experts.0.up_proj': 'W4A16_NVFP4',
...
```
This means that the lookup will never match for experts and fall back to
_nvfp4_global_algo. If the global algo is W4A16_NVFP4 then this would
go unnoticed, but for a model like
nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4 which is
MIXED_PRECISION, the expert tensors will not get prec_a4=false so they
would use the native W4A4 path on Blackwell even though the original
model tensors declare them as W4A16_NVFP4.
* cuda: gather batched tensor reads into mapped host memory
Assisted-by: Codex
* add comments
Assisted-by: Codex
* simplify the implementation based on checks that exist earlier in the flow
* simplify batched tensor gather indexing
* add conditional dispatch
* address comments - make host registration raii, better test mindful of ci time, assert instead of return false
* cuda : gather small reads and use pinned DMA for large reads
* cuda : use std::vector for host gather descriptors
* cuda : add 8-byte gather copies, generalize tensor test and log registration failures
* common buffer for descriptors
* simplify stream sync + increase to 128k descriptor pinned buf
* Add a backend interface for collectives across ranks running in separate processes, implement it in the CUDA backend using NCCL, and use it in the RPC server when available.
* Drop ggml_backend_comm_allreduce_rank_t and reuse ggml_backend_comm_allreduce_tensor for per-rank communicators
* reuse ggml_backend_cuda_comm_allreduce_nccl for per-rank communicators. Remove the helper function backend_comm_proc_address
* Change ggml_backend_comm_get_unique_id_t and ggml_backend_comm_init_rank_t to more generic multi-process backend communicator interfaces
* set maxCTAs in ncclConfig_t to 8 for dgx spark in ggml_backend_cuda_comm_init_rank
* meta: support GLM-5.3 next sm tensor
* glm5-next : cont q_absorbed for -sm tensor with non-unified KV
* meta : keep GGML_OP_ARANGE in ggml.h order
* glm5-next : split MLA heads in pairs for -sm tensor
The MLA cache has a single KV head, so the heads on a device are its GQA
ratio. The HIP tile FA kernel for D=512 needs an even GQA ratio; with 3
devices the 16 heads of the test model were split 5/5/6 and it aborted.
* meta : keep the split for views inside one slab of the split axis
With a non-unified KV cache, build_attn_mha splits q into streams with a
view. For a permuted q, e.g. the absorbed query of glm5-next, that view
aborted with "view of permuted tensor not implemented". The split is now
kept if every other dim of the source and the view stays inside one slab
of the split axis. The per-device strides are only rescaled above it.
From tetherto/qvac-fabric-llm.cpp#353.
Co-authored-by: Pratik Narola <pratik.narola@tether.io>
* glm5-next : drop the q_absorbed cont
The meta backend can now split the per-stream view of the permuted q, so
the cont is not needed for -sm tensor with a non-unified KV cache. With
the cont, the q of flash attention was not contiguous, and when FA ran on
the CPU (WebGPU) the meta get_tensor aborted.
---------
Co-authored-by: Pratik Narola <pratik.narola@tether.io>
kernel_cpy_t_t indexed the dst innermost dim with an implicit element-size
stride instead of args.nb0, so cpy into a permuted view (nb0 != type size)
wrote the wrong addresses. caught by the new test-backend-ops cpy case with
permute_dst = {1, 0, 2, 3} from #31446. addressing is unchanged when dst is
contiguous in dim 0.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* vulkan: skip maxStorageBufferRange pre-check for non-compute tensors
Pre-allocated leaf tensors and views (GGML_OP_NONE / GGML_OP_VIEW) are never
executed by the backend, but were still rejected by the descriptor-size
pre-check in ggml_backend_vk_device_supports_op when larger than
maxStorageBufferRange on devices without shader_64b_indexing. The scheduler
then aborts with 'pre-allocated tensor (...) in a buffer (Vulkan0) that cannot
run the operation (NONE)' - e.g. the MoE expert cache banks (llama-moe-cache)
of a 256-expert model with --moe-cache-mib on GCN GPUs, where a whole bank
(2.4 GB) exceeds the limit while the consuming ops only access small
per-step views.
Skip the size pre-check for non-compute ops (NONE and VIEW); the buffer type
allocation limit still bounds the actual allocation, and compute ops keep
validating their sources.
Tested on AMD Radeon Pro VII (gfx906, no shader_64b_indexing):
- before: GGML_ABORT 'pre-allocated tensor (moe_cache.0.blk.0.ffn_down_exps.weight) ... cannot run the operation (NONE)' with llama-cli -ngl 99 --cpu-moe --moe-cache-mib 6144 (Cyber-Tiel-Coder-35B-A3B, qwen35moe, 256 experts per layer)
- after: loads and generates at 8.1 t/s (vs 0.6 t/s without the cache)
- regression: -ncmoe 16 path unchanged (19.2 t/s), qwen3moe/qwen35moe without cache unchanged
Fixes#30697
* vulkan: use ggml_op_is_empty for the non-compute size-check skip
---------
Co-authored-by: FamousMonsterr <FamousMonsterr@users.noreply.github.com>
* delta_net: access conv state via views
* fix sm tensor
* delta_net: fill the conv slot index on the host
The rollback slot index of the conv windows is a constant per graph
shape. It is now an I32 input of the recurrent memory filled on the
host, instead of an arange and an argsort computed on the device, and
the meta backend no longer needs a mirrored split for ARANGE.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* vulkan : fix out-of-bounds NVFP4 UE4M3 scale lookup
On the non-OCP path, ue4m3_fp32_lut has 128 entries and is indexed with the raw scale byte. ggml_validate_row_data accepts every byte and ggml_ue4m3_to_fp32 ignores bit 7, so a byte >= 0x80 reads past the shared array. Vulkan robustness does not cover shared memory.
Size the table, both init loops, and the two host shared-memory budgets to 256 entries. ue4m3_to_fp32_build() already masks the exponent and treats only 0 and 0x7F as special, so entries 128-255 match the CPU. A 256-entry table keeps 0xFF distinct from 0x7F: the CPU decodes 0xFF as a non-zero value and 0x7F as zero.
Each NVFP4 pipeline uses 512 more bytes of shared memory. ggml_fp32_to_ue4m3 never sets bit 7, so models quantized by llama.cpp are unaffected.
Assisted-by: Claude
Assisted-by: Codex
* vulkan : fix out-of-bounds NVFP4 UE4M3 scale lookup
Assisted-by: Codex
---------
Co-authored-by: Dongsheng Luo <7137760+flyingdoog@users.noreply.github.com>
The dspark block mask is now derived at input-fill time: llm_attn_mask::block_uniform
fills every token of a sequence with the mask row of its first token, so the
graph-level reshape/repeat in dflash is gone and cnv always aliases raw.
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
* cuda : prevent premature eviction of graphs about to be used
* add last_graph_use_time for entire cuda graph map and testcase for cuda graph
Assisted-by: Codex
* delete the test
A cache that shares the cells of another one (f.ex. the gemma4 assistant) can build a graph before its sequences have any cells, f.ex. at the server warmup. get_n_kv_rows() then returned 0, which means that kv_rows are not used, so the graph fell back to the dense KQ mask and the scheduler re-planned for it. Later graphs with larger masks then aborted with GGML_SCHED_DEBUG_REALLOC=1.
* hexagon: add tiled HVX transpose path for permuted CONT/CPY
Permuted copies where dst dim 0 is strided in src, but some other src dim
is contiguous, used to fall back to HTP_COPY_KERNEL_RESHAPE. That issues
one DMA descriptor per element. The channel-first <-> channel-last
permutes around per-channel norms (e.g. [W,H,T,C] <-> [C,W,H,T] in video
VAEs) ht this path, and these copies ran at about 45 MB/s.
Add HTP_COPY_KERNEL_TRANSPOSE for same-type, same-shape f32/f16 copies
into a contiguous dst.
Host:
- drop extent-1 dims and merge dims that stay adjacent in src into a
canonical view, stored in the kernel params. There are no minimum
size limits.
- the copy is treated as a 2D transpose between dst dim 0 and the
contiguous src dim, batched over the remaining dims.
- tiles are sized to 64 KB per buffer and halved until they fit in
session VTCM. blkk is capped so the halfword vgather offsets fit in
16 bits.
Kernel, per tile:
- 2D DMA of blk0 strided src rows (blkk contiguous elements each)
into VTCM
- vgather transpose in 32x32 (f32) or 64x64 (f16) blocks
- 2D DMA of blkk contiguous rows out to dst
Staged and transposed buffers are double-buffered, so the next tile's
input DMA overlaps the current transpose and output DMA. Tiles are split
across threads and mdev sessions.
Small dims:
- packed mode: when dim k follows dim 0 and dim 0 divides the HVX
lanes (f32: 2/4/8/16, f16: 2/4/.../32), one vgather writes several
whole dst rows, and each tile goes out as one contiguous DMA.
- when the contiguous src dim fits in one vector, it is staged exactly
and the tile is made long along dim 0.
All other copies select the same kernels as before.
Tests: add test_cont cases for perms {1,2,3,0} and {3,0,1,2}, f32 and
f16, at various sizes.
* hexagon: cpy transpose cleanups
- fix vgather halfword offset comments: offsets are unsigned 16-bit, caps are kept lower for speed
- set packed directly in the packed case
- add dim_k == 3 CONT tests (full reversal)
Assisted-by: Claude Opus 5.5
* hexagon: use cpy transpose path for ne[0] == 1 permutes
- drop the nb[0] check in the transpose solver: size-1 dims are skipped, and the canonical tp.nb[0] check covers packed dim 0
- try the transpose path before SAMESHAPE_SAMETYPE, which did one DMA per element for these
- add ne[0] == 1 permute CPY tests
Assisted-by: Claude Opus 5.5
* hex-cpy: further improve cpy
Bring in improvements from aizip branch/PR.
Cleanup kparams and replace divs with fastdiv.
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
* hex-cpy: avoid vtcm bank conflicts
* hex-cpy: trim down the comments
* cpy-test: trim down new tests that cover permute scenarios
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
* feat: support PocketTTS voices with no voice cloning
The kyutai/pocket-tts-without-voice-cloning weights ship predefined voices as the
backbone KV state after the voice prompt. The converter turns each state back into
the prompt embeddings and stores them in the mmproj (voice.<name>), the zeroed voice
encoder is left out, and the "voice" field of /tts and --tts-voice select one of them.
A speaker reference is rejected when the model has no voice encoder.
Assisted-by: Claude Fable 5.1
* convert : pocket-tts voices never fail the conversion, log the drift of the state
The voice states shipped by Kyutai are not always the ones the pack's
weights give for the prompt: italian, german and spanish match within the
bf16 rounding, english_2026-04 is 2% off and french 20% (its layer 0 state
comes from other weights), while the reference loads them as is and they
still sound like the same voice. The converter now re-prompts the backbone
with the recovered embeddings, compares the whole state and logs the
distance per voice, warning above 5%. A state that does not belong to the
weights at all is skipped, so the packs of kyutai/pocket-tts convert too.
Assisted-by: Claude Fable 5.1
* mtmd : list the predefined voices, expose them in GET /props
mtmd_gen_audio_info gains n_voices and mtmd_gen_audio_voice_name() the
names, with the voices worded as a feature of any audio generation model.
The server lists them under "voices" in /props when the model ships some.
Assisted-by: Claude Fable 5.1
* server : an unknown voice is a bad request
A model with predefined voices answers 400 naming the voice instead of 500
"failed to process TTS prompt". A model without them ignores "voice" with a
warning, as the other audio generation models do.
Assisted-by: Claude Fable 5.1
* refactor api
* fix conversion
* small renaming
* fix lint
* tts : a model with predefined voices speaks with the first one by default
Without a voice prompt such a model only produces near silence. When neither
a voice nor a speaker reference is given, llama-tts and the server now take
the first predefined voice and say so in the log, as the LFM3-Audio template
does with its default voice.
Assisted-by: Claude Fable 5.1
* mtmd : one comment per voice accessor, converter states the rule only
The clip.h comment described clip_gen_get_voice but sat above all three
functions; the converter comment quoted the drift measured per pack, which
goes stale with the next upload of the weights.
Assisted-by: Claude Fable 5.1
* clean up redundant code and comments
* mtmd: match the LFM3-Audio voice name case-insensitively
The tts chat template lowers the voice name, so Male speaks with
the male voice; the voice lookup lowers it the same way and names the
bitmap after the canonical voice.
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Co-authored-by: Pascal <admin@serveurperso.com>
* llama : use llm_attn_mask and kv_rows for unified attention
Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* cont : optimize `get_seq_cells`
* llama : build the kv_rows lists from per-sequence cell bitmaps
Keep a bitmap over the cells for each sequence (the transpose of `seq`), updated in seq_pos_inc/dec, and build the cell lists by walking its set bits. The lists stay in cell order, cells shared by several sequences go to every list, and nothing has to be rebuilt when cells are removed.
Also leave out of the lists the cells that the SWA hides from every token of the slice (standard and chunked SWA).
* cont : make get_kv_rows_p_min private
* cont : minor metal
* llama : use kv_rows for every ubatch of a unified cache
A ubatch made of equal-length runs with a single sequence each (f.ex. the drafts of several slots) now splits into per-sequence slices.
Any other ubatch (interleaved, uneven or coupled sequences) becomes a single slice over the cells of all its sequences instead of falling back to the dense KQ mask, so the graph keeps the topology of the reserved graph.
* llama : merge get_kv_rows_seq_ids and get_kv_rows_p_min into get_kv_rows_slices
Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB)
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
The note under 'Run the inference' still said that the CANN backend only
supports FP16/Q4_0/Q8_0 models, while the DataType table in the same file,
and ggml-cann.cpp (BF16 is enabled for MUL_MAT, GET_ROWS, CPY and CONT for
every non-310P device), say otherwise. Spell out the per-device support and
mark BF16 as unsupported on 310P instead of leaving the cell empty.
GGML_CANN_NO_PINNED is read by ggml-cann.cpp when allocating host buffers
and when advertising the host_buffer device capability, but it was not
documented, so the escape hatch for a failing pinned host allocation was
not discoverable.
Signed-off-by: li-lizhe <147392333@qq.com>
* vulkan: add tensor-parallel all-reduce communication
* vulkan: require shared semaphores for tensor communication
Assisted-by: Codex
* vulkan: reduce directly from shared host buffers
Assisted-by: Codex
* vulkan: fix ring all-reduce races with 3+ devices
Two hazards in the tensor-parallel ring all-reduce made -sm tensor output
non-deterministic with three GPUs:
- With n = 3, step t's host-slot region starts at t * cels * 2 bytes, which
is not a multiple of 32, so consecutive steps' transfer copies share a
32-byte sector of the imported host memory. On NVIDIA the later copy's
leading bytes (up to 10 F16 values at the start of every chunk) end up
stale in host memory, even when the copies are ordered by a semaphore.
Give every step its own 256-byte-aligned region. n = 2 and n = 4 already
had aligned regions for the usual sizes, so their layout is unchanged.
- A single-shot all-reduce leaves every device's host slot being read by all
peers, but the next ring round only ordered its first transfer after nextd,
and only from round 2 on. With n >= 3 a device could overwrite its slot
while prevd was still summing it. Wait for all peers on the first ring
round after a single-shot all-reduce.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
* vulkan: fold the tensor-parallel all-reduce into the compute submissions
With -sm tensor the meta backend drives every device from one thread and calls
graph_compute once per subgraph (two per layer), so decode was bound by the CPU
cost of submitting, not by the GPUs: ~1500 vkQueueSubmit per token for a 64-layer
model on two GPUs, with the GPUs idle ~40% of the time.
- Don't split subgraphs by flops while a communicator is active. The heuristic
was tuned for whole graphs and turned each small subgraph into ~5 submits.
- Fold the small (single-shot) all-reduce into the compute submissions: the
last command buffer of a subgraph stays open, the all-reduce appends the
gather copy to it and submits it, then opens the next command buffer with the
peer waits and the reduce, which the next subgraph is recorded into. Each
subgraph now costs one submit per device. The reserved progress value is
attached by whichever path submits the open context.
- Alternate between two host slots, placed outside the ring's regions, so the
gather no longer needs a WAR wait on the peers' previous reduce.
- On synchronize, submit the deferred work of every device in the group first,
so the tail of the graph runs on all devices in parallel.
- Wait for the transfer queue as well as the compute queue before freeing or
reallocating communication buffers, and don't leave a stale cap behind when a
reallocation fails.
qwen3.5 27B Q8_0 on RTX 4090s, tg128: 2 GPUs 32.2 -> 49.7 t/s, 3 GPUs 20.9 ->
40.0, 4 GPUs 15.0 -> 27.6; pp512 on 2 GPUs 2138 -> 2320. Greedy output is
unchanged.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
* - Two-phase all-reduce over shared host memory: write own slot + signal, wait peers + add from their slots
- One exported timeline per device, no backend hooks; comm work recorded in the backend compute ctx
- F16 wire above 1 MiB, F32 below; small reductions ride with the next graph submit
- Double-buffered, grow-only host slots; zero-size splits zeroed
- Remove ring, transfer-queue leg, device staging, shrink, Windows shims
- comm_init returns nullptr when unsupported; add GGML_VK_COMM_DISABLE
- No flops-based submit splitting for graphs below 1 GFLOP
[rebased onto the ring all-reduce fix and the folded all-reduce, which this design replaces; their backend hooks are removed here]
* add minimal submit split size
* replace opaque_fd with sync_fd for spec compliance
* add precision check
* add semaphore pool
* rename env var, add logging if requirements not met
* reduce comments
* code improvements
* vulkan: use posix_memalign for the tensor-parallel host buffers
std::aligned_alloc is unavailable on Android bionic below API 28, which
upstream switched away from in #30828; do the same for the communicator's
host slots.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
* vulkan: synchronize the tensor-parallel all-reduce through the host
Exchanging SYNC_FD semaphores costs a kernel wake-up on every all-reduce:
on NVIDIA the waiting GPU starts ~35 us after its peer has finished, which
is most of the idle time in a decode step. OPAQUE_FD semaphores avoid that
but must not be imported into another physical device.
Use only device-local timeline semaphores instead. Each device signals
done[i] once its partial sum is in host memory (after a barrier that makes
it available to the host), and a bridge thread signals go[i] from the host
once all peers of i are done. The waiting side makes the host memory
visible with a host-write barrier, since its submit predates the data.
The thread spins for up to 200 us on each wait (GGML_VK_ALLREDUCE_SPIN_US)
and then blocks, so it is idle when there is no work. This also takes the
per-round fd export, dup and import calls off the main thread and drops the
VK_KHR_external_semaphore_fd requirement and the binary semaphore pool.
Graphs below 1 GFLOP, such as tensor-parallel decode slices, are no longer
split into several submits.
qwen3.5 27B Q8_0 on RTX 4090s, tg128: 2 GPUs 37.2 -> 45.4 t/s, 3 GPUs
24.1 -> 35.5, 4 GPUs 16.5 -> 27.1; pp512 unchanged. Greedy output is
unchanged.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
* use sync_fd on mesa drivers
* refactor for readability
* move sync_fd semaphore handling to thread, trigger mesa submit threads
* move allreduce functions to new cpp file
* clean comments
* smaller fixes, code improvements
---------
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Ruben Ortlam <rortlam@redhat.com>
* Update ggml-sycl.cpp
Increases -sm tensor prompt eval t/s by +19% with zero PPL change.
* Update ggml-sycl.cpp
Replace first pair of wait() calls with ext_oneapi_submit_barrier(). The wait() calls was discovered to cause MTP degradation.
* Update ggml-sycl.cpp
Adding comm_exchange function pointer to switch between 2 and 4 buffer ggml_sycl_comm_exchange versions @ initialization time (versus hot path) controlled by SYCL_ENABLE_DEFAULT_CONTEXTS=1 setting.
* Update ggml-sycl.cpp
adding back 2nd half of 4-buffer free.
* Update ggml-sycl.cpp
Removing extra whitespace.
* Update ggml-sycl.cpp
Excluding Windows until it is validated on Windows.
The router only removes models it took from the Hugging Face cache, but
the models manager offered "Delete from disk" for every model, so a model
from --models-preset or --models-dir failed with a "not removable" toast.
Read the source and can_remove fields of /models and show the delete
action disabled for those models, with a tooltip (a hint line in the
mobile drawer) saying how to remove them instead.
Assisted-by: Claude (Opus 5.5 Medium)
Use the popover background, softer border, rounded-xl corners and
shadow-md like dropdowns, hover cards and the settings dialog, and
match the dialog title line height.
Assisted-by: Claude (Opus 5.5 Medium)
A message queued while a conversation was busy was re-sent through
sendMessage, which targets the active conversation, so switching chats
before the turn ended delivered it into the wrong one. Re-send only when
the owning conversation is active; otherwise keep it queued there and
flush it when that conversation is opened again.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: show interrupted state for tool calls that never returned
The chat-level streaming flag was applied to every assistant message, so a
tool call left without a result in an earlier message kept spinning while a
newer reply streamed. Scope streaming to the last assistant message and mark
finished calls without a result as interrupted.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: mark a shell command stopped before any output as interrupted
exec_shell_command creates its result message before running, so stopping
it mid-run leaves an empty result instead of none.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: don't reload a loading model on select, prefer it in pre-selection
Assisted-by: Claude (Opus 5.5 Medium)
* ui: skip hidden models in automatic model selection
Assisted-by: Claude (Opus 5.5 Medium)
* ui: don't count a model being unloaded as loading
Assisted-by: Claude (Opus 5.5 Medium)
* webgpu: keep scratch alloc sizes monotonic in the batch length
The flash attention vec path and the MMVQ path append a scratch area
after dst, but get_alloc_size only accounted for it when the current
batch takes that path. The size was therefore not monotonic in the
batch length: a graph reserved at n_ubatch (tiled path, no scratch)
or at one token did not cover batches of 2 to 19 query rows, and
GGML_SCHED_NO_REALLOC builds aborted on an unexpected reallocation.
Pass the row count to both path predicates and size the scratch with
the row count clamped to the longest batch the short path serves.
* nit: align the mmvq pipeline key with clang-format 22
* webgpu: address review @yomaytk, name the mat-vec row threshold
* hexagon : fix DMA transfers with strides > 24 bits
On v75+ the 2D DMA descriptor stores row_size and strides in 24 bits.
dma_queue_push sent oversized strides to fallback_2d, which only splits
nrows, so strides of 16 MB or more were truncated - wrong addresses.
Fix will fall back to one transfer per row in that case.
Also added two CONT tests with source strides over 16 MB.
* hex-dma: refactor the fallback paths to remove duplication
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Limit the shared prefix to the first required output position of the parent and each child task.
This keeps decision output embeddings available in the final batch when questions share option tokens.
Fixes#30482
Signed-off-by: aiqubits <aiqubits@hotmail.com>
* chat : accept the think block in the qwen3-coder response format with reasoning_format none
Assisted-by: Claude Opus 5.5
* chat : define both qwen3-coder think block variants up front
This commit extracts all apple jobs except for macos-latest-arm64 into a
new daily workflow.
The motivation for this is to reducd the load on the macos runners.
The duplicated output looks up token_embd in its buffer context, so
token_embd must exist first. With the old order a load where input and
output share a buffer type (CPU-only, -ngl 0) held the table twice:
gemma-4-E2B-it Q4_K_M, no mmap, CPU model buffer 3163.73 -> 2947.73 MiB.
* memory: refuse rollbacks past the snapshots of the last ubatch
The recurrent memory and the DSV4 cache accepted any rollback up to
n_rs_seq, whether or not the last ubatch wrote that snapshot. Each now
records the deepest rollback its last ubatch guarantees and seq_rm
refuses anything past it:
- recurrent: n - 1 for a ubatch of n <= n_rs_seq + 1 tokens, 0 for a
longer one, which only guarantees the final state
- DSV4: min(n, n_rs_seq), the snapshots are rebuilt from the state
before the ubatch
The dead rollback block of the rollback test becomes the assertion it
asked for, and rollback_bound checks the exact bound of each memory
after ubatches of 1, 2, n_rs_seq + 1 and n_rs_seq + 2 tokens.
* tests: address review from ggerganov
Detect architectures with llama_arch_from_string in comp_padding and
arch_reserves_single_seq, as in rollback_bound.
gfx117X is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4_M and use it where it makes sense.
Assisted-by: Codex
Assisted-by: GitHub Copilot
Co-authored-by: Mario Limonciello <mario.limonciello@amd.com>
* WIP OCP FP8 E4M3 support
* Fix python type-checker
* Materialize input_scales also for FP8 when converting
* Fix conversion of https://huggingface.co/RedHatAI/Qwen3.8-27B-NVFP4
1. Per row FP8 wrongly fell to NVFP4 conversion, as NVFP4 checked on
dims alone. Also check on dtype for FP8
2. Need to reshape FP8 QKV projections scales in the same way that
weights are reshaped
* Allow remaining FP8 tensors to be dequantized during conversion
* Allow tensor-strategy for FP8-only compressed-tensor ckpts
* Fix and test gguf-py behavior for FP8_E4M3
* Disallow QKV-fusion for FP8 temporarily
* Fix conversion of fused qkvz & FP8 for Qwen3-Next
* Avoid fusing FP8 for --fuse-gate-up-exps as we may have different ws
* Emit warnings when not fusing FP8 during conversion
* More FP8 conversion fixes
* Add supported for strided copies to/from FP8 in CPU
* remove unneeded scratch buffer from FP8 cpu cpy/dup path
* add explanatory comments and TODOs
* Add dedicated encoder and decoder tests for FP8
* Whitelist ggml_op_get_rows in metal instead of blacklisting
required to fix failing CI, and white-listing should be more
future-proof than adding FP8 to the black-list
https://github.com/ggml-org/llama.cpp/actions/runs/38433916498/job/115422880633?pr=27512
* Fix more conversion bugs
* Defer weight materialization during conversion
* Blacklist fp8 in metal backend instead of white-list
* Adjust indents
* Fix lint + execute tests
* Add fast CPU kernel for F8_E4M3
* Conversion fixes
* Fix remaining gaps
* Enable tied weights with shared scales
* Make mandatory-F32 rules apply to FP8 during conversion
* Fix python type checks
https://github.com/ggml-org/llama.cpp/actions/runs/39765670052/job/119780379595?pr=27512
* Fix kaleidai unit-test failures
https://github.com/ggml-org/llama.cpp/actions/runs/39765670125/job/119780383150
* Relax numeric precision constraints for pre-check
* Leave fast-FP8 CPU on RISC-V for follow-up
We already exclude arm from the fast-path
* Allow output scale tensors to be located in any gguf shard
* Move test_conversion_fp8.py test execution to behavior-check.yaml
* Apply batched suggestions from code review
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Move mixed compressed-tensor logic into dequant_fp8 + add unit-tests for it
* Install all requirements also on Windows for behavior-check
Fixes https://github.com/ggml-org/llama.cpp/actions/runs/40535852590/job/122300758223?pr=27512
* Unmap lazy tensors before clean-up in Windows
Should fix https://github.com/ggml-org/llama.cpp/actions/runs/40539797824/job/122314782459?pr=27512
* LF line endings
https://github.com/ggml-org/llama.cpp/actions/runs/40539798168/job/122314791118?pr=27512
* Extract FP8 tests from llama-arch to its own test-fixture
See https://github.com/ggml-org/llama.cpp/pull/27512#discussion_r4400628221
* Update gguf-py/gguf/scripts/gguf_convert_endian.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
string_format("share of the MoE cache misses that the CPU computes from the host experts instead of uploading them to the GPU once the cache is full, from 0 to 1 (default: %.2f)",(double)params.moe_cache_cpu_share),
[](common_params¶ms,conststd::string&value){
constfloatshare=std::stof(value);
if(share<0.0f||share>1.0f){
throwstd::invalid_argument("invalid value");
}
params.moe_cache_cpu_share=share;
}
).set_env("LLAMA_ARG_MOE_CACHE_CPU_SHARE"));
add_opt(common_arg(
{"-ncffn","--n-cpu-ffn"},"N",
"keep the dense FFN weights of the first N layers in the CPU\n"
Some files were not shown because too many files have changed in this diff
Show More
Reference in New Issue
Block a user
Blocking a user prevents them from interacting with repositories, such as opening or commenting on pull requests or issues. Learn more about blocking a user.
Ramzi