code: def conv_approximation(size_a, size_w):
      """Conv approximation."""
      a_shape = (size_a, VRMPY_SIZE_B)
      w_shape = (size_w, VRMPY_SIZE_B)
      out_shape = (size_a, VRMPY_SIZE_INT32)
  
      @T.prim_func
      def operator(
          a_input: T.handle, b_input: T.handle, c_output: T.handle
      ) -> None:
          T.func_attr({"global_symbol": "main", "tir.noalias": True})
          a_buffer = T.match_buffer(a_input, a_shape, dtype="uint8")
          w_buffer = T.match_buffer(b_input, w_shape, dtype="uint8")
          c_buffer = T.match_buffer(c_output, out_shape, dtype="int32")
          for n, index_0 in T.grid(size_a, size_w):
              with T.block("c_buffer"):
                  vn_index, vi_index = T.axis.remap("SR", [n, index_0])
                  T.reads(
                      a_buffer[vn_index, 0:VRMPY_SIZE_B],
                      w_buffer[vi_index, 0:VRMPY_SIZE_B],
                      c_buffer[vn_index, 0:VRMPY_SIZE_INT32],
                  )
                  T.writes(c_buffer[vn_index, 0:VRMPY_SIZE_INT32])
                  with T.init():
                      for x in T.serial(VRMPY_SIZE_INT32):
                          c_buffer[vn_index, x] = 0
                  c_buffer[
                      vn_index, T.ramp(0, 1, 32)
                  ] = T.call_llvm_intrin(
                      T.llvm_lookup_intrinsic_id(
                          "llvm.hexagon.V6.vrmpyubv.acc.128B"
                      ),
                      T.uint32(3),
                      c_buffer[vn_index, T.ramp(0, 1, 32)],
                      T.reinterpret(
                          a_buffer[vn_index, T.ramp(0, 1, 128)],
                          dtype="int32x32",
                      ),
                      T.reinterpret(
                          w_buffer[vi_index, T.ramp(0, 1, 128)],
                          dtype="int32x32",
                      ),
                      dtype="int32x32",
                  )
          # Currently async DMA lowering does not add any wait to the end of schedules so
          # for timing purposes we are manually adding a wait to ensure that all copies
          # are complete when the schedule exits.
          T.evaluate(
              T.tvm_call_packed(
                  "device_api.hexagon.dma_wait",
                  0,  # QueueId
                  0,  # Wait for 0 in flight
                  dtype="int32",
              )
          )
  
      return tvm.tir.Schedule(operator)
  
file_path: tvm/tests/python/contrib/test_hexagon/test_async_dma_pipeline.py
imports: import numpy
  import pytest
  import tvm
  from tvm.script import tir
label_desc: Use image node_id 2 for calling a global function with example usage: conv_approximation(size_a, size_w) and returns: tvm