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/02/20 03:32:17 UTC
[GitHub] [tvm] MasterJH5574 opened a new pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
MasterJH5574 opened a new pull request #10327:
URL: https://github.com/apache/tvm/pull/10327
This PR introduces the support of integer BufferLoadNodes in `arith::IntervalSetEvaluator` as well as a unittest.
Prior to this PR, the evaluator always return `IntervalSet::Everything()` (that is `[-∞, +∞]`) once it encounters a BufferLoad. However, as long as 1) the buffer data _is of type integer_ and 2) the indices of the BufferLoad _contain no variable_ that needs to be relaxed during evaluation, we can keep the BufferLoad as it is.
Besides, we're considering introducing `assume_in_range` for buffers for the use of range/region analysis. With `assume_in_range`, even when the indices of BufferLoad contain some variables to be relaxed, we can just return the asserted range, which is more compact, instead of the universe.
cc @yzh119 @Hzfengsy @spectrometerHBH @junrushao1994
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810584530
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
Good catch! I never thought of this before. Is it okay by specifying `assume_in_range(A_indptr[i + 1] - A_indptr[i], [0, 127])` in the future? (not pretty sure)
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810566267
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
Due to some parser issue we're now allowed to simplify the expression for now:
```
error: Unsupported index type, expected int or tvm.tir.PrimExpr, but got <class 'tvm.script.tir.node.BufferSlice'>
--> tensorir.py:431:44
|
431 | T.reads(A_indptr[vi : vi + 1], A_data[A_indptr[vi]: A_indptr[vi + 1]])
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
note: run with `TVM_BACKTRACE=1` environment variable to display a backtrace.
```
I suppose this is some bug that can be easily fixed. However, I didn't get enough time to take a deeper look into the parser.
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810566267
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
Due to some parser issue we're now allowed to simplify the expression for now:
```
error: Unsupported index type, expected int or tvm.tir.PrimExpr, but got <class 'tvm.script.tir.node.BufferSlice'>
--> tensorir.py:431:44
|
431 | T.reads(A_indptr[vi : vi + 1], A_data[A_indptr[vi]: A_indptr[vi + 1]])
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
note: run with `TVM_BACKTRACE=1` environment variable to display a backtrace.
```
I suppose this is some bug that can be easily fixed, though I didn't get enough time to take a deeper look into the parser.
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810584318
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
Review comment:
I suppose the only way is `T.reads(A_indptr[i : i+2])` - writing read-regions like this is safe. But I have no idea how to represent `T.writes(data[i], data[i + 2])` well... Asserting an overlarge write-region is not equally safe.
--
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 change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810565801
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
```suggestion
A_data[A_indptr[i] : A_indptr[i + 1]],
```
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
Review comment:
IIRC, we don't allow annotating one buffer twice in `read` and `write` regions. Here it's better to be `A_indptr[i: i + 1]`
--
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] yzh119 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
yzh119 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810586750
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
Yes I think so, but we merge this one first.
--
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] yzh119 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
yzh119 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810586774
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
Review comment:
Okay that make sense.
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810565953
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
Review comment:
Ohhh thanks for the reminder! Will fix it soon.
--
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] yzh119 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
yzh119 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810583654
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
Review comment:
@Hzfengsy So there is no way we can do `T.reads(A_indptr[i], A_indptr[i + 2])`?
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
One problem with this analysis is that we cannot assume `A_indptr[i+1] >= A_indptr[i]`.
--
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 merged pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
Hzfengsy merged pull request #10327:
URL: https://github.com/apache/tvm/pull/10327
--
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] MasterJH5574 commented on a change in pull request #10327: [Arith] Support integer BufferLoad in IntervalSetEvaluator
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on a change in pull request #10327:
URL: https://github.com/apache/tvm/pull/10327#discussion_r810566267
##########
File path: tests/python/unittest/test_tir_transform_compact_buffer_region.py
##########
@@ -505,6 +505,74 @@ def compacted_opaque_access_annotated_func(a: T.handle) -> None:
T.store(C.data, i, T.load("float32", B.data, i))
+@T.prim_func
+def sparse_read_cache(
+ A_data: T.Buffer[(819,), "float32"],
+ B: T.Buffer[(128,), "float32"],
+ A_indptr: T.Buffer[(129,), "int32"],
+ A_indices: T.Buffer[(819,), "int32"],
+) -> None:
+ for i in T.serial(128):
+ with T.block("rowsum_outer"):
+ T.reads(
+ A_indptr[i],
+ A_indptr[i + 1],
+ A_data[A_indptr[i] + 0 : A_indptr[i] + (A_indptr[i + 1] - A_indptr[i])],
Review comment:
Due to some parser issue we're not allowed to simplify the expression for now:
```
error: Unsupported index type, expected int or tvm.tir.PrimExpr, but got <class 'tvm.script.tir.node.BufferSlice'>
--> tensorir.py:431:44
|
431 | T.reads(A_indptr[vi : vi + 1], A_data[A_indptr[vi]: A_indptr[vi + 1]])
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
note: run with `TVM_BACKTRACE=1` environment variable to display a backtrace.
```
I suppose this is some bug that can be easily fixed, though I didn't get enough time to take a deeper look into the parser.
--
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