[Tracking Issue] Padded Layout Transformations #12261

Lunderberg · 2022-08-01T14:59:03Z

) This PR adds `tvm.testing.CompareBeforeAfter`, a generalization of the `BaseBeforeAfter` utility previously used in `test_tir_transform_simplify.py`, which performs unit tests that perform a transformation on a TIR function and compare the results to an expected TIR output. This arose when minimizing the boilerplate required for unit tests in the implementation of #12261.

Implementation of API in `tvm.tir.schedule` for layout transformations with padding, as part of #12261, item "Insert pad value into generated TIR, using `tir::if_then_else`, `builtin::assume`, and `builtin::undef`". Following the RFC discussion in apache/tvm-rfcs#77 (comment) and apache/tvm-rfcs#77 (comment), this commit preferentially rewrites the loops that surround a padded transformation where possible, in order to express padding in terms of `tir::if_then_else`.

This commit adds a new sub-analyzer, `TransitiveComparisonAnalyzer`, which attempts to apply multiple known comparisons to prove an unknown. For example, `a <= b` and `b <= c` imply that `a <= c`. These simplifications are necessary for simplifying conditionals resulting from padded layout transformations (apache#12261). While some of these conditions may be proven using `ConstIntBoundAnalyzer` or `IntSetAnalyzer`, each has some limitations. `ConstIntBoundAnalyzer` can only compare against a constant, `IntSetAnalyzer` internally calls `RewriteSimplifier` which can result in infinite recursion, and neither can handle not-equal conditions because it would require tracking multiple intervals per expression. Therefore, introducing a new sub-analyzer for these simplifications.

Previously, rewrites of boolean expressions were handled directly within `RewriteSimplifier` visitors. This commit extracts the handling of `And`, `Or`, and `Not` nodes into an independent internal utility. The `RewriteBooleanOperators` utility can simplifies these nodes, assuming that the arguments to these expressions are already simplified. (e.g. Finding the negation of `(a < b)` without needing to recurse into `a` or `b`.) This will be used as part of additional simplifications required by [buffer layout padding](apache#12261). Because no external functionality is added or changed by this refactor, this commit does not introduce any additional unit tests.

Previously, `RewriteSimplifier` inspected the top branch of each `And` and `Or` node, but didn't perform simplifications that would require inspection across branches nested structures. This commit introduces `SimplifyAsAndOfOrs`, which converts to an internal representation as conjunctive normal form, performs simplifications while in that form, then converts back to a `PrimExpr`. This utility is designed for data-propagation simplifications as part of apache#12261. To avoid increasing the CPU load unnecesarily, this utility is only used on an opt-in basis, either using `RewriteSimplifer::SetEnabledFeatures` when called directly, or using `PassContext` when used as part of `tir::transform::Simplify`.

This commit adds a new sub-analyzer, `TransitiveComparisonAnalyzer`, which attempts to apply multiple known comparisons to prove an unknown. For example, `a <= b` and `b <= c` imply that `a <= c`. These simplifications are necessary for simplifying conditionals resulting from padded layout transformations (apache#12261). While some of these conditions may be proven using `ConstIntBoundAnalyzer` or `IntSetAnalyzer`, each has some limitations. `ConstIntBoundAnalyzer` can only compare against a constant, `IntSetAnalyzer` internally calls `RewriteSimplifier` which can result in infinite recursion, and neither can handle not-equal conditions because it would require tracking multiple intervals per expression. Therefore, introducing a new sub-analyzer for these simplifications.

This commit adds a new sub-analyzer, `TransitiveComparisonAnalyzer`, which attempts to apply multiple known comparisons to prove an unknown. For example, `a <= b` and `b <= c` imply that `a <= c`. These simplifications are necessary for simplifying conditionals resulting from padded layout transformations (#12261). While some of these conditions may be proven using `ConstIntBoundAnalyzer` or `IntSetAnalyzer`, each has some limitations. `ConstIntBoundAnalyzer` can only compare against a constant, `IntSetAnalyzer` internally calls `RewriteSimplifier` which can result in infinite recursion, and neither can handle not-equal conditions because it would require tracking multiple intervals per expression. Therefore, introducing a new sub-analyzer for these simplifications. * Change mutable reference to mutable pointer * Remove nullptr default on Impl unique_ptr In g++ 7, defining a default constructor attempts to define the destructor, which fails because `Impl` is an incomplete type. As far as I should tell, the destructor should only be defined at the point where `~TransitiveComparisonAnalyzer` is defined, at which point `Impl` has a full definition. This issue does not occur in g++ 10. * Require opt-in for CPU-intensive simplifications * Document the intent of using bitflags * Rename "Feature" to "Extension" * Use TVM_DLL on new public member functions * Remove duplicate BaseBeforeAfter.transform definition * Explicitly enable extension for unit tests that require it * Fix accidentally duplicate test case * Improve TryCompareFromLHS documentation * Update wording to distinguish `knowns_` and `scoped_knowns_` * Better documentation for Key enum * Document the normalization of LT/GT * Removed unused PrimExpr temp * Call out modifications of the `compared_to_x` contents * Pointed to `Comparison::Comparison` for normalization details * Updated to clarify right/RHS. * Rename TryCompareFromLHS to DFSFromLHS

Previously, `RewriteSimplifier` inspected the top branch of each `And` and `Or` node, but didn't perform simplifications that would require inspection across branches nested structures. This commit introduces `SimplifyAsAndOfOrs`, which converts to an internal representation as conjunctive normal form, performs simplifications while in that form, then converts back to a `PrimExpr`. This utility is designed for data-propagation simplifications as part of apache#12261. To avoid increasing the CPU load unnecesarily, this utility is only used on an opt-in basis, either using `RewriteSimplifer::SetEnabledFeatures` when called directly, or using `PassContext` when used as part of `tir::transform::Simplify`.

Implements `tvm::arith::NarrowPredicateExpression`, a utility that removes free parameters from a boolean expression, such that the resulting expression being true implies that the original expression is true. For example, the predicate `(0 <= i+f) && (i+f < 16)`, where `f` is a free parameter on the range `0 <= f < 2)`, can be narrowed to the expression `(0 <= i+0) && (i+2 < 16)`. In effect, `NarrowPredicateExpression` functions as a context-sentive `tvm::tir::Substitute`, where the value substituted is selected such that the resulting expression errs on the side of being false. This is an internal utility used as part of the simplifications for layout transformations ([tracking issue link](apache#12261)).

Previously, `RewriteSimplifier` inspected the top branch of each `And` and `Or` node, but didn't perform simplifications that would require inspection across branches nested structures. This commit introduces `SimplifyAsAndOfOrs`, which converts to an internal representation as conjunctive normal form, performs simplifications while in that form, then converts back to a `PrimExpr`. This utility is designed for data-propagation simplifications as part of #12261. To avoid increasing the CPU load unnecesarily, this utility is only used on an opt-in basis, either using `RewriteSimplifer::SetEnabledFeatures` when called directly, or using `PassContext` when used as part of `tir::transform::Simplify`. * [Arith][UnitTest] Unittests displaying desired behavior * Corrected example in analyzer.h * Added underscore for private members * Rename Cleanup() to RemoveTrueFalse() * Added comments for conversion to internal representation * Updated comment in test case * Preferentially order clauses according to first occurrence * Added more unit tests

Implements `tvm::arith::NarrowPredicateExpression`, a utility that removes free parameters from a boolean expression, such that the resulting expression being true implies that the original expression is true. For example, the predicate `(0 <= i+f) && (i+f < 16)`, where `f` is a free parameter on the range `0 <= f < 2)`, can be narrowed to the expression `(0 <= i+0) && (i+2 < 16)`. In effect, `NarrowPredicateExpression` functions as a context-sentive `tvm::tir::Substitute`, where the value substituted is selected such that the resulting expression errs on the side of being false. This is an internal utility used as part of the simplifications for layout transformations ([tracking issue link](apache#12261)).

Implements `tvm::arith::NarrowPredicateExpression`, a utility that removes free parameters from a boolean expression, such that the resulting expression being true implies that the original expression is true. For example, the predicate `(0 <= i+f) && (i+f < 16)`, where `f` is a free parameter on the range `0 <= f < 2)`, can be narrowed to the expression `(0 <= i+0) && (i+2 < 16)`. In effect, `NarrowPredicateExpression` functions as a context-sentive `tvm::tir::Substitute`, where the value substituted is selected such that the resulting expression errs on the side of being false. This is an internal utility used as part of the simplifications for layout transformations ([tracking issue link](#12261)).

Implements `tvm::arith::NarrowPredicateExpression`, a utility that removes free parameters from a boolean expression, such that the resulting expression being true implies that the original expression is true. For example, the predicate `(0 <= i+f) && (i+f < 16)`, where `f` is a free parameter on the range `0 <= f < 2)`, can be narrowed to the expression `(0 <= i+0) && (i+2 < 16)`. In effect, `NarrowPredicateExpression` functions as a context-sentive `tvm::tir::Substitute`, where the value substituted is selected such that the resulting expression errs on the side of being false. This is an internal utility used as part of the simplifications for layout transformations ([tracking issue link](apache#12261)).

…che#12264) This PR adds `tvm.testing.CompareBeforeAfter`, a generalization of the `BaseBeforeAfter` utility previously used in `test_tir_transform_simplify.py`, which performs unit tests that perform a transformation on a TIR function and compare the results to an expected TIR output. This arose when minimizing the boilerplate required for unit tests in the implementation of apache#12261.

Implementation of API in `tvm.tir.schedule` for layout transformations with padding, as part of apache#12261, item "Insert pad value into generated TIR, using `tir::if_then_else`, `builtin::assume`, and `builtin::undef`". Following the RFC discussion in apache/tvm-rfcs#77 (comment) and apache/tvm-rfcs#77 (comment), this commit preferentially rewrites the loops that surround a padded transformation where possible, in order to express padding in terms of `tir::if_then_else`.

…he#12863) This commit adds a new sub-analyzer, `TransitiveComparisonAnalyzer`, which attempts to apply multiple known comparisons to prove an unknown. For example, `a <= b` and `b <= c` imply that `a <= c`. These simplifications are necessary for simplifying conditionals resulting from padded layout transformations (apache#12261). While some of these conditions may be proven using `ConstIntBoundAnalyzer` or `IntSetAnalyzer`, each has some limitations. `ConstIntBoundAnalyzer` can only compare against a constant, `IntSetAnalyzer` internally calls `RewriteSimplifier` which can result in infinite recursion, and neither can handle not-equal conditions because it would require tracking multiple intervals per expression. Therefore, introducing a new sub-analyzer for these simplifications. * Change mutable reference to mutable pointer * Remove nullptr default on Impl unique_ptr In g++ 7, defining a default constructor attempts to define the destructor, which fails because `Impl` is an incomplete type. As far as I should tell, the destructor should only be defined at the point where `~TransitiveComparisonAnalyzer` is defined, at which point `Impl` has a full definition. This issue does not occur in g++ 10. * Require opt-in for CPU-intensive simplifications * Document the intent of using bitflags * Rename "Feature" to "Extension" * Use TVM_DLL on new public member functions * Remove duplicate BaseBeforeAfter.transform definition * Explicitly enable extension for unit tests that require it * Fix accidentally duplicate test case * Improve TryCompareFromLHS documentation * Update wording to distinguish `knowns_` and `scoped_knowns_` * Better documentation for Key enum * Document the normalization of LT/GT * Removed unused PrimExpr temp * Call out modifications of the `compared_to_x` contents * Pointed to `Comparison::Comparison` for normalization details * Updated to clarify right/RHS. * Rename TryCompareFromLHS to DFSFromLHS

…#12972) Previously, `RewriteSimplifier` inspected the top branch of each `And` and `Or` node, but didn't perform simplifications that would require inspection across branches nested structures. This commit introduces `SimplifyAsAndOfOrs`, which converts to an internal representation as conjunctive normal form, performs simplifications while in that form, then converts back to a `PrimExpr`. This utility is designed for data-propagation simplifications as part of apache#12261. To avoid increasing the CPU load unnecesarily, this utility is only used on an opt-in basis, either using `RewriteSimplifer::SetEnabledFeatures` when called directly, or using `PassContext` when used as part of `tir::transform::Simplify`. * [Arith][UnitTest] Unittests displaying desired behavior * Corrected example in analyzer.h * Added underscore for private members * Rename Cleanup() to RemoveTrueFalse() * Added comments for conversion to internal representation * Updated comment in test case * Preferentially order clauses according to first occurrence * Added more unit tests

Implements `tvm::arith::NarrowPredicateExpression`, a utility that removes free parameters from a boolean expression, such that the resulting expression being true implies that the original expression is true. For example, the predicate `(0 <= i+f) && (i+f < 16)`, where `f` is a free parameter on the range `0 <= f < 2)`, can be narrowed to the expression `(0 <= i+0) && (i+2 < 16)`. In effect, `NarrowPredicateExpression` functions as a context-sentive `tvm::tir::Substitute`, where the value substituted is selected such that the resulting expression errs on the side of being false. This is an internal utility used as part of the simplifications for layout transformations ([tracking issue link](apache#12261)).

Lunderberg added the type:rfc-tracking RFC progress tracking. Ref: https://github.com/apache/tvm-rfcs label Aug 1, 2022

This was referenced Aug 1, 2022

[UnitTest][TIR] Testing utility for before/after transform tests #12264

Merged

[TIR] Add tir::builtin::undef #12266

Merged

[TIR] Add tir::builtin::assume #12267

Merged

Lunderberg mentioned this issue Sep 6, 2022

[TIR] Implement API for padded layout transformations #12720

Merged

Lunderberg mentioned this issue Sep 21, 2022

[TIR][Arith] Prove conditionals by transitively applying knowns #12863

Merged

Lunderberg mentioned this issue Sep 29, 2022

[Arith][Refactor] Extract And/Or/Not handling from RewriteSimplifier #12942

Closed

Lunderberg mentioned this issue Oct 3, 2022

[Arith] Optional rewriting and simplification into AND of ORs #12972

Merged

This was referenced Oct 10, 2022

[Arith][TIR] Check for constant offsets of known literal constraints #13023

Merged

[TIR][Arith] Use TryCompare to narrow inequalities if possible #13024

Merged

Lunderberg mentioned this issue Oct 11, 2022

[Arith] Add internal NarrowPredicateExpression utility #13041

Merged

areusch added the needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it label Oct 19, 2022

driazati removed the needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it label Nov 16, 2022

driazati added the tir any TIR core issues which don’t fit into other tir: labels (include/tvm/tir, src/tir) label Nov 16, 2022

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Tracking Issue] Padded Layout Transformations #12261

[Tracking Issue] Padded Layout Transformations #12261

Lunderberg commented Aug 1, 2022 •

edited

Loading

[Tracking Issue] Padded Layout Transformations #12261

[Tracking Issue] Padded Layout Transformations #12261

Comments

Lunderberg commented Aug 1, 2022 • edited Loading

This issue is to track progress for padded layout transformations

Lunderberg commented Aug 1, 2022 •

edited

Loading