Commit Graph
10764 Commits
Author SHA1 Message Date
copilot-swe-agent[bot] 7cce890466 Initial plan 2026-09-02 16:46:05 +00:00
Georgi GerganovandXuan-Son Nguyen e750b887a8 common, server : enable preserve_reasoning kwarg by default, log its effective state (#28174)
* common, server : enable preserve_reasoning kwarg by default, log its effective state

If the preserve_reasoning chat template kwarg is not specified explicitly
via --reasoning-preserve / --no-reasoning-preserve, it is enabled by
default after argument processing. The server logs the effective state of
the kwarg, warns that it is enabled by default when the template supports
it, and only warns "has no effect" when it was enabled explicitly on a
template that does not support it. Setting the kwarg via
--chat-template-kwargs is deprecated.

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* cont : update comment

Co-authored-by: Xuan-Son Nguyen <[email protected]>

---------

Co-authored-by: Xuan-Son Nguyen <[email protected]>
b10763
2026-09-02 19:19:54 +03:00
Xuan-Son Nguyen 7798007a29 mtmd: support DeepSeek-V4-Flash-Vision-Exp (#28133)
* mtmd: support DeepSeek-V4-Flash-Vision-Exp

* handle min/max token counts from CLI

* rm debugging

* use GGML_ROPE_TYPE_VISION

* nits

* apply review comments

* correct token count
b10762
2026-09-02 16:43:43 +02:00
Aman Gupta 8e93a9773b CUDA + ggml: add sparse-fa for DSV4/GLM (#27970) 2026-09-02 17:27:37 +03:00
Pascal 0f3a71be15 mtmd: Fix Qwen3-tts-0.6b (#28231)
* mtmd: load the qwen3-tts code predictor proj_in as optional

The talker and the code predictor share the hidden size on the 0.6B
checkpoints, so the reference builds no small_to_mtp_projection and
the conversion emits no tensor for it. The graph already falls back
to identity when the weight is missing, the loader now agrees.

* mtmd: keep the qwen3-tts code predictor ffn_down in F32

The code predictor carries a massive activation: its layer 2 FFN
intermediate peaks around 1.5e5, well past the 65504 ceiling of F16.
mul_mat casts its input to the weight type, so an F16 ffn_down turns
that peak into inf, the residual follows, and the next rms_norm yields
NaN. Reference forward in float32 gives 145109 against 145396 measured
in the graph.
b10760
2026-09-02 12:46:16 +02:00
Aman Chadha(IVIXMMI)andAcmmi b81c99b479 ggml: avoid KleidiAI buffer type init on dispatch (#27891)
Co-authored-by: Acmmi <[email protected]>
b10759
2026-09-02 09:16:15 +03:00
Max Krasnyansky 960dffab05 hexagon: MUL_MAT and MUL_MAT_ID fusion and fixes (#28202)
* hex-mm: fuse QKV and FFN matmuls that land on HMX

* hex-mm: remove hardcoded ne[1] < 32K restriction

* hex-get-rows: explicitly reject repacked Q8_0 just in case somebody decided to add an override

* hex-mm: correct overhead sizing to make sure we dont exceed vtcm budget for large dims

* hex-mm: fuse MUL_MAT_ID into MUL_MAT_ID_NX (2x,3x,...) where possible

* hex-fusion: update opbatch and opqueue sizing to acount for new fusion and reduce overhead for trace buffer alloc

* hex-bufs: sort buffers while finalizing opbatch, helps avoid va space fragmentation

* hex-bufs: add simple va defrag to make sure we dont abort just because the va space is fragmented

* hex-mm: replaced more scalar divs with fastdiv and minor cleanup

* hex-mm: tighten up supported fusion checks to exactly match supported kernels
b10758
2026-09-02 09:15:21 +03:00
Laurent ZuijdwijkandMarshall ba8818cbf3 vulkan: handle larger batch sizes (>4) efficiently for IQ3_S mat-vec (#27449)
* vulkan: handle larger batch sizes (>4) efficiently for IQ3_S mat-vec when NUM_COLS > 4. 5x perf at n=8

Assisted-by: Claude Opus 5

* adds 2 cases per quant type at `k=16*256` to the `all_types` mat-vec sweep

---------

Co-authored-by: Marshall <[email protected]>
b10757
2026-09-02 09:14:52 +03:00
Mads Marquart 56dd8150cc vulkan : only request VK_KHR_shader_bfloat16 extension if supported (#28155) b10756 2026-09-02 09:13:25 +03:00
Alan Tseng 2637dfe373 ggml-cpu : conditionally add SpacemiT IME kernel sources (#27961)
When building with gcc < 15, CMakeLists.txt unconditionally adds
ime2_kernels.cpp, which fails to compile. FindSMTIME.cmake only defines
RISCV64_SPACEMIT_IME2 when the IME2 instructions are detected, and gcc 14
only has IME1, so ime2_kernels.cpp hits its #error.

This PR fixes it by using IN_LIST to add each kernel source according to
the spec that was actually detected.
2026-09-02 09:12:28 +03:00
Hongqiang Wang 43d87ff2dd opencl: fix out‐of‐bound reads in the Adreno image kernels (#27632)
* opencl: clamp the q4_K decode GEMV's fetch row on a padded x-grid

* opencl: enforce the tiling contract of the image KQ/KQV GEMMs

* opencl: decide the image KQ/KQV split at the dispatch, not from strides
b10754
2026-09-01 22:28:45 -07:00
Trivikram Reddy 69320fef12 hexagon: add missing FARF logs for cpy/get_rows/set_rows/gdn ops (#28217)
* hexagon: fix bug ne[2] printed in proc_op_req prep-src log

* hexagon: add shape/VTCM farf logs to cpy, get/set rows, gdn
b10753
2026-09-01 22:20:29 -07:00
Jhen-Jie Hong b96806d960 metal : add metallib build support for xcframework (#28163) b10752 2026-09-02 07:45:56 +08:00
anujj 3466812d1f cuda: fuse MoE weighted expert reduction (#25952)
* cuda : fuse MoE weighted reduction (mul + view + add)

The MoE combine tail currently writes weighted expert outputs to
global memory before reducing them. That intermediate global-memory
traffic is the main cost. The production baseline generally runs two
physical fused kernels; this path runs one.

This change matches the full expert-weighting plus ordered-reduction
subgraph and replaces it with one weighted-reduction kernel.

Supported graphs:
- unscaled: experts * router_weights
- scaled:   (experts * expert_scale) * router_weights

k = 2..15 is handled by one runtime-k kernel.

Matching is structural: op sequence, shapes, strides, expert views,
and the left-to-right ADD chain. The fused kernel keeps that same
reduction order. Results are not claimed bit-identical; CUDA FP32
contraction can change rounding slightly.

Allocator integration uses add_alloc_dep from the graph-optimizer
API so experts, router weights, and optional expert scales stay live
until the fused destination is written. Memory ranges are rechecked
before the fused kernel runs.

Unrecognized or unsafe graphs are left alone and keep the existing
per-op path. Set GGML_CUDA_MOE_WEIGHTED_REDUCTION=0 to disable the
fusion.

test-backend-ops covers scaled/unscaled, aligned/unaligned, and
representative values across k=2..15, plus a k=16 case that must
stay on the per-op path.

* Pruned the test matrix from 15 to 6

* Addressed the aman and olivers review comments
b10751
2026-09-01 21:48:47 +02:00
Pascal b356fa2624 kv-cells: look up the n-gram history in the sequence position index (#28040)
get_prev_tokens() rebuilt a (seq, pos) -> token hash map on every
ubatch by walking all used cells, while llama_kv_cells already keeps
an ordered index of the positions of each sequence in seq_pos, updated
on every cell mutation to serve seq_pos_min() and seq_pos_max().

The index now stores (pos, cell) pairs in a std::set instead of a
position -> count map, so a repeated position (cache reuse via rm + add,
vision inputs with shared positions) yields distinct entries and the
removal of a cell erases its own pair. The new seq_pos_tok_le() returns
the token of the cell at the largest position <= p in logarithmic time,
which is exactly what the old window lookup and its M-RoPE gap fallback
computed together.

get_prev_tokens() shrinks to a direct lookup per (token, offset) and
for_each_token_in() goes away with its only caller. The kv-cache keeps
no n-gram logic of its own.

Measured on Qwen3.8-Flash-Next UD-Q4_K_XL at 71k context, alternating
two binaries with the first run discarded: tg 69.3 -> 72.7 t/s (+4.9%),
pp unchanged at ~2720 t/s, greedy output identical, needle retrieved.
b10750
2026-09-01 20:16:07 +02:00
Sigbjørn Skjæret dfc29b64eb context : autoscale n_ctx_train when yarn scaling specified (#28030) b10749 2026-09-01 19:59:54 +03:00
Sigbjørn Skjæret f28493c783 models : appropriately flag noscan ssm_a tensors (#28121) 2026-09-01 19:59:15 +03:00
Sigbjørn Skjæret 73159c3039 model : fix gemma4-assistant (#28183) 2026-09-01 19:58:44 +03:00
Sigbjørn Skjæret d11b3cc7ed model : load relevant arrays with n_layer_all (#28173) 2026-09-01 19:58:29 +03:00
Titaniumtown c845263f8b Revert "sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 12…" (#28184)
This reverts commit 1f3d318734.
2026-09-01 19:04:31 +03:00
Jingxin (Philip) Li 1f3d318734 sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 1280 (#28016) 2026-09-01 11:47:08 -04:00
Lukasz Stolcman 8887a48f05 metal : add fa-vec tuning for M2 Pro (#28122)
* metal: add fa-vec tuning for M2 Pro

* metal : update fa-vec tuning for M2 Pro with new dtypes
b10743
2026-09-01 21:24:44 +08:00
Jhen-Jie Hong be789c3448 metal : add fa-vec tunings for A18 Pro (MacBook Neo) (#28152) b10742 2026-09-01 21:15:59 +08:00
Sigbjørn Skjæret 9d817213a0 model : load hparams.n_layer_nextn before n_layer() calls (#28159)
* load hparams.n_layer_nextn before n_layer() calls

* remove duplicate loads
b10741
2026-09-01 13:55:45 +02:00
Niklas WenzelandYiChen Lv fe2120bc9d metal : fix more leaks due to missing autoreleasepools (#27883)
* metal : fix more leaks due to missing autoreleasepools

* metal : rename variable

* metal : fix another missing pool warning

Co-authored-by: YiChen Lv <[email protected]>

---------

Co-authored-by: YiChen Lv <[email protected]>
b10740
2026-09-01 13:50:47 +02:00
Georgi Gerganov d08c7872d6 metal : add fa-vec tuning for M2 Max (#28015)
Rows for M2 Max (30 GPU cores) collected with 'ggml-metal-tuning fa-vec
--dtype f16,q8_0', pasted into fa_vec_tuned_table.

ref: https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18205786

Assisted-by: pi:llama.cpp/Qwen3.8-27B
b10739
2026-09-01 13:37:40 +03:00
Neo Zhang 5eec3ad017 sycl : support limit max alloc memory within 2GB for host-pinned memory (#27559) b10738 2026-09-01 13:35:47 +03:00
Daniel Han 36b1015438 qwen4exp: fix seq_cp, block position keying, mtmd input, cuda abort, add tests (#27941)
* qwen4exp: follow up fixes

* -kvu NaN collapse fix

Assisted-by: Claude

* indexer cache ext.x/ext.y restore fix

Assisted-by: Claude

* kv-cells: rename seq_set to seq_get_all

seq_get is already taken by the single-id getter, so the suggested name
cannot be overloaded on return type alone.

Assisted-by: Claude

* memory-hybrid-idx: implement set_input_qsa on the memory class

The context held the whole implementation, where the pattern elsewhere is a
thin context forwarding to the memory class, as llama_kv_cache_context does
for set_input_kq_mask. The body reads no context state, so it moves unchanged
and the context keeps a forwarder.

Also shortens the seq_get_all comment as suggested.

* tests: check that a sequence state survives a save/restore round-trip

Saves seq 0, erases it, restores the blob and saves again, requiring the two
blobs to match. Compares blobs rather than generated text, which cannot see a
field dropped on the way back in.

Note this passes on master for qwen4exp, so it does not demonstrate the
ext.x/ext.y drop this PR fixes; reaching that needs 2D mrope content.

* tests: give the synthetic qwen4exp a PLE so the state test bites

has_cell_ext() is n_pos_per_embd() > 1 || ple_n_heads > 0, and the indexer
cache sets rope_type = NONE, so without a PLE it serializes no cell ext at
all and the round-trip test cannot see a dropped ext.x/ext.y. With one,
removing the ext_set restore in state_read_meta fails the test: 198 of
335692 bytes differ, first at offset 282092.

Loading such a model needed two fixes:

- the row count of per_layer_token_embd came from require_weight(), which a
  model synthesised from metadata alone has no file to answer. Derive it
  from the head ranges and prefer the file's padded count where there is one.
- the PLE conv history is a row of the recurrent cache, so a PLE on a full
  attention layer dereferenced a null p_l. Reject it at load time instead.

The meta mirror is skipped for qwen4exp. It returned NaN logits before this
fixture carried a PLE, which the nmse check passes since a NaN comparison is
false, and aborts with one. -sm tensor on real devices works.

Assisted-by: Claude

* llama: disable -sm tensor for qwen4exp

test-llama-archs skipped the tensor split for this arch from inside the
test, so the arch still advertised support it does not have. Declare it in
llm_arch_supports_sm_tensor instead and drop the test-side exception; the
existing llm_arch_supports_sm_tensor branch then does the skipping.

Assisted-by: Claude
b10737
2026-09-01 13:22:04 +03:00
Georgi Gerganov d086dbb348 tests : fix log verbosity for test-llama-archs (#28147)
* tests : fix log verbosity for test-llama-archs

* cont : naming

* cont : add note
b10736
2026-09-01 13:07:12 +03:00
Xuan-Son Nguyen 1b89a43e38 quantize: row-slab stream to avoid thread starvation (#27830) 2026-09-01 11:18:54 +02:00
James Francis d5d993a093 metal: enable Metal 4.0 tensor API on M5+/A19+ (#27461)
* metal : request Metal 4.0 language version for the tensor API

* metal : load the tensor API kernels from a separate metallib

* tests : add external-metallib tensor API regression test

* metal : fix metallib build order for the tensor API kernels
b10734
2026-09-01 12:02:42 +03:00
Ludovic Henry 234a6ebaa0 ci: Bump ggml-org/ccache-action to v1.2.24 (#28083) b10733 2026-09-01 12:00:17 +03:00
Jonathan Clohessy 518b76236b kleidiai : Update KleidiAI Documentation (#26078)
Signed-off-by: Jonathan Clohessy <[email protected]>
2026-09-01 10:45:13 +02:00
Pascal 0eadefebd3 qwen4exp: support recurrent state rollback (#28123)
MTP speculative decoding needs the target state to move back by the
number of rejected draft tokens. Without rollback support the context
is classified as SEQ_RM_TYPE_FULL and the server serializes the whole
recurrent state to host memory on every round, which costs more than
the drafting saves.

The recurrent cache already holds n_rs_seq + 1 snapshot planes and the
delta net writes its SSM state into them, but build_conv_state_at wrote
a single plane, so a rollback restored a convolution history that was
never captured. It now writes one snapshot per slot, each ending one
token earlier, for the delta net QKV convolution and for the PLE
convolution alike.

Measured on Qwen3.8-Flash-Next UD-Q4_K_XL with the standalone MTP
draft, n-max 3 and a single slot: decoding reaches 183 tok/s on code
and 144 tok/s on prose. The same branch before this change, where the
server falls back to checkpointing the state to host memory, reaches
123 and 83 tok/s, for 108 tok/s without a draft.
b10731
2026-09-01 06:24:49 +02:00
Pascal 09412af38a qwen4exp: sum the indexer heads by slices (#28023)
* qwen4exp: sum the indexer heads by slices

The head reduction went through a transpose and a sum_rows over ne[1],
which left sum_rows with ne0 = 4, one block per row for a four element
reduction, and the transpose copied the whole block by token surface
twice on the way in.

The heads are adjacent on ne[1], so each one is a strided view and the
sum is a short chain of adds.

RTX PRO 6000, Qwen3.8-Flash-Next UD-Q4_K_XL, fa on, 55k context, warm
runs on top of #28011:

  prompt processing   2170 -> 2366 t/s

Generation is unaffected. The removed work scales with n_blocks by
n_tokens, so the gain grows with context and with ubatch size.

* qwen4exp: drop the redundant cont on the indexer query

rope returns a freshly allocated, contiguous tensor, so the reshape that
feeds the matmul does not need a copy. ggml_reshape_3d asserts
contiguity, so a layout that would need the cont cannot slip through
silently.

Greedy output is unchanged token for token.

Address review from @ggerganov
b10730
2026-09-01 06:23:59 +02:00
Buğra Özgürsoy 458681e1d5 metal : add fa-vec tunings for M1 Ultra (#28088)
* metal : add fa-vec tunings for M1 Ultra

* metal : move M1 Ultra tunings after M1 Max section

* metal : remove duplicate blank line
b10729
2026-08-31 23:47:27 +02:00
ynankani e4b9af007b CUDA: XOR swizzle flash attn K,V smem fp16 tiles (#25635)
* CUDA: XOR swizzle flash attn  K,V smem fp16 tiles

Signed-off-by: ynankani <[email protected]>

* Fix use 64bit generic pointer instead of 32bit shared pointer

Signed-off-by: ynankani <[email protected]>

* fix shared memory race in FA on DGX Spark

* Handle corener case

Signed-off-by: ynankani <[email protected]>

* Add swizzle test cases and gate sync for swizzled path only

Signed-off-by: ynankani <[email protected]>

* gate CUDA PTX

Signed-off-by: ynankani <[email protected]>

* offset calculation specific for swizzle branch

Signed-off-by: ynankani <[email protected]>

* Reafctor code

Signed-off-by: ynankani <[email protected]>

* Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)

Signed-off-by: ynankani <[email protected]>

* rebase and update test case args

Signed-off-by: ynankani <[email protected]>

* Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0

Signed-off-by: ynankani <[email protected]>

---------

Signed-off-by: ynankani <[email protected]>
b10728
2026-08-31 22:18:01 +02:00
Georgi Gerganov ab0b3bd3c8 metal : add concat support for quantized types (#28116)
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
b10727
2026-08-31 23:16:04 +03:00
BartowskiandGeorgi Gerganov 85c55223ca AVX2: Speed up large batch size prompt processing of IQ models (#27402)
* Batched gemm for grid IQ quants

Style updates and a bit more performance

Clean up comments

Move code around

Vectorize IQ panel decode, lower threshold for speedup

IQ panel: single-source gather layout, gate bias, vectorize interleave

Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer

Move IQ panel code out of repack into iqp.cpp, clean up comments

Another comment sweep

* Add myself as iqp.* codeownder

* Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size

* Renaming and moving

* The other half of renaming and moving

* Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition

* Update ggml/src/ggml-cpu/iqp.h

Co-authored-by: Georgi Gerganov <[email protected]>

* Add iqp_rows work buffer

* Revert "Add iqp_rows work buffer"

This reverts commit 425542991e.

* Add NUMA fallback

* Add 10 row batch tests for IQP coverage on all grid IQ types

* Swap assert for return false in support check

* Move IQP mul_mat_id test

---------

Co-authored-by: Georgi Gerganov <[email protected]>
b10726
2026-08-31 14:33:50 -04:00
Georgi Gerganov 2a74817f93 metal : add top-k radix implementation (#28073)
Assisted-by: DeepSeek-v4-Flash-0731
2026-08-31 21:31:53 +03:00
itsnotoger 2d8d612e4c kv-cache : optimize restoring non-contiguous cells (#27991)
* kv cache : batch state restore scatter reads per contiguous run

When restoring state into non-contiguous destination cells (e.g. a
prompt-cache snapshot into a fragmented ring), state_read_data issued
one small copy per KV cell - ~1.4M copies of a few KiB each for a
40k+ token restore, taking 25-63 s on the CUDA backend.

The snapshot stores cell rows in cell order, so a maximal run of
consecutive destination indices maps to one contiguous block and can
be restored with a single copy. Precompute the runs once and use them
in all three scatter loops (K, V, transposed V). Byte-identical.

The on-device reader copies with a byte cursor when the read and
write chunking differs, so the batched reads are safe for it as well.
Batching makes equal tensor counts with a different split reachable
(save ranges [2,1] vs restore runs [1,2]); the next commit teaches the
reader's 1:1 path to fall back to the byte cursor in that case.

Verified in a production setup: 1,363,616 copies / 25-63 s -> 224
copies / 221-424 ms for the same restores (42,603 cells, 4 runs).

Assisted-by: Claude Code (unsloth/qwen3.8-27b)

* context : fall back to the byte cursor when read and write chunking differ

the on-device reader copies saved state back with a 1:1 copy by tensor
index whenever the write and read sides recorded the same number of
tensors, guarded by a per-tensor size assert.

equal tensor counts do not imply equal chunking: a state restore may
batch its reads per contiguous run of destination cells while the save
used per-range reads, so both sides can record two tensors that split
the same data differently, and the assert aborts in all builds.

compare the per-tensor sizes and only take the 1:1 path when the
chunking actually matches, otherwise fall through to the existing
byte-cursor copy. both sides enumerate the same logical data in the
same order, so the cursor copy is well-defined across tensor
boundaries.

Assisted-by: Claude Code (unsloth/qwen3.8-27b)

* tests : cover state restore scatter reads on host and on-device paths

decode the same prefix on two sequences, interleaving the seq 0 cells
between the seq 1 cells, so the seq 1 cells are isolated from each
other in the kv cache (three cells, two saved ranges). save the seq 1
state, free the interleaved seq 0 cells, and restore: the destination
is then non-contiguous (two runs), and the restore-side chunking has
the same tensor count as the save-side with a different split, so the
scatter path is batched per contiguous run and the on-device reader's
byte-cursor fallback is exercised.

the restored state is saved again on the host and compared byte for
byte with the first save: the blob is serialized in sequence cell
order, so the two saves are identical if and only if the scatter
restore wrote exactly the same KV content. this documents the
byte-identical guarantee of the run-batched scatter reads.

one test per io backend: the host (CPU) path and the on-device path.

Assisted-by: Claude Code (unsloth/qwen3.8-27b)
b10724
2026-08-31 19:49:58 +03:00
Hongqiang Wang 010be9683a opencl: tune the quant paths for Intel Xe-LP GPUs to improve its TG and PP performance (#26438)
* opencl: Q4_K/Q5_K mul_mv N_DST 4->8 on Intel for 2x activation reuse

* opencl: Q4_K mul_mm 8x8 tile fot Intel

* opencl: Q5_K mul_mm 8x8 tile for Intel

* opencl: Q4_K mul_mv N_DST 8->16 for Intel
b10723
2026-08-31 08:56:22 -07:00
Pascal 774ee0e200 ui: copy the displayed text of grouped agentic responses (#27832)
* ui: copy the displayed text of grouped agentic responses

Agentic sessions render as a single entry anchored on the first
assistant turn, whose content is typically just the first tool call,
so the copy button wrote an empty string to the clipboard. Derive the
text sections of the whole session and copy them joined, matching the
visible response. Plain messages keep the previous behavior.

* const
2026-08-31 17:48:43 +02:00
8e53fcefd2 webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation (#28045)
* webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation

* chore : improve code readability

Co-authored-by: Sigbjørn Skjæret <[email protected]>

---------

Co-authored-by: Stanisław Szymczyk <[email protected]>
Co-authored-by: Sigbjørn Skjæret <[email protected]>
b10721
2026-08-31 16:04:38 +02:00
Jaden_Mach f8dbcd6189 ROCm: add radix TOP_K for long rows (#27466)
* ROCm: add radix TOP_K for long rows
b10720
2026-08-31 15:00:04 +02:00
Niklas Wenzel 5d4a3be26d metal : add fa-vec tunings for M1 (#28078) b10719 2026-08-31 13:58:55 +02:00
ynankani 41ef91f7c8 CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were restricted to 1 token (#27621)
* CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were resticted to 1 token

Signed-off-by: ynankani <[email protected]>

* Address review comments

Signed-off-by: ynankani <[email protected]>

* Add SWIGLU_CLAMP case to multi-token moe fusion

Signed-off-by: ynankani <[email protected]>

---------

Signed-off-by: ynankani <[email protected]>
b10718
2026-08-31 19:22:28 +08:00
Neo Zhang a32af33de2 sycl : Enhance to get the free memory of Intel GPU (#27968)
* enhance get mem info by l0 an SYCL API

* remove debug code, format the code

* update SYCL.md for GGML_SYCL_GET_MEM_API
b10717
2026-08-31 13:33:02 +03:00
Sigbjørn Skjæret 580e88d8b7 ci : add check for unzip (#28082) 2026-08-31 12:17:51 +02:00
662a0b0121 spec : fuse the DFlash encoder into the KV cache injection (#27310)
* dflash : fuse the encoder into the KV injection decode

The encoder is a single fc + norm, but running it as a separate
llama_encode forced a device-to-host round trip of its output before the
injection decode could re-upload it, plus a second graph build per
round. Fold the encoder into the decoder's embd branch and feed the
target features directly to one llama_decode.

Assisted-by: Claude Fable

* nit

* Apply batched suggestions from code review

Co-authored-by: Ruixiang Wang <[email protected]>

* Fix missing references from renaming

---------

Co-authored-by: Sigbjørn Skjæret <[email protected]>
Co-authored-by: Ruixiang Wang <[email protected]>
b10715
2026-08-31 11:19:20 +02:00