[BugFix][Transform] Preserve nested cast-store guards - #2820
Conversation
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
📝 WalkthroughWalkthroughThe change updates ChangesNested guard preservation
Estimated code review effort: 3 (Moderate) | ~20 minutes Suggested reviewers: Sequence Diagram(s)sequenceDiagram
participant VectorizedLoopTransformation
participant extract_if_condition
participant GeneratedCopyLoops
VectorizedLoopTransformation->>extract_if_condition: Validate and combine nested predicates
extract_if_condition-->>VectorizedLoopTransformation: Return condition or no result
VectorizedLoopTransformation->>GeneratedCopyLoops: Create guarded copy-from and copy-to loops
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
e279bfc to
986dc8e
Compare
Conjoin root-level nested if predicates when DecoupleTypeCast builds copy loops, and conservatively skip branch shapes that cannot be represented by one common predicate. This prevents generated copies from writing lanes left untouched by the source. Co-authored-by: dingsg <shengge.ding@enflame-tech.com>
986dc8e to
6b23e09
Compare
Fixes #2616
Summary
DecoupleTypeCastmoves cast stores into generated copy loops. When a source store was nested under severalifstatements, the transform kept only one guard. The generated copy could therefore overwrite lanes that the source program left untouched.The pass now combines nested no-
elseguards. If the surrounding control flow cannot be represented by one predicate, including anif/else, it leaves that store undecoupled. The pass may move a conversion, but it must preserve which lanes are written; that check belongs in the TileLang transform rather than in individual kernels.Changes
ifpredicates around cast stores.if/else, as ineligible for decoupling.Review Notes
elseguards are combined with logical conjunction.Validation
testing/python/issue/test_tilelang_issue_decouple_type_cast_nested_guards.pypreserves sentinel values in unwritten FP8 lanes on H100.git diff --checkpassed.