code: def _schedule_spatial_pack(
      cfg, s, output, conv, data_vec, kernel_vec
  ):
      """schedule the spatial packing for conv2d"""
      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)
      BC, TC, VC = cfg["tile_co"].size
      BH, TH, VH = cfg["tile_oh"].size
      BW, 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 (
          isinstance(data_vec.op, te.tensor.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
      tile_and_bind3d(s, data_vec, h, w, ci, 1)
      if vh.dom.extent.value < max_unroll:
          s[data_vec].unroll(vh)
      if vw.dom.extent.value < max_unroll:
          s[data_vec].unroll(vw)
  
      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
              co, ci, kh, kw, vc = s[kernel_vec].op.axis
              fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
              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
      n, c, h, w, vh, vw, vc = s[conv].op.axis
      kc, kh, kw = s[conv].op.reduce_axis
  
      cfg["reorder_0"].apply(
          s, conv, [n, c, h, w, kc, kh, kw, vh, vw, vc]
      )
      tile_and_bind3d(s, conv, c, h, w, TC, TH, TW)
  
      cfg["ann_reduce"].apply(
          s,
          conv,
          [kh, kw],
          axis_lens=[
              get_const_int(kernel_vec.shape[2]),
              get_const_int(kernel_vec.shape[3]),
          ],
          max_unroll=max_unroll,
      )
  
      cfg["ann_spatial"].apply(
          s,
          conv,
          [vh, vw, vc],
          axis_lens=[VH, VW, VC],
          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]
  
      _, co, oh, ow = s[output].op.axis
      tile_and_bind3d(s, output, co, oh, ow, TC, TH, TW)
  
      return s
  
file_path: tvm/python/tvm/topi/bifrost/conv2d.py
imports: import tvm
  from tvm import te
  from tvm import relay
  from tvm import autotvm
  from .gemm import decl_winograd_gemm, schedule_gemm
  from .transforms import tile_and_bind, tile_and_bind3d
  from ..utils import traverse_inline, get_const_int, get_const_tuple
  from ..None import nn
  from ..nn.winograd_util import winograd_transform_matrices
  from ..arm_cpu.conv2d_spatial_pack import conv2d_spatial_pack_nchw
label_desc: Use image node_id 3 for calling a global function with example usage: _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec) and returns: s