code: def schedule_conv2d_winograd(cfg, s, output, pre_computed):
      """Schedule winograd template"""
      inverse = s[output].op.input_tensors[0]
      bgemm, A = s[inverse].op.input_tensors
      kernel_pack, data_pack_trans = s[bgemm].op.input_tensors
      data_pack = s[data_pack_trans].op.input_tensors[0]
      input_tile, B = s[data_pack].op.input_tensors
      pad_data = s[input_tile].op.input_tensors[0]
  
      # data transform
      s[B].compute_inline()
      s[A].compute_inline()
  
      # probably will improve real topology execution
      if autotvm.GLOBAL_SCOPE.in_tuning:
          # Padding to texture
          AA = s.cache_read(
              pad_data,
              get_texture_storage(pad_data.shape),
              [input_tile],
          )
          bind_data_copy(s[AA])
  
      s[input_tile].compute_inline()
  
      OL = s.cache_write(data_pack, "local")
      c, p, eps, nu, cb = s[data_pack].op.axis
      fused = s[data_pack].fuse(c, p, eps, nu)
      bx, tx = s[data_pack].split(fused, 128)
      s[data_pack].vectorize(cb)
      s[data_pack].bind(bx, te.thread_axis("blockIdx.x"))
      s[data_pack].bind(tx, te.thread_axis("threadIdx.x"))
  
      _, _, eps, nu, cb = s[OL].op.axis
      r_a, r_b = s[OL].op.reduce_axis
      s[OL].unroll(eps)
      s[OL].unroll(nu)
      s[OL].unroll(r_a)
      s[OL].unroll(r_b)
      s[OL].vectorize(cb)
      s[OL].compute_at(s[data_pack], tx)
      s[data_pack].set_scope(get_texture_storage(data_pack.shape))
  
      s[data_pack_trans].compute_inline()
  
      # transform kernel
      if not pre_computed:
          kernel, G = s[kernel_pack].op.input_tensors
          eps, nu, ci, co, cob = s[kernel_pack].op.axis
          if autotvm.GLOBAL_SCOPE.in_tuning:
              # skip this part during tuning to make recrods accurate
              # this part will be pre-computed during pre-compute optimization pass
              s[G].pragma(s[G].op.axis[0], "debug_skip_region")
              s[kernel_pack].pragma(eps, "debug_skip_region")
          else:
              s[G].compute_inline()
              r_a, r_b = s[kernel_pack].op.reduce_axis
              for axis in [eps, nu, r_a, r_b]:
                  s[kernel_pack].unroll(axis)
  
              fused = s[kernel_pack].fuse(ci, co)
              bb, tt = s[kernel_pack].split(fused, 128)
              s[kernel_pack].reorder(bb, tt, eps, nu, r_a, r_b, cob)
              s[kernel_pack].vectorize(cob)
              s[kernel_pack].bind(bb, te.thread_axis("blockIdx.x"))
              s[kernel_pack].bind(tt, te.thread_axis("threadIdx.x"))
      else:
          kernel = kernel_pack
  
      if (
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "filter_pack" in kernel.op.tag
      ):
          # manage scheduling of datacopy
          pack_data = pad_data.op.input_tensors[0]
          bind_data_copy(s[pack_data])
          bind_data_copy(s[kernel])
      elif (
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "dilate" in kernel.op.tag
      ):
          s[kernel].compute_inline()
      s[pad_data].compute_inline()
  
      ##### space definition begin #####
      cfg.define_knob("auto_unroll_max_step", [0, 4, 16])
      b1, b2, y, x, cb = s[bgemm].op.axis
      rcc = s[bgemm].op.reduce_axis[0]
      alpha = get_const_int(b1.dom.extent)
  
      cfg.define_split(
          "tile_y",
          y,
          num_outputs=3,
          filter=lambda entry: entry.size[2] <= 64
          and entry.size[1] <= 16,
      )
  
      min_x_div = 1
      for bn in range(4, 0, -1):
          if bgemm.shape[3] % bn == 0:
              min_x_div = bn
              break
  
      cfg.define_split(
          "tile_x",
          x,
          num_outputs=3,
          filter=lambda entry: entry.size[2] <= 64
          and entry.size[1] >= min_x_div
          and entry.size[1] <= 16,
      )
      cfg.define_split("tile_rc", rcc, num_outputs=2)
      cfg.multi_filter(
          filter=lambda entity: 32
          <= (entity["tile_y"].size[2] * entity["tile_x"].size[2])
          < 1024
      )
      ##### space definition end #####
  
      # batch gemm
      OL = s.cache_write(bgemm, "local")
      if (
          autotvm.GLOBAL_SCOPE.in_tuning
          or isinstance(kernel.op, tvm.te.ComputeOp)
          and "filter_pack" in kernel.op.tag
          and kernel.shape[2] == 1
          and kernel.shape[3] == 1
      ):
          BB = s.cache_read(
              kernel_pack, get_texture_storage(kernel_pack.shape), [OL]
          )
          bind_data_copy(s[BB])
  
      by = s[bgemm].fuse(b1, b2, y)
  
      # tile and bind spatial axes
      bgemm_scope, by = s[bgemm].split(by, nparts=1)
      by, vy, ty = cfg["tile_y"].apply(s, bgemm, by)
      bx, vx, tx = cfg["tile_x"].apply(s, bgemm, x)
      s[bgemm].bind(by, te.thread_axis("blockIdx.y"))
      s[bgemm].bind(bx, te.thread_axis("blockIdx.x"))
      s[bgemm].bind(vy, te.thread_axis("vthread"))
      s[bgemm].bind(vx, te.thread_axis("vthread"))
      s[bgemm].bind(ty, te.thread_axis("threadIdx.y"))
      s[bgemm].bind(tx, te.thread_axis("threadIdx.x"))
      s[bgemm].reorder(bgemm_scope, by, bx, vy, vx, ty, tx, cb)
      s[bgemm].vectorize(cb)
      s[bgemm].set_scope(get_texture_storage(bgemm.shape))
  
      # tile reduction axes
      s[OL].compute_at(s[bgemm], tx)
      b1, b2, y, x, cb = s[OL].op.axis
      (rcc, rcb) = s[OL].op.reduce_axis
      b = s[OL].fuse(b1, b2)
      s[OL].reorder(b, y, x, rcc, rcb, cb)
      # s[OL].unroll(rcb)
      s[OL].pragma(
          rcb, "auto_unroll_max_step", cfg["auto_unroll_max_step"].val
      )
      s[OL].pragma(rcb, "unroll_explicit", True)
      s[OL].vectorize(cb)
  
      # schedule inverse, output and fusion
      if output.op in s.outputs:
          OL = None
      else:
          OL = output
          s[OL].set_scope("local")
          output = s.outputs[0]
  
      if len(s[output].op.axis) == 4:
          n, co, h, w = s[output].op.axis
          cb = None
      else:
          n, co, h, w, cb = s[output].op.axis
      inverse_scope, n = s[output].split(n, nparts=1)
  
      fused = s[output].fuse(n, co, h, w)
      bb, tt = s[output].split(fused, 128)
      if cb is not None:
          s[output].reorder(bb, tt, cb)
          s[output].vectorize(cb)
  
      s[output].bind(bb, te.thread_axis("blockIdx.x"))
      s[output].bind(tt, te.thread_axis("threadIdx.x"))
  
      if OL is not None:
          s[OL].compute_at(s[output], tt)
  
      co, p, vh, vw, cb = s[inverse].op.axis
      r_a, r_b = s[inverse].op.reduce_axis
      for axis in [vh, vw, r_a, r_b]:
          s[inverse].unroll(axis)
      s[inverse].vectorize(cb)
      s[inverse].compute_at(s[output], tt)
  
      return s
  
file_path: tvm/python/tvm/topi/adreno/conv2d_winograd_common.py
imports: import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.topi import nn
  from tvm.topi.utils import get_const_int, get_const_tuple, traverse_inline
  from ..nn.winograd_util import winograd_transform_matrices
  from .utils import split_to_chunks, pack_input, pack_filter, bind_data_copy, get_texture_storage, infer_tile_size
label_desc: Use image node_id 3 for calling a global function with example usage: schedule_conv2d_winograd(cfg, s, output, pre_computed) and returns: s