[Apache TVM Discuss] [Questions] Problem in support Hexagon HVX instruction such as vrmpyz

[dingyongchao via Apache TVM Discuss <no...@discuss.tvm.ai>]

Hi, 
    Recently, I try to develop hexagon dsp inference from TVM, And I run simple demo from: https://developer.qualcomm.com/blog/tvm-open-source-compiler-now-includes-initial-support-qualcomm-hexagon-dsp. And I found that it do not use HVX instruction. So I modify '

> target = tvm.target.hexagon('v66', hvx=128)

', meanwhile, for HVX support, I modify compute '

> C = te.compute((N,N), lambda i, j: te.sum(A[i][k] * B[j][k], axis=k), name='C')

'. Then, The TVM IR generated like:

    Tensor(shape=[128, 128], op.name=C)
    schedule(0x16b6b80)
    input_mod:  primfn(A_1: handle, B_1: handle, C_1: handle) -> ()
      attr = {"global_symbol": "mmult", "tir.noalias": True}
      buffers = {C: Buffer(C_2: Pointer(int8), int8, [128, 128], []),
                 A: Buffer(A_2: Pointer(int8), int8, [128, 128], []),
                 B: Buffer(B_2: Pointer(int8), int8, [128, 128], [])}
      buffer_map = {A_1: A, B_1: B, C_1: C} {
      attr [IterVar(pipeline: int32, (nullptr), "ThreadIndex", "pipeline")] "pipeline_exec_scope" = 1;
      for (i.inner: int32, 0, 128) {
        for (j: int32, 0, 128) {
          C_2[((i.inner*128) + j)] = 0i8
          for (k: int32, 0, 128) {
            C_2[((i.inner*128) + j)] = ((int8*)C_2[((i.inner*128) + j)] + ((int8*)A_2[((i.inner*128) + k)]*(int8*)B_2[((j*128) + k)]))
          }
        }
      }
    }

I thought that the inner loop k can run in HVX vrmpy, However, when I see the asm generated by LLVM-IR, I found it still use vmpy for multi-sum. So I wonder how i can use insts in HVX, such as vrmpy, vrmpyz and vmem.





---
[Visit Topic](https://discuss.tvm.apache.org/t/problem-in-support-hexagon-hvx-instruction-such-as-vrmpyz/11381/1) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/256c4a28089c260bd857dfcb444098806ae6c65c2e39e191202e859bdbc9f36f).