The two T4 runners in the ggml-ci group (ggml-6-x86-vulkan-t4,
ggml-8-x86-nvidia-a10) share the same labels and can both pick up the
same jobs, so extend the per-runner results dir from #31046 to them.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
Add test-speculative-rejection-sampling to the ctest -E exclusion list of the macos-latest-arm64 CI job.
Assisted-by: dsh:llama.cpp/ggml-org/Flint-GGUF (RTX 5090 32GB, DDR5 128GB)
* ci : fix data race between runners sharing the model cache
Two runners per mac mini share HOME, so concurrent jobs corrupt each
other downloads, converted models and venv under the shared MNT cache.
- gg_wget: download in a temp dir, publish with mv, so readers only
ever see complete files (-N cannot be combined with -O)
- convert, quantize and wikitext unzip now produce under a per-runner
scratch dir in OUT, published to MNT via hardlink when missing
- gguf-split and quantize tests get a per-runner TMP_DIR
- python venv moved from MNT to OUT
- metal, vulkan and webgpu mac jobs pass a per-runner OUT path
Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)
* ci : clean up scratch models after tests
Assisted-by: dsh:llama.cpp/agent-qwen4 (RTX 5090 machine)
* ci : fix q4_0 publish when a cached model exists
cp -n skips an existing file with exit code 1, so the ln fallback
aborted the qwen3 job under set -e. Publish through a unique temp
file with an atomic mv: readers never see a partial file and a
stale cached q4_0 gets replaced.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)
* ci : cache-first gg_wget to avoid re-downloads
HF sends no Last-Modified, so wget -N re-downloads the full files on
every run. Skip the download when a non-empty cache file exists: the
HF CI artifacts are static. Downloads still go through a temp dir and
an atomic mv, so concurrent runners never read partial files.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 machine)
* hex-build: bump toolchain, enable v85
* hex-build: bump toolchain to v0.8 in the CI and build scripts
* hex-build: add a little python script for swiv needed on some iot devices
* hex-build: fix flake8 warnings
* hex-arch: update version cap to v85
* hex-inspect: flexible TOOLS_ROOT
* hex-build: update to snake_case naming
* hex-build: update leftover old toolchain refs
* hex-target: update build and run scripts to support ubuntu with adb transport
* init
* wip
* wire it up
* make it clear that it's experimental
* add props.connect_code
* nits
* fix windows build
* server: do not advertise a connect code without the tunnel
--connect-code on its own had no effect other than being served in
/props, so the Web UI showed a remote access code that no tunnel was
serving. Clear it when --connect is off, and say so once at startup.
The teardown comment claimed that terminate() closes the child stdout.
It does not: the log thread leaves fgets() only once every writer of
that pipe is gone, so a descendant of llama-connect would block the
join forever. Describe the real contract instead.
Also restore the --reasoning-preserve default in the generated docs.
* Update tools/connect/CMakeLists.txt
Co-authored-by: Adrien Gallouët <angt@huggingface.co>
* server: dial llama-connect on the bound port and a reachable host
Pass ctx_http.port so --port 0 and the GCP override work, pick the
dial host from --host (wildcards map to loopback, IPv6 is bracketed)
and refuse --connect when only unix sockets are bound.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
Co-authored-by: Adrien Gallouët <angt@huggingface.co>
* enable OPENSSL=ON as default in build script, update it in CI and dockerfile
* change to LLAMA_BUILD_BORINGSSL on windows
* refactor
* rm Download UI build
* Apply suggestion from @CISC
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Apply suggestion from @CISC
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Apply suggestion from @CISC
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* replace oneAPI package by DLE+oneDNN package
* support to install oneDNN in windows, in release and docker
* for internal test
* remove test code
* Update intel.Dockerfile
* tests : consolidate the context state tests into test-llama-context
test-save-load-state, test-recurrent-state-rollback and
test-state-restore-fragmented check the same object from three places:
the KV cache and the recurrent state of a llama_context, saved, loaded,
copied, rolled back. Each had its own main, its own loop over the
models, its own copy of the token and the comparison helpers, and a run
of all three loaded the same models several times.
They are one binary now, and the case is the unit of the test:
tests/llama-context/test-case.h the interface a case implements
tests/llama-context/cases-kv.cpp the cases of the KV cache
tests/llama-context/cases-rs.cpp the cases of the recurrent state
tests/llama-context/test-util.cpp tokens, generation, comparison
tests/llama-context/test-runner.cpp the models of a case, on threads
tests/llama-context/test-log.cpp the log file of a model
A run is a table: a row per model, a column per case, a cell holds the
verdict of that pair. The models load on several threads (-j, --threads,
min(cores, 4) by default), the table is printed when it is complete and
the log of a model that failed is printed after it, in the order of the
rows. A case returns a verdict and the reason for it, and the runner is
the only one who reports - the test does not use tests/testing.h, a run
over 130 models is a matrix of verdicts and not a tree of subtests.
Staging no longer travels through the environment: llama_model_set_attn_rot
replaces the LLAMA_ATTN_ROT_DISABLE variable, which a model loaded in one
thread used to change for the models the other threads were loading, and
llama_memory_get_buffers is added for the cases that copy the buffers of
a memory. common_log_set_stdo keeps the console quiet while the table is
being built.
ggml_aligned_malloc allocates with posix_memalign when the thread
sanitizer is on. On macOS it allocates with the Mach VM, which the
sanitizer does not intercept, so it never sees a page range change hands
and reports the first write into recycled memory as a race - 30 of them
in this test, and every one of them false.
Assisted-by: pi:llama.cpp/Qwen3.8-Flash-Next
* cont : run single-threaded by default
* cont : try to build the tests on all platforms
* cont : trim down some body fat
* cont : fix
* cont : fix the fix
* ggml-openvino : Qwen3.5 MoE perf (#312)
Squash of ravi9/llama.cpp#312:
- ggml-openvino: add detailed inference profiling (Yu, Zijun)
- ggml-openvino: use remote output tensors by default (Yu, Zijun)
- ggml-openvino: optimize single-sequence recurrent state (Yu, Zijun)
- opt1: remove recurrent reset for single sequence, opt2: direct gdn outputs (break parallel sequence) (Yu, Zijun)
- fix parallel sequences (Yu, Zijun)
- ggml-openvino: simplify graph cache key (ynimmaga)
- enable stateful for qwen35 single sequence (Yu, Zijun)
- Fix after rebasing (Yu, Zijun)
- Add k-requant option q4_asym64 (Yu, Zijun)
- Fix qwen35 llama-bench -p 0 (Yu, Zijun)
- Simplify RESHAPE translation (Yu, Zijun)
- openvino: fuse MoE routing (Yu, Zijun)
- openvino: fuse GDN qk normalization (Yu, Zijun)
- openvino: enable GPU MoE fusion by default (Yu, Zijun)
- ggml-openvino: add cache_only mode to import cached compiled model on disk directly (Yu, Zijun)
- openvino : report the device allocation limit to ggml (Łukasz Ślusarczyk)
- Fix windows build (Yu, Zijun)
Co-authored-by: ynimmaga <ynimmaga@users.noreply.github.com>
Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
* ggml-openvino: Update doc of compiled model cache
* openvino: implement PRD-compliant device enumeration and memory reporting
* openvino: fix multi-device listing issues from review
- Only the device selected by GGML_OPENVINO_DEVICE reports as GPU; the
other OpenVINO devices report as IGPU so llama.cpp does not offload to
them. Initializing a non-selected device logs a warning.
- Name devices OPENVINO<i> again and show the OpenVINO id in the
description. Raw "CPU" names shadowed the ggml CPU backend.
- Support GPU.N: create the OpenCL queue on OpenVINO's own context for
the selected device, and replace "GPU"/"NPU" string comparisons with
ggml_openvino_is_gpu()/ggml_openvino_is_npu().
- An unavailable GGML_OPENVINO_DEVICE is now an error that lists the
available devices, instead of silently falling back to CPU.
- Memory: cap iGPU/NPU free memory at system available memory, fall back
to system memory instead of 0/0 when the plugin lacks memory
properties, and ignore host USM allocations in GPU usage.
- Initialize the device config once under a lock, even if OpenCL setup
fails.
- Fix supports_op return type for non-selected devices (build error).
* openvino : take USM entry points from the selected device platform
clGetExtensionFunctionAddressForPlatform was called on the first platform
returned by clGetPlatformIDs. The address it returns is only valid for the
platform it was queried on, and the first platform is not always the one that
holds the device OpenVINO selected.
On a host whose first platform comes from another vendor the lookup returns
null, and then every read, write and memset on a GPU buffer fails with
"clEnqueueMemcpyINTEL not available".
Look both entry points up in init(), on the platform of the device OpenVINO
picked, and keep them in the device config next to the command queue.
Assisted-by: Claude Opus 5
* openvino: fuse MoE experts for models with a fused gate_up weight
FuseMoeCompressed only matches models whose gate and up projections are
separate GatherMatmul ops. gemma-4 packs both into one expert weight and
splits the result after the GEMM, so its MoE block stayed unfused and ran
the expert GEMMs as per-token GEMVs.
Add FuseMoeCompressedFusedGateUp, which matches that shape
(one GatherMatmul -> Slice/Slice -> Gelu(ERF) -> Multiply) and folds it into
the same MOECompressed op, using GEMM3_SWIGLU with GEGLU_ERF. The fused
weight, scale and zero point are split into gate/up halves by copying raw
bytes, since a graph Slice would be rewritten to StridedSlice and constant
folded, whose reference evaluator crashes on sub-byte types.
gemma-4 also applies a per-expert output scale to the down projection before
the router weights. MOECompressed takes only one per-expert weight, so that
scale is folded into the routing weights, which is exact.
The op reads the zero point straight off a weight port and needs an integer
Constant there, so the matcher requires one and leaves natively quantized
experts (exact f16 zp) to the unfused path.
gemma-4-26B-A4B on Arc B390, GGML_OPENVINO_REQUANT_KQUANT=q4_asym64_all,
llama-bench -p 512 -n 128 -r 2, against a GGML_OPENVINO_MOE_OP=0 baseline:
pp512 66.16 -> 1608.73 t/s, tg128 25.94 -> 26.46 t/s. Perplexity over 12
chunks is unchanged (1451.3 +/- 177.9 unfused vs 1427.6 +/- 175.1 fused).
No effect without that requant option, on models with separate gate/up
weights, or on CPU. test-backend-ops -b OPENVINO0 is unchanged by this
commit: two MUL_MAT_ID m_v cases fail, the same two on the unmodified base.
* openvino: fix rank-3 axis handling so MoE works under stateful execution
Stateful execution drops the leading size-1 batch dim, so OV tensors are rank
3 while GgmlOvDecoder::get_shape/get_stride still report GGML_MAX_DIMS=4
reversed entries. Several MoE ops derive OV axis indices straight from that
metadata, so they picked the wrong axis. A MoE model with
GGML_OPENVINO_STATEFUL_EXECUTION=1 aborts while building the graph:
Check 'is_axis_valid(axis, r)' failed at src/core/src/validation_util.cpp:336
While validating node 'opset11::TopK ... _ffn_moe_probs ...'
Axis 3 out of the tensor rank range [-3, 2].
Fix idiom throughout: take the axis from the real OV rank, or shift a
metadata-derived axis down by metadata_rank - actual_rank.
argsort.cpp the router top-k axis is 2 on rank 3, not 3. This is the
abort quoted above.
add.cpp the MoE expert-sum bypass collapses the 8-ADD chain into one
ReduceSum on hardcoded axis 2, which on rank 3 reduces n_embd
instead of the expert axis. Now rank-2, with the following
Unsqueeze at rank-3.
get_rows.cpp squeezing a hardcoded {0,1} also strips the batch dim
whenever it is 1, which is every decode step. Squeeze down to
the trailing two dims instead.
mul_mat_id.cpp pick the reshape dims by actual rank, and skip the trailing
Unsqueeze that re-adds the batch dim.
view.cpp the expert-plane slice had the Slice axis, dst_ov_axis, the
ShapeOf+Gather index and the Reshape target all rank-4.
utils.cpp process_view_input_new's "translate_view already resolved
this VIEW, skip re-slicing" shortcut required equal ranks. 4
vs 3 never matched, so every resolved expert plane got
re-sliced. Now compares the common trailing dims. Same axis
shift for the Slice in the view-chain walker.
Stateless is unchanged by construction: every edit is gated on the actual
rank, so axis_shift == 0 reproduces the previous code exactly. Checked on
OV-CPU by diffing greedy output against the unmodified base for dense
gemma-4-E2B, granite-1b-a400m and gemma-4-26B-A4B; all identical.
granite-1b-a400m on OV-CPU aborts with the error above before this change;
after it, it generates and is byte-identical to stateless. Dense gemma-4-E2B
is identical stateless vs stateful both before and after. test-backend-ops
-b OPENVINO0 is unchanged: two pre-existing MUL_MAT_ID m_v cases fail, the
same two on the unmodified base.
gemma-4-26B-A4B is a poor correctness vehicle here. On OV it already drifts
into degenerate repetition a few tokens in, in stateless as much as stateful,
and the two modes diverge somewhere inside that degenerate region instead of
matching token for token. Each mode is self-reproducible across runs.
Known limitation: FuseMoeCompressedFusedGateUp does not match the rank-3
graph, so a MoE model run with GGML_OPENVINO_STATEFUL_EXECUTION=1 loses the
prefill fusion while gaining decode. gemma-4-26B-A4B on Arc B390,
GGML_OPENVINO_REQUANT_KQUANT=q4_asym64_all, llama-bench -p 512 -n 128 -r 2:
unfused (GGML_OPENVINO_MOE_OP=0) pp512 66.16 tg128 25.94
fused, stateless (default) pp512 1608.73 tg128 26.46
fused, stateful pp512 66.18 tg128 29.91
Stateful is opt-in and off by default, and MoE did not run there at all
before this, so nothing that previously worked regresses. Making the pass
match rank 3 is the follow-up.
* OpenVINO Backend: Upgrade graph cache to use node_idx, src_idx, node type
* ggml-openvino : enable more comprehensive conv fusion
* enable conv ops
* Reject kernel size 0 and support IM2COL_3D
* openvino : abort when the GPU remote context cannot be created
init() logged the error and returned, which left the device name a GPU but
remote_context empty. The remote buffer and tensor paths assert only on the
device being a GPU and then dereference that empty optional.
Those paths have no host fallback, and a device that OpenVINO listed should
have a working OpenCL context, so stop instead of continuing. An OpenCL stack
that is broken as a whole is still caught earlier by the device availability
check, which falls back to CPU.
Assisted-by: Claude Opus 5
* openvino : fix build warnings
The single-argument form of the OpenVINO RTTI macros is the intended one, but
their selector macro leaves __VA_ARGS__ empty, which -Wpedantic reports on
every pass and op header. Turn that warning off for this backend only, the
way ggml-cuda and ggml-sycl already do for their own third-party warnings.
Also drop a break and a dead assignment around a GGML_ABORT, which is noreturn.
Assisted-by: Claude Opus 5
* OpenVINO Backend: Support common MTMD ops
* ggml-openvino: give a reshaping view its own ov::Tensor
* ggml-openvino : compute HARDSIGMOID and EXPM1 in f32
HARDSIGMOID used a 1/6 constant in the input type, which is not exact
in bf16, and EXPM1 lost precision for small inputs in f16. Both now
compute in f32 and convert back, except on NPU where the f32 path
gives wrong results.
Fixes the HARDSIGMOID/EXPM1 test-backend-ops failures on GPU.
* ggml-openvino : update device selection and --list-devices
Show the selecting GGML_OPENVINO_DEVICE value and active device in
--list-devices, startup logs, and backend tests.
Clarify OpenVINO selection uses GGML_OPENVINO_DEVICE, not -dev.
* openvino : remove unreachable OpenCL queue checks
A remote buffer exists only on a GPU device, and init() aborts there if the
queue cannot be created, so the queue is never null at these call sites.
Assisted-by: Claude Opus 5
* openvino : update OpenVINO to 2026.4.1 and GPU drivers to 26.35.39758.10
* docs : update OpenVINO validated models and GPU driver version
* ggml-openvino : skip empty views when giving a reshaping view its own tensor
A zero-size view can sit at the end of a GPU USM buffer (Qwen3.5 recurrent cache). Wrapping it as a remote tensor throws "shared USM buffer has smaller size (0)".
Assisted-by: Claude
* ggml-openvino : rebind the cached decoder when llama passes a different graph
llama keeps separate graphs for batches with and without outputs. llama-server splits the prompt into chunks for context checkpoints, so a cached decoder could be reused with a graph built in other memory and bind the previous chunk's input tensors. SWA and recurrent models then lost most of the prompt in llama-cli and llama-server.
Assisted-by: Claude
* docs : update OpenVINO validated models
Smoke test on Lunar Lake (32 GB) with the two fixes above. Re-add the Qwen3.5 and gemma models.
Assisted-by: Claude
---------
Co-authored-by: Yu, Zijun <zijun.yu@intel.com>
Co-authored-by: ynimmaga <ynimmaga@users.noreply.github.com>
Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
Co-authored-by: haarika-madaka <haarika.madaka@intel.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Pinning to >= 3.4.3 is required to enable DeviceTopK, which was affected by
a race condition https://github.com/NVIDIA/cccl/pull/10627.
We will relax this for future CTK versions which will bundle CCCL >
3.4.X (CTK 13.5 will bundle CCCL 3.5.0 for example)
* musa: build the docker image and CI container from the MUSA SDK images
Use registry.mthreads.com/mcconline/musa_sdk:5.2.0-{devel,runtime}-ubuntu22.04-s5000
instead of registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64
for the MUSA docker image and the MUSA CI container, and let the runtime stage use the
runtime image instead of reusing the devel one, which drops the MUSA toolchain from the
published images.
* musa: install the MUSA headers and loader path the SDK images omit
musa_sdk:5.2.0-*-s5000 does not ship the cub and thrust headers that the MUSA
backend builds against, and its runtime image does not register
/usr/local/musa/lib with the dynamic loader.
Install both header packages in the build stage and in the MUSA CI container,
and write the loader path in the runtime stage.
* musa: install libmthreads-compute for the MUSA runtime library
The MUSA SDK images do not install libmthreads-compute, which provides
libmusa.so.1 in /usr/lib/x86_64-linux-gnu, so linking anything against the
MUSA backend fails.
* musa: install libmthreads-compute in the runtime stages
The MUSA runtime image does not install libmthreads-compute, so the published
images would have no libmusa.so.1 at run time.
---------
Co-authored-by: yeahdongcn <yeahdongcn@users.noreply.github.com>
* ci : update the oneAPI toolkit to 2026.1
oneDNN is removed from Intel Deep Learning Essentials in 2026.0, so
staying on the deep-learning-essentials path would silently lose oneDNN
support when the toolkit version is updated. Switch both the Ubuntu and
Windows CI jobs to the new unified Intel oneAPI Toolkit installer,
which still includes oneDNN (until 2027.0) and keeps the component IDs
unchanged for the Windows install script.
Measured with the same code (b10899) built with oneAPI 2026.1 vs the
2025.3-based release build on Arc B570: prompt processing 1331 vs 434
t/s (3.1x), token generation 50.1 vs 45.3-48.0 t/s.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* docs : update the SYCL backend build requirements for oneAPI 2026.1
With the 2026.0 release the Base toolkit and the HPC toolkit are
combined into the oneAPI Toolkit, and oneDNN is removed from the Deep
Learning Essentials package. Update the install instructions, the
verified release table and the news section accordingly.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* ci : update the release workflow for oneAPI 2026.1 and Level Zero SDK 1.33.1
Align the release package build with the CI build update:
- oneAPI toolkit 2025.3.3 -> 2026.1 (the unified oneAPI Toolkit)
- Level Zero SDK 1.28.2 -> 1.33.1, and the Debian package names
(level-zero/level-zero-devel -> libze1/libze-dev)
- The Windows DLL copy list for the 2026.1 runtime: sycl9.dll and the
.6/.3 MKL library versions
Assisted-by: GLM (z-ai/glm-5.3-flash)
* ci : remove the removed .spv fallback files from the Windows DLL copy list
oneAPI 2026.1 no longer ships libsycl-fallback-bfloat16.spv and
libsycl-native-bfloat16.spv (the OpenCL fallback mechanism changed), so
the copy step failed with exit 1.
Assisted-by: GLM (z-ai/glm-5.3-flash)
* devops : update the oneAPI toolkit image in the Intel Dockerfile
Assisted-by: GLM (z-ai/glm-5.3-flash)
---------
Co-authored-by: Asahi-Prv <Asahi-Prv@users.noreply.github.com>
* musa: use 16-byte copies for MUSA like sm_70+
ggml_cuda_get_max_cpy_bytes() derives the copy width from __CUDA_ARCH__. mcc
never defines it, so MUSA fell into the generic branch and returned 8 bytes
instead of the 16 bytes that every sm_70+ target gets. The value sizes the
per-thread copy unit of the FlashAttention K/V staging code (fattn-common,
fattn-vec, fattn-tile, fattn-mma-f16 shared-memory loads) and of mmq-vec-dot,
so every MUSA FlashAttention kernel moved half as many bytes per instruction.
On an MTT S5000 (mp_31, MUSA SDK 5.2.0) with Qwen3.8-27B-UD-Q4_K_M, -ngl 999,
-p 512 -n 64, -fa on: 751.15 -> 794.73 t/s prefill and 15.59 -> 15.69 t/s
decode. -fa off is unchanged (1050.05 -> 1052.86 t/s prefill), FLASH_ATTN_EXT
is unchanged (3984 ok / 0 fail / 1323 unsupported) and perplexity is
unchanged.
* musa: enable the CUB paths on MUSA
GGML_CUDA_USE_CUB and USE_CUB are selected by "CUDART_VERSION >= 11070", which
the MUSA SDK never satisfies: CUDART_VERSION is not defined anywhere under
/usr/local/musa/include, so the condition is always false and every CUB-based
path stayed compiled out on MUSA even though the SDK ships CUB and the kernels
build for mp_31. Select them from GGML_USE_MUSA as well. The device-wide
algorithms are usable too: cub::DeviceSegmentedSort compiles and produces
correct results on mp_31.
This lifts the ne[0] <= 1024 limit that ggml_backend_cuda_device_supports_op
applied to ARGSORT and TOP_K on MUSA. On an MTT S5000 (S5000, mcc 5.2.0):
ARGSORT 48 ok / 52 not supported -> 100 ok / 0 (CUDA parity), TOP_K 0 ok /
354 not supported -> 527 ok / 0. The other 20 per-op suites are unchanged, the
Qwen3-0.6B f16 (14.4679) and Qwen3.8-27B iq4_nl (5.1724) perplexities are
unchanged, and the 0.6B graph keeps the same nodes and splits (18 CPU + 18
MUSA0, SET_ROWS 1008) as before.
* musa: take the upstream code path where the toolkit supports it
Several guards were written for an older MUSA toolkit. Verified against MUSA SDK
5.2.0 and on an MTT S5000 (mp_31):
- device init: query cudaDevAttrCooperativeLaunch instead of hardcoding false.
The device reports cooperativeLaunch=1 and musaLaunchCooperativeKernel works
(verified with a kernel whose result was checked).
- device init: keep prop.warpSize instead of overriding it with 32. The device
reports 32 anyway, so this only removes the divergence.
- CUDA_SET_SHARED_MEMORY_LIMIT and the FA shared-memory raise: musaFuncSetAttribute
returns success and sharedMemPerBlockOptin is 192 KiB, so the kernels can use
more than the default 48 KiB.
- vendors/musa.h: add the cudaDeviceGetAttribute and cudaDevAttrCooperativeLaunch
mappings the device-init change needs.
Measured on one S5000 with Qwen3.8-27B Q4_K_M (-ngl 999, -r 3): pp512 968.27 ->
957.09 t/s, tg64 10.09 -> 10.23 t/s, FLASH_ATTN_EXT sweep identical (3975/3982
both), perplexity identical (80.2841 +/- 7.26772 both).
* musa: drop compile-time guards that MUSA's runtime gates already cover
mcc never defines __CUDA_ARCH__, so the arch-gated fallbacks in this group
were already taken on MUSA and the GGML_USE_MUSA guards on top of them only
kept the upstream text from being compiled:
- wkv.cu: the "#pragma unroll" suppression has no effect on the generated
code that is not already covered by the surrounding guards
- common.cuh: the MUSA-only __builtin_unreachable() in no_device_code() is
not needed to silence the compiler
- ssm-scan.cu: the SSD (Mamba-2 prefill) block and its dispatch are gated at
runtime by GGML_CUDA_CC_IS_NVIDIA(cc) and turing_mma_available(cc), which
are both false for PH1 (cc 0x100310), so compiling them changes nothing
- common.cuh: warp_reduce_max(half2) is guarded the same way as
warp_reduce_sum(half2) (FP16_AVAILABLE); the MUSA-only guard left the
function with no return statement. It has no caller today.
MTT S5000 (mp_31, MUSA SDK 5.2.0), MUSA_ARCHITECTURES=31: build rc=0. Against
an unmodified build of the same tree on the same card, FLASH_ATTN_EXT
(3984 ok / 0 fail / 1323 unsupported), SSM_SCAN (15/0), RWKV_WKV6 (6/0),
GATED_DELTA_NET (38/0) and MUL_MAT (1299/0/385 unsupported) are identical, and
perplexity with -fa on is bit-identical (5.1639 +/- 0.36673, 4 chunks).
* musa: do not use MMQ on PH1
test-backend-ops on an MTT S5000 (mp_31, MUSA SDK 5.2.0) fails 260 cases and every
one of them goes through the MMQ path:
- MUL_MAT with a batched src1 (any bs/nr != [1,1]): 109 cases across all
quantized types, e.g. 12 of 13 cases at n=16, while the plain [1,1] layout
passes
- every quantized MUL_MAT_ID: 147 cases, while the f16/f32 variants of the same
shapes pass
- MUL_MAT with more than ~512 tokens: 4 cases (n=509..4096); the small-n cases pass
The cuBLAS/dequant path is correct for all of them and the MMVQ path used for
small batches is unaffected, so quantized matmuls now take that path on PH1
instead of returning wrong values. 27B perplexity with default flags goes from
nan to finite, and the full suite reports 0 failures out of 22237 cases.
The MMQ defect itself (fastdiv, __umulhi, uint3 kernel parameters and
__CUDA_ARCH__-based MMA availability were all checked and are correct on this
part) is not addressed here.
* musa: keep the block barrier of the fused TOPK_MOE kernel reachable
topk_moe_cuda returns early for the rows past the end of the graph, but one block
covers TOPK_MOE_ROWS_PER_BLOCK (8) rows, so the last block is only partially filled
whenever n_rows is not a multiple of 8. On MUSA a warp that has already returned
blocks the block wide __syncthreads() below, which makes the kernel hang and the
launch time out. CUDA tolerates the exited warps, which is why the CUDA numbers
never showed it.
For MUSA, clamp the row index of those warps to the last row so that every warp of
the block reaches the barrier; they recompute the last row and write the same
values. The CUDA code path is unchanged.
On an MTT S5000 (mp_31) the fused TOPK_MOE cases change from a launch timeout with
no completed case to 418 ok / 0 not supported / 0 failed, i.e. the CUDA result, and
the other 101 per op suites are unchanged (0 failed, no count changes).
* musa: enable GATED_DELTA_NET
The op was turned off for every MUSA target because mcc could not build the kernel
at the time. The current toolkit builds it: with mp_31 and MUSA SDK 5.2.0 the file
compiles with zero errors and all 36 test-backend-ops GATED_DELTA_NET cases pass
against the CPU reference. 27B perplexity is unchanged.
While the op is refused, the scheduler has no choice but to run it on the CPU: 48
GATED_DELTA_NET nodes per forward pass. On an MTT S5000 (Qwen3.8-27B Q4_K_M, -ngl
999, one container, -r 3):
pp512 (FA off) 964.51 -> 2119.26 t/s
tg64 (FA off) 10.15 -> 15.50 t/s
* musa: name the stream capture query API for the graph aware kernels
argsort.cu and mean.cu call cudaStreamCaptureStatus, cudaStreamIsCapturing and
cudaStreamCaptureStatusNone inside their USE_CUDA_GRAPH blocks, but the MUSA
compatibility headers do not alias those names, so building with the experimental
GGML_MUSA_GRAPHS option fails with 7 errors in those two files. Map the three
names to their musa* counterparts, under the same guard that enables the graph
code, so the default build is untouched.
The option stays off by default: on an MTT S5000 the captured path measured
slower (pp512 693 vs 772 t/s, tg128 15.20 vs 15.39 t/s over two sessions) and the
borderline MUL_MAT cases are not reproducible between runs.
* musa: build the CI and docs for PH1 (MTT S5000)
The MUSA CI job and the documented default still targeted the first generation
(MTT S80, MUSA_ARCHITECTURES=21) while the current MUSA SDK targets PH1
(MTT S5000, 31). Move the job, ci/run.sh's default and the build docs to 31,
and run the job in the PH1 MUSA SDK devel image:
registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64
That image needs two things the previous one did not: python3-venv for the
ccache-buckets step, which builds a virtual environment for the Hugging Face
CLI, and no time prefix on the build command, because container jobs run their
steps with sh and the image ships no time binary.
Sreeram