Skip to content

[BugFix] Keep loop-invariant stores during vectorize planning - #2922

Closed
penguin-wwy wants to merge 1 commit into
tile-ai:mainfrom
penguin-wwy:fix_vectorize
Closed

[BugFix] Keep loop-invariant stores during vectorize planning#2922
penguin-wwy wants to merge 1 commit into
tile-ai:mainfrom
penguin-wwy:fix_vectorize

Conversation

@penguin-wwy

@penguin-wwy penguin-wwy commented Aug 8, 2026

Copy link
Copy Markdown
Contributor

VectorizePlanner::ComputeBufferVectorSize returns vector_size=1 for stores whose address is invariant within the vector boundary, because vectorizing them produces a broadcast store: every lane reads the same old value and
stores to the same address, so the last lane silently wins (accumulation pattern).
However, Plan() then bucketed such entries by address invariance alone and moved them to the local/fragment bucket, which the simple-case strategy (GCD(memory_min, non_cast_call_node_min)) ignores — the constraint
was dropped and the loop was vectorized anyway.

Summary

  • Fixed vectorization planning to preserve loop-invariant global and shared-buffer stores as memory constraints.
  • Added a CUDA regression test for invariant-store accumulation.
  • Verified results against Init + a.sum(dim=1) using float16 tensors and toleranced comparison.

C++ style / lint notes

  • The PR changes C++ code but does not modify the rules in docs/developer_guide/cpp_style.md.
  • The C++ API Style Audit is not relevant because the PR adds no public C++ API.
  • No correctness or build issues are identified. Any audit findings would be warning-only.

Tests

  • Added vectorize_invariant_store_accumulate(M, K).
  • Added test_vectorize_invariant_store_accumulate().

@github-actions

github-actions Bot commented Aug 8, 2026

Copy link
Copy Markdown

👋 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! 🚀

@penguin-wwy
penguin-wwy marked this pull request as draft August 8, 2026 16:50
@coderabbitai

coderabbitai Bot commented Aug 8, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Walkthrough

Loop-vectorization planning now treats loop-invariant global and shared-buffer stores as memory constraints. A CUDA kernel and float16 regression test cover accumulation into an invariant output location.

Changes

Vectorize Accumulation

Layer / File(s) Summary
Invariant store constraint and regression test
src/transform/loop_vectorize.cc, testing/python/language/test_tilelang_language_vectorize.py
Keeps invariant global and shared-buffer stores in the memory constraint bucket. Adds a kernel that accumulates each matrix row into one output element and tests the expected float16 results.

Estimated code review effort: 2 (Simple) | ~10 minutes

Possibly related PRs

Suggested reviewers: leiwang1999, ljc00118, siriusneo

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ 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%. 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.
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.
Title check ✅ Passed The title clearly and concisely describes the loop-vectorization fix implemented by the 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

Comment @coderabbitai help to get the list of available commands.

@penguin-wwy
penguin-wwy marked this pull request as ready for review August 8, 2026 17:09
@SiriusNEO SiriusNEO self-assigned this Aug 10, 2026
elem_offset, inner_for_->loop_var, vector_size_, analyzer_);
}
if (depends_on_loop_var) {
if (depends_on_loop_var || info.is_store) {

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This check is only in the global/shared scope. Consider the following case:

acc = T.alloc_local((1,), T.float32)
  for i in T.vectorized(4):
      acc[0] += A[row, i]

This loop-invarient store to a local/fragment seems to have the same problem. Maybe we should fix them too.

# store: every lane reads the same old B[row], and the last
# lane's store wins, dropping all addends but A[row, K-1].
for k in T.vectorized(K):
B[row] = B[row] + A[row, k]

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In my opinion, this should either report an error, or throw a warning and don't vectorize it. Vectorization should not be together with an accumulation statement.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants