[Denverjin]

Manual T.Stage pipelines can reject valid configuration set/use groups when their stages are nonadjacent or decrease in source order. Hardware configuration registers have one physical value per core, so their dependencies must preserve the reaching setter and the next overwrite within the pipeline's issue phases.

Track HF32 mode, MMAD direction, padding, and atomic mode per core, and construct phase constraints for each setter/consumer group. Preserve conservative logical dependencies for conditional setters and opaque effects, along with the existing scalar-buffer and control/pipe-fence dependencies. Document the stage semantics for configuration changes.

Summary

• Preserve phase dependencies for HF32, MMAD direction, padding, and atomic-mode registers in manual staged pipelines when offsetting is enabled and configuration effects are known.
• Retain logical dependencies for unresolved registers, conditional setters, opaque effects, and loop-control fences. Preserve existing scalar-buffer dependency behavior.
• Pass register phase dependencies through the Python scheduler. Validate them and use them in feasibility checks, verbose output, and saved schedule metadata.
• Document configuration-register behavior across stages, including setter overwrites, reader ordering, and startup constraints.
• Add tests for staged setter/consumer ordering, nested and guarded effects, register generations, infeasible schedules, and scalar-buffer dependencies.
• Add a .clang-tidy configuration.

C++ style / lint notes

The PR changes C++ code and adds .clang-tidy. It does not change rules in docs/developer_guide/cpp_style.md. CI includes the “C++ API Style Audit (warning only)” step. No audit findings or lint results were supplied, so introduced warning counts are unavailable. No build or test results were supplied.

[coderabbitai]

No actionable comments were generated in the recent review.

Recent review info

Run configuration

• Source mode: Auto
• Repository: tile-ai/tilelang
• Branch: main
• File: .coderabbit.yaml
• Commit: baceed6
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Advanced
• Run ID: 1ac70a75-1274-4608-8f6c-27d4af20f3e7

Commits

Reviewing files that changed from the base of the PR and between d1e4d49 and 9730b5e.

Files selected for processing: 4

• src/ascend/transform/auto_schedule.cc
• testing/ascend/auto_schedule/test_schedule_integration.py
• testing/ascend/auto_schedule/test_special_registers.py
• tilelang/ascend/language/schedule_hint.py

Files skipped from review as they are similar to previous changes: 1

• tilelang/ascend/language/schedule_hint.py

Included review availability: This review used your included allowance. Your plan provides up to 8 included reviews per hour; 7 remain after this review.

Walkthrough

Manual offset scheduling now derives physical-phase dependencies for analyzable HF32, MMAD direction, padding, and atomic-mode effects. The Python scheduler applies these constraints. Unresolved configuration effects and ordinary scalar-buffer dependencies retain logical ordering.

Changes: Register Phase Scheduling

Layer / File(s)	Summary
Python scheduler phase constraints tilelang/ascend/transform/z3_scheduler.py	The Python scheduler accepts register phase dependencies, validates their inputs, applies modulo-phase constraints during feasibility checks, and includes them in verbose output and saved metadata.
Register effect analysis and scheduler wiring src/ascend/transform/auto_schedule.cc	The loop scheduler analyzes configuration-register effects and passes physical-phase dependencies to the Python scheduler. Unresolved effects and loop-control ordering retain logical dependencies.
Scheduling tests and Stage documentation testing/ascend/auto_schedule/test_special_registers.py, testing/ascend/auto_schedule/test_schedule_integration.py, tilelang/ascend/language/schedule_hint.py	Tests cover staged register ordering, conservative cases, runtime padding behavior, infeasible schedules, and scalar-buffer dependencies. The Stage documentation describes shared register values and dependency rules.

Priority: Normal

Estimated code review effort: 4 (Complex) | ~45 minutes

Change: Bug fix

Sequence Diagram(s)

sequenceDiagram
  participant LoopScheduler
  participant RegisterEffectAnalysis
  participant z3_schedule_loop_ffi
  participant z3_schedule_loop_python
  LoopScheduler->>RegisterEffectAnalysis: Analyze configuration-register effects
  RegisterEffectAnalysis-->>LoopScheduler: Return register phase dependencies
  LoopScheduler->>z3_schedule_loop_ffi: Pass serialized dependency tuples
  z3_schedule_loop_ffi->>z3_schedule_loop_python: Forward register_phase_deps

Loading

Possibly related PRs

• tile-ai/tilelang#3678: Its FlashMLA kernels use the manual T.Stage pipeline behavior affected by these configuration-register phase dependencies.

Suggested labels: bug, python, ascend, documentation

Merge Risk: Minimal · up to 9730b

No merge-blocking risk was identified. Ascend runtime validation remains advisable.

Pre-merge checks: 5

Check name	Status	Explanation
Description Check		Check skipped - CodeRabbit’s high-level summary is enabled.
Title check		The title clearly and concisely describes the main change: preserving configuration register phases in Ascend manual pipelines.
Docstring Coverage		Docstring coverage is inconclusive: files from earlier reviews still need to be rechecked within the check's analysis limits. Not blocking.
Linked Issues check		Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check		Check skipped because no linked issues were found for this pull request.

Finishing Touches

🧪 Generate unit tests (beta)

• Create a new PR

• Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[github-actions]

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀