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/09 19:46:06 UTC
[GitHub] [tvm] nverke commented on a diff in pull request #13301: [TIR] Update ReductionIterNotIndexOutputBuffer to check BlockRealizeN…
nverke commented on code in PR #13301:
URL: https://github.com/apache/tvm/pull/13301#discussion_r1018349185
##########
src/tir/schedule/analysis/reducer.cc:
##########
@@ -572,9 +572,25 @@ bool ReductionIterNotIndexOutputBuffer(const Block& block) {
if (!store) {
return true;
}
- ICHECK(buffer_written.count(store->buffer.get()))
- << "ValueError: The buffer \"" << store->buffer
- << "\" is written in the block but is not in the block's signature";
+ const auto* body_block = block->body.as<BlockRealizeNode>();
Review Comment:
Hmm I am not sure if I understand. Are you talking about a situation like this?
```
@T.prim_func
def nested_reduction_loop_with_match_buffers(
in0: T.Buffer[(4, 4, 4), "int8"],
in1: T.Buffer[(4, 4, 4), "int8"],
out: T.Buffer[(4, 4, 4), "int8"],
) -> None:
# body
# with T.block("root")
for y in T.serial(4):
with T.block("C"):
T.reads(in0[y, 0:4, 0:4], in1[y, 0:4, 0:4])
T.writes(out[y, 0:4, 0:4])
for x in T.serial(4):
with T.block("C"):
T.reads(in0[y, x, 0:4], in1[y, x, 0:4])
T.writes(out[y, x, 0:4])
A = T.match_buffer(in0[y, x, 0:4], [4], dtype="int8", offset_factor=1)
B = T.match_buffer(in1[y, x, 0:4], [4], dtype="int8", offset_factor=1)
C = T.match_buffer(out[y, x, 0:4], [4], dtype="int8", offset_factor=1)
A_i8x4: T.int8x4 = A[0:4]
A_i32: T.int32 = T.reinterpret(A_i8x4, dtype="int32")
B_i8x4: T.int8x4 = B[0:4]
B_i32: T.int32 = T.reinterpret(B_i8x4, dtype="int32")
C[0:4] = T.reinterpret(A_i32 + B_i32, dtype="int8x4")
```
My understanding is that this check is on the level of the store statement and already has collected all of the write regions for the parent loops so just adding the regions that the match buffers check should be enough.
Alternatively are you referring to something like this?
```
@T.prim_func
def nested_reduction_loop_with_match_buffers(
in0: T.Buffer[(4, 4, 4), "int8"],
in1: T.Buffer[(4, 4, 4), "int8"],
out: T.Buffer[(4, 4, 4), "int8"],
) -> None:
# body
# with T.block("root")
for y in T.serial(4):
with T.block("C"):
T.reads(in0[y, 0:4, 0:4], in1[y, 0:4, 0:4])
T.writes(out[y, 0:4, 0:4])
A = T.match_buffer(in0[y, 0:4, 0:4], [4, 4], dtype="int8", offset_factor=1)
B = T.match_buffer(in1[y, 0:4, 0:4], [4, 4], dtype="int8", offset_factor=1)
C = T.match_buffer(out[y, 0:4, 0:4], [4, 4], dtype="int8", offset_factor=1)
for x in T.serial(4):
with T.block("C"):
T.reads(in0[y, x, 0:4], in1[y, x, 0:4])
T.writes(out[y, x, 0:4])
A_i8x4: T.int8x4 = A[x, 0:4]
A_i32: T.int32 = T.reinterpret(A_i8x4, dtype="int32")
B_i8x4: T.int8x4 = B[x, 0:4]
B_i32: T.int32 = T.reinterpret(B_i8x4, dtype="int32")
C[x, 0:4] = T.reinterpret(A_i32 + B_i32, dtype="int8x4")
```
Here I believe we are still able to pickup the match buffers from the body block.
--
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