[BugFix][Transform] Detect cross-thread WAW hazards

[JayceSu98]

Fixes #2698

Summary

ThreadSync modeled both sides of a same-iteration shared-memory WAW check with the same thread variable. That only tests same-thread overlap and misses aliases between different threads, such as one phase writing S[tx] and another writing S[tx + 10].

Same-iteration WAW analysis now uses two thread instances, requires at least one thread dimension to differ, and renames each side's flat-Bind constraints independently. Barrier placement is owned by ThreadSync, so the fix models the cross-thread conflict there without changing kernels or inserting unconditional barriers.

Changes

• Use distinct thread variables for same-iteration WAW conflict analysis.
• Add the constraint that at least one thread dimension differs.
• Rename each side's derived constraint definitions with ConstrSet::RenameFrom.
• Preserve the existing loop-carried WAW policy boundary.
• Add deterministic barrier-placement regressions and the original CUDA runtime reproducer.

Review Notes

• The fix is deliberately limited to same-iteration WAW.
• Loop-carried WAW has separate constraint-shifting and barrier-placement concerns and is unchanged.
• Tests cover offset aliasing, provably disjoint regions, thread-private writes, and flat-Bind-derived indices.

Validation

• testing/python/issue/test_tilelang_issue_thread_sync_cross_thread_waw.py passed 10 consecutive iterations on NVIDIA A100.
• testing/python/issue/test_tilelang_issue_thread_sync_cross_thread_waw.py passed 10 consecutive iterations on NVIDIA H100.
• Full ThreadSync transform suite: 43 passed.
• clang-format and git diff --check passed.

[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

ThreadSync conflict analysis now detects same-iteration cross-thread write/write hazards while preserving same-thread loop-carried behavior. CUDA regressions cover barrier insertion, disjoint and thread-private writes, bind-derived indices, and copy/fill ordering.

Changes

ThreadSync WAW analysis

Layer / File(s)	Summary
Cross-thread write conflict modeling src/transform/thread_storage_sync.cc	Same-iteration write/write accesses use distinct-thread modeling, with updated substitutions, constraints, and variable renaming while loop-carried analysis remains supported.
Shared-memory hazard regression coverage testing/python/transform/test_tilelang_transform_thread_sync.py, testing/python/issue/test_tilelang_issue_2698.py	CUDA tests validate required barriers, omitted barriers for disjoint or thread-private writes, bind-derived indices, and copy/fill output ordering.

Estimated code review effort: 3 (Moderate) | ~20 minutes

Possibly related PRs

• tile-ai/tilelang#2700: Both PRs modify shared-memory ThreadSync barrier detection for write hazards.
• tile-ai/tilelang#2805: Both PRs update ThreadSync cross-thread write conflict modeling and regression coverage.

Suggested reviewers: leiwang1999

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Linked Issues check	✅ Passed	The changes model distinct thread writes, preserve existing cases, insert required synchronization, and add regression coverage for the linked issue [#2698].
Out of Scope Changes check	✅ Passed	All production and test changes directly support cross-thread WAW detection, synchronization, and regression coverage described in the linked issue.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly and concisely describes the main change: detecting cross-thread write-after-write hazards in the transform logic.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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.}

[coderabbitai]

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@testing/python/issue/test_tilelang_issue_2698.py`:
- Around line 16-29: Add a companion shared-copy kernel and CUDA regression test
alongside _copy_then_offset_fill and
test_copy_then_offset_fill_orders_cross_thread_writes that performs a
nonzero-offset T.clear after copying into shared memory, then copies back and
verifies the targeted slice is zero. Repeat the test across multiple executions
and compare against an expected tensor to preserve coverage of cross-thread
write ordering.

🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: f8f604aa-fe31-4302-af1e-53be6e838d8d

📥 Commits

Reviewing files that changed from the base of the PR and between 26ce2ff and bb4b86c.

📒 Files selected for processing (3)

• src/transform/thread_storage_sync.cc
• testing/python/issue/test_tilelang_issue_2698.py
• testing/python/transform/test_tilelang_transform_thread_sync.py

[coderabbitai]

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Add the required offset-clear regression.

This reproducer covers T.fill only; Issue #2698 also requires a nonzero-offset T.clear after shared copy-in. Add a companion kernel/test that verifies the cleared slice is zero across repeated CUDA executions.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@testing/python/issue/test_tilelang_issue_2698.py` around lines 16 - 29, Add a
companion shared-copy kernel and CUDA regression test alongside
_copy_then_offset_fill and test_copy_then_offset_fill_orders_cross_thread_writes
that performs a nonzero-offset T.clear after copying into shared memory, then
copies back and verifies the targeted slice is zero. Repeat the test across
multiple executions and compare against an expected tensor to preserve coverage
of cross-thread write ordering.