fix: validate memory_order per-operation for atomic_load/store (Fixes #2574)#2589
fix: validate memory_order per-operation for atomic_load/store (Fixes #2574)#2589muhamedfazalps wants to merge 2 commits into
Conversation
…ile-ai#2574) T.atomic_load previously accepted memory_order='release'/'acq_rel' and T.atomic_store accepted 'consume'/'acquire'/'acq_rel' — memory orders that are illegal for those operations in C++/libcu++. The kernel compiled cleanly but died with a device-side assert inside libcu++ at runtime. Add per-operation legality checks in atomic_load and atomic_store that validate the memory_order against the C++ standard rules before emitting the call. Illegal orders now raise a clear ValueError at compile time. Also updated docstrings to list only the per-operation legal orders.
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Pro Run ID: 📒 Files selected for processing (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughWalkthroughAtomic loads and stores now validate operation-specific memory orders before code generation. A regression test module covers rejected orders and legal orders for CUDA compilation. ChangesAtomic memory-order validation
Estimated code review effort: 2 (Simple) | ~10 minutes Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
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. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (1)
testing/python/language/test_tilelang_issue_2574.py (1)
45-53: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueConsider adding
"consume"to the legal load orders test.
_LOAD_ORDERSincludes"consume", buttest_atomic_load_accepts_legal_ordersonly parametrizes over["relaxed", "acquire", "seq_cst"]. Adding"consume"would close the coverage gap and ensure the legal-order path is fully exercised.🧪 Suggested addition
-@pytest.mark.parametrize("order", ["relaxed", "acquire", "seq_cst"]) +@pytest.mark.parametrize("order", ["relaxed", "consume", "acquire", "seq_cst"]) def test_atomic_load_accepts_legal_orders(order):🤖 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/language/test_tilelang_issue_2574.py` around lines 45 - 53, Expand the parameterization of test_atomic_load_accepts_legal_orders to include "consume", matching the legal values defined in _LOAD_ORDERS and covering that accepted load-order path.
🤖 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.
Nitpick comments:
In `@testing/python/language/test_tilelang_issue_2574.py`:
- Around line 45-53: Expand the parameterization of
test_atomic_load_accepts_legal_orders to include "consume", matching the legal
values defined in _LOAD_ORDERS and covering that accepted load-order path.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
Run ID: a0fc93e0-c5f4-40d6-a442-9e067f57e5f0
📒 Files selected for processing (2)
testing/python/language/test_tilelang_issue_2574.pytilelang/language/atomic.py
1 participant
Fixes #2574
Summary
T.atomic_loadpreviously acceptedmemory_order="release"/"acq_rel"andT.atomic_storeacceptedmemory_order="consume"/"acquire"/"acq_rel"— memory orders that are illegal for those operations in C++/libcu++. The kernel compiled cleanly but died with a device-side assert inside libcu++ at runtime:C++ restricts which memory orders are legal per atomic operation:
relaxed/consume/acquire/seq_cst— neverreleaseoracq_relrelaxed/release/seq_cst— neverconsume,acquire, oracq_relThis PR adds per-operation legality checks in
atomic_loadandatomic_storethat validate thememory_orderagainst these rules before emitting the call. Illegal orders now raise a clearValueErrorat compile time:Changes
tilelang/language/atomic.py: Added_LOAD_ORDERSand_STORE_ORDERSsets defining per-operation legal memory orders. Added validation inatomic_load()andatomic_store()before the_MEMORY_ORDER_ID_MAPlookup. Updated docstrings to list only the per-operation legal orders.testing/python/language/test_tilelang_issue_2574.py: Added pytest tests that verify:atomic_loadrejectsreleaseandacq_relatomic_storerejectsconsume,acquire, andacq_relTesting
Summary
memory_ordervalidation forT.atomic_loadandT.atomic_store(only the C++/libcu++ legal orders per operation are accepted).ValueErrors during compilation/code-generation setup instead of failing with device-side assertions.atomic_load/atomic_storedocstrings to list only supported memory orders and reflect the new error behavior.#2574covering illegal and legalmemory_ordercombinations for CUDA targets.C++ style / lint notes
docs/developer_guide/cpp_style.md) or any C++ sources.