code: def conv2d_spatial_pack_nhwc(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      out_dtype,
      num_tile=2,
  ):
      """Spatial pack compute for Conv2d NHWC"""
      out_dtype = out_dtype or data.dtype
  
      N, IH, IW, IC = get_const_tuple(data.shape)
      assert (
          len(kernel.shape) == 4
      ), "AlterOpLayout not enabled for NHWC yet"
      KH, KW, _, OC = get_const_tuple(kernel.shape)
  
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
  
      dilated_kernel_h = (KH - 1) * dilation_h + 1
      dilated_kernel_w = (KW - 1) * dilation_w + 1
  
      pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
          padding, (dilated_kernel_h, dilated_kernel_w)
      )
      HSTR, WSTR = (
          strides
          if isinstance(strides, (tuple, list))
          else (strides, strides)
      )
  
      OH = (IH + pad_top + pad_down - dilated_kernel_h) // HSTR + 1
      OW = (IW + pad_left + pad_right - dilated_kernel_w) // WSTR + 1
      data_pad = nn.pad(
          data, [0, pad_top, pad_left, 0], [0, pad_down, pad_right, 0]
      )
  
      # ==================== define configuration space ====================
      # If it has dynamic shape in batch, we fix the split factor to 1
      n = cfg.axis(N) if isinstance(N, int) else cfg.axis(1)
      oc, oh, ow = cfg.axis(OC), cfg.axis(OH), cfg.axis(OW)
      ic, kh, kw = (
          cfg.reduce_axis(IC),
          cfg.reduce_axis(KH),
          cfg.reduce_axis(KW),
      )
  
      if num_tile == 2:  # for arm cpu
          oco, oci = cfg.define_split("tile_co", oc, num_outputs=2)
          oho, ohi = cfg.define_split("tile_oh", oh, num_outputs=2)
          owo, owi = cfg.define_split("tile_ow", ow, num_outputs=2)
      elif num_tile == 3:  # for mali gpu
          oco, _, oci = cfg.define_split("tile_co", oc, num_outputs=3)
          oho, _, ohi = cfg.define_split("tile_oh", oh, num_outputs=3)
          owo, _, owi = cfg.define_split("tile_ow", ow, num_outputs=3)
      else:
          raise RuntimeError("Invalid num_tile")
  
      cfg.define_reorder(
          "reorder_conv",
          [n, oho, owo, oco, kh, kw, ic, ohi, owi, oci],
          policy="candidate",
          candidate=[
              [n, oho, owo, oco, kh, kw, ic, ohi, owi, oci],
              [n, oho, owo, oco, ohi, kh, kw, ic, owi, oci],
              [n, oho, owo, oco, ohi, kh, kw, owi, ic, oci],
              [n, oho, owo, ohi, oco, kh, kw, owi, ic, oci],
          ],
      )
  
      cfg.define_annotate("ann_reduce", [kh, kw], policy="try_unroll")
      cfg.define_annotate(
          "ann_spatial", [owi, oci], policy="try_unroll_vec"
      )
      # ====================================================================
  
      # If there are no tuning records, use this config
      if cfg.is_fallback:
  
          def _tile_size(axis, candidates):
              for candidate in candidates:
                  tiles_divisible_by_candidate = axis % candidate == 0
                  if tiles_divisible_by_candidate:
                      return candidate
              return 1
  
          # For data tensors with unity height and width we can leave it to the
          # backend to vectorize the inner loop. This has been observed to be more
          # performant on SVE targets with a vector width > 128bits.
          target = Target.current(allow_none=False)
          if target.features.has_sve and OW == OH and OW == 1:
              tile_size = [OC]
              vectorize = "none"
          else:
              # Tile size 8 results in efficient vectorization for these schedules.
              # If the axis is not divisible by 8, try 4
              tile_size = [8, 4]
              vectorize = "vec"
  
          cfg["tile_oh"] = SplitEntity([-1, 1])
          cfg["tile_ow"] = SplitEntity([-1, _tile_size(OW, [8, 4])])
          cfg["tile_co"] = SplitEntity([-1, _tile_size(OC, tile_size)])
          cfg["ann_spatial"] = AnnotateEntity(["none", vectorize])
          cfg["ann_reduce"] = AnnotateEntity(["none", "none"])
          cfg["reorder_conv"] = ReorderEntity(
              [0, 1, 2, 3, 4, 5, 6, 7, 8, 9]
          )
          cfg["compat"] = OtherOptionEntity(0)
  
      OCI = cfg["tile_co"].size[-1]
      OHI = cfg["tile_oh"].size[-1]
      OWI = cfg["tile_ow"].size[-1]
      OCO = OC // OCI
      OHO = OH // OHI
      OWO = OW // OWI
  
      kvshape = (OCO, KH, KW, IC, OCI)
      ovshape = (N, OHO, OWO, OCO, OHI, OWI, OCI)
      oshape = (N, OH, OW, OC)
  
      if dilation_h != 1 or dilation_w != 1:
          # undilate input data
          dvshape = (N, OHO, OWO, KH, KW, IC, OHI, OWI)
          data_vec = te.compute(
              dvshape,
              lambda n, oho, owo, kh, kw, ic, ohi, owi: data_pad[n][
                  (oho * OHI + ohi) * HSTR + kh * dilation_h
              ][(owo * OWI + owi) * WSTR + kw * dilation_w][ic],
              name="data_vec_undilated",
          )
      else:
          dvshape = (
              N,
              OHO,
              OWO,
              KH + (OHI - 1) * HSTR,
              KW + (OWI - 1) * WSTR,
              IC,
          )
          data_vec = te.compute(
              dvshape,
              lambda n, oho, owo, ohi, owi, ic: data_pad[n][
                  oho * OHI * HSTR + ohi
              ][owo * OWI * WSTR + owi][ic],
              name="data_vec",
          )
  
      if autotvm.GLOBAL_SCOPE.in_tuning:
          kernel_vec = tvm.te.placeholder(
              kvshape, kernel.dtype, name="kernel"
          )
      else:
          kernel_vec = te.compute(
              kvshape,
              lambda oco, kh, kw, ic, oci: kernel[kh][kw][ic][
                  oco * OCI + oci
              ],
              name="kernel_vec",
          )
  
      ic = te.reduce_axis((0, IC), name="ic")
      kh = te.reduce_axis((0, KH), name="kh")
      kw = te.reduce_axis((0, KW), name="kw")
  
      if dilation_h != 1 or dilation_w != 1:
          conv = te.compute(
              ovshape,
              lambda n, oho, owo, oco, ohi, owi, oci: te.sum(
                  data_vec[n, oho, owo, kh, kw, ic, ohi, owi].astype(
                      out_dtype
                  )
                  * kernel_vec[oco, kh, kw, ic, oci].astype(out_dtype),
                  axis=[ic, kh, kw],
              ),
              name="conv",
          )
      else:
          conv = te.compute(
              ovshape,
              lambda n, oho, owo, oco, ohi, owi, oci: te.sum(
                  data_vec[
                      n, oho, owo, ohi * HSTR + kh, owi * WSTR + kw, ic
                  ].astype(out_dtype)
                  * kernel_vec[oco, kh, kw, ic, oci].astype(out_dtype),
                  axis=[ic, kh, kw],
              ),
              name="conv",
          )
  
      idiv = tvm.tir.indexdiv
      imod = tvm.tir.indexmod
      output = te.compute(
          oshape,
          lambda n, oho, owo, oc: conv[n][idiv(oho, OHI)][
              idiv(owo, OWI)
          ][idiv(oc, OCI)][imod(oho, OHI)][imod(owo, OWI)][
              imod(oc, OCI)
          ],
          name="output_unpack",
          tag="spatial_conv_output_NHWC",
      )
      return output
  
file_path: tvm/python/tvm/topi/arm_cpu/conv2d_spatial_pack.py
imports: from __future__ import absolute_import
  import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.target import Target
  from tvm.autotvm.task.space import SplitEntity, OtherOptionEntity, AnnotateEntity, ReorderEntity
  from ..None import nn
  from ..utils import get_const_tuple
  from ..nn.utils import get_const_int, get_pad_tuple
label_desc: Use image node_id 3 for calling a global function with example usage: conv2d_spatial_pack_nhwc(cfg, data, kernel, strides, padding, dilation, out_dtype, num_tile) and returns: output, int, candidate