[TIR] Add structural error printing for TensorIR by shingjan · Pull Request #9306 · apache/tvm

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan · 2021-10-18T07:09:35Z

This PR intends to improve the error rendering by annotating regions of interest in TIR like synr following PR #8121.
Previously when there is an error:

The IR is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
with tir.block([128, 128, 128, 128], "B") as [vi, vj, vk, vl]:
tir.bind(vi, i)
tir.bind(vj, j)
tir.bind(vk, k)
tir.bind(vl, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)
Regions of interest:
tir.Block#0
block B(iter_var(vi, range(min=0, ext=128)), iter_var(vj, range(min=0, ext=128)), iter_var(vk, range(min=0, ext=128)), iter_var(vl, range(min=0, ext=128))) {
reads([A[vi, vj, vk, vl]])
writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = (A[vi, vj, vk, vl]*2f)
}

With this PR, the error will be rendered like below:

The IR with diagnostic is:
@tvm.script.ir_module
class Module:
@tir.prim_func
def main(a: tir.handle, b: tir.handle) -> None:
A = tir.match_buffer(a, [128, 128, 128, 128], dtype="float32")
B = tir.match_buffer(b, [128, 128, 128, 128], dtype="float32")
# body
# with tir.block("root")
for i, j, k, l in tir.grid(128, 128, 128, 8):
tir.Block#0
with tir.block("B"):
^^^^^^^^^^^^^^^^^^^^
vi, vj, vk = tir.axis.remap("SSS", [i, j, k])
vl = tir.axis.spatial(128, l * 16)
tir.reads([A[vi, vj, vk, vl]])
tir.writes([B[vi, vj, vk, vl]])
B[vi, vj, vk, vl] = A[vi, vj, vk, vl] * tir.float32(2)

cc: @vinx13 @junrushao1994

vinx13

Thanks for the PR. I left some comments

junrushao · 2021-10-18T17:41:25Z

Hey thanks for the PR! It's a pretty nice POC, and I would like to deliberate on the format :-)

Shall we update the syntax in the description to reflect the latest change on mainline? Also please add a few unittests
If the error location is at a particular loop, let's avoid fusing it with other loops with tir.grid
Also we should consider printing the names of the error locations under the annotations, so that the error message could be clearly referring to these blocks/loops

junrushao · 2021-10-18T17:45:06Z

Hmm I just went through the code, ~~but figured that the snippet in the description isn't the actual output of the PR~~ (Sorry I was wrong)? Shall we update the description with a real-world example?

We have plenty of these errors in tests/python/unittest/test_tir_schedule_*.py, guarded by:

withpytest.raises(tvm.tir.ScheduleError, ...):`
...

vinx13

left some comments

Hzfengsy · 2021-10-19T00:45:08Z

Looks good to me. One thing I'm not sure is that will it compatible with the new block syntax.

tqchen · 2021-10-19T13:56:33Z

@shingjan you may need to update the test cases and code now that the new block syntax is introduced

shingjan · 2021-10-21T23:03:51Z

Printing nested loop is fixed as well as some comments addressed. Should be good for another good before we merge this in. @vinx13 @junrushao1994

vinx13

otherwise LGTM

vinx13

LGTM, @Hzfengsy would you like to take a second look?

Hzfengsy

LGTM

* add structural error printing * remove old code * address comments * address comments * add test * fix test case * fix nested loop * rm print * change simple loop cond * address comments * fix test * address comments * remove msg * add override * address comments * address comments

shingjan requested review from Hzfengsy, ZihengJiang, comaniac, jroesch, junrushao, kparzysz-quic, masahi, tqchen, vinx13, yzhliu and zhiics as code owners October 18, 2021 07:09

shingjan force-pushed the tir_structural_error branch from 621db6c to a656ba9Compare October 18, 2021 07:12

vinx13 reviewed Oct 18, 2021
View reviewed changes

Comment threadsrc/tir/schedule/error.cc Outdated
Comment threadsrc/tir/schedule/error.cc Outdated

vinx13 approved these changes Oct 18, 2021
View reviewed changes

Comment threadsrc/printer/tvmscript_printer.cc Outdated
Comment threadsrc/printer/tvmscript_printer.cc Outdated

vinx13 requested changes Oct 18, 2021
View reviewed changes

vinx13 requested changes Oct 19, 2021
View reviewed changes

Comment threadsrc/tir/schedule/error.cc Outdated
Comment threadsrc/printer/tvmscript_printer.cc Outdated

shingjan force-pushed the tir_structural_error branch from 45aa4e9 to f6da444Compare October 19, 2021 04:27

shingjan requested review from areusch and merrymercy as code owners October 19, 2021 04:27

shingjan requested a review from vinx13 October 19, 2021 04:31

vinx13 reviewed Oct 19, 2021
View reviewed changes

Comment threadsrc/printer/tvmscript_printer.cc Outdated

shingjan force-pushed the tir_structural_error branch 2 times, most recently from 5191293 to 6e0e9c4Compare October 21, 2021 23:00

shingjan requested a review from vinx13 October 21, 2021 23:02

vinx13 self-assigned this Oct 21, 2021

shingjan requested a review from vinx13 October 22, 2021 00:25

vinx13 reviewed Oct 22, 2021
View reviewed changes

Comment threadtests/python/unittest/test_tir_schedule_reorder.py Outdated

shingjan requested a review from vinx13 October 22, 2021 00:43

vinx13 reviewed Oct 22, 2021
View reviewed changes

Comment threadsrc/printer/tvmscript_printer.cc Outdated
Comment threadsrc/printer/tvmscript_printer.cc Outdated
Comment threadsrc/printer/tvmscript_printer.cc Outdated

shingjan requested a review from vinx13 October 22, 2021 01:02

vinx13 reviewed Oct 22, 2021
View reviewed changes

Comment threadtests/python/unittest/test_tvmscript_error_report.py Outdated
Comment threadsrc/tir/schedule/error.cc Outdated
Comment threadsrc/printer/tvmscript_printer.cc Outdated

vinx13 reviewed Oct 22, 2021
View reviewed changes

Comment threadsrc/printer/tvmscript_printer.cc Outdated

shingjan added 15 commits October 22, 2021 11:24

add structural error printing
d7fa520

remove old code
d0ce41d

address comments
1070e1f

address comments
2eb3f1f

add test
d115c3b

fix test case
6cc7a03

fix nested loop
47da01e

rm print
edda6a3

change simple loop cond
aa48129

address comments
6bf45cf

fix test
92a114e

address comments
3ff431a

remove msg
b3c232c

add override
05bfc31

address comments
6388858

shingjan force-pushed the tir_structural_error branch from ac89341 to 6388858Compare October 22, 2021 18:24

address comments
f690a33

vinx13 approved these changes Oct 22, 2021
View reviewed changes

Hzfengsy approved these changes Oct 23, 2021
View reviewed changes

Hzfengsy merged commit e9a66a1 into apache:mainOct 23, 2021

shingjan deleted the tir_structural_error branch October 26, 2021 00:37