* Avoid copying the per-step SSM state (CUDA)
* Avoid copying the per-step SSM state (CPU)
* Allocate only what is necessary for per-step SSM state
* Cleanup
* Use AVX version VNNI intrinsic when AVX512VNNI not available.
* remove changes under HAVE_FANCY_SIMD
---------
Co-authored-by: XZiar <xziar@xziar.xziar>
The default of 0x602 (Windows 8) causes a build failure on any toolchain
where _WIN32_WINNT propagates into vendored cpp-httplib (notably MinGW with
the bundled w64devkit GCC). cpp-httplib's httplib.h has, for some time
now, contained:
#ifdef _WIN32
#if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00
#error "cpp-httplib doesn't support Windows 8 or lower. Please use
Windows 10 or later."
#endif
#endif
so the entire llama-server target fails to compile on Windows + MinGW
unless the user passes -DGGML_WIN_VER=0x0A00 manually.
Bumping the default to 0x0A00 (Windows 10) keeps Windows 8 reachable for
anyone who explicitly requests it (-DGGML_WIN_VER=0x602) while letting the
default Windows + MinGW build succeed end-to-end. Windows 8 / 8.1 reached
end of support in January 2023, and Windows 10 is a strict superset of the
Win8 surface used elsewhere (PrefetchVirtualMemory etc.), so this is
strictly additive on the API side.
Verified by building with w64devkit 2.8.0 (gcc 16.1.0) on Windows 11
without any -DGGML_WIN_VER override: all 266 ninja targets link cleanly,
including bin/llama-server.exe, and llama-cli runs Qwen3-4B-Thinking-2507
IQ4_XS at ~6.2 tok/s with q8_0 KV at 4096 context.
* 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>
Forkoz
Kawrakow
Jun Yamog
Samuel Oliveira Alves
XZiar
Alex
dmaivel
markaalonzo
Horacio Vico
Nexes the Elder
Adam Caldwell
Alec Koumjian