NicolaLancellotti commented on a change in pull request #10344:
URL: https://github.com/apache/tvm/pull/10344#discussion_r816961600
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/dma.py
##########
@@ -287,31 +321,69 @@ def get_ifm_params(pointer, producers):
The serializable padding.
"""
- pad = producers[pointer]
+ pad = producers_consumers.get_producer(pointer, stmt)
serial_padding, input_pointer, _ = get_pad_params(pad)
- upscale = producers[input_pointer]
+ upscale = producers_consumers.get_producer(input_pointer, pad)
input_pointer, _ = get_upscale_params(upscale)
- convert_to_nhwc = producers[input_pointer]
+ convert_to_nhwc = producers_consumers.get_producer(input_pointer, upscale)
in_channels, input_pointer, _ = get_convert_to_nhwc_params(convert_to_nhwc)
- read = producers[input_pointer]
+ read = producers_consumers.get_producer(input_pointer, convert_to_nhwc)
serial_ifm, _, _ = get_read_params(read)
serial_ifm.channels = in_channels
+
+ floor_mod_stmt = None
+ for_stmt = None
+
+ def _get_buffer_var(stmt):
+ nonlocal for_stmt
+ nonlocal floor_mod_stmt
+ if isinstance(stmt, tvm.tir.For):
+ for_stmt = stmt
+ if isinstance(stmt, tvm.tir.FloorMod):
+ floor_mod_stmt = stmt
+
+ tvm.tir.stmt_functor.post_order_visit(stmt, _get_buffer_var)
+
+ if floor_mod_stmt is not None:
+ layout = get_op_attrs(read)[0]["layout"]
+ channels = serial_ifm.channels
+ if for_stmt.body.loop_var == floor_mod_stmt.a.a.a:
Review comment:
With rolling buffers the `floor_mod_stmt` should be always like that:
`floormod(((HEIGHT_OR_WIDTH_INDEX + KERNEL_INDEX) + BUFFER_OFFSET),
BUFFER_SIZE)`
I cannot think of any other case.
--
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]
