Reject sparse MMA K tiles with unhandled tails - #3367
dajiaohuang wants to merge 1 commit into
Conversation
|
👋 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! 🚀 |
|
Navigate logical layers of code changes, visualize relationships, and explore their blast radius. No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Repository: tile-ai/tilelang/.coderabbit.yaml Review profile: CHILL Plan: Advanced Run ID: 📒 Files selected for processing (2)
Included review availability: This review used your included allowance. Your plan provides up to 8 included reviews per hour; 4 remain after this review. 📝 WalkthroughWalkthroughThe sparse MMA lowering now rejects K tile sizes that are not divisible by the sparse MMA K atom. A CUDA regression test checks that a kernel with K=192 and block_K=96 raises the expected ChangesSparse GEMM K validation
Priority: ➖ Normal Estimated code review effort: 2 (Simple) | ~10 minutes Change: Bug fix · Severity of issue fixed: Medium Merge Risk: ⚪ Minimal · up to Sparse GEMM kernels whose K tile is not a multiple of the MMA K atom now fail at compile time with a clear error instead of silently dropping the K tail. The change is small and well-scoped. The new test could not be run locally without a CUDA environment, so it should be confirmed in CI. 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 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 |
Summary
Fixes #2605.
Validation
Summary
GemmSPMMA.lowernow raisesValueErrorwhen a sparse MMA K tile is not divisible by its dtype-specific K atom. This prevents the fallback loop from silently dropping the tile’s K tail.K=192andblock_K=96. It checks that compilation rejects the tile, although total K is divisible by the 64-element atom.Validation
git diff --checkreportedly pass.fcntl, which is unavailable in the local Windows Python 3.13 environment.