manupa-arm commented on a change in pull request #10344:
URL: https://github.com/apache/tvm/pull/10344#discussion_r829857975
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/dma.py
##########
@@ -197,18 +197,72 @@ def get_read_params(stmt):
base_address = [get_base_address(index) for index in inner.value.indices]
data_type = inner.buffer.data.type_annotation.element_type.dtype
+
+ def check_rolling_buffer():
Review comment:
Lets make this a function of dma.py and make it accept stmt which is a
read loop nest.
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/dma.py
##########
@@ -197,18 +197,72 @@ def get_read_params(stmt):
base_address = [get_base_address(index) for index in inner.value.indices]
data_type = inner.buffer.data.type_annotation.element_type.dtype
+
+ def check_rolling_buffer():
+ rolling_buffer = True
+ floor_mod = None
+
+ def _get_rolling_var(stmt):
+ nonlocal rolling_buffer, floor_mod
+
+ if isinstance(stmt, tvm.tir.FloorMod):
+ if floor_mod is not None:
+ rolling_buffer = False
+ elif (
+ isinstance(stmt.b, tvm.tir.expr.IntImm)
+ and isinstance(stmt.a, tvm.tir.expr.Add)
+ and isinstance(stmt.a.a, tvm.tir.expr.Var)
+ and isinstance(stmt.a.b, tvm.tir.expr.IntImm)
+ ):
+ floor_mod = stmt
+ else:
+ rolling_buffer = False
+ elif isinstance(stmt, tvm.tir.FloorDiv):
+ rolling_buffer = False
+
+ tvm.tir.stmt_functor.post_order_visit(inner.value, _get_rolling_var)
+
+ if rolling_buffer and floor_mod is not None:
+ rolling_var = floor_mod.a.a
+ tile_length = floor_mod.b - floor_mod.a.b
+ return rolling_var, tile_length
Review comment:
I think it is clear if we return a NamedTuple that contains
tile_height_0, tile_height_1, tile_width_0, tile_address_0, tile_address_1,
tile_address_2.
If you agree, then we might need to change the function to be something like
CreateTiles(...)
Then, I think we should tests for CreateTiles function with a set of test
cases that include different Stmt types.
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/dma.py
##########
@@ -197,18 +197,72 @@ def get_read_params(stmt):
base_address = [get_base_address(index) for index in inner.value.indices]
data_type = inner.buffer.data.type_annotation.element_type.dtype
+
+ def check_rolling_buffer():
+ rolling_buffer = True
+ floor_mod = None
+
+ def _get_rolling_var(stmt):
+ nonlocal rolling_buffer, floor_mod
+
+ if isinstance(stmt, tvm.tir.FloorMod):
+ if floor_mod is not None:
+ rolling_buffer = False
+ elif (
+ isinstance(stmt.b, tvm.tir.expr.IntImm)
+ and isinstance(stmt.a, tvm.tir.expr.Add)
+ and isinstance(stmt.a.a, tvm.tir.expr.Var)
+ and isinstance(stmt.a.b, tvm.tir.expr.IntImm)
Review comment:
Out of curiosity, would it be possible to use this ?
https://github.com/apache/tvm/blob/b01e3fc4d21bba898a5ea17d526013c52ea720eb/python/tvm/ir/base.py#L160-L209
--
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]
