NicolaLancellotti commented on a change in pull request #10344:
URL: https://github.com/apache/tvm/pull/10344#discussion_r829413144
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -398,11 +398,11 @@ def assign_addresses(buffer_info, npu_ops,
scratch_region_map):
def replace_npu_fm_with_address(npu_fm):
assert isinstance(npu_fm.tiles.addresses[0], tvm.tir.Load)
- # We currently does not support tiles
- # Change this when tiles are needed
- # (i.e. when using rolling buffers)
- assert npu_fm.tiles.addresses[1:] == [0, 0, 0]
- npu_fm.tiles.addresses[1:] = [0, 0, 0]
+ for i in range(1, 4):
+ address = npu_fm.tiles.addresses[i]
+ if isinstance(address, tvm.tir.expr.Load):
+ address = address.index
+ npu_fm.tiles.addresses[i] = int(address)
Review comment:
You are right, that logic (the integer conversion) worked before the
USMP was enabled, I forgot to remove it during a rebase.
Now we have the following logic:
```
npu_fm.tiles.addresses[0] = address + int(index)
npu_fm.tiles.addresses[1] = address if isinstance(npu_fm.tiles.addresses[1],
tvm.tir.BufferLoad) else 0
npu_fm.tiles.addresses[2] = address if isinstance(npu_fm.tiles.addresses[2],
tvm.tir.BufferLoad) else 0
npu_fm.tiles.addresses[3] = 0
```
Where the index of tile1 and tile2, when used, are always 0.
--
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]
