perf(gemma4-cuda): TC flash-attention prefill (#146/#147) + SoA Q8_0 weight layout (#149) — prefill 2853→4240

README.md

@@ -51,7 +51,7 @@ coherent (`scripts/bench-all.ps1`); top-1 parity vs llama.cpp b8585 verified on
 | Qwen3.6-35B-A3B-MTP (GDN+MoE) | (same) | 22 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **65.0** | **22.9** | Requires `SHARPI_CPU_MOE=1`: 30 GDN + 10 attn + shared expert on GPU, routed experts CPU mmap. 100% acceptance. Fused GDN scan + batched SDPA (#114-B/#118), bit-identical, grows with ctx |
 | Carnice (Qwen3.6-35B-A3B-MTP finetune) | [mudler](https://huggingface.co/mudler/Carnice-Qwen3.6-MoE-35B-A3B-APEX-MTP-GGUF) | 17 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **43.6** | **25.0** | agentic finetune of 35B-A3B-MTP; 77% acceptance (`bench-carnice.ps1` — the default prompt 1-token-EOSes on this terser tune). APEX mixed-precision (Q3_K + Q8_0 experts); Q8_KS per-32 int dots auto-enable at load (#99/#101/#107), +4.6% decode at ~4× tighter parity vs plain Q8_K (`SHARPI_Q3K_Q8K=0`/`SHARPI_Q8_0_Q8K=0` to disable). Fused GDN scan + wave SDPA (#114-B/#118) bit-identical past 4096 |
 | Gemma 4 E4B-it Q8 | [unsloth](https://huggingface.co/unsloth/gemma-4-E4B-it-GGUF) | 8 GB | CPU | 4.9 | 5.0 | dense 42-layer gemma4: per-layer head_dim (256 SWA / 512 global), dual-RoPE, KV-share tail (18 layers), 5:1 SWA:global, logit softcap 30, PLE-256 injection (~4.2 GB mmap-resident) |
-| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g -1 -c 2048` | **2853** | **59** | all 42 layers fit at `-c 2048`. KV-share alias + SWA/global split per layer; PLE projections (~215 MB) upload at construction. **Prefill (#141):** int8 **tensor-core MMQ** matmul (`mma.m16n8k32.s8`, each Q8_0 weight read once as int8 — beats the dequant→fp16→cuBLAS GEMM, drops its fp16 HBM temp; `SHARPI_PREFILL_MMQ=0` reverts) + a memory-efficient **flash-attention** prefill (shared K/V tiles + online softmax + half2 fp16x2 QK dot — replaces the scalar O(n²) per-query attention that re-streamed each query's K/V window up to ~512×; `SHARPI_PREFILL_FLASH=0` reverts) + a batched Q8_0 embedding lookup. **~1.8× at ~1K ctx (1564→2853), ~2.05× at 1.8K** — profiling showed *attention*, not the matmul, was the dominant prefill cost at realistic prompt lengths, so the win grows with prompt length. **Decode (#142):** dp4a/Q8_1 int8 matvec (`SHARPI_Q80_DP4A=0` to bisect) + CUDA-graph capture/replay default-on (`SHARPI_CUDA_GRAPH=0` to bisect). All prefill/decode fast paths are argmax-stable vs the fp32 path, not bit-exact. Remaining gap to llama.cpp (~8000 prefill / ~78 decode): full tensor-core flash at d=512 + decode matvec work |
+| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g -1 -c 2048` | **3698** | **59** | all 42 layers fit at `-c 2048`. KV-share alias + SWA/global split per layer; PLE projections (~215 MB) upload at construction. **Prefill (#141):** int8 **tensor-core MMQ** matmul (`mma.m16n8k32.s8`, each Q8_0 weight read once as int8 — beats the dequant→fp16→cuBLAS GEMM, drops its fp16 HBM temp; `SHARPI_PREFILL_MMQ=0` reverts) + a **tensor-core flash-attention** prefill (#146/#147): both QK^T and P·V on the mma cores (`mma.m16n8k16.f16`), multi-warp **d-split** so the O tile stays register-resident — replaces the scalar O(n²) per-query attention (which re-streamed each query's K/V window up to ~512×) and is **+27% at ~1K / +40% at 1.8K** over the earlier half2 flash kernel (`SHARPI_PREFILL_FLASH_TC=0` reverts to half2, `=…_FLASH=0` to scalar) + a **SoA Q8_0 weight repack** (#149): all Q8_0 readers (MMQ, dp4a, fp32 matvec, GEMM-N, dequant) read the quants 16-byte-aligned with the fp16 scales split out, killing the `qs` 2-byte-misalignment funnelshift tax — **+10-12% prefill, bit-identical**; `SHARPI_MMQ_SOA=0` reverts) + a batched Q8_0 embedding lookup. **109→3698 at ~1K ctx, →4240 at 1.8K** — profiling showed *attention*, then the matmul inner-loop efficiency, were the dominant prefill costs at realistic prompt lengths. **Decode (#142):** dp4a/Q8_1 int8 matvec (`SHARPI_Q80_DP4A=0` to bisect) + CUDA-graph capture/replay default-on (`SHARPI_CUDA_GRAPH=0` to bisect). All prefill/decode fast paths are argmax-stable vs the fp32 path, not bit-exact (the SoA repack is bit-identical). Remaining gap to llama.cpp (~8475 prefill / ~78 decode): cp.async-pipelined MMQ on the SoA layout + decode matvec work |
 | Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g 22 -c 2048` (hybrid) | 6.6 | 6.8 | 22 GPU + 20 CPU layers. `-g ≤ 22` required so the CPU shared-KV tail can read its own-KV source layers; CPU dense-FFN dominates decode (bandwidth-bound). `SHARPI_CUDA_PROFILE=1` for per-phase breakdown |
 
 _Numbers re-measured across every on-disk row at ~1K ctx so the prefill column is comparable; per-issue

scripts/bench-gemma4-flashtc-146.ps1

@@ -0,0 +1,55 @@
+# A/B: half2 flash prefill (default) vs tensor-core flash prefill (#146,
+# SHARPI_PREFILL_FLASH_TC=1). Same ~1K prompt + warm-discard as bench-gemma4-136.ps1,
+# all-GPU (-g -1). Reports prefill t/s for each so the TC kernel can be judged
+# against the half2 baseline (re-baselined at 2589 t/s on merged master).
+param(
+    [string]$ModelDir = "E:\models",
+    [int]$Sections = 6
+)
+$ErrorActionPreference = "Continue"
+$model = Join-Path $ModelDir "gemma-4-E4B-it-Q8_0.gguf"
+
+$para = @(
+"Modern large language model inference is dominated by memory bandwidth rather than raw compute.",
+"Each decode step streams the full weight matrix for every layer, so quantization formats like Q4_K and Q5_K trade a small accuracy loss for a large reduction in bytes moved per token.",
+"Mixture-of-experts models complicate this: only a handful of the hundreds of experts fire per token, and which fire varies token to token, defeating simple weight caching.",
+"Gated DeltaNet layers replace quadratic attention with a linear recurrent state update, bounding per-token cost as context grows, at the price of a strictly sequential scan.",
+"Hybrid placement keeps the attention and recurrent trunk on the accelerator while streaming routed-expert weights from host memory, overlapping the two so neither stalls."
+) -join " "
+$sb = [System.Text.StringBuilder]::new()
+[void]$sb.Append("Read the following engineering notes and then write a concise technical summary.`n`n")
+for ($i = 1; $i -le $Sections; $i++) { [void]$sb.Append("Section $i. $para`n`n") }
+[void]$sb.Append("Summarize the main performance trade-offs across the sections above.")
+$prompt = $sb.ToString()
+
+$cudaArgs = @("-g","-1","--backend","cuda","--ctx-size","2048")
+
+Write-Host "--- warming model ---" -ForegroundColor DarkGray
+$null = .\scripts\bench-textgen.ps1 -Model $model -Tag "ftc-warm" -NTokens 8 -Prompt "Hello, world." -TimeoutSec 900 -ExtraArgs $cudaArgs
+
+$rows = @()
+# variant: half2 | tc1 (single-warp #146) | tc2 (multi-warp #147)
+function Run($tag, $variant) {
+    try {
+        Remove-Item env:SHARPI_PREFILL_FLASH_TC  -ErrorAction SilentlyContinue
+        Remove-Item env:SHARPI_PREFILL_FLASH_TC1 -ErrorAction SilentlyContinue
+        if ($variant -eq "tc1") { $env:SHARPI_PREFILL_FLASH_TC = "1"; $env:SHARPI_PREFILL_FLASH_TC1 = "1" }
+        elseif ($variant -eq "tc2") { $env:SHARPI_PREFILL_FLASH_TC = "1" }
+        $null = .\scripts\bench-textgen.ps1 -Model $model -Tag "$tag-w" -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+        $r = .\scripts\bench-textgen.ps1 -Model $model -Tag $tag -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+        $script:rows += [PSCustomObject]@{ Tag=$tag; PrefTok=$r.PrefillTok; PrefillTps=$r.PrefillTps; DecodeTps=$r.DecodeTps; Wall=$r.WallSec; TO=$r.TimedOut; Sample=$r.Sample }
+        Write-Host ("  {0,-18} pref={1,7} t/s  dec={2,6} t/s  ({3} tok)" -f $tag,$r.PrefillTps,$r.DecodeTps,$r.PrefillTok) -ForegroundColor Green
+    } finally {
+        Remove-Item env:SHARPI_PREFILL_FLASH_TC  -ErrorAction SilentlyContinue
+        Remove-Item env:SHARPI_PREFILL_FLASH_TC1 -ErrorAction SilentlyContinue
+    }
+}
+
+Run "flash-half2" "half2"
+Run "flash-tc1"   "tc1"
+Run "flash-tc2"   "tc2"
+
+Write-Host ""
+Write-Host "=== Gemma 4 flash half2 vs TC (#146, ~1K ctx, warm) ===" -ForegroundColor Cyan
+$rows | Format-Table Tag,PrefTok,PrefillTps,DecodeTps,Wall,TO -AutoSize
+$rows | ForEach-Object { Write-Host ("  {0}: {1}" -f $_.Tag, $_.Sample) -ForegroundColor DarkGray }

scripts/bench-gemma4-mmq-soa-149.ps1

@@ -0,0 +1,35 @@
+# A/B: Q8_0 SoA weight layout (#149, SHARPI_MMQ_SOA) off vs on, both with the TC
+# flash default. Same ~1K/1.8K prompt + warm-discard, all-GPU (-g -1).
+param([string]$ModelDir = "E:\models", [int]$Sections = 13)
+$ErrorActionPreference = "Continue"
+$model = Join-Path $ModelDir "gemma-4-E4B-it-Q8_0.gguf"
+$para = @(
+"Modern large language model inference is dominated by memory bandwidth rather than raw compute.",
+"Each decode step streams the full weight matrix for every layer, so quantization formats like Q4_K and Q5_K trade a small accuracy loss for a large reduction in bytes moved per token.",
+"Mixture-of-experts models complicate this: only a handful of the hundreds of experts fire per token, and which fire varies token to token, defeating simple weight caching.",
+"Gated DeltaNet layers replace quadratic attention with a linear recurrent state update, bounding per-token cost as context grows, at the price of a strictly sequential scan.",
+"Hybrid placement keeps the attention and recurrent trunk on the accelerator while streaming routed-expert weights from host memory, overlapping the two so neither stalls."
+) -join " "
+$sb = [System.Text.StringBuilder]::new()
+[void]$sb.Append("Read the following engineering notes and then write a concise technical summary.`n`n")
+for ($i = 1; $i -le $Sections; $i++) { [void]$sb.Append("Section $i. $para`n`n") }
+[void]$sb.Append("Summarize the main performance trade-offs across the sections above.")
+$prompt = $sb.ToString()
+$cudaArgs = @("-g","-1","--backend","cuda","--ctx-size","2048")
+
+Write-Host "--- warming model ---" -ForegroundColor DarkGray
+$null = .\scripts\bench-textgen.ps1 -Model $model -Tag "soa-warm" -NTokens 8 -Prompt "Hello, world." -TimeoutSec 900 -ExtraArgs $cudaArgs
+$rows = @()
+function Run($tag, $on) {
+    try {
+        if ($on) { $env:SHARPI_MMQ_SOA = "1" } else { Remove-Item env:SHARPI_MMQ_SOA -ErrorAction SilentlyContinue }
+        $null = .\scripts\bench-textgen.ps1 -Model $model -Tag "$tag-w" -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+        $r = .\scripts\bench-textgen.ps1 -Model $model -Tag $tag -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+        $script:rows += [PSCustomObject]@{ Tag=$tag; PrefTok=$r.PrefillTok; PrefillTps=$r.PrefillTps; DecodeTps=$r.DecodeTps; TO=$r.TimedOut }
+        Write-Host ("  {0,-14} pref={1,7} t/s  dec={2,6} t/s  ({3} tok)" -f $tag,$r.PrefillTps,$r.DecodeTps,$r.PrefillTok) -ForegroundColor Green
+    } finally { Remove-Item env:SHARPI_MMQ_SOA -ErrorAction SilentlyContinue }
+}
+Run "soa-off" $false
+Run "soa-on"  $true
+Write-Host "`n=== Gemma 4 MMQ SoA #149 (~$([int]($Sections*158))ctx, warm) ===" -ForegroundColor Cyan
+$rows | Format-Table Tag,PrefTok,PrefillTps,DecodeTps,TO -AutoSize

scripts/bench-gemma4-rebaseline.ps1

@@ -0,0 +1,42 @@
+# Focused re-baseline of the all-GPU Gemma 4 prefill/decode on merged master,
+# before the #146 full-TC-flash work. Same ~1K prompt + warm-discard methodology
+# as bench-gemma4-136.ps1, but only the all-GPU (-g -1) cell, run twice (keep 2nd).
+param(
+    [string]$ModelDir = "E:\models"
+)
+$ErrorActionPreference = "Continue"
+$model = Join-Path $ModelDir "gemma-4-E4B-it-Q8_0.gguf"
+
+# Identical ~1K-token prompt to bench-allrows-1k.ps1 / bench-gemma4-136.ps1.
+$para = @(
+"Modern large language model inference is dominated by memory bandwidth rather than raw compute.",
+"Each decode step streams the full weight matrix for every layer, so quantization formats like Q4_K and Q5_K trade a small accuracy loss for a large reduction in bytes moved per token.",
+"Mixture-of-experts models complicate this: only a handful of the hundreds of experts fire per token, and which fire varies token to token, defeating simple weight caching.",
+"Gated DeltaNet layers replace quadratic attention with a linear recurrent state update, bounding per-token cost as context grows, at the price of a strictly sequential scan.",
+"Hybrid placement keeps the attention and recurrent trunk on the accelerator while streaming routed-expert weights from host memory, overlapping the two so neither stalls."
+) -join " "
+$sb = [System.Text.StringBuilder]::new()
+[void]$sb.Append("Read the following engineering notes and then write a concise technical summary.`n`n")
+for ($i = 1; $i -le 6; $i++) { [void]$sb.Append("Section $i. $para`n`n") }
+[void]$sb.Append("Summarize the main performance trade-offs across the sections above.")
+$prompt = $sb.ToString()
+
+$cudaArgs = @("-g","-1","--backend","cuda","--ctx-size","2048")
+
+Write-Host "--- warming model ---" -ForegroundColor DarkGray
+$null = .\scripts\bench-textgen.ps1 -Model $model -Tag "g4-rb-warm" -NTokens 8 -Prompt "Hello, world." -TimeoutSec 900 -ExtraArgs $cudaArgs
+
+# Run twice, keep the second (warm).
+$null = .\scripts\bench-textgen.ps1 -Model $model -Tag "g4-rb-w" -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+$r = .\scripts\bench-textgen.ps1 -Model $model -Tag "g4-rb" -NTokens 60 -Prompt $prompt -TimeoutSec 900 -ExtraArgs $cudaArgs
+
+Write-Host ""
+Write-Host "=== Gemma 4 all-GPU re-baseline (merged master, ~1K ctx, warm) ===" -ForegroundColor Cyan
+[PSCustomObject]@{
+    PrefillTok = $r.PrefillTok
+    PrefillTps = $r.PrefillTps
+    DecodeTps  = $r.DecodeTps
+    WallSec    = $r.WallSec
+    TimedOut   = $r.TimedOut
+    Sample     = $r.Sample
+} | Format-List