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>
a child process writes in the OEM code page, which is not UTF-8 on a
western Windows install, so accented output reaches the JSON layer as
invalid bytes and gets replaced there, silently losing the characters
run() spawns without a console, so the child never inherits the console
code page and GetOEMCP is the one that applies
decode with MB_ERR_INVALID_CHARS so a wrong code page returns the text
untouched instead of emitting replacement characters, and pass text that
already decodes as UTF-8 through so a child emitting UTF-8 is never
decoded twice
the check drops an incomplete trailing sequence before validating, since
a streamed chunk can end in the middle of a multi-byte character
* Resolve -1 to 1024 instead of ctx-len for samplers
Because of backend-sampling we initialize samplers before the complete
llama_context is there. Therefore, we cannot infer the resolved context
length yet at the time we construct the samplers.
* Shared default of 64 for history-based samplers, remove context_size
