code: def _schedule_winograd(cfg, s, op):
      """schedule winograd fast convolution F(2x2, 3x3) for conv2d"""
      # get ops and tensors
      output = op.output(0)
  
      Y = op.input_tensors[0]
      M, A = s[Y].op.input_tensors
      U, V = s[M].op.input_tensors
      d, B = s[V].op.input_tensors
      data_pad = s[d].op.input_tensors[0]
  
      # padding
      s[data_pad].compute_inline()
  
      # transform kernel
      if isinstance(U.op, tvm.te.ComputeOp):
          kernel, G = s[U].op.input_tensors
          s[G].compute_inline()
          (eps, nu, co, ci, vco) = s[U].op.axis
          if not autotvm.GLOBAL_SCOPE.in_tuning:
              r_kh, r_kw = s[U].op.reduce_axis
              s[U].reorder(co, ci, eps, nu, r_kh, r_kw, vco)
              _ = [s[U].unroll(x) for x in [eps, nu, r_kh, r_kw]]
              s[U].vectorize(vco)
              tile_and_bind(s, U, co, ci, 1, 256)
  
          # dilation
          if (
              isinstance(kernel.op, tvm.te.ComputeOp)
              and "dilate" in kernel.op.tag
          ):
              s[kernel].compute_inline()
  
      # transform image
      s[B].compute_inline()
      VL = s.cache_write(V, "local")
  
      eps, nu, p, ci, vp = s[V].op.axis
      s[V].reorder(p, ci, eps, nu, vp)
      for axis in [eps, nu]:
          s[V].unroll(axis)
      s[V].vectorize(vp)
      fused = s[V].fuse(p, ci)
  
      bb, tt = cfg["tile_t1"].apply(s, V, fused)
      s[V].bind(bb, te.thread_axis("blockIdx.x"))
      s[V].bind(tt, te.thread_axis("threadIdx.x"))
  
      eps, nu, p, ci, vp = s[VL].op.axis
      r_a, r_b = s[VL].op.reduce_axis
      for axis in [eps, nu, r_a, r_b]:
          s[VL].unroll(axis)
      s[VL].vectorize(vp)
      s[d].compute_at(s[V], tt)
      s[VL].compute_at(s[V], tt)
  
      # batch gemm
      bna = cfg["tile_bna"].val
      bnb = cfg["tile_bnb"].val
  
      eps, nu, k, b = s[M].op.axis
      alpha = eps.dom.extent
      c = s[M].op.reduce_axis[0]
      yo, xo, yi, xi = s[M].tile(k, b, bna, bnb)
      c, c_unroll = cfg["c_unroll"].apply(s, M, c)
      s[M].reorder(yo, xo, c, c_unroll, yi, xi)
      s[M].unroll(c_unroll)
      s[M].unroll(yi)
      s[M].vectorize(xi)
      z = s[M].fuse(eps, nu)
      tile_and_bind3d(s, M, z, yo, xo, 1, cfg["yt"].val, 1)
  
      # inverse transform
      s[A].compute_inline()
      k, b, vh, vw = s[Y].op.axis
      r_a, r_b = s[Y].op.reduce_axis
      for axis in [vh, vw, r_a, r_b]:
          s[Y].unroll(axis)
  
      # schedule output and fusion
      if output.op not in s.outputs:
          s[output].compute_inline()
          output = s.outputs[0]
  
      n, co, h, w = s[output].op.axis
      m = alpha - 3 + 1
      h, w, hi, wi = s[output].tile(h, w, m, m)
      s[output].unroll(hi)
      s[output].unroll(wi)
      fused = s[output].fuse(n, co, h, w)
      bb, tt = cfg["tile_t2"].apply(s, output, fused)
      s[output].bind(bb, te.thread_axis("blockIdx.x"))
      s[output].bind(tt, te.thread_axis("threadIdx.x"))
  
      s[Y].compute_at(s[output], tt)
  
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 10 for calling a global function with example usage: _schedule_winograd(cfg, s, op) without return types