* server: report the isolate working directory from get_info
Without an explicit cwd, get_info fell back to the server process
working directory even when a tools runtime was configured. That named a
host path no tool would ever run in, since an isolate starts in a
directory of its own.
It now asks the isolate for its working directory in that case, and
keeps the process one only when the tools run on the host.
* remove redundant comment
---------
Co-authored-by: Xuan-Son Nguyen <thichthat@gmail.com>
The working directory chip showed up as soon as the server exposed any
builtin tool, so a server started with just get_datetime, or a user who
turned every filesystem tool off in the settings, still got a control
that nothing would read.
Tools now declare whether they resolve their paths and run against the
working directory, next to the write permission they already publish in
the /tools listing. The WebUI shows the chip and enables the /cwd
command only when at least one such tool is both served and left
enabled.
get_output runs the waveform work the pipeline defers to it, from a
single trailing window to a full pass depending on the model. Measuring
it keeps the reported total and the audio to process ratio honest.
* feat: Add contenteditable tokenizer for badge/code-chip chat input
* feat: Add source-space undo/redo history for the rich input
* feat: Split text glued to a closing code fence onto its own line
* feat: Add ChatFormContenteditable rich input renderer
* feat : wire the contenteditable into ChatForm with auto-switch gating
* base : slash-command/misc foundation - model icon and focus-selector constants
* feat : slash-command picker and command parsing helpers
* refactor : wire command and @-mention pickers into the chat form
* ui : improve model selector keyboard navigation and load/dismiss
* feat: Unify markdown/raw-text rendering under one setting with migration
* fix: Misc fixes - tool-call subtitle, assistant wrap, progress guards
* feat: Clamp and style numeric settings inputs from registry bounds
ggml_metal_op_norm sized the threadgroup with
`nth = std::min(nth, args.ne00_t)`, which can leave nth not a multiple of
the simdgroup size. The kernels finish their row reduction with a
cross-simdgroup step where each lane of the last simdgroup reads one
per-simdgroup partial sum out of shmem_f32:
if (tiisg == 0) { shmem_f32[sgitg] = sumf; }
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
When the last simdgroup is partial it has fewer lanes than the
threadgroup has simdgroups, so the tail of the partial sums is never
read and the row sum is too small. For ne00_t = 33 nth becomes 33: two
simdgroups, but only one lane in the second, so one of the two partial
sums is dropped. The mean and variance are then wrong for the whole row.
Round ne00_t up to a whole number of simdgroups instead. Rounding up
rather than dropping the clamp keeps the threadgroup as small as
possible: deleting the line would raise nth to the next power of two
(ne00_t = 544 -> 1024 instead of 544), which costs idle lanes on 26 row
lengths below 8192 that were already correct, including 1536 and 3584.
GGML_OP_NORM is affected as well as GGML_OP_RMS_NORM - both dispatch
through ggml_metal_op_norm.
No mainstream LLM hidden size hits this: ne00_t is ne00/4 on the
vectorized path, so 4096, 8192, 2048 and friends all give a multiple of
32. It is reachable from other norm shapes, e.g. 320-channel norms.
Add NORM and RMS_NORM cases for ne0 = 33, 132 and 260 across the
existing eps values. 33 exercises the scalar path and 132/260 the
vectorized one, since only those divide by 4.
Before, on M3 Pro:
test-backend-ops test -b MTL0 -o NORM 25/50
test-backend-ops test -b MTL0 -o RMS_NORM 26/51
After:
test-backend-ops test -b MTL0 -o NORM 50/50
test-backend-ops test -b MTL0 -o RMS_NORM 51/51
test-backend-ops test -b MTL0 13943/13943
* base : @-mention picker foundation - glob search, picker nav, highlight
* feat : @-mention file/folder picker and mention badges in message bubbles
* fix: Imports
* feat : wire the @-mention picker into the chat form
* fix: Bound the glob-search result cache key and prune stale entries
* webui: load the model selected via ?model=
Opening the WebUI with ?model= selects the model but doesn't load it. The load only starts when you send your first message, so you wait for it then.
This loads it as soon as the page opens, while you're still typing your prompt. It's what the model dropdown already does, and it isn't awaited, so the UI still works while the model loads.
This is the path the Llama macOS app uses to open the WebUI, so it's a common way in.
* webui: gate the load behind ?load=true
Loading on landing is opt-in, so a plain ?model= link behaves as before and doesn't allocate memory on its own.
* webui: name the chat URL params
Collects the query params the chat routes read into a URL_PARAMS constant, instead of repeating the literals across three files. NEW_CHAT_PARAM folds into it.
The reference implementation has two mutually exclusive prompt layouts.
In non streaming mode the prefill carries the whole utterance text plus
tts_eos summed with codec_pad, and the trailing text hidden collapses to
a single tts_pad row. In streaming mode the prefill carries only the
first text token and the trailing rows stream the rest of the text
followed by tts_eos.
The pipeline built the non streaming prefill but the streaming overlay,
so the talker saw the utterance a second time during generation and read
it twice before emitting codec_eos.
The overlay is now the single tts_pad row that matches the prefill.
* ggml : add fallback definitions for missing aarch64 HWCAP bits
* ggml : require HWCAP_ASIMDHP for the aarch64 fp16 cpu variants
Also rename has_fp16_va to has_fp16, the field gates the whole FEAT_FP16
extension, scalar and vector half-precision arithmetic together.
* ui: read model modalities from the router model list
The router advertises input modalities for every model, loaded or not.
Reading them at list build time lets the UI accept image and audio
uploads for a model selected through ?model=, which has no /props yet.
* enum
The NVFP4 quantization format stores a scaling factor for every group of
16 weights, packed into a single UE4M3 byte.
The SYCL GPU code was converting these scale values using the E4M3 path,
but that's *signed*, and these are unsigned values.
* tests: add SWIGLU perf cases
perf mode had no GLU coverage. Adds SWIGLU at 17408 columns, 512 and
2048 tokens, f16 and f32, with the operands both fused and split.
* sycl: consolidate fused-GLU kernels
They differed only in which op_* they called, so take the op as an argument and share a common launcher.
Their block sizes were all 256, so launch geometry is unchanged;
SYCL_GELU_BLOCK_SIZE and SYCL_SILU_BLOCK_SIZE lose their last users so are dropped.
* sycl: contiguous fast path for the fused GLU ops
o0 == n and o1 == n collapse the de-interleave index math to the
identity, so dispatch a flat kernel in that case. It fires for
ggml_glu_split with packed operands; a fused [gate|up] tensor keeps the
strided path. test-backend-ops perf -o SWIGLU on an Arc Pro B70: split
+14% f16 and +4% f32, fused unchanged.
- In MSL, declaring an array of matrix types like `threadgroup half4x4` causes
a 'no matching constructor' compilation error because MSL matrix types do not
have zero-argument default constructors and threadgroup variables cannot have
initializers.
- Fix this by declaring a POD `threadgroup half` array instead and casting
to `threadgroup half4x4 *` for matrix indexing.
Signed-off-by: JamePeng <jame_peng@sina.com>
* ci: prepare for amd rocm ci
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: fix editorconfig-checker
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: fix device not recognised
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: rename gpu-amd to gpu-hip
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: gpu-hip to gpu-rocm
haha
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* CUDA: allow integrated-GPU host output buffer in debug assert
On integrated GPUs (APUs), the scheduler can legitimately place a graph
node's output on the host-visible buffer, which ggml_cuda_compute_forward
already handles. The debug assert in ggml_cuda_graph_evaluate_and_capture
required every node output to be on the device buffer, so a debug build
aborts on such a node (e.g. attn_residual ADD -> ROCm_Host on RDNA3.5).
The source-tensor assert directly below already permits this via the
integrated + cuda_host exception; apply the same exception to the node's
own output buffer. Debug-only; no effect on release/compute.
Fixes test-recurrent-state-rollback on gfx1151 (Strix Halo).
* ci: enable unified memory for ROCm gfx1151 job
Work around a coherence issue on integrated RDNA3.5 (gfx1151) where GPU
kernels reading mmap-loaded weights can return incorrect output, which
makes test-llama-archs (and real inference) intermittently wrong.
GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 uses managed memory, which restores
coherence. Remove once the underlying ROCm/HIP issue is fixed.
* test-llama-archs: skip jamba on HIP backend
jamba produces incorrect output (~0.55 NMSE vs CPU) on the HIP backend on
RDNA3.5 (gfx1151); the SSM kernels need separate investigation. Skip it
for now, matching the existing per-backend carve-outs (WebGPU), so the
ROCm CI can run the test for the remaining architectures.
* ci: use HIP_LAUNCH_BLOCKING for ROCm gfx1151 job
The gfx1151 ROCm CI job produced incorrect inference output (qwen3 perplexity ~88 vs ~9.4) due to an async-execution correctness issue in the HIP path. Serializing kernel launches with HIP_LAUNCH_BLOCKING=1 restores correctness. This replaces the earlier GGML_CUDA_ENABLE_UNIFIED_MEMORY workaround, which did not fix batched inference.
* test-backend-sampler: skip top-k subtests on HIP backend
The ROCm backend does not support the TOP_K/ARGSORT op at vocab scale (no CUB; bitonic argsort is capped at ncols <= 1024), so top-k/top-p backend samplers cannot be offloaded. The penalties, set_sampler, mixed, and top_p subtests assert that offload happened, so they fail on HIP. Skip them until TOP_K is supported on the ROCm backend.
* Update tests/test-backend-sampler.cpp
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
* Update tests/test-backend-sampler.cpp
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Jim Wu <ywu@xilinx.com>
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
This commit adds the --model-name flag to the causual and embedding
model conversion scripts.
The motivation for this is that this is the name used for the metadata
field general.name and it can be useful to specify this explicitely if
the default (the basename of the model path) is not what we want.
* vulkan: add debug tooling to get more information about a DeviceLost error
* fix submission threshold applied too late
* use logging macros, throw instead of aborting
* clean up circular dependency
* ggml: add ggml_build_forward_order
ggml_build_forward_expand marks the tensor and all its ancestors for
compute, so using it as a pure ordering hint (keeping q, k and v
together) defeats ggml_build_forward_select: the unselected branch is
forced to run with inputs that were never uploaded. In the mtmd audio
graph this makes GEN_WAV calls execute the GEN_CODE branch with a
stale inp_code0, hitting the get_rows bound assert on CPU.
Add ggml_build_forward_order, which inserts nodes without the compute
flag; the flag is restored when the branch is actually selected.
Switch the q/k/v hints in clip_graph::build_attn to it.
* nit: reduce comments (AGENTS.md)
* server: don't walk Windows junctions in file_glob_search
std::filesystem reports a junction as a plain directory, so the symlink
guard misses it and a junction pointing back at an ancestor is walked
until the path length gives out
read the reparse tag and treat a symlink and a mount point as links,
leaving any other reparse point walkable so cloud placeholders and dedup
stubs still get searched
look junk directory names up case insensitively on Windows, where NTFS
makes Build the same directory as build
test that a junk directory stays selectable while its contents stay out
of search results
* server: report a directory the walk could not read
a directory that fails to open or to iterate was skipped in silence, so
a caller got a listing that looked complete while a whole subtree was
missing: a path over the platform limit, a volume going away, a name the
filesystem rejects
skip_permission_denied never reaches this path, so an error here is an
incomplete answer rather than a deliberate omission, and it now sets the
truncated flag
* server: simplify the file_glob_search listing plumbing
return a small result struct instead of two out params and a caller path
that only fed an error string, taking list_entries from six parameters
down to three
scope the error code to the directory being read, act on the status code
the entry lookups already returned, and treat an unreadable link state as
a link so the walk never descends on a guess
check the deadline when a directory is popped, not only per entry, so a
tree of empty directories cannot outlive the budget
read the path parameter once, and reject an invalid limit the way an
invalid type is already rejected, instead of silently falling back
normalize the resolved path, so a "." or ".." a caller typed reaches
neither git nor the client, and return the generic path form with '/'
separators on every platform, so the base sent to clients no longer needs
a local fixup
* ui: expire cached picker searches
the cache grew for the lifetime of the component: entries went stale
after the TTL but were never removed, so every distinct query typed in a
session stayed in memory
drop expired entries when a new result is stored
* server: address review from @ngxson
trim comments to one line each, and drop two that restate the code
rename junk_lookup_name to get_effective_name, and move it and the link
check to private static members next to junk_dir_names
merge the Windows and Linux link checks into one is_link, so symlinks are
checked everywhere and junctions only add to it on Windows
* server: convert tool paths as UTF-8 on Windows
a narrow path uses the active code page there, so a file name came back
mangled and a path with an accent could not be opened at all
convert explicitly at every crossing between a std::string, which always
carries UTF-8 here, and fs::path
read the home directory through the wide environment, since the narrow
one returns the profile path in the active code page too
the walker no longer normalizes separators by hand, since paths now come
back in generic form
* server: fold the platform branch inside console_output_to_utf8
match the shape of the other helpers, one definition with the #if inside,
instead of two definitions wrapped in #if and #else
inline the single caller helper and trim the comment
* mtmd: support multi-row batching for deepseek-ocr
* mtmd: weave deepseek-ocr rows in one shot instead of per row (#26615)
---------
Co-authored-by: Saba Fallah <sabafallah@gmail.com>
