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>
Replace maddubs_epi16 + madd_epi16 with dpbusd_epi32 in the
mul_mat_q8_0_r8_q8_2 dot product lambda when HAVE_VNNI256 is
defined. Same sign trick operands (abs(x), sign(y,x)), just
fewer instructions per sub-block.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Ungate the VNNI path in mul_mat_q8_1_r8_q8_2 by changing the
guard from HAVE_FANCY_SIMD to HAVE_VNNI256. This block only uses
256-bit intrinsics so it is safe for AVX-VNNI (non-512) CPUs.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
* Allow using -rtr and -muge together
* Various Qwen-3.5 tweaks
* No need to make v, g, beta contiguous
* Adjust NEON delta-net to non-contiguous v, g, b
* Cleanup
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
Add new CPU macro HAVE_VNNI256 for CPUs with 256-bit VNNI
(AVX-VNNI) support or better (AVX512-VNNI+VL), separate from
HAVE_FANCY_SIMD which requires the full AVX-512 set. Relax four
#ifdef guards in mul_mat_q4_k_r4_q8_k and mul_mat_q5_k_r4_q8_k
to use HAVE_VNNI256 instead of HAVE_FANCY_SIMD, enabling vpdpbusd
and cvtepi8_epi32 on Alder Lake, Raptor Lake, and similar CPUs.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Kawrakow
markaalonzo
Horacio Vico
Nexes the Elder
Adam Caldwell
Alec Koumjian
Max Homilius