Conversation
…validate_node (NVIDIA#703) The SDPA forward graph factory creates O with output_tensor(), so it carries no dim/stride unless the caller declares them, and infer_properties_node() materializes a packed BHSD layout for an undeclared one -- the same contract the Stats, Max and Sum_exp outputs of this node already follow. pre_validate_node() checked that layout before inference ran, so a legal omission was rejected with ATTRIBUTE_NOT_SET "The dim for output_names::O is invalid". Only a caller-declared O layout is judged in pre now (dim and stride must be declared together); the materialized layout is checked in post_validate_node(), whose rank/last-stride check also replaces the old get_stride().back() -- UB on an empty stride -- and keeps the original error codes and messages. A node-author rule next to INode::pre_validate_node() records the phase contract. Related to NVIDIA#703.
|
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)
Included review availability: Your plan provides up to 12 included reviews per hour; 8 remain after this review. 📝 WalkthroughWalkthroughThe SDPA node now defers validation for undeclared output properties, infers packed BHSD dimensions and strides, and validates the materialized output. The node interface documents this validation order. Tests cover declared, inferred, partial, invalid, repeated, and expanded layouts. ChangesSDPA output validation
Priority: ⬇️ Low Estimated code review effort: 3 (Moderate) | ~25 minutes Change: Bug fix Suggested reviewers: Merge Risk: ⚪ Minimal · up to The deferred output-validation and inference paths have no identified merge-blocking risk. 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1🛠️ Fix failing CI checks 💡
🧪 Generate unit tests (beta)
Warning Some tools did not complete. Review the errors below. 🔧 Cppcheck (2.21.0)test/cpp/validate.cppinclude/cudnn_frontend_utils.h:193:9:performance:useInitializationList:Variable 'error_status' is assigned in constructor body. Consider performing initialization in initialization list. ... [truncated 8653 characters] ... nction parameter 'x' should be passed by const reference. Comment |
YangXu1990uiuc
left a comment
YangXu1990uiuc
left a comment
There was a problem hiding this comment.
Codex bot review · model: gpt-6-astra
Reviewed 33353d1f1584 against base ebe3bba998b8.
Thank you for your contribution! The validation/inference ordering change makes sense: undeclared O gets a complete layout before dependent checks, explicit layouts remain intact, and partial or unsupported declarations still fail. I found no actionable correctness or compatibility issue in this change. The shared implementation also passed independent composite/unified GQA checks with different Q and V embedding sizes.
Validation: Compiled the exact-head C++ suite with C++17 and warning-as-error flags against CUDA 13.2/cuDNN 9.25.1. Repository suite: 37 passed, 3 expected skips requiring cuDNN >=9.27; the additional independent numerical probe passed too (combined 38 passed,3 skipped,100039 assertions). Built the two new validation_order tests against exact-base headers: both fail before the fix, including the omitted/repeated/expand paths. Both pass on the reviewed head; explicit layouts and invalid declaration controls retain their behavior. On SM100, independent analytic O and Stats checks passed all16 configurations across cuDNN 9.19.1/9.25.1, COMPOSITE/UNIFIED implementations, inferred BHSD/explicit BSHD output, and Dqk=64 with Dv=64/128. B=2,Hq=4,Hkv=2,Sq=16,Skv=32; output buffers were initialized to a nonzero sentinel. Measured max O/Stats error was zero in these uniform-attention cases. Rechecked the exact head/base, mergeability and current reviews. The fork Style workflow is pending maintainer authorization, not a failed code check.
Limitations: Numerical probes cover dense GQA on SM100, not a full paged/ragged/FP8 matrix. No separate GPU performance sweep: the change is in graph validation/layout inference and does not modify execution or kernels. Warm validation measurements were noisy and are not used as a performance claim.
Approved: no P0 or high-risk P1 found. Remaining findings stay with the owner; merge timing stays with the owner.
|
I am requesting test CI for Codex bot review — model |
|
@cudnn-ci-bot run python_tests |
|
🏁 Pipeline finished SHA: 18 passed, 7 failed, 6 manual
|
3 participants
Problem
CompositeSDPANode's forwardpre_validate_node()validates the shape and layout of an output that its owninfer_properties_node()is responsible for producing.scaled_dot_product_flash_attention.h(base L285) runsCUDNN_FE_SDPA_VALIDATE_DIM_STRIDE(output_names::O, ...), which requires rank-4 dim AND stride onO. ButOis created by theGraph::sdpafactory as a virtual output (graph_interface.h:379), and the very same node already treats that contract as "absent means inferred" forStats,MaxandSum_exp. A caller who requestsOexactly as the factory hands it back is therefore rejected on a property that has not been materialized yet:Change
Only the derived-output check moves; nothing else about ordering changes.
ATTRIBUTE_NOT_SET.infer_properties_node()materializes the packed BHSD layout.INode::pre_validate_node()gains the rule a node author needs: a derivable output attribute can only be validated after its materialization, and pre must not read "not yet derived" as "unsupported".Regression tests cover both
validate_subtree()andexpand_subtree(), and were seen RED before the change and GREEN after.Verification
Run on an L20 with two real cuDNN installs,
cudnnGetVersion()self-reporting each, because the check itself carries no version gate:An explicitly-declared
Ois not rewritten: the caller's layout survives inference byte for byte. On GPU, an omitted-Ograph and an explicitly-declared one both validate, build and execute, agreeing element-wise (max |A-B| = 0 over 32768 elements) even when the two declare different layouts.C++ suite, reading the runner's own summary line:
38 cases | 35 passed | 3 skipped | 620 assertions, exit 0[validation_order]against the unfixed tree:2 cases | 0 passed | 2 failed, exit 3All tests passed (27 assertions in 2 test cases), exit 040 cases | 37 passed | 3 skipped | 647 assertions, exit 0The three skips are pre-existing and identical on both sides:
KernelCache revision()/size() require cuDNN >= 9.27.Cost
Omitted versus explicit spelling of the same graph in one process, AB/BA interleaved: first
validate()947.5 us versus 943.3 us (+0.45%, inside the inter-quartile spread), warmvalidate()3.4 us on both, fullGraph::build398.6 ms versus 399.6 ms with the sign flipping between runs. This is a correctness fix, not a speedup, and the measurements are reported that way. The base "omitted" path is not comparable -- it fails in pre-validation and never builds a graph.Not in scope
The backward has the identical shape:
CompositeSDPABackwardNode::pre_validate_node()(L1251-1253) validates the factory-created virtual dQ/dK/dV the same way and reproduces it (ATTRIBUTE_NOT_SET "The dim for output_names::dQ is invalid"). It is left untouched here and reported on #703 for whoever owns that surface, together with a second static instance insdpa_fp8_bwd.h:136-138that is unreachable on Ada because the same function requiresprop_major >= 9first.Related to #703.
Summary by CodeRabbit
Bug Fixes
Documentation
Tests