ggml-cpu : fix missing Q5_0 dispatch in SpaceMiT backend (#26792)

ci: rm GGML_HIP_ROCWMMA_FATTN (#26760)

Signed-off-by: Aaron Teo aaron.teo1@ibm.com

server: report the isolate working directory from get_info (#26773)

• server: report the isolate working directory from get_info

Without an explicit cwd, get_info fell back to the server process
working directory even when a tools runtime was configured. That named a
host path no tool would ever run in, since an isolate starts in a
directory of its own.

It now asks the isolate for its working directory in that case, and
keeps the process one only when the tools run on the host.

• remove redundant comment

Co-authored-by: Xuan-Son Nguyen thichthat@gmail.com

CUDA: fuse rms_norm + mul + rope (+ view + set_rows) (#26767)

• CUDA: fuse rms_norm + mul + rope (+ view + set_rows)

• tests: add broadcast weight case to rms_norm_mul_rope

• CUDA: check memory ranges before rms_norm rope fusion

• CUDA: check memory ranges in rope set_rows fusion

server, ui: only offer a working directory when a tool reads it (#26762)

The working directory chip showed up as soon as the server exposed any
builtin tool, so a server started with just get_datetime, or a user who
turned every filesystem tool off in the settings, still got a control
that nothing would read.

Tools now declare whether they resolve their paths and run against the
working directory, next to the write permission they already publish in
the /tools listing. The WebUI shows the chip and enables the /cwd
command only when at least one such tool is both served and left
enabled.

server: add initial tool isolation support (via docker) (#26507)

• server: add initial tool isolation support (via docker)

• add docs

• adapt get_info

• py: fix type check

• cont

• separate tools_io_sandbox / tools_io_docker

• rename sandbox --> isolate

• x-tool-docker --> x-tool-runtime

Co-authored-by: Pascal admin@serveurperso.com

CUDA: fix thread/block count in quantized cpy kernel launches (#26731)

• CUDA: fix thread/block count in quantized cpy kernel launches

• tests: add uneven block count cpy case

tts: account for the vocoder pass in the timings line (#26733)

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.

sycl: coalesce the ssm_conv window loads (#26612)

test-backend-ops perf -o SSM_CONV on an Arc Pro B70, interleaved A/B against
master, 6 reps, us/run:

ne_a=[515,3328,1,1] ne_b=[4,3328,1,1] n_t=512 97.68 -> 52.95 1.85x
ne_a=[937,8192,1,1] ne_b=[4,8192,1,1] n_t=934 516.16 -> 276.13 1.87x
ne_a=[4,3328,1,1] ne_b=[4,3328,1,1] n_t=1 2.73 -> 2.71 flat

llama-bench on qwen35 27B Q4_K - Medium (48 of its 64 blocks run ssm_conv),
-ngl 99 -fa 1 -ctk f16 -ctv f16, interleaved passes of r=3:

-b 2048 -ub 2048 pp2048 1045.1 / 1043.5 / 1043.7 -> 1069.5 / 1066.3 / 1065.9 +2.2%
-b 2048 -ub 512 pp2048 771.8 / 772.7 -> 785.5 / 786.6 +1.8%
-b 2048 -ub 512 tg128 23.81 / 23.88 -> 23.87 / 23.86 flat

metal : fix NORM/RMS_NORM for row lengths that leave a partial simdgroup (#26708)

ggml_metal_op_norm sized the threadgroup with
nth = std::min(nth, args.ne00_t), which can leave nth not a multiple of
the simdgroup size. The kernels finish their row reduction with a
cross-simdgroup step where each lane of the last simdgroup reads one
per-simdgroup partial sum out of shmem_f32:

if (tiisg == 0) { shmem_f32[sgitg] = sumf; }
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);

When the last simdgroup is partial it has fewer lanes than the
threadgroup has simdgroups, so the tail of the partial sums is never
read and the row sum is too small. For ne00_t = 33 nth becomes 33: two
simdgroups, but only one lane in the second, so one of the two partial
sums is dropped. The mean and variance are then wrong for the whole row.

Round ne00_t up to a whole number of simdgroups instead. Rounding up
rather than dropping the clamp keeps the threadgroup as small as
possible: deleting the line would raise nth to the next power of two
(ne00_t = 544 -> 1024 instead of 544), which costs idle lanes on 26 row
lengths below 8192 that were already correct, including 1536 and 3584.

GGML_OP_NORM is affected as well as GGML_OP_RMS_NORM - both dispatch
through ggml_metal_op_norm.

No mainstream LLM hidden size hits this: ne00_t is ne00/4 on the
vectorized path, so 4096, 8192, 2048 and friends all give a multiple of
32. It is reachable from other norm shapes, e.g. 320-channel norms.

Add NORM and RMS_NORM cases for ne0 = 33, 132 and 260 across the
existing eps values. 33 exercises the scalar path and 132/260 the
vectorized one, since only those divide by 4.

Before, on M3 Pro:

test-backend-ops test -b MTL0 -o NORM        25/50
test-backend-ops test -b MTL0 -o RMS_NORM    26/51

After:

test-backend-ops test -b MTL0 -o NORM        50/50
test-backend-ops test -b MTL0 -o RMS_NORM    51/51
test-backend-ops test -b MTL0                13943/13943