* BLAS : Document AOCL-BLAS build and label the device AOCL-BLAS
* AOCL-Blas : Add an AOCL-BLAS Quick Start and drop the fixed version path
* AOCL-BLAS doc : Note on ZenDNN
LLM-jp-4.1 uses the GPT-OSS format, but its tokenizer decodes a space
after every special token and parallel tool calls are separated by
<|end|>. The GPT-OSS handler rejects this output, so add a dedicated
handler, selected by the chat_format=llm-jp-harmony-v1 declaration in
the chat template.
Assisted-by: Claude Fable 5.1
The committed docs/ops/CPU.csv is out of sync with the current
test-backend-ops suite: 11 ops with CPU support (COL2IM_1D,
MUL_MAT_HADAMARD, SWIGLU_CLAMP, MUL_MAT_W4A4/W4A8, MUL_MAT_ID_W4A4/W4A8,
DSV4_HC_COMB/PRE/POST, LIGHTNING_INDEXER) are missing entirely, and
many other ops have fewer test cases than the suite generates now.
docs/ops.md (which CI requires to match the CSVs) therefore
understates CPU support.
Regenerated with:
test-backend-ops support -b CPU --output csv > docs/ops/CPU.csv
scripts/create_ops_docs.py
Note: ADD1 now reads unsupported on CPU because ggml_add1 is
GGML_DEPRECATED and the suite no longer generates test cases for it;
the CPU implementation itself is still present.
Assisted-by: Xing
* 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>
* hex-topk: trying to improve/cleanup the pipeline
* hex-sampling: add STEP op
* hex-sampler: add SUM op
* hex-sampler: update CPY to support sampling cases
* hex-binary: add support for chunking to handle large logits
* hex-argmax: super basic version of ARGMAX
* hex-binary: support for scalars in extended buffers
* hex-binary: fix wrong indexing for dim 1 broadcasts across dim 2 slices
* hex-argsort: fix missing header
* hex-sampler: cleanup dma usage in the sampler related ops, and binary
* hex-build: disable autovectorizer, it is better to use explicit hints for critical loops
* hex-binary: fix perf regression due to is_1d fallback
* hex-ops: update supported ops
* 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.
* hex-dma64: enable support extended buffer mappings and 64bit dma
hex-dma64: expand binary ops to support more DMA scenarios
hex-dma64: add binary-ops.h
hex-dma64: add --hex-dma64 to run.py and fix minor issues
hex-dma64: update SSM_CONV to use dma with proper support for 64bit
hex-ops: remove obsolete gate for % 128 in binary ops
hex-l2: dont check weight tensors against dirty ranges
hex-dma64: most binary ops now support dma
hex-dma: use dma_addr_t instead of plain uint64_t to avoid overhead on older targets
hex-dma: update all dma users to use dma_data (instead of pointers)
hex-dma64: simplify lazy buffer mapping and clonning
hex-fusion: factor out try_fuse_common that checks for dma64 buffers
hex-bufs: minor cleanup for mmaping logic
hex-bufs: simplify buffer clonning
hex-ssm-conv: tighten gating checks and check vtcm size in kparams
hex-binary: fix incorred mod/wrap in scalar ops
hex-binary: make sure to call precompute kparams in support checks
hex-dma64: update addr handling in mm,concat,binary
hex-dma64: fixing up leftover of dma_addr_t conversion
hex-binary: redo the kernel selection again and fix regressions in MOEs
hex-binary: specialize per-type/per-op
hex-binary: vtcm-layout and per-src dma-queue
hex-dma64: update dma_push to transparently handle 64bit/extended
* hex-cpy: fix improper rebase with the fixes for cont. tensors
* hex-dma-cpy: update CPY to use safe dma rows/size limits
* hex-mmap: bump number of mmaps to 64 to allow avoid eviction in larger models
* hex-dma: add support for the secondary ring as a fallback for too-large transactions
* hex-rope: fix freq_factors access with 64bit dma
* hex-dma: audit all ops for proper use/gards for 64bit addresses
* hex-dma64: uninline glu-compute funcs to avoid register pressure due to 64bit addr math
* hex-dma64: refactor binary ops to separate dma loops
* hex-devel: add inspect script to help with dbg and analysis
* hex-dma: refactor dma-pipelines in unary-ops
* hex-dma: rewrite softmax to use dma
* hex-dma: rewrite GDN dma loops and improve HVX register usage
* hex-gdn: fuse GDN+CPY
* hex-mm: factor out HVX solver
* hex-mm: remove hvx-flat kernels, the chunked version now handles vtcm limits much better
* hex-buffs: reject huge buffer allocations that we cannot memory map
* hex-inspect: add logic to look for float promo calls
* hex-mm: reduce HVX register spills in HVX prompt kernels
* hex-bufs: do not double count buffers from tensors in the same op
* hex-roll: fix merge conflict
* hex-dma: reroute all matmul ddr kernels to new chunked dma/vtcm kernels
* hex-dev: update developer docs to include inspection for register spils and float promos
* hex-ops: forgot to add new headers
* hex-softmax: fix gpt-oss dims
* hex-dma64: cleanup dma_addr_t casts
* hex-dma64: add support for dma/vtcm for flash-atten with sinks
* hex-mm-add: fix MUL_MAT+ADD fusion with bias.weights in extended bufs
* hex-add-id: add support for dma for src1 (exp. table)
* hex-dma: imrpove v73 fallback paths
* hex-bufs: do not drop extended mappings during va defrag
* hex-scripts: fix flake8 warnings
* hex-docs: fix editor-config warnings
* hex-inspect: fix warnings from ty
* Update to openvino-2026.4
* Update OV docs
* ggml-openvino : fix clangd and MSVC warnings
* fix int to ptr cast, more internal linkage enforcement, and avoiding duplicate switch case
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
* exclude GPU/NPU failing POOL_2D case
* Fix pool case
* ggml-openvino: fix stateful decode for Gemma-4 per-layer-type head sizes
* ggml-openvino: fix MSVC narrowing error in permute
* ggml-openvino: classify sliding-window layers structurally on interleaved-SWA models
* ggml-openvino: add GGML_OPENVINO_REQUANT_KQUANT to select a 4-bit requant target
* ggml-openvino: add GGML_OPENVINO_SPILL_DIR to spill weight buffers to disk
* Stateful Performance: Added pass::KVStateSeqAxis to change KV layout
* ggml-openvino: fix stateful decode past the sliding-window size
Assisted-by: Claude Sonnet
* ggml-openvino: refuse stateful decode that cannot resume from the KV state
The stateful path seeds its KV state from ggml's cache when the decode position
is ahead of what the state holds. That only works when ggml's cache is a plain
prefix, where cell i holds position i. A sliding-window layer keeps just the last
n_swa positions and drops the rest, so past the window cell i no longer holds
position i and the seeded state is wrong.
Slicing the state to the decode position also had no bounds check, so a position
past the end surfaced as a bare ov::Exception from the ROI constructor
(llama_decode ret = -3, with no reason given at default verbosity).
Refuse both cases with a clear message instead, and refuse on the compile path
too, where a new model starts with an empty state and so can only serve a
sequence from its beginning. Reproducible with llama-bench -d, which restores a
saved sequence state rather than recomputing the depth prefill.
Assisted-by: Claude Opus 5
* ggml-openvino: use the per-layer KV head count for the stateful KV state
The stateful path reinterprets ggml's KV buffer [1, 1, seq, n_heads_kv * head_size]
as [1, seq, n_heads_kv, head_size]. The head size is already taken from the
tensor's own combined dim, because gemma-4 varies it per layer type, but the head
count still came from a model-level scalar that compute_llm_params() overwrites
per attention node, so it ended up holding whatever the last layer said.
gemma-4 varies the head count per layer too: 12B has 8 x 256 sliding layers and
1 x 512 full layers, 31B has 16 x 256 and 4 x 512. So 40 of 12B's 48 layers were
split as 1 x 2048 instead of 8 x 256, and attention read the state with the wrong
head split - both models decoded garbage on CPU and GPU. E2B is unaffected, its
head count is 1 everywhere.
Record the count per layer instead and look it up by the cache_k_l<N> leaf name.
Key it by layer, not by layer type: the sliding/full classification comes from
cache extents, which tie at a small -c, while the head count does not.
The stateful state trim now derives its sequence axis per state for the same
reason, since pass::KVStateSeqAxis matches per state on the head count.
Assisted-by: Claude Opus 5
* ggml-openvino: apply the KV state relayout to any KV head count
pass::KVStateSeqAxis was limited to states with a single KV head, where moving
the sequence axis from dim 1 to dim 2 is a pure metadata change. The limit was
also based on a measurement showing no gain for a multi-head model, but that was
taken at depth 0, which is the one depth where this change does nothing.
With several heads the pass does more than move metadata: it drops the reader
side transpose of the whole accumulated state, which the graph otherwise redoes
every token at a cost that grows with the context length, and replaces it with a
transpose of the single new row. Measured on GPU, tg128, alternating arms:
gemma-4-12B 6.27 -> 9.11 t/s at depth 8192 (stateless is 7.69, so stateful now
wins at depth instead of losing), Llama-3.2-1B 47.8 -> 59.6 t/s. Both are within
noise at depth 0, which is why the earlier check saw nothing.
The state refill needs the rows copied rather than reinterpreted now: ggml stores
[seq][n_heads_kv * head_size], and a relayout state with several heads is a
different element order. Without that, a refill would seed wrong data - it is
reachable today through llama-bench -d.
Assisted-by: Claude Opus 5
* ggml-openvino : support ggml_rope_set_offset and simplify op support gating
* add more cpy cases
* reject BF16 cpy on NPU
* Remove mul_mat_id fallback, gate large mul_mat_id only for mxfp4
* ggml-openvino: fuse the MoE expert block into MOECompressed on GPU
* ggml-openvino: skip GPU MUL_MAT_ID for unbound expert tensors
* ggml-openvino: requantize grouped 8-bit MoE experts on GPU
* Enable special strided CPY for conv state writeback
* openvino: support cacheless encoder models on NPU
Packed QKV views used by mmBERT were rejected by the ROPE support check. This split Q/K RoPE onto CPU, prevented cacheless attention detection, and sent fragmented encoder graphs through the decoder-oriented NPUW path.
Accept packed QKV RoPE views, detect cacheless attention from its mask, and run these models as a single full-sequence prefill without NPUW or a decode graph. Also provide static mask, output index, and mean-pooling shapes and inputs.
* openvino: optimize norm and RoPE translation
Replace the decomposed mean/variance normalization graph with an opset6 MVN operation. This preserves the GGML epsilon placement while allowing OpenVINO plugins to compile normalization as one operation with fewer intermediate tensors.
Cache RoPE sine and cosine outputs in the graph-wide tensor map. Build the cache key from all RoPE parameters and the optional frequency-factor input so compatible Q/K and layer nodes share one subgraph without mixing different RoPE configurations.
Expose NodeContext::put_shared() to publish translator-created outputs for graph-level reuse.
* ggml-openvino : simplify op translators and enable IMROPE/NEOX RoPE fusion
* remove unnecessary include and clean up PAD
* fix mulmat bug
* use ov::as_type_ptr instead of std::dynamic_pointer_cast
* ggml-openvino: fix mixed-dtype ADD/SWIGLU_CLAMP, gate unsupported ROPE/SOFTPLUS cases
- translate_add: upcast mismatched operand types (e.g. f16/f32 in fused
ADD_ADD) to f32, add, then cast once to the output type. opset1::Add
requires matching input types and downcasting first lost precision.
- translate_glu_swiglu_clamp: same fix, f16 Swish/Clamp rounding was
drifting past the test tolerance.
- supports_op: reject ROPE with ne[3] > 1 (multi-sequence) since the
cos/sin tables only cover one sequence, and SOFTPLUS on GPU since the
OpenVINO GPU kernel overflows to inf for large inputs (CPU is fine).
- ci/run.sh: serialize test-backend-ops on OpenVINO GPU; running two
workers concurrently crashes the GPU plugin (CL_OUT_OF_RESOURCES).
* openvino: share compiled models with per-context inference state; fix thread-safety
* ggml-openvino: gate MoE expert-sum ReduceSum shortcut past 8 experts
The ReduceSum shortcut for the MoE expert-plane-sum ADD chain drifts past
the 1e-7 test tolerance for >8 experts (f32 accumulation order vs CPU
reference), intermittently, like the existing Q4_K/Q5_K NMSE case.
Expose is_moe_expert_sum_add() so supports_op can gate on expert count
and fall back to CPU for just that reduction op.
* ggml-openvino: gate degenerate m=1,n=1 MUL_MAT on GPU
CI hit ERR=1.8e-3 (> 5e-4 tolerance) for a scalar-output f32 dot product
(m=1,n=1,k=2048); didn't reproduce locally in 8 tries, so likely an
internal fp16 accumulation path the GPU plugin picks for this tiny
shape. m=1 output dim doesn't occur in real model weights, so gate it.
* ggml-openvino: make SoftPlus decomposition opt-in native
Assisted-by: Codex
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
Co-authored-by: ravi9 <ravi.panchumarthy@intel.com>
* fix for unsupport zes API
* optimize the code
* adjust the log level
* rm unused head files
* Update docs/backend/SYCL.md
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
* fix the error to detect level zero SDK/dev package, stop build after detect the error
* update the message
* fix the build error when missed to install level zero dev package
* rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions
---------
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
Co-authored-by: Neo Zhang <NA>
* hex-row-split: add support for multi-device row spliting
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* hex-mdev: add work splitting to fused kernels
* hex-mdev: use mdev_ prefix for all multi-device state
* hex-mdev: make device configuration more expressive to support device groups
* hex-mdev: fix mdev session init
* hex-mdev: fused nx (2x,3x) matmuls must update row counts for each w/o
* hex-mdev: fix MUL_MAT work partitioning bugs introduced by mdev
* hex-cont: fix crashes with new tests due to wrong striding
* hex-mdev: move fences after l2flushes
* hex-cont: fix work splitting for mnpu -- align chunks to cachelines
* hex-mdev: fix CPY tests with multi-dev
* hex-mmid: fix work partitioning with mnpu
* hex-mm: fix test failures with mdev
* hex-binary: fix work partitioning for mdev
* hex-argsort: fix mdev partitioning
* hex-mdev: fix work partitioning and general updates for all simple ops
* hex-fa: fix mdev work splitting issues
* hex-mdev: fixing more failing ops test
* hex-mdev: update the rest of the ops
* hex-mdev: refactor all mdev splitting logic to be contained within if (mdev_count > 1) {...}
* hex-mdev: fix macros
* hex-mdev: simplify session flush logic
* hex-sync: fix recursion in session flush
* hex-mdev: factor out fence buffer and allocator
* hex-fence: make fence allocation more robust with reserved slots for mdev
* hex-mdev: keep all mdev state in htp_mdev_group
* hex-mdev: further cleanup mdev group handling at the host
* hex-mdev: update group idx in the opbatch before serializing
* hex-batch: remove separate op_pending and use batch_req/rsp_seq
* hex-async: workaround another missing tensor_init in ggml-meta
* hex-fence: cleanup and robustify fences and error handling in multi-device scenarios
* hex-ar: improve ALLREDUCE error handling
* hex-async: robust error handling for op_cpy_fence
* hex-async: use seq0 from allreduce context to allocate fence_seq
* hex-mdev: fix remaining issues with fence and barrier clearing in CPY_FENCE
* hex-misc: realign macros and fix misplaces trace events
* hex-misc: align macros
* hex-mdev: fix unclone buffer re-entrancy
* hex-glu: fix mdev partitioning logic
* hex-mdev: make buffer uncloning/cleanup work with tensor-split scenarios
* hex-mdev: tighten up the can_split check in act-ops
* hex-mdev: factor out common bits of the partitioning logic
* hex-mm: minor realignment of the macros
* hex-bufs: fix incorrectly placed assert for MAX_BUFS
* hex-pad: tighten up gating checks for PAD
* hex-kparams: make sure all kernels properly use kparams->n_threads
* hex-docs: update user and developer docs with new features and detailed guide for ops development
* hex-scripts: update run script to properly parse dev groups
* hex-misc: formatting
* hex-sess: minor cleanup for session init
* hex-ar: fix vtcm size calc in allreduce kparams
* hex-scripts: fix flake8 warnings
* hex-rope: update ROPE to support mdev work split
* hex-ops: remove redunant checks and minor reformat
* hex-dev-guide: update dev-guide to avoid redundant null checks
* hex-async: improve event_wait, event_sync and fence implementations
* hex-async: remove synchronous flush from event_sync
* hex-async: symplify fence recovery protocol and make sync more robust
* hex-async: futher simplify error recovery for fences
* hex-err: return status instead of just -1
* hex-async: print all seq nums in hex
* hex-async: make sure fences flush dirty ranges
* hex-async: add dirty ranges merging to reduce fence flushes
* hex-async: properly sync before freeing the event
* hex-async: make sure fence owner session is not overriden
* hex-async: more fence write order more robust
* hex-async: make sure not to fuse ALLREDUCE+ADD if their dsts overlap
* hex-fusion: cleanup redundant checks
---------
Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com>
* ggml-webgpu: Update to a recent version of Dawn
* No module scanning
* Accept review suggestion to update comment
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
---------
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
* scripts : add initial profiling script (wip)
* src : add precompile headers (PCH) for models.h
* common : add common.h as PCH
* ggml : add PCH for ggml-impl.h
* mtmd : use PCH for models.h
* scripts : add script to build with Server/Tools/Tests
* server : add PCH for common.h
* docs: add profiling progress notes (wip)
* ggml : add exclude for GCC + SVE on ARM
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33393906061/job/99493756214?pr=28091
* ggml : attempt to fix use of std::hardware_destructive_inference_size
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33396221677/job/99501265689?pr=28091
* squash! ggml : attempt to fix use of std::hardware_destructive_inference_size
Add a version check for GCC 12 to conditionally apply the `-Winterference-size`
pragma.
* editorconfig : exclude profiling reports dir
This directory will not be included in the merge later and this commit
can be ignore at that point. Just fixing to keep CI happy.
* ggml : skip PCH for gcc on non-x86 architectures
* tests : add PCH for peg-parser/tests.h
There are 7 peg-parser tests that can share one PCH instead of then each
parsing the full tests.h.
* common : add PCH for chat.h
* docs : update linux build profiling full results
Just updating after a number of PCH additions. These are not exact
figures and will vary a bit from run to run, but they give a general idea
of the performance impact of PCH.
* cmake : introduce unity build for models
This commit introduces a unity build for the models to improve
compilation time.
The improvements were roughly the following:
```console
+------------------------+-----+------------+------------+------------+
| Build | TUs | Frontend | Backend | Total |
+------------------------+-----+------------+------------+------------+
| Full, master | 396 | 811.0 s | 692.2 s | 1,503.2 s |
| Full, with PCH | 405 | 380.0 s | 664.7 s | 1,044.7 s |
| Full, with PCH + UB | 264 | 357.7 s | 635.7 s | 993.4 s |
+------------------------+-----+------------+------------+------------+
TU = Translation Unit.
Full = includes Server, Tools, and Tests.
PCH = precompiled headers.
UB = unity build for models.
```
* docs : update linux profiling table with unitiy build results
* docs : update mac profiling results to include unity build [no ci]
* docs: remove profiling reports
* scripts : merge build profile scripts into one script
I was lazy before and just copied the first script to enable Tests,
Server, and Tools. This now merges them into a single script.
* Revert "editorconfig : exclude profiling reports dir" [no ci]
This reverts commit 2922a12118.
* src : rename ggml_view_2d_slice to gemma3n_view_2d_slice
This is to be consistent with the rename in gemma4.cpp which was
required to avoid a name clash.
* cmake : add build profile script for windows [no ci]
This commit adds a port of the scripts/build-profile.sh script to
windows powershell.
This was developed on Windows on ARM but should work on X64 as well but
needs to be tested there as well.
define two new environment variables to better understand how much
memory is being allocated, and when. This has been invaluable in
inproving the --fit algorithm, and is likely to be useful when debugging
other memory-related issues.
`-lv 4` will be required to enable the following:
GGML_SYCL_MEMTRACE=1 will show per-site memory usage, updated whenever
it increases by more than 64MiB.
GGML_SYCL_MEMTRACE=2 will show every allocation and deallocation.
To change the default 64MiB threshold for reporting memory usage increases, use
GGML_SYCL_MEMTRACE_STEP.
A sample log line:
[SYCL-MEMTRACE] device memory query (dev): total 59493 MiB, free 4494, in use 54998; allocated 0 (buffers 0 + scratch 0), peak 0 MiB
Fuse RMS_NORM+MUL+ADD and ADD+ADD under GGML_SYCL_ENABLE_FUSION.
ADD+ADD uses the same binbcast indexing and type matrix as standalone
add() (f32, f16, f16/f32, i32, i16, bf16, including broadcast and
non-contiguous). Unsupported combinations fall back to two add() launches.
* port setup-build.ps1 to setup_sdk.py, to facilitate installation of Hexagon and OpenCL SDKs on Windows
* rename setup_sdk.py -> setup-sdk.py
* flake8 fix: print() -> logger.info()
---------
Co-authored-by: Kristopher Urquhart <kurquhar@qti.qualcom.com>
* OpenVINO Backend: Fuse IM2COL + MatMul convolution into OpenVINO convolution
* ci:ggml-ov: Skip recurrent state rollback tests
* ci:ggml-ov: Skip recurrent state rollback tests
* Update OPENVINO.md
* ggml-openvino : add env-var gated op support debugging
* Fix ggml_rope_set_offset case
* OpenVINO backend: Support Whisper.cpp
* Fix code style
* openvino : enable qwen35 on NPU
Static shapes:
- get_graph_input_shape() left the s_copy / s_copy-leaf inputs dynamic
([1,1,1,-1]) even in static mode, which propagated a dynamic slot dim through
GET_ROWS into the conv/GDN state, the state reshapes and the GDN output.
- With -np 1 the s_copy defrag remainder gathers zero rows; short-circuit that
CPY to the untouched cache instead of emitting a degenerate Slice/Concat, and
skip binding its zero-byte ggml tensor as an output (the dynamic path already
did the latter, the static path wrote the full cache over a 0-byte buffer).
Token-count independence:
- In static mode the compiled model's token count is the prefill chunk size or
1, not the captured cgraph's. Offsets derived from the captured count were
therefore wrong. Anchor the GDN state slice at the end of the packed
[attn | state] output and drop the rs_src_begin runtime inputs, and make
VIEWs over the GDN output / conv_input pass through so the consumer does the
slicing.
- CONT could not identify its token axis when the graph was captured with a
single token (every trailing dim has the same stride and size 1) and baked
the captured shape into the prefill model.
Chunked prefill:
- The last chunk is padded with fabricated tokens. Attention masks them, but
the recurrent path folded them into cache_r/cache_s permanently. Add a
chunk_valid_len runtime input, use it to zero g and beta for padded steps
(making the recurrence an exact identity) and to end the conv snapshot window
at the last valid token, and disable the recurrent-cache reset after the
first chunk so earlier chunks are not wiped.
- get_is_prefill() and the chunk loop bound read inp_pos->ne[0] directly, but
IMROPE stacks 4 position planes, so every decode step was run through the
padded prefill model and the loop ran extra out-of-bounds chunks.
cache_rs_reset_idx/len now stay runtime Parameters in static mode, since
can_reuse_statically() does not invalidate the cached model on ComputeParams
changes. Add GGML_OPENVINO_FORCE_STATIC to exercise the static path on CPU.
* Update to OpenVINO 2026.3.1
* ggml-openvino: forward NPU compilation mode parameters
Add GGML_OPENVINO_NPU_COMPILE_CONFIG to the backend's cached environment so callers can configure the NPU compiler without using the generic property escape hatch.
When the value is non-empty, pass it to OpenVINO as NPU_COMPILATION_MODE_PARAMS. This enables settings such as optimization-level=3 for NPU compilation while preserving the existing behavior when the variable is unset and leaving CPU and GPU configuration unchanged.
Document the variable, its NPU-only scope, and the optimization-level=3 example in the OpenVINO backend runtime configuration table.
* ggml-openvino : support RELU, POOL_2D, QUICK_GEGLU, and ROLL ops
* reorder op table
* exclude GPU/NPU failing POOL_2D case
* move op type detection to compute_op_case
* Relax rope supported cases
* Fix pool case
* Update openvino doc, gpu driver in ov docker
* openvino: remove unused static remote context branch
* openvino: parallelize static model build
* Apply editorconfig
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
* Add benchmark-only synthetic speculative acceptance to llama-server and llama-cli
* Address review comments
* Address review comments
* Add some comments in the code
* hexagon: use non-host bufs by default and make the backend fully async
* hex-hb: remove optional hostbuf support and fix async copy
* hex-unary: relax supported unary check
* hex-bufs: use same get_alignment for host bufs
* snapdragon: bump android_platform to 34
* hex-rows: super hacky get/set rows for q8_0
* hex-get-rows: fix q8_0
* hex-get-rows: supprot for f16 and cleanup for q8_0
* hex-get-rows: generic macros and specialized thread funcs
* hex-get-rows: add DMA pipeline, vtcm_layout and kernel params
* hex-set-rows: fix q8_0 support, add dma and tracing
* hex-tests: override nmse threshold for HTP of Q8_0 quants
* hex-fa: add support for Q8_0 with inplace dequantizers
* hex-get-rows: simplify type dispatch
* hex-rows: simplify GET/SET_ROWS DMA pipeline
* hex-async: add events, set/get-tensor-async and rest of the async api support
* hex-repack: use slice instead of expert in repack functions
* hex-cpy: update event/async-cpy logging
* hex-set-rows: optimize smaller tensors
* hex-geglu: fix perf regression with larger tensors
* hex-get-rows: add missing header
* hex-set-rows: add missing header
* hex-bufs: ressurect GGML_HEXAGON_HOSTBUF but disable it by default
* hexagon: do not reject ops with non-heaxon buffers
* hex-get-rows: apply >=32 restriction only for q8_0
* hex-res: bump vtcm acquire timeout to 10 seconds
* hex-bufs: add support for cloning buffers between sessions to speed up tensor copies
* hex-async: rework event recording and batch flushing and integrate with meta backend
* hex-bufs: improved handling of repacked tensors
* hex-repack: handle get_tensor_2d offsets
* hex-dev: add support for devices with multiple NPUs
* hex-sync: add support for sync tokens to synchronize npu devices for async splits
* hex-mmap: cleanup mmap calls and add a retry for robustness
* hex-sync: add failsafe if sync wait gets stuck
* hex-sync: use sync_seq to check for completed events
* hex-sync: rotate tokens for extra robustness
* hex-devs: add supprot for legacy device names for now
* hex-bufs: add support for auto-cloning buffers from diff sessions
* hex-fusion: simplify and optimize htp-opnode fusion handling
* hex-sync: override opnode name so that it shows up in the profiles
* hex-trace: update scripts to handle multiple devices
* hex-sync: bump the size of the opbatch queue and number of sync tokens
* hex-cpy-sync: do not explicitly flush opbatches in cpy_tensor_async and add support for cpy-dma
* hex-sync: add graph-flush threshold to avoid single op batches
* hex-sync: add sync_peer so that we can flush peers we depend on during cross-device ops
* hex-bufs: introduce tensor->extra and shadow_bufs for repacking
* hex-l2: flush tiny tensors inline
* hex-sync: use explicit l2flush for sync tokens
* hex-extra: track weight flags via tensor extra
* hex-fence: rename sync to fence
* hex-repack: proper handling of set-tensor-2d in the shadow_buf
* hex-trace: remove obsolete opstage mask that we used for profiling
* hex-env: remove obsolete use_hmx variable
* hexagon: new unified run.py and build.py and updated docs
* snapdragon: update run script to auto-escapt test-backend-op -p argument
* hex-scripts: fix trailing spaces
* hex-scripts: fix flake8 warnings
* snapdragon: cleanup dst lib/bin dirs before copying new build
* hex-ops: add support for allreduce
* hex-ar: improved allreduce with dma pipeline
* hex-ar: align macros
* hex-ar: consistent use of fence_seq
* hex-ar: add AR_SELECT env var to select ALLREDUCE kernel or fallback
* hex-ar: add proper synchronize handling for ALLREDUCE
* hex-opbatch: looks like we now just rely on backend.synchronise to flush the batches, no need to flush them by threshold
* hex-ar: bump block size to improve dma efficiency
* hex-ar: fused ALLREDUCE+ADD
* hex-ar: cleaner fence buffer management
* hex-ar: futher allreduce tweaking to remove race conditions
* hex-ar: add simple solver and remove non-dma kernels
* hex-ar: add row-broadcast to fuse with bias ADD
* hex-fence: pass seq numbers via op_params
* hex-ar: allow for both entry/exit seq for completing entry wait
* hex-ar: align macros
* hex-ar: do not refetch broadcast row
* hex-fusion: move all fusion into opbatch::add_op for consistency with ALLREDUCE and things
* hex-fusion: fix incorrect MUL_MAT reordering
* hex-mm: make fused 2x and 3x matmuls more generic
* hex-fusion: move tensor fusion tagging to graph_compute
* hexagon: make sure to copy tensor->extra by value
* hex-get-rows: fix offset calc with row-chunking
* hex-repack: get_tensor_2d fixes for non-zero offsets
* snapdragon: make profile/trace scripts more robust and donot mix stdout/stderr by default
* hex-devices: use legacy device nameing by default to ease the transition
* hex-devices: hardcode CDSP domain IDs for current devices for now
* hex-optrace: improve multi-NPU timestamp alignment and overall handling of cycle values
* hex-optrace: more robust handling of the fence events
* CI: Use LLVM's OpenMP over MSFT_DEBUG_non_redist on Windows
Currently, we ship the non-redist debug version of microsoft's libomp.
This PR changes this to official LLVM's release, also packaging
the license as needed.
* Remove LLVM SHA from job name to increase legibility
* Add temp validations to CI
* Revert "Add temp validations to CI"
This reverts commit eef97c88b5.
* Build OpenMP in CI
* Make OpenMP fetch self-contained in cmake and cache in CI
* Robustify Licens-packaging
1. Ship OpenMP license, not LLVM's.
2. Invalidate cache also on checksum of the license
* Remove stale reference in docs/build.md
* No longer package base license in release
This was scope-creep
* Add explanatory comment to OpenMP license
* Remove arm64 smoke
Forgot this during conflict resolution during rebase of
c54c0e9cf6
* Remove GGML_OPENMP_FETCH_CACHE_DIR as requested by @CISC
* whitespace changes
* update to ov-2026.3, update device drivers
* ci: skip nemotron-h rollback test on OpenVINO
The OpenVINO backend does not support SSM_SCAN, so the Nemotron-H recurrent state rollback graph is split and cannot preserve the recurrent cache output shape. Keep the test enabled for other backends and retain the qwen35 OpenVINO rollback coverage.
---------
Co-authored-by: ravi9 <ravi.panchumarthy@intel.com>
* dspark: support speculators-format checkpoints (SpecForge exports)
Speculators-format DSpark drafts (e.g. SpecForge exports for the
Gemma-4-26B-A4B target) differ from the dense DeepSpec checkpoints in
three ways:
- the config nests the backbone hparams under transformer_layer_config
and gives the extract layers as aux_hidden_state_layer_ids
- the block is the DFlash 1+N fill-in layout: the anchor slot is a bonus
token, not a prediction slot. Written as dflash.bonus_anchor; such
drafts build the block and read the mask positions exactly like
DFlash (n_max drafts from a 1+n_max block), only the Markov/confidence
sampling comes from DSpark
- the draft output vocab may be reduced (draft_vocab_size < vocab_size)
with a d2t remap table. The converter expands lm_head/markov_w2 back
to the full vocab and synthesizes an lm_head bias of -1e9 on the rows
the draft cannot produce, so the runtime needs no d2t remapping. Such
drafts ship their own (now optional) token_embd/output tensors instead
of sharing the target's
Verified against gemma4-26b-a4b-dspark: greedy outputs are byte-identical
with and without the draft; acceptance 0.46, mean draft len 3.7 (n_max 6).
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Fable 5
* dspark: fold the speculators draft class into DSparkModel
One class now covers every DSpark variant. What used to pick the class is
a single flag, because the arch name turns out to be the only thing that
separates the two families: SpecForge also exports a flat schema that
carries no speculators_* fields yet still uses the 1+N bonus-anchor block,
so keying on those fields would silently mis-read its drafts.
Also rename i0 to i_first_pred in the draft read loop and the Markov head,
and give the head a real bonus_anchor bool instead of testing i0 > 0.
Converting the Qwen3-8B DeepSpec draft and both gemma-4 speculators drafts
produces byte-identical GGUFs. The one behaviour change is that the
markov_head_type check now also covers the DeepSpec checkpoints, which
previously skipped it.
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Opus 5
* dspark: address review comments
- rename bonus_anchor to sample_from_anchor (GGUF key and code), matching
the checkpoint config field; absent key still means anchor-first
- rework the reduced draft vocab to match EAGLE3: d2t is written as I64
absolute target ids and the logits are scattered at runtime, instead of
expanding lm_head/markov_w2 and synthesizing an output bias at conversion
- move the t2d skip to modify_tensors, like EAGLE3
- drop _is_specforge: the arch name only picks the sample_from_anchor
default, embed/lm_head sharing is decided by the draft vocab size
- deduplicate the tok_embd create_tensor left behind by the rebase
Verified with the RedHat gemma-4-31b speculator draft: greedy output is
byte-identical with and without the draft; acceptance 0.26 (n_max 7).
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Fable 5
* dspark: fold the sample_from_anchor read into the block_size block
* dspark: fix flake8 continuation indent
* clean up
* dspark: key the sample_from_anchor default off the export format
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Fable
* dspark: drop t2d in filter_tensors
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Fable
* dspark: map model.lm_head instead of bypassing the dflash prefix
Co-authored-by: desovo7 <942845546@qq.com>
Assisted-by: Claude Fable
---------
Co-authored-by: desovo7 <942845546@qq.com>
Co-authored-by: ruixiang63 <wangruixiang07@outlook.com>
* ci : allow make-release to target a specific commit
The make-release workflow now accepts an optional 'commit' input. When
set, that commit is checked out and the release checks verify that it
belongs to the branch selected in the "Run workflow" dialog and is not
older than 3 days from the branch tip. The check is part of
make-release-checks.sh (driven by the RELEASE_BRANCH env), so it follows
the same dry-run semantics as the other checks.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : scan latest 100 relase workflow runs
* cont : do not check manually for release.yml success
Replace the deprecated --mmap, --no-mmap, --mlock, and --direct-io flags with
the unified --load-mode argument across scripts, examples, and documentation.
Internal warning message and env var docs updated accordingly.
Signed-off-by: Fathi Boudra <fathi.boudra@linaro.org>
Pradeep Rao