* server: spec checkpoints for recurrent models
* fix: save/restore sampler state during speculative checkpoint
When speculative decoding rejects draft tokens and restores the
recurrent state checkpoint, the sampler (RNG, grammar, prev tokens)
must also be restored to maintain consistency. Without this, the
sampler state reflects the rejected draft tokens, leading to
potential divergence.
Uses common_sampler_clone() to snapshot the sampler before the
speculative batch decode, and restores it on rejection.
* server: snapshot recurrent state in tensor
* reset ngram mod state for rejected tokens
* server: refactor checkpoint state logic
* speculative: fix sampler for checkpoints
* recurrent model: implement recurrent kernel checkpoint
* recurrent model: refactor api
* spec: free rbudget before overwriting
The token loop reads sK[] in the state update (bottom of loop) but has
no barrier before the next iteration overwrites sK[] (top of loop).
Without an explicit memory fence, hardware/compiler reordering can
cause non-deterministic reads from shared memory.
Per review: with this barrier in place, the prior __syncthreads() after
the cross-warp reduction and the one immediately after loop exit are
both redundant. The new barrier is a full block-level fence that also
orders all_sum1/all_sum2 reads vs. the next iteration's writes, and
every thread reaches it before leaving the loop. Both redundant
barriers removed.
No performance impact — GPU utilization is 31-33% during inference,
bottlenecked by CPU MoE expert computation, not the CUDA kernel.
Co-authored-by: Mark Alonzo <mark.alonzo@outlook.com>
* WIP: Gemma4 vision
Crashes on the GPU because of rms_norm requiring ne0 to be multiple
of warp_size.
Runs on the CPU, but produces garbage.
* Remove unnecessary assert in CUDA rms_norm
* GLU was not advertised as supported on CUDA
* Still not working
* This seems to work
get_mmq_x_max_host() returned 128 for NVIDIA Volta when built with
GGML_CUDA_FORCE_MMQ, but get_mmq_x_max_device() caps at
MMQ_DP4A_MAX_BATCH_SIZE (= 64) for the same configuration because the
wider mmq_x device kernels are not instantiated on sm_70. The host
dispatcher selects an mmq_x in [8, mmq_x_max] and instantiates the
matching template; any pick > 64 hits NO_DEVICE_CODE at runtime and
the kernel returns without writing output, producing zero-valued
activations and incoherent generation.
Repro on Tesla V100 (sm_70) with a Q6_K Qwen3.5-MoE model:
cmake -B build -DGGML_CUDA=ON -DGGML_CUDA_FORCE_MMQ=ON \
-DGGML_CUDA_NO_PEER_COPY=ON -DCMAKE_CUDA_ARCHITECTURES=70
./llama-cli -m model.gguf -ngl 99 -p "<~200-token prompt>"
# stderr: mmq.cuh:3941: ERROR: CUDA kernel mul_mat_q has no
# device code compatible with CUDA arch 700.
Cap the host bound at MMQ_DP4A_MAX_BATCH_SIZE on Volta regardless of
GGML_CUDA_FORCE_MMQ; this matches the device template availability and
yields correct output. Measured on a single V100-SXM2:
before patch: PP unusable (NO_DEVICE_CODE), TG ~100 t/s
after patch: PP ~630 t/s, TG ~102 t/s
The change is a no-op on architectures that have int8 MMA (Turing+),
on AMD/HIP, and on builds without GGML_CUDA_FORCE_MMQ.
* Use build_std_attention for Gemma4 when possible
It is possible for the 26b MoE and 31b dense models.
It is not possible for the E4B/E2B vaiants because they
don't have KV cache in each layer.
* Standardize Gemma4 dense ffn
* WIP: Gemma4 split mode graph
Runs but produces NaNs
* WIP: Gemma4 split mode graph
Runs but very high PPL. At least it is no longer NaN.
* WIP
* This works!
* Put attn_norm, attn_post_norm, ffn_norm, ffn_post_norm on all GPUs
* Fix crash when saving/loading KV cache
* WIP: split mode graph for Gemma4-MoE - crashes
* Split mode graph for Gemma4-MoE - this works
* Disable SWA optimization
Something goes wrong there
* Consolidate MoE and dense graph parallel
* Little maintenance
* llama-quantize : Add the missing items in the help
* Add GGML_MAX_CONTEXTS define in the general cmakelist.txt
* Make the KV cache (CPU) based warnings clearer
* Correct placement of GGML_MAX_CONTEXTS definition
* Revert wrong indents
This reverts commit d0728cbb6c6d4d6d8dc13f062e542d232647a38d.
* Moving the GGML_MAX_CONTEXTS definition to src/CMakeLists.txt
* Update warning message for unsupported KV cache types
* forgotten antislash
* Gemma4: WIP
* Gemma4: WIP - runs with totally wrong results
* Gemma4: WIP - add CPU 512, 512 FA
* Gemma4: WIP
It gives a meaningful response in llama-cli, but PPL is still much too
high. Is this due to tokenizer issues?
* Gemma4: this works
I had forgotten the softcap on the final output.
* Remove log
* Gemma4: WIP E4B/E2B
* Gemma4: Q4B/E2B appear to work now
* gemma4: tokenizer fixes
Add HAVE_FANCY_SIMD path that processes 16 rows at a time using 512-bit
operations, combining two R8 groups via _mm512_inserti32x8. Reuses the
existing qx_r8_q8_dot_product 512-bit overload for the inner dot product.
Also updates num_rows for Q8_1 to 16 under HAVE_FANCY_SIMD.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Add a separate HAVE_VNNI256 code path using _mm256_dpwssd_epi32 and
_mm256_dpbusd_epi32 for the Q3_K R4 kernel. The existing HAVE_FANCY_SIMD
(AVX-512 VNNI) path is preserved unchanged.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Add a separate HAVE_VNNI256 code path using _mm256_dpwssd_epi32 and
_mm256_dpbusd_epi32 for the Q6_K R4 kernel. The existing HAVE_FANCY_SIMD
(AVX-512 VNNI) path is preserved unchanged.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
* fix: FA vec kernels for D=256 quantized KV cache
Three bugs prevented Flash Attention from working with quantized
(q8_0) KV cache on models with head dimension 256, such as
Qwen3-Coder-Next:
1. need_f16_K/V used dimension-based logic (Dk != 128) that forced
q8_0→f16 conversion at D=256. The kernel then read f16 data as
q8_0 blocks. Changed to type-based logic (type_K == GGML_TYPE_F16)
matching mainline.
2. quantize_q8_1_to_shared output pointers were not advanced between
loop iterations, so the second iteration overwrote the first half
of Q data in shared memory. Added i0 offset to output pointers.
3. Q_i32 register array in vec_f32 kernel was sized to 1 instead of 2
for D=256 due to an errant comparison (Dk >= expr evaluates to
boolean 1, not expr). Removed the comparison.
Also adds extern template declarations for (256, Q8_0, Q8_0) so the
kernel is compiled and linked.
Tested on Pascal (GTX 1080, CC 6.1) with Qwen3-Coder-Next (D=256,
n_head_kv=2) using --cache-type-k q8_0 --cache-type-v q8_0.
* fix: use dimension-based need_f16 for D=256 vec kernels
Use the sign trick with dpbusd instead of maddubs+madd+add,
replacing 3 AVX2 instructions with 1 fused VNNI instruction.
Removes dead HAVE_FANCY_SIMD code left over from the R16 split.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
dmaivel
Kawrakow
Samuel Oliveira Alves
markaalonzo
Horacio Vico
Nexes the Elder
Adam Caldwell
Alec Koumjian
Max Homilius