[Bug] Fix hanging from reduction on sm120

[PannenetsF]

Summary by CodeRabbit

• Chores
  • Expanded reduction operation support to include SM120 hardware targets.

_{✏️ Tip: You can customize this high-level summary in your review settings.}

[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! 🚀

[coderabbitai]

📝 Walkthrough

Walkthrough

The changes extend SM120 GPU target support to existing AllReduce operation paths by updating target-checking conditions in reduce and finalize reducer operations to treat SM120 the same as Hopper and Sm100 targets, routing them through the run_hopper path.

Changes

Cohort / File(s)	Summary
AllReduce SM120 Target Extension src/op/finalize_reducer.cc, src/op/reduce.cc	Extended target-checking conditions to include TargetIsSM120() alongside Hopper and Sm100; SM120 targets now route through the run_hopper AllReduce path in both ReduceOpNode and FinalizeReducerOp lowering

Estimated code review effort

🎯 1 (Trivial) | ⏱️ ~3 minutes

Poem

🐰 ✨

SM120 joins the hopper band,
With Hopper and Sm100, hand in hand,
AllReduce paths now expand their reach,
Two simple tweaks, a lesson to teach!

🐰✨

Pre-merge checks and finishing touches

✅ Passed checks (3 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly identifies the specific issue being fixed: a hanging problem in reduction operations on SM120 targets.
Docstring Coverage	✅ Passed	Docstring coverage is 100.00% which is sufficient. The required threshold is 80.00%.

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment

---

📜 Recent review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 2e82f37 and c6b0ef3.

📒 Files selected for processing (2)

• src/op/finalize_reducer.cc
• src/op/reduce.cc

🧰 Additional context used

🧬 Code graph analysis (2)

src/op/reduce.cc (1)

src/target/utils.cc (6)

• TargetIsHopper (52-57)
• TargetIsHopper (52-52)
• TargetIsSm100 (59-64)
• TargetIsSm100 (59-59)
• TargetIsSM120 (66-71)
• TargetIsSM120 (66-66)

src/op/finalize_reducer.cc (1)

src/target/utils.cc (6)

• TargetIsHopper (52-57)
• TargetIsHopper (52-52)
• TargetIsSm100 (59-64)
• TargetIsSm100 (59-59)
• TargetIsSM120 (66-71)
• TargetIsSM120 (66-66)

⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (1)

• GitHub Check: Test for Python 3.12 with Metal (on macos-latest)

🔇 Additional comments (2)

src/op/finalize_reducer.cc (1)

101-102: LGTM! Consistent SM120 support across finalize_reducer.

The change correctly extends SM120 support to match the pattern in src/op/reduce.cc, ensuring consistent AllReduce handling across both reduce and finalize_reducer operations. This maintains architectural consistency and should resolve hanging issues on SM120 targets.

src/op/reduce.cc (1)

312-313: LGTM! SM120 target support added correctly.

The change appropriately extends the AllReduce path to include SM120 targets (Blackwell architecture), routing them through run_hopper alongside Hopper and Sm100. The underlying run_hopper implementation in src/tl_templates/cuda/reduce.h uses the bar.sync barrier synchronization instruction, which is fully supported on SM120 architecture.

---

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 and usage tips.}