class_docstr: Fixed point Multiply test class
class_docstr_tok: ["Fixed","point","Multiply","test","class"]
docstr: Test fixed point multiply with vectorization.
  Vectorization size is more than hw vector length
docstr_tok: ["Test","fixed","point","multiply","with","vectorization",".","Vectorization","size","is","more","than","hw","vector","length"]
code: def test_per_tensor_with_lanes(
      self, hexagon_session: Session, vector_size
  ):
      """Test fixed point multiply with vectorization.
      Vectorization size is more than hw vector length"""
      ishape = [2, 256, 16]
  
      def q_mul_shift(shape):
          x = te.placeholder(shape, name="X", dtype="int32")
          out = te.compute(
              shape,
              lambda i, j, k: tvm.tir.q_multiply_shift(
                  x[i, j, k],
                  tvm.tir.const(1395864320, "int32"),
                  tvm.tir.const(31, "int32"),
                  tvm.tir.const(1, "int32"),
              ),
              name="compute",
          )
          return te.create_prim_func([x, out])
  
      mod = q_mul_shift(ishape)
  
      # Schedule with vectorization
      sch = tvm.tir.Schedule(mod)
      b00 = sch.get_block(name="compute", func_name="main")
      fused = sch.fuse(*sch.get_loops(block=b00))
      _, v = sch.split(loop=fused, factors=[None, vector_size])
      sch.vectorize(v)
  
      with tvm.transform.PassContext(opt_level=3):
          hex_lib = tvm.build(
              sch.mod["main"], target=get_hexagon_target("v68")
          )
          host_lib = tvm.build(mod, target=tvm.target.Target("llvm"))
  
      asm = hex_lib.get_source("asm")
  
      # Check that 'vmpye' instruction was generated in asm file.
      vmpye_regex = re.compile(
          r"v\d{1,2}.w = vmpye\(v\d{1,2}.w,v\d{1,2}.uh\)"
      )
      assert vmpye_regex.search(asm) is not None
  
      # Check that 'vmpyo' instruction was generated in asm file.
      vmpyo_regex = re.compile(
          r"v\d{1,2}.w \+= vmpyo\(v\d{1,2}.w,v\d{1,2}.h\):<<1:rnd:sat:shift"
      )
      assert vmpyo_regex.search(asm) is not None
  
      # Verify accuracy
      a_np = (
          np.random.randint(-1000, 1000, size=np.prod(ishape))
          .reshape(ishape)
          .astype("int32")
      )
      b_np = (
          np.random.randint(-1000, 1000, size=np.prod(ishape))
          .reshape(ishape)
          .astype("int32")
      )
      hex_args = [
          tvm.runtime.ndarray.array(
              arg, device=hexagon_session.device, mem_scope="global"
          )
          for arg in [a_np, b_np]
      ]
      host_args = [
          tvm.runtime.ndarray.array(arg) for arg in [a_np, b_np]
      ]
  
      hex_rt = hexagon_session.load_module(hex_lib)
      hex_rt(*hex_args)
      host_lib(*host_args)
  
      assert np.allclose(hex_args[1].numpy(), host_args[1].numpy())
  
file_path: tvm/tests/python/contrib/test_hexagon/test_fixed_point_multiply.py
imports: import re
  import numpy
  import tvm.testing
  from tvm import relay
  from tvm import te
  from tvm.relay.backend import Executor
  from tvm.contrib.hexagon.session import Session
  from tvm.contrib.hexagon.pytest_plugin import HEXAGON_AOT_LLVM_TARGET
  from .infrastructure import get_hexagon_target
label_desc: Use image node_id 3 for calling the TestFixedPointMultiply obj's underlying member method code with example usage: obj.test_per_tensor_with_lanes(hexagon_session, vector_size) and returns: te