Split a loop into a list of consecutive loops. It requires: - The loop can't have annotation or thread binding. - The loop must start with 0. Predicates may be added to ensure the total loop numbers keeps unchanged. In `factors`, at most one of the factors can be No
(
self,
loop: LoopRV,
factors: list[int | ExprRV | None],
preserve_unit_iters: bool = True,
disable_predication: bool = False,
)
| 736 | |
| 737 | @type_checked |
| 738 | def split( |
| 739 | self, |
| 740 | loop: LoopRV, |
| 741 | factors: list[int | ExprRV | None], |
| 742 | preserve_unit_iters: bool = True, |
| 743 | disable_predication: bool = False, |
| 744 | ) -> list[LoopRV]: |
| 745 | """Split a loop into a list of consecutive loops. It requires: |
| 746 | |
| 747 | - The loop can't have annotation or thread binding. |
| 748 | - The loop must start with 0. |
| 749 | |
| 750 | Predicates may be added to ensure the total loop numbers keeps unchanged. |
| 751 | In `factors`, at most one of the factors can be None, |
| 752 | which will be automatically inferred. |
| 753 | |
| 754 | Parameters |
| 755 | ---------- |
| 756 | loop : LoopRV |
| 757 | The loop to be split |
| 758 | |
| 759 | factors: List[int | ExprRV | None] |
| 760 | The splitting factors |
| 761 | Potential inputs are: |
| 762 | |
| 763 | - None |
| 764 | - ExprRV |
| 765 | - Positive constant integers |
| 766 | |
| 767 | preserve_unit_iters : bool |
| 768 | Whether or not to preserve unit iterators in block bindings |
| 769 | |
| 770 | disable_predication : bool |
| 771 | If enabled, don't create a predicate for guarding the loop. This can |
| 772 | be useful when splitting with scalable factors that the schedule writer |
| 773 | knows are divisible by the loop bound. |
| 774 | |
| 775 | Warning: enabling this feature may result in incorrect code generation |
| 776 | if not used carefully. |
| 777 | |
| 778 | Returns |
| 779 | ------- |
| 780 | split_loops : List[LoopRV] |
| 781 | The new loops after split |
| 782 | |
| 783 | Examples |
| 784 | -------- |
| 785 | |
| 786 | Before split, in TensorIR, the IR is: |
| 787 | |
| 788 | .. code-block:: python |
| 789 | |
| 790 | @T.prim_func(s_tir=True) |
| 791 | def before_split(a: T.handle, b: T.handle) -> None: |
| 792 | A = T.match_buffer(a, (128, 128)) |
| 793 | B = T.match_buffer(b, (128, 128)) |
| 794 | for i, j in T.grid(128, 128): |
| 795 | with T.sblock("B"): |
no outgoing calls