code: def _schedule_spatial_pack(cfg, s, op, layout):
      """schedule the spatial packing for conv2d"""
  
      assert layout in ("NCHW", "NHWC")
  
      output = op.output(0)
      conv = op.input_tensors[0]
      data_vec = conv.op.input_tensors[0]
      data_pad = data_vec.op.input_tensors[0]
      s[data_pad].compute_inline()
      kernel_vec = conv.op.input_tensors[1]
      if kernel_vec.op.name == "kernel_vec":
          kernel = kernel_vec.op.input_tensors[0]
      else:
          kernel = kernel_vec
      if (
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "dilate" in kernel.op.tag
      ):
          s[kernel].compute_inline()
      data = s[data_vec].op.input_tensors[0]
  
      max_unroll = 16
      vec_size = [1, 2, 4, 8, 16]
      # get tunable parameters (they are defined in compute)
      _, TC, VC = cfg["tile_co"].size
      _, TH, VH = cfg["tile_oh"].size
      _, TW, VW = cfg["tile_ow"].size
  
      # schedule padding
      if isinstance(data.op, tvm.te.ComputeOp) and "pad" in data.op.tag:
          data_pad = data
          s[data_pad].compute_inline()
  
      # schedule data packing
      if layout == "NCHW":
          if (
              isinstance(data_vec.op, tvm.te.ComputeOp)
              and data_vec.op.name == "data_vec_undilated"
          ):
              _, h, w, ci, _, _, vh, vw = s[data_vec].op.axis
          else:
              _, h, w, ci, vh, vw = s[data_vec].op.axis
          z, y, x, unroll1, unroll2 = h, w, ci, vh, vw
      else:
          if (
              isinstance(data_vec.op, tvm.te.ComputeOp)
              and data_vec.op.name == "data_vec_undilated"
          ):
              _, oho, owo, _, _, ic, ohi, owi = s[data_vec].op.axis
          else:
              _, oho, owo, ohi, owi, ic = s[data_vec].op.axis
          z, y, x, unroll1, unroll2 = oho, owo, ohi, ic, owi
      tile_and_bind3d(s, data_vec, z, y, x, 1)
      if unroll1.dom.extent.value < max_unroll:
          s[data_vec].unroll(unroll1)
      if unroll2.dom.extent.value < max_unroll:
          s[data_vec].unroll(unroll2)
  
      if (
          isinstance(kernel_vec.op, tvm.te.ComputeOp)
          and kernel_vec.name == "kernel_vec"
      ):
          if not autotvm.GLOBAL_SCOPE.in_tuning:
              max_threads = tvm.target.Target.current(
                  allow_none=False
              ).max_num_threads
              ax1, ax2, ax3, ax4, ax5 = s[kernel_vec].op.axis
              fused = s[kernel_vec].fuse(ax1, ax2, ax3, ax4, ax5)
              fused, vec = s[kernel_vec].split(fused, VC)
              bb, tt = s[kernel_vec].split(fused, max_threads)
              s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
              s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
              if VC in vec_size:
                  s[kernel_vec].vectorize(vec)
  
      # schedule convolution
      ic, kh, kw = s[conv].op.reduce_axis
      if layout == "NCHW":
          kh_dim, kw_dim = kernel_vec.shape[2], kernel_vec.shape[3]
      else:
          kh_dim, kw_dim = kernel_vec.shape[0], kernel_vec.shape[1]
      cfg["ann_reduce"].apply(
          s,
          conv,
          [kh, kw],
          axis_lens=[get_const_int(kh_dim), get_const_int(kw_dim)],
          max_unroll=max_unroll,
      )
  
      if layout == "NCHW":
          n, c, h, w, vh, vw, vc = s[conv].op.axis
          cfg["reorder_0"].apply(
              s, conv, [n, c, h, w, ic, kh, kw, vh, vw, vc]
          )
          tile_and_bind3d(s, conv, c, h, w, TC, TH, TW)
          unroll_vec_axes = [vh, vw, vc]
          axis_lens = [VH, VW, VC]
      else:
          n, oho, owo, oco, ohi, owi, oci = s[conv].op.axis
          cfg["reorder_conv"].apply(
              s, conv, [n, oho, owo, oco, kh, kw, ic, ohi, owi, oci]
          )
          tile_and_bind3d(s, conv, oho, owo, oco, TH, TW, TC)
          unroll_vec_axes = [ohi, owi, oci]
          axis_lens = [VH, VW, VC]
  
      cfg["ann_spatial"].apply(
          s,
          conv,
          unroll_vec_axes,
          axis_lens,
          max_unroll=max_unroll,
          vec_size=vec_size,
          cfg=cfg,
      )
  
      # schedule output
      if output.op not in s.outputs:  # has bias
          s[output].compute_inline()
          output = s.outputs[0]
      if layout == "NCHW":
          _, co, oh, ow = s[output].op.axis
          tile_and_bind3d(s, output, co, oh, ow, TC, TH, TW)
      else:
          _, oh, ow, co = s[output].op.axis
          tile_and_bind3d(s, output, oh, ow, co, TH, TW, TC)
  
      return s
  
imports: import logging
  import tvm
  from tvm import te
  from tvm import relay
  from tvm import autotvm
  from tvm.autotvm.task.space import get_factors
  from ..utils import traverse_inline, get_const_int, get_const_tuple
  from ..None import nn
  from ..nn.winograd_util import winograd_transform_matrices
  from ..nn.conv2d import conv2d_winograd_nhwc, _conv2d_winograd_nhwc_impl
  from ..arm_cpu.conv2d_spatial_pack import conv2d_spatial_pack_nchw
  from ..arm_cpu.conv2d_spatial_pack import conv2d_spatial_pack_nhwc
label_desc: Use image node_id 5 for calling a global function with example usage: _schedule_spatial_pack(cfg, s, op, layout) and returns: s