wrongtest opened a new pull request #10203:
URL: https://github.com/apache/tvm/pull/10203
Hi~ The PR fix a problem when use `cache_write` after loop transformations.
It may create out of bound accesses to external buffer without compiler
warnings, which could be dangerous and hard to detect at immediate. The example
to reproduce is as below:
```python
import tvm
from tvm.script import tir as T
from tvm import tir
@T.prim_func
def f(x: T.handle)->None:
X = T.match_buffer(x, [28], "int32")
for i in range(28):
with T.block("block"):
vi, = T.axis.remap("S", [i])
X[vi] = 1
s = tir.schedule.Schedule(f)
block = s.get_block("block")
i, = s.get_loops(block)
ii, io = s.split(i, factors=[None, 16]) # 28 % 16 != 0
s.cache_write(block, 0, "global")
print(s.mod["main"].script())
```
The result is:
```python
@T.prim_func
def func(X: T.Buffer[(28,), "int32"]) -> None:
X_global = T.alloc_buffer([28], dtype="int32")
for i_0, i_1 in T.grid(2, 16):
with T.block("block"):
vi = T.axis.spatial(28, i_0 * 16 + i_1)
T.where(i_0 * 16 + i_1 < 28)
X_global[vi] = 1
for ax0 in T.serial(28):
with T.block("X_global"):
v0 = T.axis.spatial(28, ax0)
X[v0] = X_global[v0]
```
Use `AnalyzeRegionUpperBound` which is equipped with affine analysis can
eliminate the problem in reported cases.
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
