forked from tile-ai/tilelang
-
Notifications
You must be signed in to change notification settings - Fork 6
Closed
Labels
enhancementNew feature or requestNew feature or request
Description
Required prerequisites
- I have searched the Issue Tracker that this hasn't already been reported. (comment there if it has.)
Motivation
The LegalizeVectorizeLoop pass infers the vector size of a vectorized loop according to the target and the loop extent. After deciding the vector size, it rewrites the original vectorized loop into a nested outer/inner loop, with the inner loop extent matching the vector size, making the later VectorizeLoop pass easier to process.
A simiple example:
The original vectorized loop
for i in [0, 32) vectorized:with vector_size = 8, is rewritten into
for outer in [0, 4):
for inner in [0, 8) vectorized:
body[i = outer * 8 + inner]The original LegalizeVectorizeLoop pass infer the max possible vector size based on the target type:
if (TargetIsSm100 && HasGlobalAccess)
vector_load_bits_max_ = 256;
else
vector_load_bits_max_ = 128;The SUNMMIO ZPU supports a custom vector size of 8192 bits. We need to alter the pass to support this.
Solution
No response
Alternatives
No response
Additional context
No response
Reactions are currently unavailable
Metadata
Metadata
Assignees
Labels
enhancementNew feature or requestNew feature or request