[BugFix][Transform] Normalize vectorized loop domains

[JayceSu98]

Fixes #2669

Summary

The vectorization planner assumed that every T.vectorized loop had a zero minimum. A loop such as T.vectorized(2, 6) reached legalization with min=2 and tripped ICHECK(is_zero(fnode->min)). Planning from the unnormalized base could also select a width that was misaligned at the first real access.

Legalization now plans on a zero-based domain and keeps the original minimum and step when rewriting accesses. Nonzero bounds are valid frontend input, so the normalization is handled in the TileLang transform rather than imposed on kernel authors.

Changes

• Normalize vectorized loop domains before vector-width planning.
• Carry the original loop minimum and step into access rewriting.
• Select widths using the actual base alignment after normalization.
• Add transform-level checks and an end-to-end CUDA regression for a nonzero loop minimum.

Review Notes

• The loop still executes the original logical values [2, 6).
• For the int32 regression, the real base alignment selects width 2 instead of an invalid width 4.
• Zero-min vectorized loops keep their existing behavior.

Validation

• testing/python/issue/test_tilelang_issue_vectorized_nonzero_min.py passed on NVIDIA A100 and NVIDIA H100.
• testing/python/transform/test_tilelang_transform_legalize_vectorized_loop.py passed.
• Python compilation and git diff --check 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]

📝 Walkthrough

Walkthrough

Loop vectorization now normalizes nonzero-minimum or non-unit-step loops before planning and rewriting. Regression tests cover structural legalization and CUDA execution for vectorized loops beginning at a nonzero index.

Changes

Loop vectorization normalization

Layer / File(s)	Summary
Normalize loop domains and planner analysis src/transform/loop_vectorize.cc	Loop analysis and vectorization planning use a zero-based, unit-step representation while preserving the original iteration mapping.
Rewrite normalized loops and validate behavior src/transform/loop_vectorize.cc, testing/python/transform/..., testing/python/issue/...	Innermost loop rewriting validates normalized extents before tiling, with structural IR and CUDA execution regression tests for nonzero minima.

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

Possibly related PRs

• tile-ai/tilelang#2752: Related handling of non-unit-step loops in vectorization transforms.
• tile-ai/tilelang#2784: Related loop-minimum and step handling in another loop transformation.
• tile-ai/tilelang#2807: Related vectorization utilities used by layout-aware lowering.

Suggested reviewers: leiwang1999

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Linked Issues check	✅ Passed	The implementation addresses issue #2669 by supporting nonzero-start vectorized loops and adds targeted regression coverage.
Out of Scope Changes check	✅ Passed	All code and test changes directly support vectorized-loop normalization and regression coverage for issue #2669.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly and concisely describes the main change: normalizing vectorized loop domains in the transform pass.

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