[kalebbroo]

CPU FlashAttention previously suppressed a NaN/+infinite sink result on an empty row. CUDA monolithic Flash applied ALiBi before softcap, and materialized attention returned NaN (F32/BF16) or the mean value (F16) for fully excluded rows. This draft corrects those reproduced defects.

• Apply softcap before ALiBi in both F32 and F16-KV CUDA Flash kernels; preserve non-finite normalization.
• Use an attention-owned materialized/tiled softmax that returns zero probabilities for fully excluded rows and propagates NaNs. Add F32 mask bias without clamping/narrowing it, removing the F16 mask clamp/cast buffers and separate add launches.
• Preserve NaN across standalone/fused split-combine and native half-prefill epilogues; carry NaN through fused Q8 sidecar scales. Add finite-sentinel and nonzero-dot/bias controls.
• Keep generic softmax unchanged. Rebuild the applicable SM80/SM75 artifacts with their prescribed pinned NVRTC/NVCC compilers; add Windows support to that existing helper.
• Add 45 executed-route semantic cases covering CUDA F32/F16/BF16, native cuDNN F16/BF16, Vulkan scalar F32, batch/head indexing and reduction tails. The tested cuDNN cohort already handles these rows correctly, so no cuDNN patch/rejection was introduced.

This is PR 02 under the user's subsequent authorization, stacked on #905. PR 01 remains incomplete and draft. The user authorized the next PR after the current PRs are ready and merged; their exit gates remain required.

Validation on Windows 11 / RTX 3060 / driver 610.88 / cuDNN 9.21.0:

• Full CPU lane, 13 projects: 3,219 reported passes, two known B0 failures and one skip. The failures are MeshOperateRouteTests.OperationsListsEveryOperationWithItsParameters and MiniMaxH3FlatLayoutResolutionTests.ImageResolverFindsTheVaeAndEncoderInTheModelsRootLayout. A T5 missing-reference guarded pass remains unvalidated.
• Full GPU categories run sequentially: Core 19 passes; LLM 101 passes/19 skips; Diffusion 18 passes/2 skips; Engine 12 passes/10 skips. No failures. Four reported passes are guarded two/eight-GPU NCCL/Kimi checks and remain unvalidated; skips do not establish real-weight coverage.
• All 45 new GPU semantics cases pass, with executed-family assertions or direct typed kernel launches.
• CUDA builds for net8.0/net10.0 and the benchmark build pass. Changed modules assemble without spills on their applicable SM75/SM86 lanes. Only SM86 executed; SM75 hardware remains unverified.
• An unmodified SM80 Flash rebuild matched shipped PTX after line-ending normalization before arithmetic edits.
• Eight controlled operation/boundary timing cells for the earlier e4ff1d1 production candidate pass quality and route checks using the same harness, shared bytes and unchanged tolerances against pristine B0 and the candidate. Two alternating pairs per cell; verified Swarm quiet intervals and the 3060 lock; no shared service state changed. Latency results are mixed (descriptive ratios 0.880–1.317, including a slower unchanged Vulkan route). This is an unresolved performance concern, not statistical improvement/non-inferiority. The initial default-arm F32/TF32 mismatch and all slower results remain retained. Current sibling-kernel review fixes require fresh registered timing in the GPU handoff.

Still required: affected real-generation A/B correctness and timing, frozen Qwen/Krea path and primary endpoints on the 4090, PR 01's incomplete baseline/workload evidence, and review. Bfix is not frozen; no primary speed, unchanged-quality optimization, Kitchen parity, or full-validation claim is made. Keep draft.

Raw logs, failed probes, fixtures, binaries and timing results are retained under artifacts/attention-pr02 in the PR worktree.

Detailed local evidence and timing report and 4090 handoff.

Current production test pin: 57b143f004050d484dc1b667ae8d5f91fc82ac22. Complete GPU-agent assignment and copy-back report template. Review replies address all six suggestions; fresh Codex reviews are requested. The user will return the GPU agent's report before readiness/merge and the next PR.

[kalebbroo]

@claude Review Pr

[claude]

Claude finished @kalebbroo's task in 2m 37s —— View job

---

PR review

• Read AGENTS.md and the PR diff (origin/codex/attention-evidence-pr01...HEAD)
• Review the CUDA kernels (softcap/ALiBi order, attention_softmax)
• Review the CudaBackend/CudaKernels wiring
• Review the CPU fix and tests
• Post the review

Verdict: I found no blocking defects in this PR's changes. The three reproduced defects are fixed correctly, and the tests target the right contract edges. The main gap is that the same NaN masking still exists in sibling Flash kernels, so CUDA results still depend on which route runs. I reviewed the code only; I did not build or run anything (this runner has no GPU).

What I checked and found correct

• flash_attn_f32.cu ordering: both kernels now apply softcap and then ALiBi. That matches docs/ATTENTION_CONTRACT.md §3 and the CPU kernel (FlashAttentionKernels.cs:65-71).
• CUDA empty rows with a -inf sink: these still return 0 after the new isnan(l) branch. That only holds because INF_F is FLT_MAX (flash_attn_f32.cu:16), so m starts at -FLT_MAX and __expf(-inf - (-FLT_MAX)) is 0, not NaN.
  • This is a quiet dependency: if INF_F is ever changed to a true infinity, EmptyFlashRowsPreserveSinkSemantics(-inf) starts producing NaN. A one-line comment next to the final l check would protect it.
• CPU fix (FlashAttentionKernels.cs:79, :89): without the guard, a -inf sink on an empty row computes exp(-inf - -inf), which is NaN. With the guard and the IsNaN(l) check, a +inf or NaN sink gives NaN and a -inf sink gives 0. All three cases are correct.
• attention_softmax.cu:
  • A row whose maximum is -inf produces zero probabilities and an inverse of 0.
  • NaN carries through both the max reduction (explicit isnan, since fmaxf drops NaN) and the sum.
  • A +inf score gives inf - inf, which is NaN, as intended.
  • The early row >= rows return applies to the whole block, so it cannot cause a __syncthreads divergence.
  • The tree reduction assumes a power-of-two block size, and BlockSize is fixed at 256 (CudaKernels.cs:380).
• Mask wiring (CudaBackend.cs:5565-5589): maskPointer is offset in F32 elements with stride sq*skv. Inside the kernel, bias[row*rowLen + c] addresses the per-head [Sq,Skv] block. Key-only masks are expanded first, so the indexing holds.
• Memory: dropping the clamp/cast buffers removes about 1.5× the mask size in temporary allocations on the F16 path.
• Module cleanup: the new module loads through LoadOwnedModule, so it is disposed with the others.
• Commits: all five are by kalebbroo with no co-author trailers, as AGENTS.md requires.

Issues and suggestions

1. Sibling Flash routes still turn NaN into zero (medium, consistency). The monolithic kernel now propagates a NaN l, but these still write 0:

• flash_attn_f32_split.cu:193, the split combine
• lm_decode_fused.cu:272, a verbatim copy of that combine
• attention/flash_attn_f16.cu:249-250 and lm/flash_attn_causal_f16.cu:226-227 (l > 0 ? 1/l : 0)

A NaN or non-finite Q·K in a Gemma-2 split decode therefore returns zeros, while the monolithic route returns NaN. The new tests can't catch this because they only assert flash.decode. If these are out of scope for PR 02, list them in the progress report's open items.
Fix this →

2. Finite-sentinel masks now behave differently on the F16 path (medium, undocumented). The old clamp to ±32752 kept a fully masked row with a finite sentinel such as -1e30 shift-invariant. In practice that meant softmax over the raw scores.

• Now score + (-1e30) in F32 rounds to exactly -1e30 for every key, so the row becomes a uniform mean of V.
• That matches the F32 reference and is arguably more correct, but models that pad with -1e30 or -FLT_MAX will see different output on S.13/S.15.
• No test pins it, because every test uses -inf. Please add an F16 case with a finite sentinel on a fully masked row and record the change in the report, since it affects real-generation A/B results.
  Fix this →

3. Tests only ever produce scores of 0 (low). Every case fills Q and K with zeros, so the score-plus-bias addition is only tested with a zero score. At least one case with non-trivial Q·K would test real score/bias mixing and the F16 score rounding. BiasRetainsF32Range (100000 vs 100001) does test bias range, which is the key F16 check.

4. LaunchSoftmaxF16 no longer has any callers (low). See CudaKernels.cs:3281; the softmax_f16 module is still loaded at :802. Per "reuse before writing", either delete it or note why it stays.

5. Extra launches and memory traffic in the new softmax (low).

• The masked path launches the softmax once per (b,h) (totalHeads launches). That isn't a regression, since it replaces the same number of add launches, but you could fold the mask-block index into the kernel and launch once.
• The kernel reads score and bias twice and makes a third read-write pass. For rowLen ≤ k·256 you could keep values in registers. This may account for some of the 1.1–1.3× slower timing cells.

6. The NVRTC helper can fall back to another version silently (low). On Windows it falls back from nvrtc64_130_0.dll to nvrtc64_120_0.dll, and Linux does the same with the .so names. Either fallback can produce PTX from a compiler other than the pinned one without any message. Printing nvrtcVersion() to stderr would make each artifact's compiler traceable.

Agreement with the PR body

The PR body keeps the same hedges: SM75 not executed, no speed claim, timing concern left open. I didn't see any claim in the code or report that goes beyond the evidence, so keeping it draft looks right.

[kalebbroo]

Addressed the review in 57b143f; pinned GPU handoff is in 69883e4.

1. Split decode with F32/F16 KV and both combine kernels reproduced NaN-to-zero loss. Normalization now preserves NaN, and fused Q8 metadata preserves NaN for the next projection. Direct finite/empty/NaN and one-channel value-NaN controls cover the sidecar. Native half-prefill already propagated the tested NaN through its accumulator; both epilogues now explicitly preserve a NaN denominator too.
   2–3. Added F32/F16/BF16 finite-sentinel controls with nonzero scores, an F16 -float.MaxValue case, and ordinary finite bias/nonzero-dot mixing. The report explicitly records that removing the F16 sentinel clamp can change real outputs. Existing dtype tolerances are retained.
2. LaunchSoftmaxF16 stays as a public compatibility adapter for external CudaKernels consumers; the XML summary/report explain why in-tree attention no longer uses it. Removing it would break the program's public-compatibility requirement.
3. Further mask-index launch fusion/register retention is proposed optimization work. The current fusion replaces the old per-head adds plus shared softmax; the existing mixed timing concern remains open and is assigned a fixed registered 4090 campaign. No performance claim or attribution to this mechanism is made.
4. Every NVRTC compile now prints the loaded major/minor and target; exact build identity remains in the emitted PTX header. NVCC-owned modules were rebuilt with cached Linux CUDA 13.0.88, which reproduced the unmodified fused artifact exactly before edits. Compiler ownership and SM75 floors are preserved.
   The monolithic finite-sentinel dependency is also documented next to INF_F.

Local validation: all 45 GPU semantic cases pass; full GPU categories run sequentially with Core 19, LLM 101, Diffusion 18 and Engine 12 reported passes and no failures. Resource skips/guarded passes remain unvalidated. Full CPU lane: 3,219 reported passes, the same two B0 failures, one skip. net8.0 and benchmark builds pass; changed PTX assembles without spills on the applicable SM75/SM86 lanes. No 4090 generation, final parity or Bfix claim is made. Both PRs stay draft pending hardware evidence.

@codex review

[chatgpt-codex-connector]

Codex Review Summary: ✅ Completed

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-11-15T06:34:28.776561Z	69883e4	Manual request

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

[chatgpt-codex-connector]

Codex Review: Didn't find any major issues. Bravo.

Reviewed commit: 69883e4be8

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

[kalebbroo]

PR 03 support-matrix inspection found a separate existing CPU dense exclusion gap. Read-only reproduction: F32 Q=[1,1,1,2] and K=[1,1,2,2] all zero, V filled with 7, additive bias [1,2] filled with -Infinity, scale 1. CpuBackend.ScaledDotProductAttention returns NaN instead of the contract's exact zero. AttentionKernels.SoftmaxInPlace evaluates exp(-Infinity - -Infinity). The PR 03 Core signatures and normalized method bodies match the pinned parent on net8/net10, so its file split did not introduce this.

This is recorded as mandatory-common case M03 in draft #915. PR 03 does not change arithmetic or treat the current result as allowed conformance. Keep the GPU agent's existing immutable test pins unchanged during its run; the report/exit review needs to account for this additional semantic follow-up before declaring Bfix complete.

[kalebbroo]

Imported the parent fixture-API review fix without changing production code or the active P2 pin 57b143f. The corrected operation harness is independently pinned at H=4ce2c223c6abe67b3d9096276a67b40735455575. The GPU handoff now builds the same H tools against immutable B0/P1/P2 production roots and explains which earlier specialized C# fixture results are invalid attribution. Unsupported cells remain unvalidated; no K0 exception or parity claim is inferred. Current handoff: https://github.com/HartsyAI/HartsyInference/blob/ce15e0330/benchmarks/attention/gpu-agent-handoff.md . The known CPU dense excluded-row gap remains recorded separately; this harness correction does not repair arithmetic or establish Bfix.