[TIR][Transform] Handle grid barriers and unbounded pointer ranges

[LeiWang1999]

Summary

• Recognize tl.sync_grid as an existing synchronization point during shared-memory thread-sync planning, avoiding redundant block barriers.
• Keep pointer-access conflict analysis conservative when touched ranges contain TVM's symbolic unbounded sentinels, preventing integer-versus-handle comparison failures.

Changes

• Record grid-wide synchronization as a kSync access entry for the active storage scope.
• Stop disjointness proofs before ordering non-integer interval bounds.
• Cover both an unbounded data-dependent atomic range and a grid barrier separating the same hazard pattern.

Validation

• ./format.sh
• cmake --build build -j16
• python -m pytest testing/python/transform/test_tilelang_transform_thread_sync.py -x (41 passed)

Summary

• Recognize tl.sync_grid as a synchronization point for the active storage scope.
• Record grid synchronization as kSync access to avoid redundant block barriers.
• Handle symbolic unbounded pointer ranges conservatively.
• Add CUDA regression tests for atomic ranges and grid barriers.

C++ style / lint notes

• The PR changes C++ code.
• No changes affect rules documented in docs/developer_guide/cpp_style.md.
• The C++ API Style Audit remains warning-only.
• No correctness, build, or test issues are reported.
• No new blocking style issue is identified.

Validation

• Formatting completed successfully.
• Build completed successfully.
• 41 thread-sync transform tests 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]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 57e4c8a1-2e9b-42f8-811f-633032ab38a8

📥 Commits

Reviewing files that changed from the base of the PR and between 23c29fa and b38e6a1.

📒 Files selected for processing (2)

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

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

---

📝 Walkthrough

Walkthrough

The thread synchronization planner now treats tl::sync_grid() as a storage-scope barrier and handles non-integer symbolic bounds conservatively. Two CUDA regression tests cover both behaviors.

Changes

Thread synchronization fixes

Layer / File(s)	Summary
Conservative symbolic-bound analysis src/transform/thread_storage_sync.cc, testing/python/transform/test_tilelang_transform_thread_sync.py	PointerAccessIsDisjoint avoids comparisons for non-integer bounds. A regression test verifies conservative synchronization for data-dependent atomic accesses.
Grid synchronization barrier tracking src/transform/thread_storage_sync.cc, testing/python/transform/test_tilelang_transform_thread_sync.py	tl::sync_grid() is recorded as a full synchronization access. A regression test verifies that no extra shared-memory synchronization is emitted afterward.

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

Merge Risk: ⚪ Minimal · up to b38e6

This PR adds handling for grid barriers and unbounded pointer ranges with focused test coverage; no actionable merge-blocking risk remains after normal checks and review.

Possibly related PRs

• tile-ai/tilelang#2738: Modifies storage-scope synchronization handling and adds barrier regression tests.
• tile-ai/tilelang#2805: Modifies PointerAccessIsDisjoint analysis and related regression tests.
• tile-ai/tilelang#2821: Modifies synchronization and access-analysis logic in thread_storage_sync.cc.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 60.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly and concisely summarizes both main changes: grid barrier handling and unbounded pointer range analysis.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

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

[LeiWang1999]

@regression-perf

[github-actions]

Performance Regression Test Report

Triggered by: @LeiWang1999
Workflow run: https://github.com/tile-ai/tilelang/actions/runs/32221932980

Results

File	Original Latency	Current Latency	Speedup
example_warp_specialize_gemm_softpipe_stage2	0.0149955	0.0151108	0.992371
example_mha_fwd_bhsd	0.0070268	0.00708043	0.992426
example_gqa_sink_bwd_bhsd	0.0250054	0.0251879	0.992753
example_dequant_gemm_fp4_hopper	0.53323	0.536853	0.993252
example_mha_fwd_varlen	0.0204857	0.0206191	0.993528
example_tilelang_sparse_gqa_decode_varlen_mask	0.0282698	0.0283968	0.995527
fp8_lighting_indexer	0.0121523	0.0122055	0.99564
example_dequant_gemm_w4a8	2.68061	2.69224	0.995682
example_convolution	0.586377	0.588714	0.99603
example_linear_attn_fwd	0.0228784	0.0229631	0.996313
example_dequant_gemv_fp16xint4	0.0176318	0.0176875	0.996847
example_gemm	0.0143881	0.0144327	0.996911
example_tilelang_nsa_fwd	0.00398822	0.00400026	0.996992
example_mha_sink_bwd_bhsd	0.0407073	0.0408251	0.997113
example_fusedmoe_tilelang	0.0764902	0.0767048	0.997203
example_gqa_bwd_tma_reduce_varlen	0.0275526	0.0276229	0.997457
example_per_token_cast_to_fp8	0.00431809	0.0043285	0.997596
example_mhc_post	0.0654378	0.0655925	0.997641
example_dynamic	0.389241	0.390119	0.99775
example_gemm_intrinsics	0.0200323	0.0200766	0.99779
example_mha_sink_bwd_bhsd_sliding_window	0.0262533	0.0263102	0.997837
example_tilelang_nsa_decode	0.00419278	0.00420174	0.997867
example_gqa_bwd	0.0280415	0.0280959	0.998063
example_mha_bwd_bshd	0.0141058	0.0141257	0.998596
example_gqa_decode	0.0308412	0.0308817	0.998689
example_mha_fwd_bshd	0.0149156	0.0149341	0.998759
example_tilelang_gemm_splitk	0.592003	0.592735	0.998764
example_tilelang_gemm_splitk_vectorize_atomicadd	0.585109	0.58583	0.998769
sparse_mla_fwd	0.0529622	0.0530093	0.999111
example_mhc_pre	0.115243	0.115321	0.999324
sparse_mla_fwd_pipelined	0.0342715	0.034293	0.999375
example_tilelang_gemm_fp8	0.172815	0.172909	0.999457
example_warp_specialize_gemm_copy_0_gemm_1	0.0234247	0.0234371	0.999473
example_elementwise_add	0.0689941	0.0690259	0.99954
example_mha_sink_fwd_bhsd_sliding_window	0.0095205	0.00952304	0.999734
example_warp_specialize_gemm_copy_1_gemm_0	0.0150618	0.015065	0.999789
block_sparse_attn_tilelang	0.00620548	0.00620599	0.999918
example_gemv	0.144922	0.144915	1.00005
example_linear_attn_bwd	0.0973373	0.0973171	1.00021
example_convolution_autotune	0.595552	0.595414	1.00023
topk_selector	0.0271992	0.0271889	1.00038
example_dequant_gemm_bf16_fp4_hopper	0.268519	0.2684	1.00044
example_tilelang_gemm_fp8_2xAcc	0.0678928	0.067847	1.00067
example_mha_inference	0.0331413	0.0331174	1.00072
example_vertical_slash_sparse_attn	0.135363	0.135245	1.00087
example_gqa_fwd_bshd	0.0296384	0.0296067	1.00107
example_tilelang_block_sparse_attn	0.00573976	0.00573332	1.00112
example_tilelang_sparse_gqa_decode_varlen_indice	0.0108521	0.0108379	1.00131
example_dequant_gemm_bf16_mxfp4_hopper	0.258994	0.258607	1.0015
sparse_mla_bwd	0.137253	0.137007	1.00179
example_mha_sink_fwd_bhsd	0.00967507	0.00963333	1.00433
example_blocksparse_gemm	0.0117533	0.011702	1.00438
example_warp_specialize_gemm_barrierpipe_stage2	0.0247393	0.0246239	1.00469
example_mha_bwd_bhsd	0.014051	0.0139621	1.00637
example_gqa_sink_bwd_bhsd_sliding_window	0.0152212	0.0151246	1.00639
example_group_per_split_token_cast_to_fp8	0.00569321	0.00565597	1.00659
example_mla_decode	0.311464	0.309023	1.0079

Artifacts

• regression_result.png (speedup plot) is attached as a workflow artifact. Download it from the workflow run page above.