Revises all 12 issues from the self-review pass:
1. (HIGH) Soften bit-equivalence claim — single-sample helper and
batched kernel at B=1 are different CUDA kernels; FP order may
differ. Acceptance is relative_eq ≤ 1e-6, not bit-exact.
2. (HIGH) Explicit asymmetry: only cfc_step has a single-sample GPU
oracle. GRN/VSN/attn bwd rely on smoke + chain-rule preservation.
feedback_no_cpu_test_fallbacks.md forbids a CPU reference oracle.
3. (HIGH) Realistic targets — 3× floor, 5× stretch. Drops 15× claim
which was Amdahl-bounded under any reasonable assumption.
4. (MED) AdamW-after-reducer invariant stated explicitly: final grad
buffer is OVERWRITE by reducer, meaningful only after reducer ran.
5. (MED) New B=32 smoke test (stacked_trainer_loss_shrinks_at_batch_32)
actually exercises the cross-batch reduction path; existing
perception_overfit suite is all B=1.
6. (MED) Rollout commit 1 bundles reduce_axis0 + first consumer
(cfc_step refactor) to avoid feedback_wire_everything_up.md
orphan-kernel anti-pattern.
7. (LOW) Drop "merge to ml-alpha-phase-a" — user already chose direct
commits to that branch; clarify in Rollout.
8. (LOW) Add explicit scratch-sizing formula:
scratch_bytes ≈ B × Σ(param tensor sizes per kernel).
9. (LOW) Remove resolved open question (memset_zeros ordering).
10. (STRUCT) Post-refactor bottleneck analysis section — names the
next L40S floor (Mamba2 scan, launch latency, cuBLAS).
11. (STRUCT) cudaFuncSetAttribute note — refactored cfc_bwd's per-
block shared mem drops to 1 KiB; no attribute change needed.
12. (STRUCT) Explicit Rollback section — atomic-commit-per-kernel +
revert strategy; rollback baseline is the spec commit
(54aa69c10) on ml-alpha-phase-a.
Plus locks the target pool to L40S — speedup must be attributable to
the refactor, not to a hardware bump. H100 / BF16 / larger batch
become candidates for a follow-up spec once the L40S floor is known.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
ea1d8703b7