From d00fea0d57678f71767032243e2e4ad0de4892db Mon Sep 17 00:00:00 2001 From: kmbandy Date: Mon, 15 Jun 2026 15:10:41 -0400 Subject: [PATCH 1/7] cuda/paged-attn: portable non-WMMA tile kernel + fix avg_q_len grid sizing MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 93a1e6a11. - 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 --- ggml/src/ggml-cuda/mt_pagedattn.cu | 19 ++- ggml/src/ggml-cuda/mt_pagedattn_tile.cu | 165 ++++++++++++++++++++++-- 2 files changed, 172 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-cuda/mt_pagedattn.cu b/ggml/src/ggml-cuda/mt_pagedattn.cu index 9257824a2213..45a21c27b6a0 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn.cu @@ -1105,7 +1105,11 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d const int avg_q_len = num_seqs > 0 ? (total_q_tokens / num_seqs) : 0; const bool wmma_ok = amd_wmma_available(cc); const bool tile_env_on = get_paged_tile_mode() != 0; - const bool tile_gate_on = wmma_ok && tile_env_on; + // Tile kernel now has a portable FMA #else path (non-WMMA: Pascal + // sm_61, GCN gfx803) in addition to the RDNA3/4 WMMA path, so it is + // no longer gated on wmma_ok. GGML_PAGED_TILE=0 reverts to the slow + // scalar prefill kernel. (Multi-warp variant stays WMMA-only below.) + const bool tile_gate_on = tile_env_on; if (tile_gate_on && avg_q_len >= 16) { if (probe_on) { int n = probe_tile.fetch_add(1, std::memory_order_relaxed); @@ -1123,7 +1127,9 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d (const int32_t *) slot_mapping->data, (const int32_t *) q_lens->data, num_seqs, (int) k_cur->ne[2], n_kv_heads, stream); - const bool mw_on = get_paged_tile_multiwarp_mode() != 0; + // Multi-warp tile kernel is still WMMA-only (not yet ported); + // non-WMMA hardware uses the portable single-warp tile kernel. + const bool mw_on = get_paged_tile_multiwarp_mode() != 0 && wmma_ok; if (mw_on) { launch_paged_attn_tile_mw( (__half *) dst->data, @@ -1134,7 +1140,7 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d (const int32_t *) context_lens->data, (const int32_t *) q_lens->data, num_seqs, n_heads, n_kv_heads, max_bps, - avg_q_len, + total_q_tokens, // see note below: size grid.z by max, not avg, q_len scale, stream); } else { launch_paged_attn_tile( @@ -1146,7 +1152,12 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d (const int32_t *) context_lens->data, (const int32_t *) q_lens->data, num_seqs, n_heads, n_kv_heads, max_bps, - avg_q_len, // approximation; tile kernel skips q_tile_start >= q_len + total_q_tokens, // grid.z (n_q_tiles) must cover the LARGEST per-seq + // q_len. avg_q_len = total/num_seqs floors with idle + // parallel slots (1 active + N idle) and under-sizes + // the grid -> high Q-tiles never launch -> uncomputed + // rows -> garbage. total_q_tokens is a safe upper bound + // (kernel still skips tiles with q_tile_start >= q_len). scale, stream); } return; diff --git a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu index 6c41d2744b5a..38b13dd154a7 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu @@ -571,13 +571,159 @@ __global__ void mt_paged_attention_tile_kernel( } } #else - // Non-WMMA hardware should not be dispatched here; existing scalar - // kernel is the fallback. NO_DEVICE_CODE traps if we somehow are. - GGML_UNUSED(out); GGML_UNUSED(q); GGML_UNUSED(k_cache); GGML_UNUSED(v_cache); - GGML_UNUSED(block_tables); GGML_UNUSED(context_lens); GGML_UNUSED(q_lens); - GGML_UNUSED(max_blocks_per_seq); GGML_UNUSED(n_kv_heads); GGML_UNUSED(n_heads); - GGML_UNUSED(scale); - NO_DEVICE_CODE; + // ── Portable FMA tile flash-attention (non-WMMA: NVIDIA Pascal sm_61, + // AMD GCN gfx803). Mirrors the WMMA path's Q-tile + online-softmax + // structure, but the two matmuls run as explicit shared-memory GEMMs + // instead of matrix-core fragments. Lane ownership is kept identical to + // the WMMA fragment layout (row = tid%16, cols 8*(tid/16)+l), so the + // masking / softmax / writeback logic matches the #if branch line-for-line. + // One 32-lane warp per block; shuffles pass explicit WARP_SIZE so a 32-wide + // logical reduction is correct on gfx803's 64-lane waves too. + static_assert(HEAD_SIZE % K_INNER == 0, "HEAD_SIZE must be multiple of K_INNER=16"); + constexpr int N_INNER = HEAD_SIZE / K_INNER; + + const int head_idx = blockIdx.x; + const int seq_idx = blockIdx.y; + const int q_tile_idx = blockIdx.z; + const int tid = threadIdx.x; + + const int q_len = q_lens[seq_idx]; + const int q_tile_start = q_tile_idx * Q_TILE_M; + if (q_tile_start >= q_len) { + return; + } + const int q_tile_actual = (q_tile_start + Q_TILE_M <= q_len) ? Q_TILE_M : (q_len - q_tile_start); + + const int kv_head_idx = head_idx / (n_heads / n_kv_heads); + const int ctx_len_after_q = context_lens[seq_idx]; + const int * seq_block_table = block_tables + seq_idx * max_blocks_per_seq; + + size_t seq_q_offset = 0; + for (int s = 0; s < seq_idx; ++s) { + seq_q_offset += (size_t) q_lens[s]; + } + + const int q_pos_base = (ctx_len_after_q - q_len) + q_tile_start; + + // smem: Q[16,HS] + K[16,HS] + V[16,HS] half, then scores[16,16] f32. + extern __shared__ unsigned char smem_raw[]; + __half * smem_q = (__half *)(smem_raw); + __half * smem_k = smem_q + Q_TILE_M * HEAD_SIZE; + __half * smem_v = smem_k + K_TILE_N * HEAD_SIZE; + float * smem_s = (float *)(smem_v + K_TILE_N * HEAD_SIZE); // [Q_TILE_M * K_TILE_N] + + const size_t q_global_base = ((seq_q_offset + (size_t) q_tile_start) * (size_t) n_heads + (size_t) head_idx) + * (size_t) HEAD_SIZE; + stage_q_tile(smem_q, q, q_global_base, q_tile_actual, n_heads, tid); + __syncthreads(); + + // Lane ownership: row = tid%16 (Q row), half = tid/16 -> cols 8*half+l. + const int row = tid % 16; + const int half = tid / 16; + + // Online-softmax state + output accumulators (this lane's 8 cols per HS block). + float running_max = -INFINITY; + float running_sum = 0.0f; + float acc[N_INNER][8]; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) acc[n][l] = 0.0f; + } + + const int q_pos_last = q_pos_base + (q_tile_actual - 1); + const int valid_ctx = q_pos_last + 1; + const int q_pos = q_pos_base + row; + const bool row_valid = (row < q_tile_actual); + + for (int k_tile_start = 0; k_tile_start < valid_ctx; k_tile_start += K_TILE_N) { + stage_k_tile( + smem_k, k_cache, seq_block_table, k_tile_start, valid_ctx, + kv_head_idx, n_kv_heads, tid); + __syncthreads(); + + // scores[row, col] = scale * (Q[row] . K[col]) for this lane's 8 cols. + float sc[8]; + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int col = 8 * half + l; + float dot = 0.0f; + #pragma unroll + for (int d = 0; d < HEAD_SIZE; ++d) { + dot += __half2float(smem_q[row * HEAD_SIZE + d]) * __half2float(smem_k[col * HEAD_SIZE + d]); + } + const int k_pos = k_tile_start + col; + const bool visible = row_valid && (k_pos <= q_pos) && (k_pos < valid_ctx); + sc[l] = visible ? (dot * scale) : -INFINITY; + } + + // Per-row max: 8-wide local then pair-lane (tid^16) via shfl_xor. + float local_max = -INFINITY; + #pragma unroll + for (int l = 0; l < 8; ++l) local_max = max(local_max, sc[l]); + const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16, WARP_SIZE)); + const float new_max = max(running_max, row_max); + + // Rescale running state on new max. + if (running_max > -INFINITY) { + const float rescale = __expf(running_max - new_max); + running_sum *= rescale; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) acc[n][l] *= rescale; + } + } + + // exp(scores - new_max) + per-row sum. + float local_sum = 0.0f; + #pragma unroll + for (int l = 0; l < 8; ++l) { + const float e = (sc[l] == -INFINITY) ? 0.0f : __expf(sc[l] - new_max); + sc[l] = e; + local_sum += e; + smem_s[row * K_TILE_N + (8 * half + l)] = e; // publish full row for PV + } + running_sum += local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16, WARP_SIZE); + running_max = new_max; + + stage_v_tile( + smem_v, v_cache, seq_block_table, k_tile_start, valid_ctx, + kv_head_idx, n_kv_heads, tid); + __syncthreads(); // smem_s (all 16 cols) + smem_v visible to the warp + + // acc[n][l] += sum_k scores[row, k] * V[k, d], d = n*16 + 8*half + l. + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int d = n * K_INNER + 8 * half + l; + float pv = 0.0f; + #pragma unroll + for (int k = 0; k < K_TILE_N; ++k) { + pv += smem_s[row * K_TILE_N + k] * __half2float(smem_v[k * HEAD_SIZE + d]); + } + acc[n][l] += pv; + } + } + __syncthreads(); // before next iter overwrites smem_k / smem_v / smem_s + } + + // Writeback (same layout as the WMMA branch). + const float inv_sum = 1.0f / (running_sum + 1e-6f); + if (row < q_tile_actual) { + const int q_row_global = q_tile_start + row; + const size_t out_row_base = + ((seq_q_offset + (size_t) q_row_global) * (size_t) n_heads + (size_t) head_idx) * (size_t) HEAD_SIZE; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int d = n * K_INNER + 8 * half + l; + out[out_row_base + (size_t) d] = __float2half(acc[n][l] * inv_sum); + } + } + } #endif // AMD_WMMA_AVAILABLE } @@ -603,7 +749,10 @@ void launch_paged_attn_tile( dim3 grid(n_heads, num_seqs, n_q_tiles); dim3 block(TILE_NUM_THREADS); - const size_t smem_bytes = (size_t)(Q_TILE_M + 2 * K_TILE_N) * (size_t) HEAD_SIZE * sizeof(__half); + // half Q/K/V tiles + (non-WMMA path only) an f32 scores[Q_TILE_M*K_TILE_N] + // scratch. Allocated unconditionally; the WMMA branch simply ignores it. + const size_t smem_bytes = (size_t)(Q_TILE_M + 2 * K_TILE_N) * (size_t) HEAD_SIZE * sizeof(__half) + + (size_t) Q_TILE_M * (size_t) K_TILE_N * sizeof(float); mt_paged_attention_tile_kernel <<>>( From 1d75aee747bf3f4ab631d6762908d9bf1b21e896 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Mon, 15 Jun 2026 15:11:45 -0400 Subject: [PATCH 2/7] docs: paged-attn debug notes for 2026-06-15 session 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 --- docs/paged-attn-debug.md | 87 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 87 insertions(+) diff --git a/docs/paged-attn-debug.md b/docs/paged-attn-debug.md index 8032e96c81b2..bfa134f63b89 100644 --- a/docs/paged-attn-debug.md +++ b/docs/paged-attn-debug.md @@ -170,3 +170,90 @@ exercised). * Reverting the multi-seq commits — they aren't the cause. * Rewriting from scratch — the existing kernel is mostly right; it's one or two layout / shape assumption bugs. + +--- + +## 2026-06-15 session — non-WMMA prefill, gfx803 perf, checkpoint+paged, tool bleed + +Investigation on **mad-lab-2026** (GTX 1070 / sm_61 / CUDA `build-army`; +RX 480 / gfx803 / HIP `build-rocm-gfx803` in docker `rx480-army`). + +### A. Paged PREFILL was on the slow scalar kernel for ALL non-WMMA GPUs (fixed) + +`nsys` on a 2501-tok paged prefill (1070): `mt_paged_attention_kernel` += **84.7% of GPU time, 623 ms/launch**, grid `n_heads×num_seqs` (NO query +parallelism). Root cause: the fast query-tiled prefill kernel +(`mt_pagedattn_tile.cu`, grid `n_heads×num_seqs×n_q_tiles`) is gated at +`mt_pagedattn.cu` behind `tile_gate_on = wmma_ok && tile_env_on`, and +`wmma_ok = amd_wmma_available(cc)` is **RDNA3/4-only**. So Pascal (1070) +and GCN/gfx803 (480) both fall through to the scalar O(n_q·n_kv) kernel. +Decode was unaffected (`launch_paged_attn_decode` is not WMMA-gated). + +**Fix (uncommitted, working tree):** +1. Added a portable FMA `#else` body to `mt_paged_attention_tile_kernel` + (was `NO_DEVICE_CODE`): QK^T as smem dot-products, scores→`smem_s`, + scores·V as smem GEMM. Same lane layout as the WMMA path; shfl + reductions pass explicit `WARP_SIZE` (correct on gfx803 64-lane waves). + Footprint ~13 KiB @ HS128 (fits Pascal's 48 KiB). +2. Dispatch: `tile_gate_on = tile_env_on` (drop `wmma_ok`); force + `mw_on=false` on non-WMMA (multi-warp tile still WMMA-only). +3. **Also fixed a latent grid-sizing bug** (same class as the decode + `93a1e6a11` fix): the tile launches sized `n_q_tiles` from + `avg_q_len = total_q_tokens/num_seqs`, which floors with idle + `--parallel` slots (1 active + N idle) → high Q-tiles never launch → + uncomputed rows → token-salad. Now use `total_q_tokens`. **This also + fixes the WMMA path and is very likely the core of MAD-288.** + +**Verified:** 1070 (CUDA) paged prefill **160 → 328 t/s** (non-paged ref +395), output byte-matches the scalar oracle (`GGML_PAGED_TILE=0`) for +Qwen3.5-9B (HS128) and LFM2.5-8B-A1B (HS64 **padded to 128** in the paged +cache → uses the tile path; no separate HS64 work needed — clarifies +MAD-298). 480 (gfx803): output correct, no crash. Revert: `GGML_PAGED_TILE=0`. + +### B. gfx803 (RX 480) ROCm is ~10× slower than Vulkan — and it's NOT the attn kernel + +Same card, Qwen3.5-9B, `-p512 -n128`: + +| Path | Prefill | Decode | +|---|---|---| +| gfx803 ROCm/HIP, paged tile | 16.6 t/s | — | +| gfx803 ROCm/HIP, paged scalar | 13.8 t/s | — | +| **480 via Vulkan0 (RADV, f16 KV)** | **162 t/s** | 20.4 t/s | +| 1070 via Vulkan | 181 t/s | — | + +Tile ≈ scalar on gfx803 (both ~16) → the bottleneck is **gfx803 ROCm/HIP +itself**, not the attention kernel and not the Polaris silicon. Suspects: +turbo4 per-element dequant in `stage_k/v_tile` (`ops::k_load`), paged-block +gather, or `gated_delta_net` hybrid layers under HIP. **Decision (Kurtis): +keep the custom paged/turbo4/tiered features on gfx803, find the ROCm perf +bug — do NOT port to Vulkan** (Vulkan lacks all custom kernels; only viable +for standard dense models that fit VRAM). NEXT: profile the gfx803 paged +prefill (rocprof not installed in `rx480-army` — needs a profiling path). + +### C. Checkpoint reuse vs paged cache — block-alignment crash (NOT fixed) + +`--ctx-checkpoints>0` would let hybrid/recurrent models reuse the prompt +prefix (the restore path at `server-context.cpp` ~3397 already exists, and +with paged the checkpoints are tiny — **0.282 MiB**, blocks referenced not +copied). But restore sets `n_past` to the checkpoint pos (e.g. 1364), and +`llama_kv_cache_paged::seq_rm` requires **block-aligned (×16)** ranges → +unaligned trim leaves stale partial blocks → `GGML_ASSERT compute_slot_mapping +failed` abort. This is why LFM2.5 runs `--ctx-checkpoints 0` and full- +reprocesses every turn (~2.5 min for ~3700 tok). **Two fix attempts failed:** +round-at-restore is wrong (recurrent state can't be partially rolled back — +must reuse exactly `pos_max+1`); block-aligning the checkpoint batch-break at +creation crashed turn-1 prefill. **Reverted.** Needs a deeper design (paged +`seq_rm`/`compute_slot_mapping` tolerant of the boundary, or a hybrid+paged- +aware checkpoint path). Upstream unmerged references: PRs #20955, #21099, #20428. + +### D. Tool-schema bleed (LFM2.5) + +"hello!" → model regurgitates tool-schema JSON fragments then hallucinates. +The LFM2.5 jinja template renders the tool list as a **plain-text system +message** `List of tools: [{json}]` (faithful to the official template; +`common/chat.cpp:2226` detects this variant, vs LFM2's native +`<|tool_list_start|>` tokens at :622). Tool *calls* are correct +(grammar-constrained `<|tool_call_start|>`). With `thinking=1` the model +also reasons verbosely about the tools. Fix TBD (limit/disable reasoning on +tool turns, or revisit tool presentation). Single/multi-turn tool calls +otherwise correct on 1070. From 7e5d23d9abce6575cbc91dbcfb8f9dec911d80e8 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Mon, 15 Jun 2026 20:42:28 -0400 Subject: [PATCH 3/7] cuda/paged-attn: fix gfx803 decode under-utilization (CHUNK_KV + size-by-actual) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- ggml/src/ggml-cuda/mt_pagedattn_decode.cu | 44 ++++++++++++----------- 1 file changed, 24 insertions(+), 20 deletions(-) diff --git a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu index 2b876dc65038..b8562dfacdb2 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu @@ -70,7 +70,7 @@ static constexpr int DECODE_K_TILE_N = 16; // 1024 = 64 sub-chunks/block. At 400K ctx → ~400 chunks × n_heads = 12800 // blocks (well above RDNA4 R9700's ~240 concurrent block ceiling, so we // run in waves and amortize scratch). -static constexpr int CHUNK_KV = 1024; +static constexpr int CHUNK_KV = 256; // MAD-301A: was 1024 (tuned for 400K ctx / RDNA4 240-block ceiling); too coarse for gfx803's 36 CUs at mid ctx (ctx 1751 -> only 2 chunks -> few blocks -> CU under-utilization). 256 -> 4x more parallel blocks. TODO: make arch-tunable. // Threads/block. 4 warps (32 lanes each) — fits LDS budget for 2× tile // (smem_k + smem_v at HEAD_SIZE=128 = 4 KiB each = 8 KiB) plus Q+logits. @@ -519,21 +519,12 @@ __global__ void mt_paged_attention_decode_kernel( }; if (chunk_start >= valid_ctx_max) { - // No visible tokens for any query — write neutral partials so the - // reducer's pass over all chunks doesn't see uninitialized memory. - // Loop over (qh, qi) so we cover every head this block owns. - for (int qh = 0; qh < num_queries_per_kv; ++qh) { - const int head_idx = head_base + qh; - const size_t base = partial_chunk_base_for_head(head_idx); - for (int qi = 0; qi < q_len; ++qi) { - const size_t off = base + (size_t) qi * (HEAD_SIZE + 2); - for (int d = tid; d < HEAD_SIZE; d += DECODE_NUM_THREADS) partials[off + d] = 0.0f; - if (tid == 0) { - partials[off + HEAD_SIZE] = -INFINITY; - partials[off + HEAD_SIZE + 1] = 0.0f; - } - } - } + // MAD-301A: this chunk is entirely beyond the seq's real context + // (num_chunks is sized by ALLOCATED ctx, not actual — and with a small + // CHUNK_KV + large --ctx-size there are thousands of these). The reduce + // kernel now bounds its loop to ceil(context_lens/CHUNK_KV), so it never + // reads partials from these chunks — skip the neutral-partial writes and + // just retire the block. Makes decode cost scale with actual depth. return; } const int chunk_end = min(chunk_start + CHUNK_KV, valid_ctx_max); @@ -1058,9 +1049,11 @@ __global__ void mt_paged_attention_decode_reduce_kernel( __half * __restrict__ out, const float * __restrict__ partials, const int32_t * __restrict__ q_lens, + const int32_t * __restrict__ context_lens, + int num_chunks, + int chunk_kv, int n_heads, int n_seqs, - int num_chunks, int max_q_len) { const int head_idx = blockIdx.x; const int seq_idx = blockIdx.y; @@ -1071,6 +1064,16 @@ __global__ void mt_paged_attention_decode_reduce_kernel( for (int s = 0; s < seq_idx; ++s) seq_q_offset += (size_t) q_lens[s]; const int q_len = q_lens[seq_idx]; + // MAD-301A: num_chunks is sized by ALLOCATED ctx (max_blocks_per_seq*block_size). + // With small CHUNK_KV + a large --ctx-size, that's thousands of chunks while a + // decode step only fills context_lens[seq] tokens. Bound the reduction to the + // chunks that actually hold data so cost scales with real depth, not allocated + // capacity. Pass-1 chunk blocks beyond this range early-return without writing + // neutral partials, so they must never be read here. + const int ctx_len = context_lens[seq_idx]; + const int chunks_full = (ctx_len + chunk_kv - 1) / chunk_kv; + const int valid_chunks = chunks_full < num_chunks ? chunks_full : num_chunks; + // Stride from chunk c → chunk c+1 in the partials buffer. const size_t chunk_stride_q = (size_t) max_q_len * (size_t) (HEAD_SIZE + 2); // (head, seq) base. @@ -1082,7 +1085,7 @@ __global__ void mt_paged_attention_decode_reduce_kernel( // Pass 1: global max across chunks for this query position. float global_max = -INFINITY; - for (int c = 0; c < num_chunks; ++c) { + for (int c = 0; c < valid_chunks; ++c) { const float m = partials[partial_seq_base + (size_t) c * chunk_stride_q + qi_stride + HEAD_SIZE]; global_max = max(global_max, m); } @@ -1090,7 +1093,7 @@ __global__ void mt_paged_attention_decode_reduce_kernel( // Pass 2: merge across chunks. float global_sum = 0.0f; float v_d = 0.0f; - for (int c = 0; c < num_chunks; ++c) { + for (int c = 0; c < valid_chunks; ++c) { const size_t cbase = partial_seq_base + (size_t) c * chunk_stride_q + qi_stride; const float c_max = partials[cbase + HEAD_SIZE]; if (c_max == -INFINITY) continue; @@ -1187,7 +1190,8 @@ void launch_paged_attn_decode( dim3 block2(HEAD_SIZE); mt_paged_attention_decode_reduce_kernel <<>>( - out, partials_scratch, q_lens, n_heads, num_seqs, num_chunks, max_q_len); + out, partials_scratch, q_lens, context_lens, num_chunks, CHUNK_KV, + n_heads, num_seqs, max_q_len); } // ── explicit instantiations ──────────────────────────────────────────── From 9fbeba2304664dc36d68e1ab9a98a7dabef13bc6 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Mon, 15 Jun 2026 21:51:56 -0400 Subject: [PATCH 4/7] cuda/mmq: route gfx80x large-batch GEMM to dequant+hipBLAS (MAD-301B) 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 --- ggml/src/ggml-cuda/ggml-cuda.cu | 7 ++++++- ggml/src/ggml-cuda/mmq.cu | 8 ++++++++ ggml/src/ggml-cuda/mt_pagedattn_decode.cu | 2 +- 3 files changed, 15 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 789fc4fdb671..8ec8ef63b671 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2930,7 +2930,12 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * } } - if (ggml_cuda_should_use_mmq(src0->type, cc, ne12, /*n_experts=*/ne02)) { + // routing_active forces MMQ: it is the only mul_mat_id path that honors the + // weight-pager expert pointers here (MMVQ handled above for small batch). On + // gfx80x ggml_cuda_should_use_mmq() now returns false for large batches (no + // hardware dp4a -> route to dequant+hipBLAS), so without this OR the routing + // case would fall through to the GGML_ABORT below. + if (routing_active || ggml_cuda_should_use_mmq(src0->type, cc, ne12, /*n_experts=*/ne02)) { ggml_cuda_mul_mat_q(ctx, src0, src1, ids, dst); return; } diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 8ef5e6a68588..789f7d5295f8 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -477,5 +477,13 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return true; } + // gfx80x (Tonga/Fiji/Polaris, GGML_CUDA_CC_GCN4 and earlier) have no hardware dp4a + // (v_dot4): ggml_cuda_dp4a falls back to fully-scalar byte emulation, which makes MMQ + // ~8x slower than dequantization + hipBLAS for prefill-sized batches (rocBLAS ships + // gfx803 fp16/fp32 Tensile kernels). Restrict MMQ to small (decode) batches only. + if (GGML_CUDA_CC_IS_GCN(cc) && cc <= GGML_CUDA_CC_GCN4) { + return ne11 < MMQ_DP4A_MAX_BATCH_SIZE; + } + return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; } diff --git a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu index b8562dfacdb2..579d2ecb2876 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu @@ -70,7 +70,7 @@ static constexpr int DECODE_K_TILE_N = 16; // 1024 = 64 sub-chunks/block. At 400K ctx → ~400 chunks × n_heads = 12800 // blocks (well above RDNA4 R9700's ~240 concurrent block ceiling, so we // run in waves and amortize scratch). -static constexpr int CHUNK_KV = 256; // MAD-301A: was 1024 (tuned for 400K ctx / RDNA4 240-block ceiling); too coarse for gfx803's 36 CUs at mid ctx (ctx 1751 -> only 2 chunks -> few blocks -> CU under-utilization). 256 -> 4x more parallel blocks. TODO: make arch-tunable. +static constexpr int CHUNK_KV = 128; // MAD-301A: was 1024 (tuned for 400K ctx / RDNA4 240-block ceiling); too coarse for gfx803's 36 CUs at mid ctx (ctx 1751 -> only 2 chunks -> few blocks -> CU under-utilization). CHUNK_KV sweep on gfx803: 128 (56/38 t/s) > 256 (52/37) > 512 (45/23) @ctx 1751/10501. TODO: make arch-tunable. // Threads/block. 4 warps (32 lanes each) — fits LDS budget for 2× tile // (smem_k + smem_v at HEAD_SIZE=128 = 4 KiB each = 8 KiB) plus Q+logits. From 60dc49a1b0a4c4834f94106f928ea05dd7a9fd24 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Tue, 16 Jun 2026 10:06:08 -0400 Subject: [PATCH 5/7] cuda/paged-attn: portable FMA multi-warp tile kernel for gfx803 (MAD-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 --- ggml/src/ggml-cuda/mt_pagedattn.cu | 9 +- ggml/src/ggml-cuda/mt_pagedattn_tile.cu | 255 +++++++++++++++++++++++- 2 files changed, 255 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/mt_pagedattn.cu b/ggml/src/ggml-cuda/mt_pagedattn.cu index 45a21c27b6a0..a4e1c12a1f57 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn.cu @@ -1127,9 +1127,12 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d (const int32_t *) slot_mapping->data, (const int32_t *) q_lens->data, num_seqs, (int) k_cur->ne[2], n_kv_heads, stream); - // Multi-warp tile kernel is still WMMA-only (not yet ported); - // non-WMMA hardware uses the portable single-warp tile kernel. - const bool mw_on = get_paged_tile_multiwarp_mode() != 0 && wmma_ok; + // Multi-warp tile kernel now has a portable FMA #else path + // (MAD-301C), so non-WMMA hardware (gfx803, Pascal) uses it too: + // K/V staged once per block + shared across warps (HBM-reuse win) + // and N_WARPS*32 threads fill gfx803's 64-lane wave (occupancy win). + // GGML_PAGED_TILE_MULTIWARP=0 reverts to the single-warp tile kernel. + const bool mw_on = get_paged_tile_multiwarp_mode() != 0; if (mw_on) { launch_paged_attn_tile_mw( (__half *) dst->data, diff --git a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu index 38b13dd154a7..ac1805b62bea 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu @@ -1114,12 +1114,255 @@ __global__ void mt_paged_attention_tile_mw_kernel( } } #else - GGML_UNUSED(out); GGML_UNUSED(q); GGML_UNUSED(k_cache); GGML_UNUSED(v_cache); - GGML_UNUSED(block_tables); GGML_UNUSED(context_lens); GGML_UNUSED(q_lens); - GGML_UNUSED(max_blocks_per_seq); GGML_UNUSED(n_kv_heads); GGML_UNUSED(n_heads); - GGML_UNUSED(scale); - NO_DEVICE_CODE; -#endif + // ── Portable FMA multi-warp tile flash-attention (non-WMMA: NVIDIA + // Pascal sm_61, AMD GCN gfx803). Mirrors the WMMA mw path's block layout + // (Q_TILES warps/block, K/V staged once into smem and shared across warps + // = the HBM-reuse win) but runs the two matmuls as explicit shared-memory + // GEMMs per warp, identical math to the single-warp FMA #else above. On + // gfx803 the N_WARPS*32 threads fill the 64-lane wave (2 logical warps per + // physical wave), recovering the ~2x SIMD utilization the 32-thread + // single-warp tile kernel left on the table. No extra LDS over the WMMA + // layout: the PV step rebuilds each Q row's full 16-wide score vector from + // the pair lane via __shfl_xor(.,16) instead of an smem scratch. + static_assert(HEAD_SIZE % K_INNER == 0, "HEAD_SIZE must be multiple of K_INNER=16"); + constexpr int Q_TILES = TileConfig::Q_TILES_PER_BLOCK; + constexpr int N_WARPS = Q_TILES; + constexpr int N_THREADS = N_WARPS * 32; + constexpr int N_INNER = HEAD_SIZE / K_INNER; + + using ops = paged_cache_ops; + + const int head_idx = blockIdx.x; + const int seq_idx = blockIdx.y; + const int q_tile_grp_idx = blockIdx.z; + const int tid = threadIdx.x; + const int warp_id = tid >> 5; // tid / 32 + const int lane_id = tid & 31; // tid % 32 + + const int q_len = q_lens[seq_idx]; + const int q_tile_base = q_tile_grp_idx * Q_TILES; + + if (q_tile_base * Q_TILE_M >= q_len) { + return; + } + + const int my_q_tile_idx = q_tile_base + warp_id; + const int my_q_tile_start = my_q_tile_idx * Q_TILE_M; + const bool warp_active = (my_q_tile_start < q_len); + const int my_q_tile_actual = warp_active + ? ((my_q_tile_start + Q_TILE_M <= q_len) ? Q_TILE_M : (q_len - my_q_tile_start)) + : 0; + + const int kv_head_idx = head_idx / (n_heads / n_kv_heads); + const int ctx_len_after_q = context_lens[seq_idx]; + const int * seq_block_table = block_tables + seq_idx * max_blocks_per_seq; + + size_t seq_q_offset = 0; + for (int s = 0; s < seq_idx; ++s) { + seq_q_offset += (size_t) q_lens[s]; + } + + const int my_q_pos_base = (ctx_len_after_q - q_len) + my_q_tile_start; + + // Uniform K-loop bound across the block (all warps must hit the same + // cooperative __syncthreads); per-warp causal masking handles the rest. + const int block_last_q_row = min(q_tile_base * Q_TILE_M + Q_TILES * Q_TILE_M, q_len) - 1; + const int block_last_q_pos = (ctx_len_after_q - q_len) + block_last_q_row; + const int block_valid_ctx = block_last_q_pos + 1; + + extern __shared__ unsigned char smem_raw[]; + __half * smem_q = (__half *)(smem_raw); + __half * smem_k = smem_q + Q_TILES * Q_TILE_M * HEAD_SIZE; + __half * smem_v = smem_k + K_TILE_N * HEAD_SIZE; + + // ── cooperative Q load (all threads, all Q_TILES tiles) ── + { + constexpr int TOTAL_Q_ELEMS = Q_TILES * Q_TILE_M * HEAD_SIZE; + for (int idx = tid; idx < TOTAL_Q_ELEMS; idx += N_THREADS) { + const int qt = idx / (Q_TILE_M * HEAD_SIZE); + const int qt_off = idx % (Q_TILE_M * HEAD_SIZE); + const int qrow = qt_off / HEAD_SIZE; + const int qcol = qt_off % HEAD_SIZE; + + const int q_tile_idx_local = q_tile_base + qt; + const int q_row_global = q_tile_idx_local * Q_TILE_M + qrow; + + __half val = __float2half(0.0f); + if (q_row_global < q_len) { + const size_t base = ((seq_q_offset + (size_t) q_row_global) * (size_t) n_heads + + (size_t) head_idx) * (size_t) HEAD_SIZE; + val = q[base + (size_t) qcol]; + } + smem_q[idx] = val; + } + } + __syncthreads(); + + // Lane ownership within a warp: row = lane%16 (Q row), half = lane/16 -> + // this lane owns cols 8*half + l, l in [0,8). + const __half * my_smem_q = smem_q + warp_id * Q_TILE_M * HEAD_SIZE; + const int row = lane_id & 15; + const int half_id = lane_id >> 4; + + float running_max = -INFINITY; + float running_sum = 0.0f; + float acc[N_INNER][8]; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) acc[n][l] = 0.0f; + } + + const int q_pos = my_q_pos_base + row; + const bool row_valid = (row < my_q_tile_actual); + + for (int k_tile_start = 0; k_tile_start < block_valid_ctx; k_tile_start += K_TILE_N) { + // Cooperative K load (shared across all warps in the block). + if constexpr (CACHE_TYPE == GGML_TYPE_TURBO4_0) { + coop_stage_turbo4_tile( + smem_k, k_cache, seq_block_table, k_tile_start, block_valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); + } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO3_0) { + coop_stage_turbo3_tile( + smem_k, k_cache, seq_block_table, k_tile_start, block_valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); + } else { + for (int idx = tid; idx < K_TILE_N * HEAD_SIZE; idx += N_THREADS) { + const int krow = idx / HEAD_SIZE; + const int kcol = idx % HEAD_SIZE; + const int token = k_tile_start + krow; + float val = 0.0f; + if (token < block_valid_ctx) { + const int logical_block = token / BLOCK_SIZE; + const int tok_in_block = token % BLOCK_SIZE; + const int physical = seq_block_table[logical_block]; + if (physical != kInvalidBlockTableEntry) { + val = ops::k_load(k_cache, physical, kv_head_idx, n_kv_heads, tok_in_block, kcol); + } + } + smem_k[idx] = __float2half(val); + } + } + __syncthreads(); + + // QK: scores[row, col] = scale * (Q[row] . K[col]) for this lane's 8 cols. + float sc[8]; + if (warp_active) { + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int col = 8 * half_id + l; + float dot = 0.0f; + #pragma unroll + for (int d = 0; d < HEAD_SIZE; ++d) { + dot += __half2float(my_smem_q[row * HEAD_SIZE + d]) * __half2float(smem_k[col * HEAD_SIZE + d]); + } + const int k_pos = k_tile_start + col; + const bool visible = row_valid && (k_pos <= q_pos) && (k_pos < block_valid_ctx); + sc[l] = visible ? (dot * scale) : -INFINITY; + } + + // Per-row max: 8-wide local then pair-lane (lane^16) via shfl_xor. + float local_max = -INFINITY; + #pragma unroll + for (int l = 0; l < 8; ++l) local_max = max(local_max, sc[l]); + const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16, WARP_SIZE)); + const float new_max = max(running_max, row_max); + + // Rescale running state on new max. + if (running_max > -INFINITY) { + const float rescale = __expf(running_max - new_max); + running_sum *= rescale; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) acc[n][l] *= rescale; + } + } + + // exp(scores - new_max) + per-row sum. + float local_sum = 0.0f; + #pragma unroll + for (int l = 0; l < 8; ++l) { + const float e = (sc[l] == -INFINITY) ? 0.0f : __expf(sc[l] - new_max); + sc[l] = e; + local_sum += e; + } + running_sum += local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16, WARP_SIZE); + running_max = new_max; + } + + // Cooperative V load (same dispatch as K; turbo* V shares K layout). + if constexpr (CACHE_TYPE == GGML_TYPE_TURBO4_0) { + coop_stage_turbo4_tile( + smem_v, v_cache, seq_block_table, k_tile_start, block_valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); + } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO3_0) { + coop_stage_turbo3_tile( + smem_v, v_cache, seq_block_table, k_tile_start, block_valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); + } else { + for (int idx = tid; idx < K_TILE_N * HEAD_SIZE; idx += N_THREADS) { + const int vrow = idx / HEAD_SIZE; + const int vcol = idx % HEAD_SIZE; + const int token = k_tile_start + vrow; + float val = 0.0f; + if (token < block_valid_ctx) { + const int logical_block = token / BLOCK_SIZE; + const int tok_in_block = token % BLOCK_SIZE; + const int physical = seq_block_table[logical_block]; + if (physical != kInvalidBlockTableEntry) { + val = ops::v_load(v_cache, physical, kv_head_idx, n_kv_heads, tok_in_block, vcol); + } + } + smem_v[idx] = __float2half(val); + } + } + __syncthreads(); // smem_v visible to the warp + + // PV: acc[n][l] += sum_c scores[row, c] * V[c, d], d = n*16 + 8*half + l. + // This lane owns only its 8 exp-scores (cols 8*half+l); fetch the pair + // lane's 8 via shfl_xor(.,16) to rebuild the full 16-wide row. + if (warp_active) { + float s16[K_TILE_N]; + #pragma unroll + for (int l = 0; l < 8; ++l) { + const float partner = __shfl_xor_sync(0xFFFFFFFF, sc[l], 16, WARP_SIZE); + s16[8 * half_id + l] = sc[l]; + s16[8 * (1 - half_id) + l] = partner; + } + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int d = n * K_INNER + 8 * half_id + l; + float pv = 0.0f; + #pragma unroll + for (int c = 0; c < K_TILE_N; ++c) { + pv += s16[c] * __half2float(smem_v[c * HEAD_SIZE + d]); + } + acc[n][l] += pv; + } + } + } + __syncthreads(); // before next iter overwrites smem_k / smem_v + } + + // ── per-warp output writeback ─────────────────────────────────────── + if (warp_active && row < my_q_tile_actual) { + const float inv_sum = 1.0f / (running_sum + 1e-6f); + const int q_row_global = my_q_tile_start + row; + const size_t out_row_base = + ((seq_q_offset + (size_t) q_row_global) * (size_t) n_heads + (size_t) head_idx) * (size_t) HEAD_SIZE; + #pragma unroll + for (int n = 0; n < N_INNER; ++n) { + #pragma unroll + for (int l = 0; l < 8; ++l) { + const int d = n * K_INNER + 8 * half_id + l; + out[out_row_base + (size_t) d] = __float2half(acc[n][l] * inv_sum); + } + } + } +#endif // AMD_WMMA_AVAILABLE } template From 58504209c6a6f0e5923aa60957dec865cc152df7 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Tue, 16 Jun 2026 11:39:20 -0400 Subject: [PATCH 6/7] cuda/paged-attn: native head_dim-64 turbo4 KV (GGML_TYPE_TURBO4_64, MAD-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, 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 --- ggml/include/ggml.h | 3 +- ggml/src/ggml-common.h | 13 ++ ggml/src/ggml-cuda/ggml-cuda.cu | 3 +- ggml/src/ggml-cuda/mt_pagedattn.cu | 143 +++++++++++++++++++++- ggml/src/ggml-cuda/mt_pagedattn_decode.cu | 77 ++++++++++++ ggml/src/ggml-cuda/mt_pagedattn_ops.cuh | 35 ++++++ ggml/src/ggml-cuda/mt_pagedattn_tile.cu | 17 +++ ggml/src/ggml-cuda/turbo-quant.cuh | 8 ++ ggml/src/ggml-quants.h | 5 + ggml/src/ggml-turbo-quant.c | 61 +++++++++ ggml/src/ggml.c | 8 ++ src/llama-kv-cache-paged.cpp | 20 ++- 12 files changed, 387 insertions(+), 6 deletions(-) diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index ae6523a9f270..d1c5b3dbe5de 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -439,7 +439,8 @@ extern "C" { GGML_TYPE_F8_E4M3 = 49, // MAD-223 Phase G: fp8 e4m3 1-byte storage, used as sidecar dtype for ml8 centroids GGML_TYPE_ML8_4_SOA = 50, // MAD-223 Phase G.7 / MAD-244: ml8-4 with stored-as-repacked SOA layout (b_packed bytes followed by b_scale fp32 per expert row). Same numerics as ML8_4; eliminates the runtime MoE repack cache. GGML_TYPE_ML8_FP8 = 51, // MAD Task 9: ml8 fp8 weight quant, 32-element blocks, fp16 per-block scale + 32 e4m3 bytes (34 bytes/block). On-disk id=51; matches gguf-py GGML_QUANT_SIZES[ML8_FP8]=(32,34). - GGML_TYPE_COUNT = 52, + GGML_TYPE_TURBO4_64 = 52, // MAD-301C Lever B: native head_dim-64 turbo4 KV cache, 64-element block (34 bytes), no 64->128 pad. Runtime KV-cache-only type (never serialized to GGUF). + GGML_TYPE_COUNT = 53, }; // precision diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 5ad2f2a5f209..282e5d852676 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -329,6 +329,19 @@ static_assert(sizeof(block_turbo4_0) == 2*sizeof(ggml_half) + QK_TURBO4*3/8 + QK static_assert(QK_TURBO4 == 128, "turbo4 kernels assume QK_TURBO4 == 128"); +// MAD-301C Lever B: native head_dim-64 turbo4 KV cache block. +// Identical 4-bit PolarQuant + per-block norm as block_turbo4_0, but a 64-element +// block so a head_dim-64 head (LFM2.5, gpt-oss) is exactly one block with NO +// 64->128 zero-padding. The paged path skips RHT, so dequant returns +// centroid*norm and dot products match the padded-128 path bit-for-bit at ~half +// the storage (34 vs 68 bytes/head). 4-bit centroid layout only. +#define QK_TURBO4_64 64 +typedef struct { + ggml_half norm; // 2 bytes: per-block scale + uint8_t qs[QK_TURBO4_64 / 2]; // 32 bytes: 4-bit PolarQuant indices (nibble packed) +} block_turbo4_64; // 34 bytes total +static_assert(sizeof(block_turbo4_64) == 2 + QK_TURBO4_64/2, "wrong turbo4_64 block size"); + // TurboQuant 2-bit: 2-bit PolarQuant indices only (no QJL) // Per block: norm(fp16) + 2-bit indices (8 bytes) = 10 bytes per 32 values // = 2.5 bits/value → 6.4× compression vs fp16 diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 8ec8ef63b671..6d4f58fd22db 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5945,7 +5945,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g || op->src[1]->type == GGML_TYPE_Q8_0 || op->src[1]->type == GGML_TYPE_TURBO4_0 || op->src[1]->type == GGML_TYPE_TURBO3_0 - || op->src[1]->type == GGML_TYPE_TURBO4_FP8_BS256) // MAD-214 + || op->src[1]->type == GGML_TYPE_TURBO4_FP8_BS256 // MAD-214 + || op->src[1]->type == GGML_TYPE_TURBO4_64) // MAD-301C Lever B && op->src[6] // k_cur (fused scatter) && op->src[6]->type == GGML_TYPE_F16 && op->src[7] // v_cur (fused scatter) diff --git a/ggml/src/ggml-cuda/mt_pagedattn.cu b/ggml/src/ggml-cuda/mt_pagedattn.cu index a4e1c12a1f57..fbf33325d414 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn.cu @@ -454,6 +454,113 @@ __global__ void mt_scatter_kv_turbo4_0_kernel( } } +// Turbo4_64 scatter (MAD-301C Lever B): native head_dim-64 cooperative quantize. +// Identical pipeline to mt_scatter_kv_turbo4_0_kernel but a 64-element block +// (64 threads / 2 warps), writing block_turbo4_64 (no rnorm field). No RHT. +template +__launch_bounds__(QK_TURBO4_64) +__global__ void mt_scatter_kv_turbo4_64_kernel( + void * __restrict__ k_cache, + void * __restrict__ v_cache, + const __half * __restrict__ k_cur, + const __half * __restrict__ v_cur, + const int32_t * __restrict__ slot_mapping, + int n_kv_heads) { + + constexpr int Q_BLOCK = QK_TURBO4_64; // 64 + constexpr int N_QBLOCKS_PER_TOKEN = HEAD_SIZE / Q_BLOCK; + constexpr int N_WARPS = Q_BLOCK / WARP_SIZE; // 2 + static_assert(HEAD_SIZE % Q_BLOCK == 0, "HEAD_SIZE must be divisible by QK_TURBO4_64"); + + const int j = threadIdx.x; // 0..63 + const int global_token_idx = blockIdx.x; + const int y_idx = blockIdx.y; + const int kv_select = blockIdx.z; // 0 = K, 1 = V + const int kv_head_idx = y_idx / N_QBLOCKS_PER_TOKEN; + const int qb_idx = y_idx % N_QBLOCKS_PER_TOKEN; + + const int slot = slot_mapping[global_token_idx]; + if (slot < 0) return; // padding token + + const int paged_block = slot / BLOCK_SIZE; + const int slot_in_block = slot % BLOCK_SIZE; + + const int d = qb_idx * Q_BLOCK + j; + const __half * src = (kv_select == 0) ? k_cur : v_cur; + const size_t src_off = (size_t) global_token_idx * n_kv_heads * HEAD_SIZE + + (size_t) kv_head_idx * HEAD_SIZE + + (size_t) d; + + void * dst_buf = (kv_select == 0) ? k_cache : v_cache; + const int64_t block_ib = ((int64_t) paged_block * n_kv_heads + kv_head_idx) * BLOCK_SIZE * N_QBLOCKS_PER_TOKEN + + (int64_t) slot_in_block * N_QBLOCKS_PER_TOKEN + + (int64_t) qb_idx; + block_turbo4_64 * blk = (block_turbo4_64 *) dst_buf + block_ib; + + __shared__ float x[Q_BLOCK]; + x[j] = __half2float(src[src_off]); + __syncthreads(); + + __shared__ float warp_accum[N_WARPS]; + { + float v_sq = x[j] * x[j]; + for (int offset = WARP_SIZE / 2; offset > 0; offset >>= 1) { + v_sq += __shfl_xor_sync(0xffffffffu, v_sq, offset); + } + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = v_sq; + } + __syncthreads(); + + __shared__ float s_norm_sq; + if (j == 0) { + float total = 0.0f; + for (int w = 0; w < N_WARPS; ++w) total += warp_accum[w]; + s_norm_sq = total; + } + __syncthreads(); + const float grp_norm = sqrtf(s_norm_sq); + const float inv_norm = (grp_norm > 1e-10f) ? (1.0f / grp_norm) : 0.0f; + + x[j] *= inv_norm; + __syncthreads(); + + // No RHT (see mt_scatter_kv_turbo4_0_kernel step-4 comment): centroid-quant + // normalized K directly; dequant returns centroid*norm ~= K. + const float rv = x[j]; + const uint8_t idx = turbo_nearest_centroid_4bit(rv); + + const int lane = j % WARP_SIZE; + const uint8_t my_nibble = idx & 0xF; + const uint8_t partner_nibble = __shfl_sync(0xffffffffu, my_nibble, lane ^ 1, WARP_SIZE); + if ((j & 1) == 0) { + blk->qs[j / 2] = my_nibble | (partner_nibble << 4); + } + + { + const float c = TURBO_CENTROIDS_4BIT[idx]; + float rc = c * c; + for (int offset = WARP_SIZE / 2; offset > 0; offset >>= 1) { + rc += __shfl_xor_sync(0xffffffffu, rc, offset); + } + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = rc; + } + __syncthreads(); + + __shared__ float s_recon_sq; + if (j == 0) { + float total = 0.0f; + for (int w = 0; w < N_WARPS; ++w) total += warp_accum[w]; + s_recon_sq = total; + } + __syncthreads(); + const float recon_norm = sqrtf(s_recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? (grp_norm / recon_norm) : grp_norm; + + if (j == 0) { + blk->norm = __float2half(corrected_norm); + } +} + // Turbo3_0 scatter: per-128-element-block cooperative quantize. // // Same threading and pipeline as turbo4 above, but packs 3-bit indices into @@ -641,6 +748,18 @@ static void launch_scatter_kv( mt_scatter_kv_turbo4_0_kernel <<>>( k_cache, v_cache, k_cur, v_cur, slot_mapping, n_kv_heads); + } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO4_64) { + // MAD-301C Lever B: native head_dim-64 turbo4. Same grid topology as + // TURBO4_0 but 64 threads/block (one 64-element block per token/head). + constexpr int Q_BLOCK = QK_TURBO4_64; + constexpr int N_QBLOCKS_PER_TOKEN = HEAD_SIZE / Q_BLOCK; + dim3 grid(num_tokens_total, n_kv_heads * N_QBLOCKS_PER_TOKEN, 2); + dim3 block(Q_BLOCK); + GGML_UNUSED(q_lens); + GGML_UNUSED(num_seqs); + mt_scatter_kv_turbo4_64_kernel + <<>>( + k_cache, v_cache, k_cur, v_cur, slot_mapping, n_kv_heads); } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO3_0) { // Same grid topology as TURBO4_0 — one CUDA block per // (token, (kv_head, qb_idx), K-or-V); 128 threads per block. @@ -1097,9 +1216,14 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d HS, (int) QK_TURBO3); } else { - // MAD-180: WMMA tile FA gate. - if constexpr (((HS == 128) || (HS == 256)) && (BS == 16) - && (CT == GGML_TYPE_F16 || CT == GGML_TYPE_TURBO4_0 || CT == GGML_TYPE_TURBO3_0)) { + // MAD-180: WMMA tile FA gate. HS=128/256 keep F16/TURBO4_0/TURBO3_0. + // MAD-301C Lever B adds HS=64 paired ONLY with GGML_TYPE_TURBO4_64 + // (native head_dim-64 turbo4) — HS=64 for other cache types stays on + // the scalar path, so no <64,*> tile instances are needed for them. + if constexpr ((BS == 16) + && ((((HS == 128) || (HS == 256)) + && (CT == GGML_TYPE_F16 || CT == GGML_TYPE_TURBO4_0 || CT == GGML_TYPE_TURBO3_0)) + || (HS == 64 && CT == GGML_TYPE_TURBO4_64))) { const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int total_q_tokens = (int) k_cur->ne[2]; const int avg_q_len = num_seqs > 0 ? (total_q_tokens / num_seqs) : 0; @@ -1179,6 +1303,9 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d // (rare) heterogeneous case we use it as max_q_len so // the partials buffer's inner stride fits the largest // seq's queries. + // MAD-301C Lever B: GGML_TYPE_TURBO4_64 now has a flash-decode + // path (decode_coop_stage_turbo4_64), so it uses the same gate as + // the other turbo types — no exclusion. const bool decode_env_on = get_paged_decode_mode() != 0; // **Decode regression root cause (2026-05-18 rocprof hunt)**: // the original gate `avg_q_len >= 1` used integer division @@ -1307,6 +1434,16 @@ void ggml_cuda_op_paged_attn_mt(ggml_backend_cuda_context & ctx, ggml_tensor * d case GGML_TYPE_TURBO4_0: run_typed(std::integral_constant{}); break; + case GGML_TYPE_TURBO4_64: + // MAD-301C Lever B: native head_dim-64 turbo4 — only valid at HS=64. + // Guard so run_typed (and its kernel instances) is only + // instantiated for HS=64, not 128/256. + if constexpr (HS == 64) { + run_typed(std::integral_constant{}); + } else { + GGML_ABORT("mt_paged_attn: GGML_TYPE_TURBO4_64 only supports head_size=64 (got %d)", HS); + } + break; case GGML_TYPE_TURBO3_0: run_typed(std::integral_constant{}); break; diff --git a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu index 579d2ecb2876..278f0f355028 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_decode.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_decode.cu @@ -257,6 +257,70 @@ static __device__ __forceinline__ void decode_coop_stage_turbo4( } } +// TURBO4_64 cooperative dequant for the decode tile (MAD-301C Lever B). +// Native head_dim-64: 64-element block => 32 lanes × 2 elements (1 qs byte/lane, +// 2 nibbles). Matches mt_scatter_kv_turbo4_64_kernel packing. +template +static __device__ __forceinline__ void decode_coop_stage_turbo4_64( + __half * __restrict__ smem_dst, + const void * __restrict__ cache, + const int * __restrict__ seq_block_table, + int tile_start, + int valid_ctx, + int kv_head_idx, + int n_kv_heads, + int warp_id, + int lane_id) { + constexpr int Q_BLOCK = QK_TURBO4_64; // 64 + constexpr int QBLOCKS_PER_TOKEN = HEAD_SIZE / Q_BLOCK; + constexpr int N_QBLOCKS_PER_TILE = DECODE_K_TILE_N * QBLOCKS_PER_TOKEN; + static_assert(HEAD_SIZE % Q_BLOCK == 0, "HEAD_SIZE must be multiple of QK_TURBO4_64=64"); + static_assert(Q_BLOCK == 64, "cooperative dequant expects QK_TURBO4_64=64 (32 lanes × 2 elements)"); + + const block_turbo4_64 * blocks = (const block_turbo4_64 *) cache; + + #pragma unroll + for (int qb = warp_id; qb < N_QBLOCKS_PER_TILE; qb += DECODE_NUM_WARPS) { + const int row = qb / QBLOCKS_PER_TOKEN; + const int qb_in_token = qb % QBLOCKS_PER_TOKEN; + const int token = tile_start + row; + + const block_turbo4_64 * blk = nullptr; + float norm_f = 0.0f; + + if (token < valid_ctx) { + const int logical_block = token / BLOCK_SIZE; + const int tok_in_block = token % BLOCK_SIZE; + const int physical = seq_block_table[logical_block]; + if (physical != kInvalidBlockTableEntry) { + const int64_t ib = ((int64_t) physical * n_kv_heads + kv_head_idx) * BLOCK_SIZE * QBLOCKS_PER_TOKEN + + (int64_t) tok_in_block * QBLOCKS_PER_TOKEN + + (int64_t) qb_in_token; + blk = &blocks[ib]; + if (lane_id == 0) { + norm_f = __half2float(blk->norm); + } + } + } + norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0, WARP_SIZE); + + uint8_t packed = 0; + if (blk != nullptr) { + packed = blk->qs[lane_id]; // 1 byte = 2 nibbles (elements 2*lane, 2*lane+1) + } + + const int smem_row_base = row * HEAD_SIZE; + const int smem_col_base = qb_in_token * Q_BLOCK + lane_id * 2; + + #pragma unroll + for (int l = 0; l < 2; ++l) { + const uint8_t idx_nib = (packed >> (l * 4)) & 0xF; + const float val = TURBO_CENTROIDS_4BIT[idx_nib] * norm_f; + smem_dst[smem_row_base + smem_col_base + l] = __float2half(val); + } + } +} + // TURBO3_0 cooperative dequant for the decode tile. Same threading shape as // decode_coop_stage_turbo4 (32 lanes × 4 elements per qblock), but unpacks // the 3-bit index as (qs low-2 | signs hi-1). @@ -346,6 +410,10 @@ static __device__ __forceinline__ void decode_stage_k( decode_coop_stage_turbo3( smem_dst, cache, seq_block_table, tile_start, valid_ctx, kv_head_idx, n_kv_heads, warp_id, lane_id); + } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO4_64) { + decode_coop_stage_turbo4_64( + smem_dst, cache, seq_block_table, tile_start, valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); } else { decode_stage_kv_f16( smem_dst, (const __half *) cache, seq_block_table, tile_start, valid_ctx, @@ -373,6 +441,10 @@ static __device__ __forceinline__ void decode_stage_v( decode_coop_stage_turbo3( smem_dst, cache, seq_block_table, tile_start, valid_ctx, kv_head_idx, n_kv_heads, warp_id, lane_id); + } else if constexpr (CACHE_TYPE == GGML_TYPE_TURBO4_64) { + decode_coop_stage_turbo4_64( + smem_dst, cache, seq_block_table, tile_start, valid_ctx, + kv_head_idx, n_kv_heads, warp_id, lane_id); } else { decode_stage_kv_f16( smem_dst, (const __half *) cache, seq_block_table, tile_start, valid_ctx, @@ -1228,5 +1300,10 @@ template void launch_paged_attn_decode<256, 16, GGML_TYPE_TURBO3_0>( __half *, const __half *, const void *, const void *, const int32_t *, const int32_t *, const int32_t *, float *, int, int, int, int, int, int, float, cudaStream_t); +// HEAD_SIZE=64 — MAD-301C Lever B native head_dim-64 turbo4 flash-decode. +template void launch_paged_attn_decode<64, 16, GGML_TYPE_TURBO4_64>( + __half *, const __half *, const void *, const void *, + const int32_t *, const int32_t *, const int32_t *, + float *, int, int, int, int, int, int, float, cudaStream_t); } // namespace mt diff --git a/ggml/src/ggml-cuda/mt_pagedattn_ops.cuh b/ggml/src/ggml-cuda/mt_pagedattn_ops.cuh index 096ca18527ca..915757ecd1d9 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_ops.cuh +++ b/ggml/src/ggml-cuda/mt_pagedattn_ops.cuh @@ -152,6 +152,41 @@ struct paged_cache_ops { } }; +// Turbo4_64 specialization (MAD-301C Lever B): native head_dim-64 4-bit PolarQuant, +// 64-element blocks. Same layout shape as TURBO4_0 (1 block/token for head_dim 64, +// since N_QBLOCKS_PER_TOKEN = 64/64 = 1) and same no-RHT dequant; only the block +// width (64 vs 128) and struct (block_turbo4_64, 34 bytes) differ. No 64->128 pad. +// kv_store omitted — handled by mt_scatter_kv_turbo4_64_kernel. +template +struct paged_cache_ops { + static constexpr int Q_BLOCK = QK_TURBO4_64; // 64 + static constexpr int N_QBLOCKS_PER_TOKEN = HEAD_SIZE / Q_BLOCK; + static_assert(HEAD_SIZE % Q_BLOCK == 0, "HEAD_SIZE must be divisible by QK_TURBO4_64"); + + __device__ __forceinline__ static int64_t element_block_index( + int paged_block, int kv_head, int n_kv_heads, int token_in_block, int d) { + return ((int64_t) paged_block * n_kv_heads + kv_head) * BLOCK_SIZE * N_QBLOCKS_PER_TOKEN + + (int64_t) token_in_block * N_QBLOCKS_PER_TOKEN + + (int64_t) (d / Q_BLOCK); + } + + __device__ __forceinline__ static float k_load( + const void * buf, int paged_block, int kv_head, int n_kv_heads, + int token_in_block, int d) { + const block_turbo4_64 * blocks = (const block_turbo4_64 *) buf; + const int64_t ib = element_block_index(paged_block, kv_head, n_kv_heads, token_in_block, d); + const int iqs = d % Q_BLOCK; + const float norm = __half2float(blocks[ib].norm); + return turbo4_64_dequant_element(&blocks[ib], iqs, norm); + } + + __device__ __forceinline__ static float v_load( + const void * buf, int paged_block, int kv_head, int n_kv_heads, + int token_in_block, int d) { + return k_load(buf, paged_block, kv_head, n_kv_heads, token_in_block, d); + } +}; + // Turbo3_0 specialization: 3-bit PolarQuant with WHT rotation, 128-element blocks. // Same N_QBLOCKS_PER_TOKEN layout as TURBO4_0; differs in block payload size // (14 vs 68 bytes) and dequant centroid table (3-bit Lloyd-Max). diff --git a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu index ac1805b62bea..e269f2246d6a 100644 --- a/ggml/src/ggml-cuda/mt_pagedattn_tile.cu +++ b/ggml/src/ggml-cuda/mt_pagedattn_tile.cu @@ -64,6 +64,13 @@ static constexpr int K_INNER = 16; template struct TileConfig; +template <> +struct TileConfig<64> { + // MAD-301C Lever B: native head_dim-64 turbo4. LDS at 8 tiles: + // 8*16*64*2 + 2*16*64*2 = 16 KiB + 4 KiB = 20 KiB — well under the 64 KiB cap. + static constexpr int Q_TILES_PER_BLOCK = 8; +}; + template <> struct TileConfig<128> { static constexpr int Q_TILES_PER_BLOCK = 8; @@ -1428,6 +1435,11 @@ template void launch_paged_attn_tile<256, 16, GGML_TYPE_TURBO3_0>( __half *, const __half *, const void *, const void *, const int32_t *, const int32_t *, const int32_t *, int, int, int, int, int, float, cudaStream_t); +// HEAD_SIZE=64 — MAD-301C Lever B native head_dim-64 turbo4 (LFM2.5, gpt-oss). +template void launch_paged_attn_tile<64, 16, GGML_TYPE_TURBO4_64>( + __half *, const __half *, const void *, const void *, + const int32_t *, const int32_t *, const int32_t *, + int, int, int, int, int, float, cudaStream_t); // Multi-warp launcher instantiations — same (HEAD_SIZE, BLOCK_SIZE, CACHE_TYPE) // matrix as the single-warp path. TileConfig picks Q_TILES_PER_BLOCK at @@ -1456,5 +1468,10 @@ template void launch_paged_attn_tile_mw<256, 16, GGML_TYPE_TURBO3_0>( __half *, const __half *, const void *, const void *, const int32_t *, const int32_t *, const int32_t *, int, int, int, int, int, float, cudaStream_t); +// HEAD_SIZE=64 — MAD-301C Lever B native head_dim-64 turbo4. +template void launch_paged_attn_tile_mw<64, 16, GGML_TYPE_TURBO4_64>( + __half *, const __half *, const void *, const void *, + const int32_t *, const int32_t *, const int32_t *, + int, int, int, int, int, float, cudaStream_t); } // namespace mt diff --git a/ggml/src/ggml-cuda/turbo-quant.cuh b/ggml/src/ggml-cuda/turbo-quant.cuh index 9edd913ae6a1..cdbc3075116f 100644 --- a/ggml/src/ggml-cuda/turbo-quant.cuh +++ b/ggml/src/ggml-cuda/turbo-quant.cuh @@ -351,6 +351,14 @@ static __device__ __forceinline__ float turbo4_dequant_element( return TURBO_CENTROIDS_4BIT[idx] * norm; } +// MAD-301C Lever B: native head_dim-64 turbo4 dequant. Same nibble layout as +// block_turbo4_0 (j/2 byte, (j%2)*4 shift); only the block struct width differs. +static __device__ __forceinline__ float turbo4_64_dequant_element( + const block_turbo4_64 * __restrict__ x, int j, float norm) { + uint8_t idx = (x->qs[j / 2] >> ((j % 2) * 4)) & 0xF; + return TURBO_CENTROIDS_4BIT[idx] * norm; +} + // ---- Nearest 3-bit centroid index ---- static __device__ __forceinline__ uint8_t turbo_nearest_centroid_3bit(float val) { diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index db271e517931..adb0440d3f28 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -107,6 +107,11 @@ GGML_API void quantize_row_turbo3_0_ref(const float * GGML_RESTRICT x, block_tur GGML_API void quantize_row_turbo4_0_ref(const float * GGML_RESTRICT x, block_turbo4_0 * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_turbo3_0(const block_turbo3_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_turbo4_0(const block_turbo4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +// MAD-301C Lever B: native head_dim-64 turbo4 (64-element block). CPU refs exist to +// satisfy ggml type-traits; the hot path is CUDA paged (mt_pagedattn). No-RHT, matching +// the paged scatter convention (dequant = centroid*norm). +GGML_API void quantize_row_turbo4_64_ref(const float * GGML_RESTRICT x, block_turbo4_64 * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_turbo4_64(const block_turbo4_64 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API size_t quantize_turbo3_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_turbo4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API void quantize_row_turbo2_0_ref(const float * GGML_RESTRICT x, block_turbo2_0 * GGML_RESTRICT y, int64_t k); diff --git a/ggml/src/ggml-turbo-quant.c b/ggml/src/ggml-turbo-quant.c index 992c6aa1249f..bba63572c649 100644 --- a/ggml/src/ggml-turbo-quant.c +++ b/ggml/src/ggml-turbo-quant.c @@ -609,6 +609,67 @@ void dequantize_row_turbo4_0(const block_turbo4_0 * GGML_RESTRICT x, float * GGM #endif } +/* MAD-301C Lever B: native head_dim-64 turbo4 (64-element block, no RHT). + * Mirrors the CUDA paged scatter convention (mt_scatter_kv_turbo4_64_kernel): + * L2-norm -> normalize -> 4-bit centroid quant (NO WHT) -> recon-norm correction; + * dequant returns centroid*norm. These CPU refs exist only to satisfy ggml + * type-traits for GGML_TYPE_TURBO4_64; the hot path is CUDA paged attention. */ +void quantize_row_turbo4_64_ref(const float * GGML_RESTRICT x, block_turbo4_64 * GGML_RESTRICT y, int64_t k) { + static const float CENTROIDS_4BIT[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, + -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, + 0.068756f, 0.089527f, 0.117195f, 0.173926f + }; + assert(k % QK_TURBO4_64 == 0); + const int nb = k / QK_TURBO4_64; + const int d = QK_TURBO4_64; + for (int block = 0; block < nb; block++) { + const float * src = x + block * d; + float norm_sq = 0.0f; + for (int i = 0; i < d; i++) norm_sq += src[i] * src[i]; + const float norm = sqrtf(norm_sq); + float normalized[QK_TURBO4_64]; + if (norm > 1e-10f) { + const float inv = 1.0f / norm; + for (int i = 0; i < d; i++) normalized[i] = src[i] * inv; + } else { + memset(normalized, 0, d * sizeof(float)); + } + uint8_t indices[QK_TURBO4_64]; + for (int i = 0; i < d; i++) indices[i] = (uint8_t)nearest_centroid_4bit(normalized[i]); + float recon_sq = 0.0f; + for (int i = 0; i < d; i++) recon_sq += CENTROIDS_4BIT[indices[i]] * CENTROIDS_4BIT[indices[i]]; + const float recon_norm = sqrtf(recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? norm / recon_norm : norm; + y[block].norm = GGML_FP32_TO_FP16(corrected_norm); + memset(y[block].qs, 0, d / 2); + for (int i = 0; i < d; i++) { + y[block].qs[i / 2] |= (uint8_t)((indices[i] & 0xF) << ((i % 2) * 4)); + } + } +} + +void dequantize_row_turbo4_64(const block_turbo4_64 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + static const float CENTROIDS_4BIT[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, + -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, + 0.068756f, 0.089527f, 0.117195f, 0.173926f + }; + assert(k % QK_TURBO4_64 == 0); + const int nb = k / QK_TURBO4_64; + const int d = QK_TURBO4_64; + for (int block = 0; block < nb; block++) { + const float norm = GGML_FP16_TO_FP32(x[block].norm); + float * dst = y + block * d; + for (int i = 0; i < d; i++) { + uint8_t idx = (x[block].qs[i / 2] >> ((i % 2) * 4)) & 0xF; + dst[i] = CENTROIDS_4BIT[idx] * norm; + } + } +} + size_t quantize_turbo4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { GGML_UNUSED(imatrix); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 4f44aa726b31..1606630e42ef 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -765,6 +765,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_turbo4_0, .from_float_ref = (ggml_from_float_t) quantize_row_turbo4_0_ref, }, + [GGML_TYPE_TURBO4_64] = { + .type_name = "turbo4_64", + .blck_size = QK_TURBO4_64, + .type_size = sizeof(block_turbo4_64), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_turbo4_64, + .from_float_ref = (ggml_from_float_t) quantize_row_turbo4_64_ref, + }, [GGML_TYPE_TURBO2_0] = { .type_name = "turbo2", .blck_size = QK_TURBO2, diff --git a/src/llama-kv-cache-paged.cpp b/src/llama-kv-cache-paged.cpp index 450c3075f098..4ce2c0fd994d 100644 --- a/src/llama-kv-cache-paged.cpp +++ b/src/llama-kv-cache-paged.cpp @@ -16,6 +16,7 @@ #endif #include +#include #include #include #include @@ -120,9 +121,26 @@ llama_kv_cache_paged::llama_kv_cache_paged( } } + // MAD-301C Lever B: native head_dim-64 turbo4. Remap TURBO4_0 -> TURBO4_64 + // (64-element block) for head_dim-64 models (LFM2.5, gpt-oss). Each head is + // then exactly one 64-wide block with NO 64->128 zero-pad, ~halving the + // turbo4 KV footprint. Must run before the pad check + sizing below; the + // graph-level pad in build_attn keys off layer.k->type so it auto-skips for + // TURBO4_64 (not in its TURBO{2,3,4}_0 list). + // GGML_PAGED_TURBO4_64=0 forces the legacy padded-128 path (A/B + rollback). + { + const char * t464_env = std::getenv("GGML_PAGED_TURBO4_64"); + const bool t464_on = (t464_env == nullptr || t464_env[0] != '0'); + if (t464_on && head_dim == 64) { + if (type_k == GGML_TYPE_TURBO4_0) { type_k = GGML_TYPE_TURBO4_64; type_k_ = GGML_TYPE_TURBO4_64; } + if (type_v == GGML_TYPE_TURBO4_0) { type_v = GGML_TYPE_TURBO4_64; type_v_ = GGML_TYPE_TURBO4_64; } + } + } + // Turbo cache quantizes 128-element blocks; pad a sub-128 head_dim up to // 128 so each head is exactly one turbo block (matches the graph-level - // padding in llm_graph_context::build_attn for the paged path). + // padding in llm_graph_context::build_attn for the paged path). TURBO4_64 + // is excluded (its block is exactly head_dim 64 — no pad needed). const bool paged_cache_is_turbo = (type_k == GGML_TYPE_TURBO2_0 || type_k == GGML_TYPE_TURBO3_0 || type_k == GGML_TYPE_TURBO4_0); if (paged_cache_is_turbo && head_dim % 128 != 0) { From bf90040d7569c0a3eb5c5738db3e0be58b124569 Mon Sep 17 00:00:00 2001 From: kmbandy Date: Sat, 20 Jun 2026 10:29:14 -0400 Subject: [PATCH 7/7] chat: fix LFM2.5 reasoning bleed when opening is omitted LFM2.5's chat template does not prefill an opening 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 ), the peg-native lfm2 parser -- which required a literal opening and had no forced-open fallback -- failed to enter the reasoning rule, so the entire think block plus the stray bled into content. Make the opening optional: capture a leading reasoning block terminated by whether or not the opening tag was emitted. A pure-content turn has no , 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 --- common/chat.cpp | 10 +++++++++- tests/test-chat.cpp | 18 ++++++++++++++++++ 2 files changed, 27 insertions(+), 1 deletion(-) diff --git a/common/chat.cpp b/common/chat.cpp index 24e58ab06400..a585b0c6e8c7 100644 --- a/common/chat.cpp +++ b/common/chat.cpp @@ -1670,7 +1670,15 @@ static common_chat_params common_chat_params_init_lfm2(const common_chat_templat auto reasoning = p.eps(); if (extract_reasoning) { - reasoning = p.optional(THINK_START + p.reasoning(p.until(THINK_END)) + THINK_END); + // LFM2.5's template does not prefill an opening into the generation prompt, + // so the model must emit it itself each turn. It sometimes skips the opening tag and + // dives straight into reasoning, closing only with . Without forced-open + // handling that raw reasoning bleeds into content. Make the opening tag + // optional so a leading reasoning block terminated by is captured whether or + // not was emitted. A pure-content turn has no , so until() fails to + // find the terminator, the sequence fails, and the outer optional collapses -> the + // whole output falls through to content exactly as before. + reasoning = p.optional(p.optional(p.literal(THINK_START)) + p.reasoning(p.until(THINK_END)) + p.literal(THINK_END)); } if (!has_tools || inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_NONE) { diff --git a/tests/test-chat.cpp b/tests/test-chat.cpp index c1be9eb5a99f..b450dc18e919 100644 --- a/tests/test-chat.cpp +++ b/tests/test-chat.cpp @@ -4167,6 +4167,24 @@ static void test_template_output_peg_parsers(bool detailed_debug) { )) .run(); + // Missing opening : the model dives straight into reasoning and only closes + // with . Must still be captured as reasoning, not bled into content + // (regression test for the n_thinking=0 bleed on plain chat). + tst.test("I'm\nthinkingHello, world!\nWhat's up?") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .expect(message_assist_thoughts) + .run(); + + // Missing opening with a trailing tool call after the reasoning block. + tst.test("I need to call a function" + "Let me check the time.<|tool_call_start|>[get_time(city=\"Paris\")]<|tool_call_end|>") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .tools({ get_time_tool }) + .expect(message_with_reasoning_content_and_multiple_tool_calls( + "I need to call a function", "Let me check the time.", { { "get_time", "{\"city\":\"Paris\"}" } } + )) + .run(); + // Fake tool call marker inside reasoning is not parsed as a call tst.test("Let me think about <|tool_call_start|>[special_function(arg1=1)]<|tool_call_end|> hmm" "<|tool_call_start|>[special_function(arg1=1)]<|tool_call_end|>")