[RaceChecker] RaceChecker report warning rather than error for backward compatibility

[kurisu6912]

as title

Summary by CodeRabbit

• New Features

  • Added an option to disable data race checks during compilation.

• Improvements

  • Data race detection now logs warnings instead of terminating compilation.
  • Verification updated to correctly handle local buffers and reduce false positives.

_{✏️ 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

Adds a pass config option and attribute key to disable data-race checks, and relaxes/softens the parallel-loop data-race verifier by expanding safe buffer scopes and turning a fatal error into a warning.

Changes

Cohort / File(s)	Summary
Configuration & Attribute Declarations src/op/builtin.h, src/op/builtin.cc	Add public attribute key tl.disable_data_race_check = "tl.disable_data_race_check" and register pass config option TVM_REGISTER_PASS_CONFIG_OPTION(kDisableDataRaceCheck, Bool).
Data Race Verification Logic src/transform/verify_parallel_loop.cc	Treat buffers with scopes "local.var" and "local" as non-racy in BufferStore checks; change detected data race from a fatal log to a warning with updated message text.

Sequence Diagram(s)

(omitted — changes are small and do not introduce a multi-component sequential flow that requires visualization)

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

Possibly related PRs

• [Feat] Add strong checker to detect data racing in T.Parallel #1615: Edits the same data-race verification logic and introduces the tl.disable_data_race_check pass-config key.

Poem

🐰
I nibbled through warnings, soft and light,
Local nests now safe from frantic plight,
A switch to hush the race that once would roar,
Configured calm upon the compiler floor,
Hop, hop — peace in parallel delight.

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 20.00% which is insufficient. The required threshold is 80.00%.	You can run @coderabbitai generate docstrings to improve docstring coverage.

✅ Passed checks (2 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly and specifically summarizes the main change: converting RaceChecker from reporting errors to warnings for backward compatibility.

✨ Finishing touches

• 📝 Generate docstrings

---

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

[coderabbitai]

Actionable comments posted: 1

🤖 Fix all issues with AI agents

In @src/transform/verify_parallel_loop.cc:
- Around line 231-237: The pass currently always runs the data-race verifier;
update the pass entry (where you access PassContext) to read
PassKey.TL_DISABLE_DATA_RACE_CHECK from the PassContext and skip the verifier
when the flag is set: query the PassContext for the TL_DISABLE_DATA_RACE_CHECK
option at the start of the verifier routine and, if true, return early (before
invoking the analyzer/z3_prover logic and before emitting the LOG(WARNING)
message referencing op->buffer, op->indices, failed_vars, and failed_var_expr)
so the disable flag actually disables the check.

📜 Review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between a756074 and 00ed9be.

📒 Files selected for processing (3)

• src/op/builtin.cc
• src/op/builtin.h
• src/transform/verify_parallel_loop.cc

🧰 Additional context used

🧠 Learnings (1)

📚 Learning: 2025-09-12T09:47:46.474Z

Learnt from: kurisu6912
Repo: tile-ai/tilelang PR: 794
File: tilelang/transform/add_bufstore_wrapper.py:30-33
Timestamp: 2025-09-12T09:47:46.474Z
Learning: In TVM's PyStmtExprMutator, visit_block_ methods typically call super().visit_block_(op) to process child nodes and update internal state, but return the original op when the block itself doesn't need transformation. The pattern `return op` is correct for blocks that serve as containers where mutations happen at deeper levels.

Applied to files:

• src/transform/verify_parallel_loop.cc

🧬 Code graph analysis (1)

src/transform/verify_parallel_loop.cc (1)

src/transform/simplify.cc (16)

• op (95-104)
• op (95-95)
• op (105-108)
• op (105-105)
• op (109-112)
• op (109-109)
• op (113-136)
• op (113-113)
• op (177-180)
• op (177-177)
• op (181-184)
• op (181-181)
• op (314-324)
• op (314-314)
• op (326-337)
• op (326-326)

🔇 Additional comments (3)

src/transform/verify_parallel_loop.cc (1)

197-198: LGTM: Thread-local buffer scope exclusion is correct.

Expanding the skip condition to treat local.var and local scoped buffers as non-racy is appropriate, as thread-local buffers cannot have data races between parallel threads.

src/op/builtin.h (1)

64-65: LGTM: Constant declaration follows codebase conventions.

The new attribute key kDisableDataRaceCheck is correctly declared and follows the same pattern as other configuration constants in the file.

src/op/builtin.cc (1)

39-39: LGTM: Pass config option registration is correct.

The registration of kDisableDataRaceCheck as a Bool config option follows the correct pattern and is consistent with other option registrations in this file.

Note: The critical issue that this option is not actually checked in the VerifyParallelLoop pass has been flagged in the review of src/transform/verify_parallel_loop.cc.

[coderabbitai]

Actionable comments posted: 1

🤖 Fix all issues with AI agents

In @src/transform/verify_parallel_loop.cc:
- Around line 231-238: The warning message emitted in verify_parallel_loop (the
LOG(WARNING) that concatenates op->buffer, op->indices, failed_vars and
analyzer.z3_prover.GetModel(failed_var_expr)) has two text issues: add a space
between the closing backtick and "is written" so the output reads "...` is
written by multiple threads..." and rephrase the final clause "to pass key to
pass key to disable this check" to a clearer form such as "pass
PassKey.TL_DISABLE_DATA_RACE_CHECK in the pass configuration to disable this
check"; update the LOG(WARNING) string accordingly while keeping the same
concatenation of op->buffer, op->indices, failed_vars and the model output.

📜 Review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 00ed9be and ac6c4b0.

📒 Files selected for processing (1)

• src/transform/verify_parallel_loop.cc

🧰 Additional context used

🧠 Learnings (1)

📚 Learning: 2025-09-12T09:47:46.474Z

Learnt from: kurisu6912
Repo: tile-ai/tilelang PR: 794
File: tilelang/transform/add_bufstore_wrapper.py:30-33
Timestamp: 2025-09-12T09:47:46.474Z
Learning: In TVM's PyStmtExprMutator, visit_block_ methods typically call super().visit_block_(op) to process child nodes and update internal state, but return the original op when the block itself doesn't need transformation. The pattern `return op` is correct for blocks that serve as containers where mutations happen at deeper levels.

Applied to files:

• src/transform/verify_parallel_loop.cc

🧬 Code graph analysis (1)

src/transform/verify_parallel_loop.cc (1)

src/transform/simplify.cc (16)

• op (95-104)
• op (95-95)
• op (105-108)
• op (105-105)
• op (109-112)
• op (109-109)
• op (113-136)
• op (113-113)
• op (177-180)
• op (177-177)
• op (181-184)
• op (181-181)
• op (314-324)
• op (314-314)
• op (326-337)
• op (326-326)

⏰ 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). (2)

• GitHub Check: Test for Python 3.12 with Nightly-ROCm-7.1 (on self-hosted-amd)
• GitHub Check: Test for Python 3.12 with Metal (on macos-latest)

🔇 Additional comments (1)

src/transform/verify_parallel_loop.cc (1)

197-198: The scope assumptions in lines 197-198 are correct. The "local" and "local.var" scopes represent per-thread storage (registers and thread-local variables) and are inherently thread-safe. This is validated by the codebase pattern where lower_tile_op.cc explicitly documents that loops manipulating only local buffers indicate register usage, which justifies skipping thread binding and, by extension, race detection.