perf(cuda): tensor-core (mma.sync FP16) decode attention kernel — port Vulkan's coopmat design #199

Activity

1. jamesburton commented on Jul 27, 2026

   

   jamesburton

   on Jul 27, 2026

   OwnerAuthor

   More actions

   Design-ahead plan: tensor-core FP16 decode attention (DESIGN ONLY — not a green light to implement)

   This is a brainstorm/design document produced ahead of implementation, per this issue's own
   sequencing note (do this after #197/#198 land and re-profiling confirms the ceiling is worth
   chasing). At the time of writing, #197 and #198 are both still open with no design/implementation
   comments — there is nothing yet to build on top of from those issues; this design is
   self-contained and notes where it should be revisited once they land.

   0. A directly-relevant precedent this repo already has, that the issue description doesn't mention

   Before designing anything new, I found the codebase has already tried and rejected a full-FP16
   decode-attention path once, for this exact model family. From
   src/DotLLM.Cuda/Architectures/CudaQwen3MoeHybridTransformerModel.cs (lines ~110-151, ~2041-2051):

   Rolled forward to "F32 attention with F16-stored KV that's dequanted on read" (Option A) after
   Option B (all-F16 internal attention: Q→F16, attn out F16→F32) breached the synthetic Q5_K and
   F16 parity bars by ~0.5-1x on a hidden=32 / headDim=8 fixture — "too much drift introduced by
   the Q-cast on a 16-element headDim." Restricting F16 precision loss to cache storage only keeps
   Q and the attention accumulator in F32, which is what stays inside the existing
   AbsTol(1e-4)+RelTol(1e-3) synthetic-fixture bars.

   This is the single most load-bearing fact for this design. Two things follow from it:

   1. The tight AbsTol(1e-4)+RelTol(1e-3) synthetic-fixture bar used elsewhere in this repo
      (CudaQwen3MoeHybridParityTests, etc.) is already known to be unreachable with a Q-cast-to-FP16
      attention path on this architecture, at small headDim. Don't target it for this kernel. Target
      the same 5e-3 abs OR rel bar CudaTensorCoreAttentionParityTests already uses for the
      prefill flash-mma kernel — that's the established, working precedent for "attention with FP16
      Q/K/V" in this repo, not the tight bar.
   2. The failure mode may not directly transfer. The rejected "Option B" almost certainly did its
      QK dot product as a CUDA-core scalar reduction over FP16-cast operands (mirroring
      attention.cu's attention_f16 shape) — precision loss compounds from both the FP16 operand
      rounding and however the accumulation was done. mma.sync.m16n8k16.f16.f16.f32 (the PTX this
      issue proposes reusing from attention_flash_mma.cu) is architecturally guaranteed to accumulate
      in FP32 registers regardless of input dtype — the accumulator was never the weak link for
      tensor cores the way it may have been for Option B's scalar path. So this design is not simply
      repeating a known failure, but it does inherit the Q-cast-to-FP16 risk Option B hit, on a
      small headDim, which is exactly the failure mode reported. This needs to be re-characterized,
      not assumed safe by analogy to attention_flash_mma.cu's prefill parity (which was tested at
      headDim=64, not headDim=8/16).

   Also unresolved and worth flagging as a prerequisite fact-check: attention_flash_mma.cu is
   hardcoded HEAD_DIM=64 (#define HEAD_DIM 64, unrolled k-steps HEAD_DIM/16 == 4,
   n-subtiles HEAD_DIM/8 == 8). I could not confirm Bonsai-27B's (qwen35moe) actual headDim in
   this pass — if it isn't 64, the existing flash-mma helpers can be reused as PTX primitives but the
   kernel's tiling constants cannot be reused unmodified, and Vulkan's coopmat kernel (which loops in
   D_TILE=16 chunks up to MAX_HEAD_DIM=256 with zero-padding, not a hardcoded head dim) is the more
   relevant structural reference, not attention_flash_mma.cu's fixed-64 layout. Verify Bonsai's
   headDim before writing any kernel code.

   1. Current decode path (baseline this kernel replaces)

   Traced end-to-end for the Bonsai-27B full-attention layers
   (CudaQwen3MoeHybridTransformerModel.ForwardFullAttnBody, ~line 2050-2085):

   • KV cache is stored FP16 (_f16KCache/_f16VCache, model-private per-layer sidecar).
   • Each decode step: write the new K/V row(s) into the FP16 cache, then
     LaunchConvertF16ToF32 dequants the entire live [0, seqKv) prefix K and V into a shared F32
     scratch (_f32KvReadStagingK/V), then LaunchAttentionF32 (the attention_f32 kernel in
     native/kernels/attention_f32.cu, Schraudolph fast-exp, grid = numHeads, one block per query
     head) runs entirely in F32.
   • This is exactly the kernel perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197's ncu profiling flagged: grid=24 for Bonsai's 24 heads,
     16.5-16.8% achieved occupancy, latency-bound (~42 stall cycles/instruction, ~95% cycles with no
     eligible warp).

   A tensor-core decode kernel that reads the FP16 KV cache directly (skipping the
   LaunchConvertF16ToF32 dequant step entirely) is a real secondary win this design should capture —
   not just the matmul throughput, but removing a whole kernel launch + HBM round-trip per layer per
   decode step. Q would still need a cast to FP16 (currently produced as F32 from the preceding
   RMSNorm/projection in this model's F32 pipeline) — that cast is exactly Option B's risk vector
   above, so it needs to be the primary focus of precision characterization, not the KV-cache read
   (cache read is already lossy today by design and known to be within tolerance).

   2. Kernel structure recommendation

   Scope narrow initially: one query head per block/warp-group, tensor cores replacing the CUDA-core
   matmul only. Do NOT bundle GQA-group batching (#198) into this kernel's first version.

   Reasoning:

   • perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's GQA-group batching is an orthogonal axis (grid/block restructuring + KV read
     amortization, explicitly scoped to stay F32, no new precision risk, per perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's own text). If
     this kernel adopts perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's block-per-KV-head-group structure and changes dtype at the same time,
     any precision regression found in validation is ambiguous — is it the regrid or the dtype? Keep
     them separable so a failure in the new precision axis doesn't also indict perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's already-designed,
     lower-risk regrid.
   • attention_flash_mma.cu already has a proven group warps/block pattern (one warp per query
     head sharing a KV head, KV loaded once into shared, reused across the group's warps) that IS a
     form of GQA-group batching, at the warp level, for prefill. It is reasonable — and probably the
     right move once perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 lands — to compose the two, but the sequencing should be: (a) land this
     kernel scoped to decode's seqQ=1 shape with the SAME warp-per-query-head / KV-shared-once
     pattern attention_flash_mma.cu already has (this is not "GQA-group batching" in the perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 sense
     of restructuring the grid to fix the latency-bound occupancy problem; it's just reusing the
     existing prefill kernel's grouping, which is already group warps cooperatively loading one
     block's K/V) — (b) once perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's regrid lands and is validated, evaluate whether this kernel's grid
     needs to change to match it, as a small follow-up, with the dtype axis already proven safe by
     then.

   Row-padding scheme (from attention_f32_coopmat.comp's decode handling, ported to mma.sync
   tiling): attention_flash_mma.cu tiles the query axis in Q_TILE=16 rows via ldmatrix.x4
   (an 8x8x4 = 16x16 fragment load). For decode, seqQ=1: allocate the same Q_TILE=16-row shared
   sQ buffer, fill query row 0 with the real (single) query vector and rows 1-15 with zero
   (matching Vulkan's validRows/zero-pad approach exactly — "the wasted lanes are not a correctness
   issue"), run the identical QK/softmax/PV mma.sync pipeline unmodified, and on write-back only
   emit row 0 (gr0 < seq check in attention_flash_mma.cu's final loop already does exactly this
   kind of bounds check — for decode it becomes "only write if this is the real row", the same
   validRows idea Vulkan names explicitly). This means the existing PTX helpers
   (ldmatrix_x4/ldmatrix_x2/mma_m16n8k16) and the QK/softmax/PV structure need zero changes —
   only the grid/launch config (one query tile total instead of ceil(seq/16)), the KV loop bound
   (no causal triangle — decode attends to the full [0, seqKv) cache, replacing
   kv_end = min(seq, q0 + Q_TILE) with kv_end = seqKv), and the positionOffset handling (decode
   needs position_offset support, which the current prefill kernel hardcodes to 0 per
   CudaFlashAttention.CanUse's gate — this is a real gap to close, not just a formality: the causal
   mask inside the online-softmax loop needs gq = positionOffset instead of assuming prefill's
   q0 + lane).

   GPU utilization cost of the 16-row pad: 15/16 of the tensor-core MMA throughput is wasted on
   zero rows for decode. This is fine — the whole point (per the issue's own estimate) is that even a
   "wasteful" tensor-core kernel doing 1/16 useful MMA work can still beat a latency-bound CUDA-core
   kernel that's achieving ~17% occupancy, if the launch is no longer latency-bound (which is why
   #198 needs to land first — more resident warps per block from GQA-group sharing is what fixes the
   latency-bound diagnosis; tensor cores alone don't fix "no eligible warp to issue"). This is worth
   stating explicitly in case a future implementer is tempted to skip #198: without more resident
   warps/block, a decode-shaped mma.sync kernel launched at grid=numHeads (today's grid) risks being
   just as latency-bound as attention_f32 is today, wasting the entire tensor-core throughput
   advantage on stalls, not matmul work.

   3. Dtype recommendation

   FP16 (not BF16) for Q/K/V storage and mma.sync inputs; FP32 accumulator — matching
   attention_flash_mma.cu's existing prefill kernel exactly, not introducing a third variant.

   Justification:

   • FP32 accumulator is not really a "choice" — mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32
     (the exact instruction attention_flash_mma.cu already uses) accumulates in FP32 by hardware
     definition. This matches "common flash-attention practice" the issue asks about, and it is
     strictly more precise than the rejected Option B's likely FP16-adjacent accumulation.
   • FP16 vs BF16 for the inputs: BF16 (mma.sync...f32.bf16.bf16.f32, same PTX family, also
     Ampere-supported) trades FP16's extra mantissa bits for FP32-matching exponent range. The
     rejected Option B failure was attributed to "drift from the Q-cast," not overflow/underflow — a
     mantissa-precision problem, not a range problem — which argues for FP16 (more mantissa bits, same
     8-bit-exponent range is not the reported failure mode) over BF16. Also: reusing FP16 keeps this
     kernel numerically comparable to attention_flash_mma.cu and attention_f16 (both already FP16
     in this codebase), so any parity-test infrastructure and tolerance precedent
     (CudaTensorCoreAttentionParityTests's 5e-3 abs/rel bar) transfers directly. Introducing BF16
     would mean this is the first BF16 CUDA kernel in the repo (grep found no existing
     bf16/__nv_bfloat16 CUDA kernel) — a second new axis of risk/infrastructure this design should
     not add unless FP16 is characterized and found insufficient first.
   • Recommendation: build FP16 first. If precision characterization (below) finds the Q-cast risk
     reproduces at Bonsai's real headDim, BF16 is the documented fallback to try next — same PTX
     shape, same helper functions, only the fragment element type and conversion intrinsics change.

   4. Precision-tolerance characterization plan (must happen BEFORE writing kernel code)

   Given the prior Option B failure, this needs more rigor than "port the prefill kernel's 5e-3 bar
   and hope it holds":

   1. Reproduce the Option B failure mode in isolation first, cheaply, without CUDA: write (or
      reuse if it still exists) a small script/unit-test-style harness that does Q→FP16 cast → FP16
      dot product accumulated in FP32 (i.e. exactly what mma.sync would do) vs the current F32
      reference, on the same hidden=32/headDim=8 synthetic fixture that broke Option B. If FP32
      accumulation alone (no accumulation-precision loss, only the FP16 input rounding) already
      breaches the tight 1e-4/1e-3 bar at that headDim, that confirms the Q-cast itself (not
      accumulation) was the dominant error source, and the same risk fully carries into this design.
      If FP32-accumulated FP16 inputs pass at that fixture size where FP16-accumulated inputs failed,
      that's evidence the accumulator precision, not the input cast, was Option B's problem — good
      news for this design, since mma.sync's accumulator is FP32. This single check is cheap
      (no GPU kernel needed, pure host-side float math) and should be the first thing done, before any
      CUDA code.
   2. Reference for parity: two-tier, matching this repo's existing pattern —
      • Tier 1 (unit, tight-ish bar): CPU F32 oracle (DotLLM.Cpu, same fast-exp convention as
        attention_f32.cu) vs the new FP16 tensor-core kernel, small synthetic fixture (mirror the
        hidden=32/headDim=8 Bonsai fixture specifically, since that's the one with known history —
        add the real Bonsai headDim as a second fixture size once known). Target the loose bar (5e-3
        abs OR rel), NOT the tight synthetic bar the model-level tests use — per finding #0 above, the
        tight bar is already known unreachable for FP16-cast attention on this architecture.
      • Tier 2 (kernel parity, like CudaTensorCoreAttentionParityTests): new kernel vs
        attention_f16 (existing CUDA-core FP16 kernel) at realistic magnitude (DOTLLM_CUDA_ATTN_MAG
        pattern already exists — reuse it), decode shape (seqQ=1, varying seqKv — 1, 256, 1024,
        4096+ to cover the depths perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197/perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's investigation already profiled at). This is an
        apples-to-apples FP16-vs-FP16 check (isolates the tensor-core-vs-CUDA-core question from the
        F32-vs-FP16 question, which Tier 1 covers).
   3. Validate across realistic Bonsai-27B decode sequences, not just synthetic fixtures: run a
      real generation (reuse whatever harness produced the "300-step non-compounding drift"
      characterization CudaAttentionF32SplitKvTests.cs already did for perf(cuda): attention_f32 flash-decoding split-KV for depth-dependent decode cost #183's split-KV kernel) with
      the new kernel swapped in, opt-in, and confirm output token sequences match the F32 baseline for
      a meaningful number of steps at realistic depth (the perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197 issue text already establishes decode
      depths of interest: 256, 768, 1024, 2048+). Per this kernel's own math: attention over an
      unperturbed KV cache is not persistent-state-compounding the way GDN's scan state is (same
      argument attention_f32.cu's header already makes for why split-KV's reassociation is safe over
      many steps) — a numeric perturbation in this layer's attention output at step N does not feed
      back into what's cached for this layer at step N. The only cross-step pathway is discrete
      (a different sampled token branches the whole generation). So the right validation isn't "does
      drift compound" (it structurally can't, the same way it doesn't for perf(cuda): attention_f32 flash-decoding split-KV for depth-dependent decode cost #183) — it's "does any
      single decode step produce a large enough output delta to flip the argmax/sampled token", which
      is a distribution-level check (logit-diff histograms across many real decode steps), not a
      drift-accumulation check.
   4. Concrete target tolerance: 5e-3 abs OR rel vs attention_f16 (Tier 2, matches
      CudaTensorCoreAttentionParityTests precedent for FP16 attention already in this repo). If step
      1's isolated repro shows the FP32-accumulator does NOT save the Q-cast precision problem at
      Bonsai's real headDim, this tolerance target needs to be revisited (either loosen further and
      accept degraded quality with explicit measurement of downstream eval impact — e.g. perplexity or
      task-accuracy delta, not just tensor-diff — or fall back to the BF16 variant from section 3, or
      conclude the kernel isn't viable for this architecture and scope it to models/head-shapes where
      it does hold).

   5. Default-off / opt-in

   Yes, ship opt-in / default-off initially, following the #180 (GDN split) / #183 (attention
   split-KV) precedent, for stronger reasons than those two kernels had:

   • perf(cuda): gdn_scan_step_f32 grid-too-small occupancy — cooperative-groups row-split (ncu-informed follow-up to #173) #180/perf(cuda): attention_f32 flash-decoding split-KV for depth-dependent decode cost #183 are reassociation-only (same F32 math, different summation order) — a pure
     numerical-noise question with well-understood error bounds. This kernel is a dtype change
     with a documented prior failure on this exact architecture. That's categorically higher risk
     than either precedent.
   • Gate it the same way CudaFlashAttention.Enabled/CanUse and the GDN split's
     DOTLLM_GDN_SCAN_APPROX_SPLIT4 env var work: a named env var (something like
     DOTLLM_CUDA_ATTN_TENSORCORE_DECODE), default OFF, device-eligibility-gated to GeForce Ampere
     (same sm_86-only constraint CudaFlashAttention already has, since the PTX is Ampere-only
     mma.sync), and a CanUse-style per-call gate mirroring CudaFlashAttention.CanUse but for the
     decode shape (seqQ==1, headDim matching whatever the kernel is built for, positionOffset
     now supported rather than required to be 0, GQA group ≤ whatever cap the tuned kernel uses).
   • Criteria to flip the default later: (a) Tier 1+2 precision checks above pass at the loose
     bar across Bonsai's real headDim, not just the small synthetic fixture; (b) the real-generation
     distribution check (section 4.3) shows no measurable argmax-flip rate increase over a
     statistically meaningful number of real decode steps; (c) ncu re-profiling (post-perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198, per this
     issue's own sequencing) confirms the kernel is no longer latency-bound and the measured wall-clock
     decode speedup is a real, reproducible win over the perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197+perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198-improved F32 baseline — not just an
     ncu "Est. Speedup" number (this investigation's own stated lesson: those overstate real
     wall-clock impact).

   6. Prerequisites from #197/#198

   Neither #197 nor #198 has a design or implementation yet (checked at time of writing — both open,
   zero comments). So today, nothing concrete to build on top of. When this is picked up:

   • From perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 (if landed): the GQA-group-batched grid/block structure for decode. Section 2 above
     recommends NOT taking a hard dependency on perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's specific regrid for this kernel's first version
     (compose later, once both are independently validated) — but perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's occupancy fix is very likely
     a prerequisite for this kernel's throughput win to materialize at all (section 2's "latency-bound"
     warning). Practically: this kernel can be designed and written independently of perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198, but should
     not be benchmarked/evaluated for go/no-go until perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's occupancy fix is in, per the issue's own
     reasoning about the 8-10x ceiling only being realized outside the latency-bound regime.
   • From perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197 (if landed): independent — perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197 is a heuristic tune of the existing F32 split-KV
     kernel, not a structural change this kernel depends on. No blocking relationship either direction.
   • Re-profile after both land: confirm the diagnosed ncu numbers (16.5% occupancy, 42 stall
     cycles/instruction) actually change post-perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197/perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 before sinking implementation time into this
     kernel — if perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198 alone gets decode attention to, say, 60-70% occupancy and comfortably
     memory/compute-bound, the marginal win from also switching to tensor cores may be much smaller
     than the "training-knowledge estimate" 8-10x this issue cites (which was explicitly flagged as
     "not yet re-verified").

   7. Implementation scope/complexity and riskiest unknown

   Scope: Medium-large. The PTX plumbing (ldmatrix/mma.sync helpers, QK/softmax/PV structure)
   is genuinely reusable from attention_flash_mma.cu with only the row-padding, KV-loop-bound, and
   positionOffset changes described in section 2 — that part is closer to "port + adapt" than
   "design from scratch," which de-risks the mechanical implementation. What's NOT reusable /
   genuinely new work: (a) the CanUse gate and dispatch wiring for the decode call site
   (CudaQwen3MoeHybridTransformerModel.ForwardFullAttnBody, currently hardcoded to
   LaunchAttentionF32), (b) reading the FP16 KV cache directly instead of through the F32 dequant
   staging buffer (a real code path change, not just a kernel swap), (c) the entire precision
   characterization work in section 4, which is the actual gating cost, not the kernel code itself.

   Single riskiest technical unknown: whether the Q→FP16 cast precision loss that broke "Option B"
   once already, on this exact model family, reproduces for a tensor-core (FP32-accumulated) kernel
   at Bonsai's real (currently unconfirmed) headDim. Section 4.1's cheap host-side isolation check
   should be the very first thing done — before any CUDA code — because it's a fast, low-cost way to
   get an early signal on whether this entire design direction is viable for Bonsai-27B specifically,
   or whether it needs to be scoped to a different, less precision-sensitive model family (or shipped
   with BF16 inputs / a hybrid F32-Q/FP16-KV scheme instead of full FP16 inputs).

   ---

   What should happen before any code is written

   1. Confirm Bonsai-27B's (qwen35moe) actual headDim — determines whether
      attention_flash_mma.cu's fixed HEAD_DIM=64 tiling can be reused directly or needs
      generalizing (Vulkan coopmat's D_TILE-chunked approach as the fallback reference).
   2. Run the cheap host-side Q-cast/FP32-accumulator isolation check (section 4.1) against the
      known-failing hidden=32/headDim=8 Bonsai fixture, before writing any CUDA. This is the fastest
      possible go/no-go signal for the whole design.
   3. Land and re-profile after perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197 and perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198, per this issue's own stated sequencing — confirm
      with real ncu + real dotnet run bench throughput (not just "Est. Speedup") that decode
      attention is no longer latency-bound before sinking implementation time into a kernel whose
      entire value proposition depends on that being true.
   4. Only then: write the kernel (adapt attention_flash_mma.cu per section 2), wire it opt-in per
      section 5, and run the full Tier 1/Tier 2/real-generation validation from section 4 before
      flipping any default.

   🤖 Generated with Claude Code
2. jamesburton commented on Jul 28, 2026

   

   jamesburton

   on Jul 28, 2026

   OwnerAuthor

   More actions

   Precondition check result (before any kernel implementation): confirmed NOT met — recommend #199 stay parked

   Per the coordinator's request, before writing any tensor-core kernel code I re-verified the actual
   precondition this issue names for going ahead: "this kernel's ceiling is only realized once the
   kernel is out of the latency-bound regime #198 targets." #198 merged (PR #201, attention_f32_gqa_split_kv,
   opt-in via DOTLLM_ATTN_GQA_SPLIT=1), but its own real-world A/B came back inconclusive/noise-level
   and the ncu re-profile needed to confirm occupancy actually moved was never completed
   (ERR_NVGPUCTRPERM last session). This pass fixes that gap with real ncu --set full data.

   Setup

   • Pulled prism-ml/Ternary-Bonsai-27B-gguf / Ternary-Bonsai-27B-Q2_0.gguf (7.17 GB) onto T5500
     into the HF hub cache (hf download, HF_HUB_DISABLE_XET=1 — the default Xet transfer stalled
     indefinitely at 0 bytes; classic HTTP download worked cleanly). Fresh worktree
     issue/199-cuda-tensorcore-decode-attention off dev (b76a6aa, includes perf(cuda): port Vulkan's tuned split-KV attention heuristic to replace fixed ATTN_KV_SPLIT=4 #197/perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198/perf(cuda): combined GQA-group register-blocking + tuned split-KV attention (#197, #198) #201).
     native/build_ptx.bat + dotnet build -c Release, GPU-lock held for the whole profiling session.
   • ncu itself works fine on this host now (v2025.1.1.0, no elevation needed this session — last
     session's ERR_NVGPUCTRPERM did not reproduce).
   • Methodology note for anyone reusing this: --launch-skip/--launch-count count every launch
     matching --kernel-name regex:attention_f32, including the initial prefill (seqQ=8, grid=192)
     and — important, cost me a couple of throwaway captures — the --depth warm-context in dotllm bench is applied as one big batched forward (grid = depth * numHeads, e.g. 6000 for
     depth=250), not per-token decode calls. Only the timed -n steps are true seqQ==1 decode.
     DOTLLM_ATTN_GQA_SPLIT's kernel also only engages once seqKv >= AttentionGqaSplitMinSeqKv
     (default 256) — a shallow -p 8 -n 12 probe with no --depth never crosses that gate and
     silently profiles the plain-fallback kernel, which is what my first pass here accidentally did.
     Used -p 8 --depth 250 -n 12 (seqKv 258-270) to get real engagement, confirmed via a cheap
     ncu --metrics gpu__time_duration.sum sweep (single lightweight metric, no --set full replay
     cost) before spending a full profiled capture on it.

   Results — same shape (Bonsai-27B, seqKv≈258-265), same host, --set full, 8 launches each

   Metric	attention_f32 (baseline, unchanged)	attention_f32_gqa_split_kv (#197/#198, opt-in)
   Grid	(24,1,1) = numHeads	(4,3,1) = numKvHeads × kv_split=3 → 12 blocks
   Duration	105.1-106.1 us	174.6-176.2 us (≈1.66-1.68× SLOWER)
   Waves Per SM	0.14	0.09 (worse — smaller grid underfills the 3060's 28 SMs even more)
   Theoretical Occupancy	100%	83.33% (register-blocking costs some occupancy ceiling)
   Achieved Occupancy	16.5-16.7%	16.5-16.7% (statistically identical — no improvement)
   Compute (SM) Throughput	9.2-9.3%	9.3-9.4% (flat)
   Memory Throughput	44.6-45.0%	12.4-13.2% (dropped ~3.4×)
   Warp Cycles Per Issued Instruction	25.5-25.7	13.2-13.3 (improved ~1.9×)

   (Shallow-depth, seqKv≈8-20, baseline-only sanity capture reproduced the original pre-#197/#198
   numbers from docs/CUDA.md almost exactly: grid=24, Waves/SM 0.14, Achieved Occupancy 16.6-16.8%,
   Compute 4.2-4.3%, Memory 8.6-8.7%, Warp Cycles/Instr 42.3-42.8 — confirms the baseline path is
   genuinely untouched by #197/#198 and that this session's tooling/methodology reproduces the earlier
   elevated-ncu findings faithfully.)

   Reading

   The per-instruction stall metric did improve roughly 2× (25.6 → 13.2 cycles/instruction) —
   the GQA-group register-blocking design is doing what it was meant to do at the warp level. But
   that improvement did not translate into either higher achieved occupancy (flat at ~16.6%) or a
   shorter kernel (1.66-1.68× longer, not shorter). Two things are working against the intended
   fix simultaneously:

   1. Grid got smaller, not bigger. numKvHeads × kv_split = 4 × 3 = 12 blocks is half the
      already-too-small 24-block baseline grid. Waves/SM went from bad (0.14) to worse (0.09).
   2. Memory throughput dropped ~3.4× (44.6% → 12.4%) despite the kernel's whole design premise
      being fewer, better-amortized K/V reads. This points at the cooperative-launch (grid.sync())
      and cross-block partial-result combine overhead dominating wall time, not the QK/PV compute
      itself — consistent with perf(cuda): GQA-group K/V-read amortization + more resident warps/block for decode attention #198's own doc noting split-KV's grid-sync+combine overhead "has little
      sequential work left to amortize against" at shallow per-split KV depth (here: 258/3 ≈ 86 rows
      per split, still fairly thin).

   Net: the diagnosed root cause from the original elevated ncu pass — latency-bound, not
   compute- or memory-bandwidth-bound, near-zero achieved occupancy from a badly underfilled grid —
   is unchanged, and on the raw per-launch-duration axis, worse with #197/#198's fix active. The
   explicit precondition for #199 ("out of the latency-bound regime") is not met. Tensor cores
   raise compute throughput; compute throughput here is ~9% either way and was never the
   bottleneck — a tensor-core rewrite would not address either of the two things actually broken above
   (grid fill, cooperative-launch overhead), and would add the real precision risk this issue already
   flagged as HIGH/needs-new-tolerance-characterization, for a kernel unlikely to show a measurable
   win.

   Recommendation

   #199 stays parked, not implemented. No tensor-core kernel code was written this pass — this
   was specifically the confirm-the-precondition step. Suggested next steps for whoever revisits
   decode attention perf, in priority order:

   1. File a follow-up on the GQA-split kernel itself: grid shrinking to numKvHeads × kv_split (12
      for Bonsai-27B) rather than staying ≥ the 24-block baseline looks like a tuning bug, not a
      fundamental limit — kv_split should probably be chosen to keep total blocks ≥ baseline grid,
      not computed independently of it. Worth a quick occupancy-heuristic fix + re-profile before
      trying anything more invasive.
   2. The memory-throughput regression (44.6% → 12.4%) suggests the grid-sync/combine overhead, not
      the K/V amortization, dominates at this depth — worth an ncu source/instruction-level pass on
      attention_f32_gqa_split_kv specifically (not just SoL-level metrics) to confirm before assuming
      the split design itself is sound and just mistimed.
   3. Only after (1)/(2) actually move Achieved Occupancy and Waves/SM meaningfully off their current
      ~16.6%/0.14 floor should a tensor-core kernel (perf(cuda): tensor-core (mma.sync FP16) decode attention kernel — port Vulkan's coopmat design #199) be reconsidered — and even then, per the
      issue's own risk note, it should start as an opt-in precision-gated spike, not a default-path
      replacement.

   No code changes made. GPU-lock held only for the profiling session on T5500, released cleanly
   (scripts/gpu-lock.sh release issue199). Scratch worktree
   E:\Development\dotLLM-worktrees\issue-199-cuda-tensorcore-decode and the raw .ncu-rep captures
   left on T5500 (not committed — same convention as the .perf-runs/ncu-2026-07-27/README.md
   precedent: this comment is the durable record) for anyone who wants to re-import them.
3. jamesburton

   added 2 commits that reference this issue on Jul 30, 2026

   docs(cuda): correct the #199 premise-check with a real depth-512 re-p…

   ...

   904e0a8

   perf(cuda): remove redundant barrier in attention_f32_gqa_split_kv; h…

   ...

   78082b9
4. jamesburton commented on Jul 30, 2026

   

   jamesburton

   on Jul 30, 2026

   OwnerAuthor

   More actions

   Merged via PR #243 (30a48ed on dev), default ON. Composed tensor-core + GQA-group decode attention: 2-2.6x faster than baseline, 2-2.5x faster than the #197/#198 kernel itself, real generation-parity validated (pre-gate bit-identical, post-gate PPL improves -0.173%). Full writeup in docs/CUDA.md and [[prismml-bonsai-model]] memory. #200 remains open separately (low priority, waits on continuous batching).
5. jamesburton

   closed this as completedon Jul 30, 2026