code: def _schedule_winograd(cfg, s, output, last):
      Y = output.op.input_tensors[0]
      M, A = Y.op.input_tensors
      U, V = M.op.input_tensors
      d, B = V.op.input_tensors
      data_pad = d.op.input_tensors[0]
  
      # padding
      s[data_pad].compute_inline()
  
      # pack input tiles
      s[d].compute_inline()
  
      # transform kernel
      if isinstance(U.op, tvm.te.ComputeOp):
          kernel, G = U.op.input_tensors
          s[G].compute_inline()
          (eps, nu, k, c, kk) = s[U].op.axis
          if autotvm.GLOBAL_SCOPE.in_tuning:
              # kernel transformation will be pre-computed during compilation, so we skip
              # this part to make tuning records correct
              s[U].pragma(eps, "debug_skip_region")
          else:
              r_kh, r_kw = s[U].op.reduce_axis
              s[U].reorder(k, c, eps, nu, r_kh, r_kw, kk)
              for axis in [eps, nu, r_kh, r_kw]:
                  s[U].unroll(axis)
              s[U].vectorize(kk)
              s[U].parallel(k)
  
          if (
              isinstance(kernel.op, tvm.te.ComputeOp)
              and "dilate" in kernel.op.tag
          ):
              s[kernel].compute_inline()
  
      # transform image
      DD = s.cache_read(d, "global", [V])
      s[B].compute_inline()
      eps, nu, b, c, bb = s[V].op.axis
      r_eps, r_nu = s[V].op.reduce_axis
      s[V].reorder(b, c, eps, nu, r_eps, r_nu, bb)
      for axis in [eps, nu, r_eps, r_nu]:
          s[V].unroll(axis)
      s[DD].compute_at(s[V], c)
      s[V].vectorize(bb)
      s[V].parallel(b)
  
      # batch gemm
      eps, nu, k, b = s[M].op.axis
      c = s[M].op.reduce_axis[0]
      cfg.define_split(
          "tile_c", c, num_outputs=2, filter=lambda x: x.size[-1] <= 16
      )
      co, ci = cfg["tile_c"].apply(s, M, c)
      xo, xi = cfg["tile_p"].apply(s, M, b)
      s[M].reorder(eps, nu, xo, co, k, ci, xi)
      cfg.define_annotate("ann_reduce", [ci], policy="try_unroll")
      cfg.define_annotate(
          "ann_spatial", [k, xi], policy="try_unroll_vec"
      )
      cfg["ann_reduce"].apply(
          s,
          M,
          [ci],
          axis_lens=[cfg["tile_c"].size[-1]],
          max_unroll=16,
          cfg=cfg,
      )
      cfg["ann_spatial"].apply(s, M, [k, xi])
  
      # inverse transform
      s[A].compute_inline()
      k, b, vh, vw = s[Y].op.axis
      r_eps, r_nu = s[Y].op.reduce_axis
      for axis in [vh, vw, r_eps, r_nu]:
          s[Y].unroll(axis)
  
      # output
      n, co, h, w = s[last].op.axis
      co, coi = cfg["tile_k"].apply(s, last, co)
      p = s[last].fuse(n, co)
      s[M].compute_at(s[last], p)
      s[last].parallel(p)
  
      MM = s.cache_read(M, "global", [Y])
      m = get_const_int(V.shape[0]) + 1 - 3
      ho, wo, hi, wi = s[last].tile(h, w, m, m)
      s[Y].compute_at(s[last], wo)
      s[MM].compute_at(s[last], wo)
  
      if output != last:
          s[output].compute_inline()
  
imports: from __future__ import absolute_import
  import tvm
  from tvm import te
  from tvm import autotvm
  import tvm.contrib.nnpack
  from ..utils import traverse_inline, get_const_tuple
  from ..None import nn
  from ..nn.utils import get_const_int, get_pad_tuple
  from ..nn.winograd_util import winograd_transform_matrices
  from .arm_utils import get_tiling_B_transformed
  from .conv2d_spatial_pack import conv2d_spatial_pack_nchw, conv2d_spatial_pack_nhwc, schedule_conv2d_spatial_pack_nchw, schedule_conv2d_spatial_pack_nhwc
  from .conv2d_gemm import compute_conv2d_gemm_without_weight_transform, schedule_conv2d_gemm_interleaved, schedule_conv2d_gemm_native
  from .mprofile.dsp.conv2d import conv2d_nhwc_dsp_compute, conv2d_nhwc_dsp_schedule
label_desc: Use image node_id 8 for calling a global function with example usage: _schedule_winograd(cfg, s, output, last) without return types