adstraw opened a new pull request, #12785: URL: https://github.com/apache/tvm/pull/12785
Creates a 2-stage pipeline (DMA, compute) on Hexagon by lowering Async TIR primitives (see [RFC](https://github.com/apache/tvm-rfcs/blob/main/rfcs/0077-async-pipeline.md)) with memory copy semantics to builtin functions for DMA `Copy` and `Wait`. These builtins can then be lowered through the device API, with example lowering on Hexagon to the `HexagonUserDMA` class. The example test case contained in this PR creates a 2 stage pipeline which lowers to the following TIR which is built and run on Hexagon. The _prologue_ contains a DMA `Copy` for the first input. The _body_ contains a DMA `Copy` for the next input, a DMA `Wait` for the current input and then compute (increment) on the current input. The _epilogue_ computes the last input. ``` # from tvm.script import tir as T @tvm.script.ir_module class Module: @T.prim_func def main(A: T.Buffer[2048, "uint8"], B: T.Buffer[2048, "uint8"]) -> None: # function attr dict T.func_attr({"global_symbol": "main", "tir.noalias": True}) T.preflattened_buffer(A, [16, 128], dtype="uint8", data=A.data) T.preflattened_buffer(B, [16, 128], dtype="uint8", data=B.data) # prologue A_global = T.allocate([256], "uint8", "global") A_global_1 = T.buffer_decl([256], dtype="uint8", data=A_global) T.evaluate(T.dma_copy(0, T.address_of(A_global_1[0], dtype="handle"), T.address_of(A[0], dtype="handle"), 128, dtype="int32")) # body for i in T.serial(15): T.evaluate(T.dma_copy(0, T.address_of(A_global_1[(i + 1) % 2 * 128], dtype="handle"), T.address_of(A[i * 128 + 128], dtype="handle"), 128, dtype="int32")) T.evaluate(T.dma_wait(0, 1, dtype="int32")) for j in T.serial(128): B[i * 128 + j] = A_global_1[i % 2 * 128 + j] + T.uint8(1) # epilogue T.evaluate(T.dma_wait(0, 0, dtype="int32")) for j in T.serial(128): B[j + 1920] = A_global_1[j + 128] + T.uint8(1) ``` -- 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]
