[GitHub] [tvm] wrongtest opened a new pull request #10203: [TIR] Fix an index out of bound problem of cache write block

[GitBox <gi...@apache.org>]

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: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org

[GitHub] [tvm] junrushao1994 commented on pull request #10203: [TIR] Fix an index out of bound problem of cache write block

[GitBox <gi...@apache.org>]

junrushao1994 commented on pull request #10203:
URL: https://github.com/apache/tvm/pull/10203#issuecomment-1035911965


   Thanks @wrongtest @Hzfengsy!


-- 
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] junrushao1994 merged pull request #10203: [TIR] Fix an index out of bound problem of cache write block

[GitBox <gi...@apache.org>]

junrushao1994 merged pull request #10203:
URL: https://github.com/apache/tvm/pull/10203


   


-- 
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