code: def _qnn_dense_legalize_hexagon(attrs, inputs, types):
      """Legalize qnn.dense op for vrmpy tensorization.
  
      N dimension of weights should be aligned on vector length. If not, then N dimension is padded to
      be a multiple of 32.
      """
      assert len(types) == 7
      assert len(inputs) == 6
  
      data_tensor, kernel_tensor = types[0], types[1]
      if (
          "int8" not in data_tensor.dtype
          or "int8" not in kernel_tensor.dtype
      ):
          return None
  
      N, _ = kernel_tensor.shape
  
      if N % OUT_CHANNEL_VECTOR_LENGTH != 0:
          N_padded = helper_align_up(N, OUT_CHANNEL_VECTOR_LENGTH)
          diff = N_padded - N
  
          (
              data,
              kernel,
              data_zp,
              kernel_zp,
              data_scale,
              kernel_scale,
          ) = inputs
  
          # Pad weights by 'diff'
          padded_kernel = relay.nn.pad(
              kernel, pad_width=((0, diff), (0, 0))
          )
  
          kernel_zp_tensor, kernel_scale_tensor = types[3], types[5]
  
          # Pad kernel zero point by 'diff' elements of 0 if it is not scalar
          if len(kernel_zp_tensor.shape) != 0:
              assert isinstance(kernel_zp, relay.Constant)
              assert isinstance(diff, tvm.tir.IntImm)
              padded_kernel_zp_np = np.append(
                  kernel_zp.data.numpy(), [0] * diff.value
              )
              kernel_zp = relay.const(padded_kernel_zp_np)
  
          # Pad kernel scale by 'diff' elements of 1.0 if it is not scalar
          if len(kernel_scale_tensor.shape) != 0:
              assert isinstance(kernel_scale, relay.Constant)
              assert isinstance(diff, tvm.tir.IntImm)
              padded_kernel_scale_np = np.append(
                  kernel_scale.data.numpy(), [1.0] * diff.value
              )
              kernel_scale = relay.const(padded_kernel_scale_np)
  
          # If units is explicitly specified, it is used to compute the output shape.
          # We need to update units after padding to prevent a type error.
          new_attrs = dict(attrs)
          if attrs["units"] is not None:
              new_attrs["units"] = N + diff
  
          new_inputs = (
              data,
              padded_kernel,
              data_zp,
              kernel_zp,
              data_scale,
              kernel_scale,
          )
  
          out = relay.qnn.op.dense(*new_inputs, **new_attrs)
  
          output_tensor = types[6]
          out = relay.strided_slice(
              out, begin=[0, 0], end=list(output_tensor.shape)
          )
          return out
  
      return None
  
imports: import numpy
  from scipy import special
  import tvm
  from tvm import relay
  from tvm._ffi.base import TVMError
  from tvm.relay.qnn.op.canonicalizations import create_integer_lookup_op
  from ....target.x86 import target_has_features
  from ....topi.utils import is_target
  from ..None import op
label_desc: Use image node_id 25 for calling a global function with example usage: _qnn_dense_legalize_hexagon(attrs, inputs, types) and returns: None, None, out