code: def dot_vrmpy(x_ty, y_ty):
      """Generates vrmpy instruciton for tensorization."""
      int32_lanes = 32
      num_int8_elements = 4  # 4 int8 elements in int32
      data = te.placeholder(
          (num_int8_elements,), dtype=x_ty, name="data"
      )
      kernel = te.placeholder(
          (int32_lanes, num_int8_elements), dtype=y_ty, name="kernel"
      )
      k = te.reduce_axis((0, num_int8_elements), name="k")
      C = te.compute(
          (int32_lanes,),
          lambda i: te.sum(
              data[k].astype("int32") * kernel[i, k].astype("int32"),
              axis=k,
          ),
          name="C",
      )
  
      a_buffer = tvm.tir.decl_buffer(
          data.shape,
          dtype=x_ty,
          name="a_buffer",
          offset_factor=1,
          strides=[1],
      )
      b_buffer = tvm.tir.decl_buffer(
          kernel.shape,
          dtype=y_ty,
          name="b_buffer",
          offset_factor=1,
          strides=[te.var("ldw"), 1],
      )
  
      def _intrin_func(ins, outs):
          def _instr(index):
              ib = tvm.tir.ir_builder.create()
              if index == 1:
                  ib.emit(
                      outs[0].vstore(0, tvm.tir.const(0, "int32x32"))
                  )
                  return ib.get()
  
              vec_zero = tvm.tir.const(0, "int32x32")
  
              if x_ty == "uint8" and y_ty == "uint8":
                  a_uint8 = ins[0].vload([0], "uint8x4")
                  re_int32 = tvm.tir.call_intrin(
                      "int32", "tir.reinterpret", a_uint8
                  )
                  vec_b = ins[1].vload([0, 0], "uint8x128")
  
                  vrmpy_inst_name = "llvm.hexagon.V6.vrmpyub.acc.128B"
  
                  vec_bi32 = tvm.tir.call_intrin(
                      "int32x32", "tir.reinterpret", vec_b
                  )
  
                  quad_reduction = tvm.tir.call_llvm_pure_intrin(
                      "int32x32",
                      vrmpy_inst_name,
                      tvm.tir.const(3, "uint32"),
                      vec_zero,
                      vec_bi32,
                      re_int32,
                  )
              elif x_ty == "uint8" and y_ty == "int8":
                  a_uint8 = ins[0].vload([0], "uint8x4")
                  re_int32 = tvm.tir.call_intrin(
                      "int32", "tir.reinterpret", a_uint8
                  )
                  vec_b = ins[1].vload([0, 0], "int8x128")
  
                  vrmpy_inst_name = "llvm.hexagon.V6.vrmpybusv.acc.128B"
  
                  vec_bi32 = tvm.tir.call_intrin(
                      "int32x32", "tir.reinterpret", vec_b
                  )
  
                  quad_reduction = tvm.tir.call_llvm_pure_intrin(
                      "int32x32",
                      vrmpy_inst_name,
                      tvm.tir.const(3, "uint32"),
                      vec_zero,
                      re_int32.astype("int32x32"),
                      vec_bi32,
                  )
              else:
                  raise ValueError(
                      "Only (u8, u8) or (u8, i8) dtype pairs are supported by vrmpy."
                  )
  
              if index == 0:
                  ib.emit(outs[0].vstore(0, quad_reduction))
              else:
                  ib.emit(
                      outs[0].vstore(
                          0,
                          quad_reduction
                          + outs[0].vload([0], "int32x32"),
                      )
                  )
              return ib.get()
  
          # body, reset, update
          return _instr(0), _instr(1), _instr(2)
  
      buffer_params = {"offset_factor": 1}
      return te.decl_tensor_intrin(
          C.op,
          _intrin_func,
          binds={data: a_buffer, kernel: b_buffer},
          default_buffer_params=buffer_params,
      )
  
file_path: tvm/python/tvm/topi/hexagon/tensor_intrin.py
imports: import tvm
  from tvm.ir import register_intrin_lowering
  from tvm import te
label_desc: Use image node_id 7 for calling a global function with example usage: dot_vrmpy(x_ty, y_ty) and returns: te, _instr,  _instr,  _instr, ib, ib