Consolidate gfx803 paged/turbo4 + LFM2.5 work onto fork master#14
Merged
Conversation
…izing The flash-attn tile kernel was WMMA-gated (RDNA3/4 only); on non-WMMA archs (gfx803, GTX 1070/Pascal) it fell through to NO_DEVICE_CODE, so the paged tile path was disabled and prefill ran the slow fallback. - Add a portable FMA #else body to mt_paged_attention_tile_kernel mirroring the WMMA lane layout (row=tid%16, half=tid/16, cols 8*half+l): QK^T as smem dot products over HEAD_SIZE, exp-scores published to new smem_s[f32], scores·V as a smem GEMM. Uses __shfl_xor_sync(...,16,WARP_SIZE) (correct on 64-lane gfx803). Reuses portable stage_q/k/v_tile helpers (turbo4 dequant). - Relax tile gate to tile_env_on (drop wmma_ok); force single-warp on non-WMMA (mw_on requires wmma_ok). - Size tile grid n_q_tiles from total_q_tokens, not avg_q_len: with idle --parallel slots avg_q_len floored low, leaving rows uncomputed (garbage output). Same bug class as 93a1e6a. - Grow smem_bytes for the added Q_TILE_M*K_TILE_N f32 score tile. Verified correct vs scalar oracle (Qwen3.5-9B, LFM2.5); GTX 1070 prefill 160 -> 328 t/s. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Non-WMMA prefill fix + avg_q_len grid bug, gfx803 ROCm vs Vulkan perf table, checkpoint/paged block-align crash, LFM2.5 tool-schema bleed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…-by-actual) Paged decode on RX 480 (gfx803) ran ~1.7x slower than standard FA. rocprofv3 kernel-trace localized it to mt_paged_attention_decode_kernel (pass-1 flash-decode): 2.5 ms/call at ctx 1751, the dominant decode kernel — bigger than all decode matvecs combined. Two synergistic fixes: - CHUNK_KV 1024 -> 256. 1024 was tuned for ~400K ctx on RDNA4's ~240-block ceiling; on gfx803 (36 CUs) at mid ctx it yields only ~2 chunks -> a handful of grid blocks -> severe CU under-utilization. 256 gives ~4x more parallel blocks. - Bound the reduce kernel (and skip pass-1 neutral-partial writes) to the chunks that actually hold data: valid_chunks = ceil(context_lens/CHUNK_KV). num_chunks is sized by ALLOCATED ctx, so with the smaller CHUNK_KV and a large --ctx-size (524288) there are thousands of empty chunks; without this bound the reduce loop + neutral writes would erase the CHUNK_KV gain. The reducer takes context_lens; empty pass-1 blocks now early-return. Mathematically exact (split-K granularity + bounding chunks-with-data): A/B at the full production config (parallel 4, tiered 25/25/50, ctx 524288) produces byte-identical output vs baseline on varied prompts. Decode: 33 -> 52 t/s at ctx 1751, 33 -> 37 at ctx 10501; equal at shallow ctx. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
gfx803 (GGML_CUDA_CC_GCN4) has no hardware dp4a (v_dot4); ggml_cuda_dp4a falls back to fully-scalar byte emulation, so MMQ ran ~8x slower than dequantization + hipBLAS for prefill-sized batches. ggml_cuda_should_use_mmq previously returned true for all non-CDNA AMD, keeping every GEMM on MMQ. - mmq.cu: add a gfx80x branch (GGML_CUDA_CC_IS_GCN && cc <= GCN4) that keeps MMQ only for small (decode) batches and routes large batches to dequant+hipBLAS. rocBLAS ships gfx803 fp16/fp32 Tensile kernels. - ggml-cuda.cu (mul_mat_id): gate MMQ with `routing_active ||` so the MAD-88 weight-pager routing-active MoE path always uses MMQ (the only routing-aware kernel here) instead of falling through to GGML_ABORT. No-op for configs without the weight pager. - mt_pagedattn_decode.cu: CHUNK_KV 256->128 (gfx803 sweep: 128 > 256 > 512). Clean same-revision llama-bench A/B (LFM2.5-8B-A1B Q5_K_M, gfx803): pp512 72.3->600.7 t/s (8.3x), pp2048 74.1->583.9 (7.9x), tg64 unchanged. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…301C Lever A) The multi-warp tile FA kernel was WMMA-only (#if AMD_WMMA_AVAILABLE, no #else), so non-WMMA gfx803 fell back to the single-warp tile kernel: 32 threads/block = half a 64-lane wavefront, and K/V reloaded from HBM per Q-tile (no reuse). Add a portable FMA #else path mirroring the single-warp FMA math per warp, reusing the mw kernel's cooperative K/V staging (loaded once per block, shared across warps). N_WARPS*32 threads now fill the wave64; the PV step rebuilds each Q row's full 16-wide score vector from the pair lane via __shfl_xor(.,16), so no extra LDS over the WMMA layout. Un-gate the dispatch (drop && wmma_ok) so gfx803/Pascal reach it. Same per-warp FMA math as the single-warp path => bit-identical output (verified vs baseline at 4 and 126 K-tiles). Prefill on LFM2.5-8B-A1B (gfx803, turbo4, full prod config): ~290 -> ~370 t/s (kernel-isolated 290 -> 453, 1.56x); decode unchanged. Rollback: GGML_PAGED_TILE_MULTIWARP=0 (single-warp, no rebuild). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…AD-301C Lever B) head_dim-64 models (LFM2.5, gpt-oss) padded each 64-dim head to a 128-element turbo4 block, wasting ~half the turbo4 KV cache and doing 2x the paged-attn work. Add GGML_TYPE_TURBO4_64: a 64-element turbo4 block (34 vs 68 bytes/head) used as a paged KV cache type. The paged path skips RHT, so the 64 real values quantize identically to the padded case (norm is over real values either way) -> prefill output bit-identical; ~1% recon-norm shift only (native is marginally more accurate). - ggml core: block_turbo4_64, GGML_TYPE_TURBO4_64=52, type-traits, CPU quant/dequant - paged: paged_cache_ops<TURBO4_64>, mt_scatter_kv_turbo4_64_kernel (2-warp), HS=64 tile gate + <64,16,TURBO4_64> instances, supports_op whitelist - decode: decode_coop_stage_turbo4_64 (32 lanes x 2) so HS=64 keeps the fast flash-decode path (no decode regression) - llama-kv-cache-paged: auto-remap TURBO4_0 -> TURBO4_64 for head_dim==64, skip the 64->128 pad; graph pad auto-skips (keys off layer.k->type) LFM2.5-8B-A1B (gfx803, full prod config): prefill 370 -> 432 t/s (+17%), decode ~59 t/s unchanged, turbo4 KV -50%. Output bit-identical (prefill) / coherent (decode). A/B + rollback: GGML_PAGED_TURBO4_64=0 reverts to padded-128 (no rebuild). TURBO4_64 is a runtime KV-cache-only type (never serialized to GGUF). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
LFM2.5's chat template does not prefill an opening <think> into the generation prompt, so the model emits it itself each turn. When it skips the opening tag and dives straight into reasoning (closing only with </think>), the peg-native lfm2 parser -- which required a literal opening <think> and had no forced-open fallback -- failed to enter the reasoning rule, so the entire think block plus the stray </think> bled into content. Make the opening <think> optional: capture a leading reasoning block terminated by </think> whether or not the opening tag was emitted. A pure-content turn has no </think>, so until() fails to find the terminator and the outer optional collapses, leaving output as content exactly as before (no regression -- full test-chat suite passes). Adds regression tests for the missing-opening-tag case (no-tools and with-tools) in test_lfm2_parser. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
github-actions
Bot
added
documentation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Pre-migration consolidation of local llama.cpp work onto the fork master.
Includes the gfx803 paged-attention / turbo4 prefill work and LFM2.5 fixes from
fix/lfm25-paged-turbo4-gfx803. Local master was diverged (ahead 529 / behind 37); preserved separately asbackup/local-master-20260625.🤖 Generated with Claude Code