You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "vinx13 (via GitHub)" <gi...@apache.org> on 2023/02/22 23:38:02 UTC
[GitHub] [tvm] vinx13 opened a new pull request, #14095: [TIR] Allow TransformLayout with non-inversible index map
vinx13 opened a new pull request, #14095:
URL: https://github.com/apache/tvm/pull/14095
TransformLayout requires the index map to have inverse map that can be calculated by the analyzer in order to check whether padding is added.
However, such check doesn't always work for all cases because of limitation of the affine analysis that can only handle a set of supported patterns. In some cases, even if the index map doesn't introduce padding, the schedule primitive throws `TransformationIntroducesPaddingError` because it fails to calculate the inverse index map.
It is safe to allow buffer being padded without providing pad_value because the original loop extent is not changed and the padded region is not accessed. This PR changes the behavior of `TransformLayout` to allow non-inversible index map.
cc @Lunderberg @Hzfengsy
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1453627426
Thank you, @vinx13 , for making the changes, and the updates look reasonable to me.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1441006032
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* cc @Hzfengsy, @junrushao, @quic-sanirudh, @shingjan <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 merged pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 merged PR #14095:
URL: https://github.com/apache/tvm/pull/14095
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1441044009
It looks like this would also allow index maps that reduce the number of elements in a buffer. For example, if I use the transform `lambda i, j: [i % 2, j]`, I'd expect to receive an error, because I've reduced the buffer size from `[N, M]` to `[2, M]`. I put together a quick unit test (below), and
```python
class TestNonInjectiveTransform(BasePaddingCompare):
"""Each location must map to a unique transformed location"""
index_map = tvm.testing.parameter(lambda i, j: [i % 2, j])
def before():
A = T.alloc_buffer((4, 4), "int32")
for i, j in T.grid(4, 4):
with T.block("block"):
vi, vj = T.axis.remap("SS", [i, j])
A[vi, vj] = 0
expected = tvm.tir.schedule.schedule.ScheduleError
```
No error is thrown, and I instead get the following output.
```python
@T.prim_func
def after():
A = T.alloc_buffer((2, 4), "int32")
for i, j in T.grid(4, 4):
with T.block("block"):
vi, vj = T.axis.remap("SS", [i, j])
T.reads()
T.writes(A[vi % 2, vj])
A[vi % 2, vj] = 0
```
This could be useful, as several other primitives could then be expressed in terms of a buffer transform, which may reduce complexity (e.g. This `lambda i,j: [i%2, j]` transformation could be used to implement double-buffering), but it could also have a large effect on any of the consumers of `A`.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1442193506
I think even if we don’t require the inverse exists, proving injective can’t be difficult (for example the mapping that contains xor). I agree with @quic-sanirudh that we can add an option, and document the implication that when using implicit padding it’s users responsibility to have the valid index map to ensure correctness.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Hzfengsy (via GitHub)" <gi...@apache.org>.
Hzfengsy commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1451255450
Hi @Lunderberg @quic-sanirudh, could you please take another look at this PR? Let's merge it in 24 hours if you do not have more objections :-)
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1445139232
Hmm. I think I'd still be worried on it, though mainly for naming reason. If I were encountering `implicit_padding=True` at the callsite as a first-time user, I would assume that the transformation is going to infer the shape of the transformed output, as opposed to disabling of a safety check. On the other hand, if I see something like `assume_injective_transform=True` or `unsafe_assume_injective_transform=True`, then it more clearly states that this is a precondition that must be guaranteed by the caller. Making this argument a keyword-only argument would also require the parameter name is present at the callsite, so the reader would see the parameter name and not just a lone `True`.
That said, I'd still lean toward having the precondition checked by the transformation if at all possible. Is there a specific expression that is injective, but is failing to be recognized as such? If so, could we instead expand the analysis utilities to recognize the expression?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] quic-sanirudh commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "quic-sanirudh (via GitHub)" <gi...@apache.org>.
quic-sanirudh commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1441856005
> This effectively allows any index map. When the index map creates overlapping mapped indices, the correctness is not guaranteed.
Hi,
Apologies if I'm mistaken, but if correctness is not guaranteed, that would break the assumption most people would have, which is that no matter what scheduling we do, the actual computation is guaranteed.
Could we perhaps add another argument to the `layout_transform` primitive to force implicit padding and set it to false by default so that by default implicit padding is disabled and we get an error on `None` for `pad_value`, but we can override that by setting an extra argument such as `sch.transform_layout(..., implicit_padding=True)`.
I don't know if there are other schedule primitives that break the correctness assumption, so please correct me if I'm wrong about that.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1442114841
I was thinking on it last night, and am leaning the same as @quic-sanirudh . Since all other scheduling changes validate their pre-conditions, hanving a scheduling change that requires the user (or the automatic schedule optimizer) to validate a pre-condition in order to produce correct code would be rather unexpected.
> I don't know if there are other schedule primitives that break the correctness assumption, so please correct me if I'm wrong about that.
As far as I know, the only current scheduling primitive that may break the correctness of an operator is `sch.tensorize(...)`, and even then that would only occur if the tensor intrinsic's definition doesn't match the implementation.
> Could we perhaps add another argument to the layout_transform primitive to force implicit padding and set it to false by default so that by default implicit padding is disabled and we get an error on None for pad_value
That's effectively what the current behavior is. If there is a defined `pad_value`, even something like `T.undef()` which will be removed during lowering and has no runtime equivalent, then implicit padding is allowed. If `pad_value` is None, then implicit padding is not allowed.
> This effectively allows any index map.
What if we were to weaken the check instead of removing it altogether? To maintain correctness, we don't need to find an inverse, only to prove that the transformation is injective.
* Updating `IndexMap::NonSurjectiveInverse` to return a `std::optional<std::pair<IndexMap, PrimExpr>>`, with `std::nullopt` returned when the `DetectIterMap` can't recognize the transformation.
* Adding a helper method `bool IndexMap::IsProvablyInjective(Array<Range> initial_ranges)`, which returns true if the function maps each input location maps to a unique output location. This could delegate out to either `DetectIterMap`, `arith::Analyzer`, or some combination of the two.
* In `tir/schedule/primitive/layout_transformation.cc`, start by validating the `IndexMap` and raising a schedule error if `index_map->IsProvablyInjective()` returns false.
That way, we would still maintain correctness, but could have a looser condition.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1447007522
The option `assume_injective_transform=True` sounds good. I agree that ideally we should expand analysis to cover our use cases, however as I mentioned above, the current analysis is limited and there are no good solution to extend it to more complicated use cases in the short term. I think it'd proceed with adding `assume_injective_transform` if we can agree on this
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] quic-sanirudh commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "quic-sanirudh (via GitHub)" <gi...@apache.org>.
quic-sanirudh commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1451294043
> Hi @Lunderberg @quic-sanirudh, could you please take another look at this PR? Let's merge it in 24 hours if you do not have more objections :-)
Looks good to me. I don't have any objections, thanks for considering my feedback.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on a diff in pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "Hzfengsy (via GitHub)" <gi...@apache.org>.
Hzfengsy commented on code in PR #14095:
URL: https://github.com/apache/tvm/pull/14095#discussion_r1121078571
##########
include/tvm/tir/schedule/schedule.h:
##########
@@ -642,10 +642,17 @@ class ScheduleNode : public runtime::Object {
* Algebraic symplifications, branch elimination, and other
* optimizations may assume that this precondition is met, and
* may result in incorrect results being returned.
+ *
+ * \param assume_injective_transform If set to true, the schedule primitive will assume the
Review Comment:
```suggestion
* \param assume_injective_transform If set to true, the schedule primitive will assume the
```
##########
python/tvm/tir/schedule/schedule.py:
##########
@@ -2645,6 +2644,8 @@ def transform_layout(
buffer: Union[Tuple[str, int], str, Buffer],
index_map: Union[IndexMap, Callable],
pad_value: Optional[Union[int, float, PrimExpr, IndexMap, Callable]] = None,
+ *,
+ assume_injective_transform=False,
Review Comment:
```suggestion
assume_injective_transform: bool = False,
```
##########
src/tir/schedule/primitive.h:
##########
@@ -501,10 +501,15 @@ TVM_DLL void Unannotate(ScheduleState self, const StmtSRef& sref, const String&
* \param buffer_index_type The type of the buffer index, kRead or kWrite.
* \param index_map The transformation to apply.
* \param pad_value The value to write into padding introduced by the transformation.
+ * \param assume_injective_transform If set to true, the schedule primitive will assume the
Review Comment:
```suggestion
* \param assume_injective_transform If set to true, the schedule primitive will assume the
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14095: [TIR] Allow TransformLayout with non-inversible index map
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14095:
URL: https://github.com/apache/tvm/pull/14095#issuecomment-1441078118
This effectively allows any index map. When the index map creates overlapping mapped indices, the correctness is not guaranteed.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org