You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/11/22 08:55:17 UTC
[GitHub] [tvm] masahi opened a new pull request, #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
masahi opened a new pull request, #13463:
URL: https://github.com/apache/tvm/pull/13463
All indices in `IndexMap` are int32. The PR https://github.com/apache/tvm/pull/13327 made the dtypes of buffer shapes int64, which conflicts with the dtype s in `IndexMap`.
--
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] masahi commented on pull request #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13463:
URL: https://github.com/apache/tvm/pull/13463#issuecomment-1324239516
@Lunderberg Thanks for finding the simple test case! I hit this error when running this Hexagon test https://github.com/apache/tvm/blob/b29ab5c6baea8884a5b9c0702c1c93da6a830866/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py#L397. The original test case in this PR, with its `cache_read`, index map etc, was directly taken from this Hexagon test.
The bug is triggered when we hit the code path https://github.com/apache/tvm/blob/b6fae9b35eff4ad1f7cc2e83d8d7da5d701d8e44/src/tir/schedule/primitive/layout_transformation.cc#L282
I was not sure what makes this code path hit and why existing tests didn't hit it.
--
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] masahi commented on pull request #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
masahi commented on PR #13463:
URL: https://github.com/apache/tvm/pull/13463#issuecomment-1328727642
@Lunderberg @vinx13 PTAL, thanks.
--
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 #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on PR #13463:
URL: https://github.com/apache/tvm/pull/13463#issuecomment-1324213823
Regarding the readability of the unit test, I did some poking around, and the following PrimFunc triggers the same error on main when the layout is transformed using `lambda h: [h//8, h%8]`.
```python
@T.prim_func
def func(A: T.Buffer[T.int64(58), "int32"]):
for i in T.serial(T.int64(58)):
with T.block("block"):
vi = T.axis.remap("S", [i])
A[vi] = 0
```
--
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 #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
tvm-bot commented on PR #13463:
URL: https://github.com/apache/tvm/pull/13463#issuecomment-1323318975
<!---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 <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] Lunderberg commented on a diff in pull request #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on code in PR #13463:
URL: https://github.com/apache/tvm/pull/13463#discussion_r1029680983
##########
tests/python/unittest/test_tir_schedule_transform_layout.py:
##########
@@ -173,6 +173,35 @@ def two_elementwise_unit_dim(A: T.Buffer[(1, 128), "float32"], C: T.Buffer[(1, 1
vi, vj = T.axis.remap("SS", [i, j])
C[vi, vj] = B[vi, vj] + 1.0
+
+
+@tvm.script.ir_module
Review Comment:
Is this entire definition required for the test case? As a reader, it's hard to tell which parts of this PrimFunc are needed to trigger the bug.
##########
src/tir/schedule/primitive/layout_transformation.cc:
##########
@@ -1055,13 +1055,43 @@ class TransformationIntroducesPaddingError : public ScheduleError {
PrimExpr padding_predicate_;
};
+// Make the dtypes of indices in IndexMap be the same as the dtype of the buffer shape, to avoid
+// dtype-mismatch issues later.
+IndexMap LegalizeIndexMapDType(const IndexMap& index_map, const Buffer& buf) {
+ auto initial_indices_orig = index_map->initial_indices;
+ ICHECK(buf->shape.size() == initial_indices_orig.size());
+
+ Array<Var> initial_indices;
+ Map<Var, PrimExpr> var_map;
+
+ for (size_t i = 0; i < buf->shape.size(); ++i) {
+ if (buf->shape[i]->dtype != initial_indices_orig[i].dtype()) {
Review Comment:
I think this would have an error if only some of the index dtypes have a mismatch. In that case, `initial_indices` would only be filled with variables that have a mismatched dtype, when it should have the same size as `initial_indices_orig`.
```c++
if (buf->shape[i]->dtype == initial_indices_orig[i].dtype()) {
initial_indices.push_back(initial_indices_orig[i]);
} else {
auto new_idx = ...
}
```
##########
tests/python/unittest/test_tir_schedule_transform_layout.py:
##########
@@ -925,5 +954,24 @@ def expected(a: T.handle):
A[i, j] = T.if_then_else(i == 3 and 2 <= j, 0, 42, dtype="int32")
+def test_index_map_dtype_legalize():
+ """Test dtype legalization of the index map indices."""
+
+ def index_map_nchw32c_nchw8h8w32c(n_batch, channel, height, width, channel_32):
+ return [n_batch, channel, height // 8, width // 8, height % 8, width % 8, channel_32]
+
+ sch = tir.Schedule(Conv2dNCHW32c)
+
+ conv2d_block = sch.get_block("conv2d_NCHWc_int8")
+ sch.cache_read(conv2d_block, 0, "global.vtcm")
+
+ # The following error is raised from the IterVar constructor without the dtype legalization.
+ # TVMError: Check failed: dom->extent.dtype() == var.dtype() (int64 vs. int32) :
+ # The dtype of the extent of an IterVar (int64) must match its associated Var's dtype (int32)
+ sch.transform_layout(
+ conv2d_block, ("read", 0), index_map=index_map_nchw32c_nchw8h8w32c, pad_value=0
Review Comment:
Instead of `("read", 0)`, the buffer can be specified by name, `buffer = "data_pad_global_vtcm"`.
##########
src/tir/schedule/primitive/layout_transformation.cc:
##########
@@ -1055,13 +1055,43 @@ class TransformationIntroducesPaddingError : public ScheduleError {
PrimExpr padding_predicate_;
};
+// Make the dtypes of indices in IndexMap be the same as the dtype of the buffer shape, to avoid
+// dtype-mismatch issues later.
+IndexMap LegalizeIndexMapDType(const IndexMap& index_map, const Buffer& buf) {
+ auto initial_indices_orig = index_map->initial_indices;
Review Comment:
Nit: `const auto&` instead of `auto`.
##########
tests/python/unittest/test_tir_schedule_transform_layout.py:
##########
@@ -173,6 +173,35 @@ def two_elementwise_unit_dim(A: T.Buffer[(1, 128), "float32"], C: T.Buffer[(1, 1
vi, vj = T.axis.remap("SS", [i, j])
C[vi, vj] = B[vi, vj] + 1.0
+
+
+@tvm.script.ir_module
+class Conv2dNCHW32c:
+ @T.prim_func
Review Comment:
This location for `@T.prim_func` would run while the test is being collected. This can make it difficult to trouble-shoot, since a failure when parsing/constructing would prevent any unit test from running, not just the unit tests that make use of it. It would be better to have the `@T.prim_func` be in the unit test itself, so a failure to construct the primfunc would only cause a failure in a single test.
##########
tests/python/unittest/test_tir_schedule_transform_layout.py:
##########
@@ -925,5 +954,24 @@ def expected(a: T.handle):
A[i, j] = T.if_then_else(i == 3 and 2 <= j, 0, 42, dtype="int32")
+def test_index_map_dtype_legalize():
+ """Test dtype legalization of the index map indices."""
+
+ def index_map_nchw32c_nchw8h8w32c(n_batch, channel, height, width, channel_32):
+ return [n_batch, channel, height // 8, width // 8, height % 8, width % 8, channel_32]
+
+ sch = tir.Schedule(Conv2dNCHW32c)
+
+ conv2d_block = sch.get_block("conv2d_NCHWc_int8")
+ sch.cache_read(conv2d_block, 0, "global.vtcm")
Review Comment:
Why does the test case call `cache_read`, instead of starting with the input provided to `transform_layout`?
--
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 #13463: [TIR] Fix buffer shape and IndexMap indices dtype mismatch
Posted by GitBox <gi...@apache.org>.
vinx13 merged PR #13463:
URL: https://github.com/apache/tvm/pull/13463
--
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