Repository navigation
Conversation
cc10.3 has had the fused LDTM.STAT row-max (tcgen05.ld.red.f32.max) since bring-up, but only the MXFP8 kernel read the fused_ldtm_stat flag - the per-tensor FP8 kernel ran the manual tcgen05_ld + software reduction on every part. This wires the same specialization into the per-tensor kernel: the unmasked softmax path loads S_acc and reduces the row max in one op per 64-column chunk; masked iters keep the software path (the fused max reduces before a mask could apply); cc10.0 folds to the manual path unchanged (the flag is per-specialization state in TemplateParams, so each device class gets its own traced variant, as with MXFP8). The branch is byte-for-byte the one the SM107 sibling bakes (board-proven on cc10.7, where the LDTM lever carried most of a -7.5% step at 65k), and the helper + instruction are the same ones the MXFP8 kernel exercises on cc10.3 in production. SM100 (cc10.0) suite: 34 passed, behaviorally identical. Direct cc10.3 hardware validation is queued (B300); the fused path's only new surface here is this kernel's integration, identical in shape to both existing users. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
vedaanta
added
orig-nv-eng
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (3)
🚧 Files skipped from review as they are similar to previous changes (1)
Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review. 📝 WalkthroughWalkthroughSM100 compilation now enables fused LDTM statistics for MXFP8 and per-tensor FP8 on compute capability 10.3. The FP8 prefill kernel uses fused row-max reduction for unmasked chunks and retains software reduction otherwise. ChangesSM100 FP8 reduction path
Priority: ⬇️ Low Estimated code review effort: 3 (Moderate) | ~20 minutes Change: Feature Sequence Diagram(s)sequenceDiagram
participant compile
participant PARAMS
participant prefill_d128_fp8
participant tmem_load_max_reduction_x64
compile->>PARAMS: set fused_ldtm_stat for FP8 on cc10.3
PARAMS->>prefill_d128_fp8: provide capability flag
prefill_d128_fp8->>tmem_load_max_reduction_x64: load unmasked chunk and reduce row maxima
tmem_load_max_reduction_x64-->>prefill_d128_fp8: return values and row maxima
Suggested reviewers: Merge Risk: ⚪ Minimal · up to The FP8 fused reduction change has no identified merge-blocking correctness or availability risk. 🚥 Pre-merge checks | ✅ 3 | ❌ 2❌ Failed checks (2 warnings)
✅ Passed checks (3 passed)
Full details: Description checkExplanation The description explains the implementation and includes validation evidence, but it omits most required template sections, including checklist items, affected area, explicit Why, Related issues, API and compatibility impact, and a Testing section with exact commands and results. Resolution Complete the repository template. Add the required checklist, select the affected area, provide explicit Why, state Related issues or None, document API and compatibility impact, and move testing details into a Testing section with exact commands and results.
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
|
@cudnn-ci-bot run frost |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-586-182d3a7 |
|
B300 (sm103, cc 10.3) hardware validation done — the merge gate from the PR description is cleared. Setup: B300 SXM6 node, NGC pytorch:26.03 container, FE source-built from this branch, cuDNN backend 9.25.0, nvidia-cutlass-dsl 4.7.0. The per-tensor fp8 e2e suite runs the fused |
YangXu1990uiuc
left a comment
YangXu1990uiuc
left a comment
There was a problem hiding this comment.
Codex bot review · model: gpt-6-astra
Reviewed 182d3a732a57 against base c61a5e4eb5c5.
The specialization makes sense: use the fused load/reduction only on cc10.3 and only for unmasked iterations, retain software reduction after masking, and include the flag in the existing frozen specialization key. I found no actionable correctness defect in the reviewed change.
Validation: Built a matching extension from this head. The per-tensor FP8 suite passed independently on SM103/B300 (32 passed) and SM100 (32 passed), explicitly selecting L0 and L1. B300 baseline/head/head/baseline spot check against branch point 61d6506: B=1, Hq=8, Hkv=2, D=128, E4M3 inputs and FP16 output. GPU kernel time: S=128 dense 7.49 -> 7.44 us; S=2048 dense 23.19 -> 21.77 us; S=2048 causal 25.61 -> 24.09 us; S=2048 left-window 64 19.97 -> 19.84 us; S=4096 dense 39.72 -> 36.84 us. Prepared CPU launch medians were 4.96–5.14 us baseline and 4.89–5.26 us head. Every timed arm passed O/amax reference checks, an independent LSE reference, exact replay checks after poisoning output buffers, and a changed-V zero-output control. Verified the actual specialization flag and emitted sm_103a PTX: the head contains tcgen05.ld.red; the baseline does not. A private integration overlay on develop c61a5e4 also passed 16 B300 mask, padding, sink and GQA cases. It resolves the one api_dsl.py conflict by changing only fused_ldtm_stat while preserving the current scheduler policy logic and the renamed kernel path.
Limitations: Timing covers the prepared compiled kernel, including its amax writes. It excludes graph.execute scalar reads, amax reset/divide launches, compilation, and full-application overhead; these are focused measurements, not a general end-to-end claim. The existing exact-head CI has SM120 ragged-FP8 worker crashes and guardword failures. I have not attributed those to this change. There is no pending fork test workflow to approve, and I did not launch a duplicate pipeline. Please rebase and run the current CI before merging. When rebasing, preserve develop's sched_policy=None derivation rather than restoring the older natural/LPT selection. The adjacent comment saying FP8 does not read this flag also needs to be updated.
Approved under the trial's no-P0 threshold; any P1/P2 findings above remain for the owner to address. Merge timing stays with the owner.
|
@cudnn-ci-bot run frost |
|
Could not launch CI The bot failed before the pipeline started. The details are in the bot's log; ask a maintainer to look, then try |
# Conflicts: # python/cudnn/sdpa/fwd/api_dsl.py
|
Codex bot review · model The previous CI request failed before a pipeline started. I am retrying FROST test CI for the current head ( |
|
@cudnn-ci-bot run frost |
|
🏁 Pipeline finished SHA: 22 passed, 3 failed, 6 manual
|
YangXu1990uiuc
left a comment
YangXu1990uiuc
left a comment
There was a problem hiding this comment.
Codex bot review · model: gpt-6-astra
Reviewed 93a91a9e774f against base 0ed2710bc6db.
Re-reviewed the rebase at 93a91a9e774f; no new actionable finding. The updated entry point preserves scheduler-policy derivation, gates the fused operation to cc10.3, and leaves masked iterations on their software reduction. Existing approvals remain in place; this is an incremental comment.
Validation: Reviewed all three changed files. The kernel and configuration are byte-identical to the B300-tested integration overlay from the previous review; the compile method has the same AST after removing the old unused local. Checked intervening dispatch changes and a clean merge onto develop 0ed2710bc6db. Fresh exact-head validation: 31 focused SM100 FP8 tests passed, covering masks, output dtypes/Amax, sinks, GQA, padding, strided/omitted LSE, rectangular bottom-right and short/empty Q. Retained target-hardware evidence from the previous review: 32 SM103 tests, 16 tests of the integration overlay, and the B300 CPU/GPU ABBA spot checks with capture/correctness controls. These are previous measurements, not a new-head B300 run. Retried the explicitly failed pre-launch CI request; FROST CI now runs the verified full head SHA.
Limitations: Current-head SM103 CI is still pending. Fresh local tests exercise the SM100 fallback; target-hardware timing was not repeated for this equivalent rebase.
Comment only: this PR already has an active approval; preserving existing approvals.
|
Codex bot review · model CI follow-up for current head The remaining numerical test failure is SM120 backward D256 sliding-window bitwise reproducibility, outside this SM103 FP8-forward change; I have not reproduced it on the base. Rubin failed during runner setup for lack of disk space, and the text scan flagged an unchanged convolution comment. Existing approvals are retained. |
This branch has not been deployed
5 participants
Summary
cc10.3 has the fused LDTM.STAT row-max, but only the MXFP8 kernel read the
fused_ldtm_statflag — the per-tensor FP8 kernel ran the manualtcgen05_ld+ software reduction on every part, leaving the win on the table on SM103. This wires the identical specialization into the per-tensor kernel (unmasked path only; masked iters keep the software reduction since the fused max reduces before a mask could apply). cc10.0 folds to the manual path unchanged.Three-line summary of the mechanism for reviewers new to it:
fused_ldtm_statis per-specialization state in the frozenTemplateParams(auto-set from the device in the adapter), so one kernel file yields distinct traced variants per device class — the same routing MXFP8 has used since its bring-up, and the SM107 sibling simply bakes.Evidence
tmem_load_max_reduction_x64) and instruction are the same ones the MXFP8 kernel exercises on cc10.3 in production.🤖 Generated with Claude Code
Summary by CodeRabbit