code: def schedule_winograd_cuda(cfg, s, output, pre_computed):
      """Schedule winograd template"""
      # get stages
      inverse = s[output].op.input_tensors[0]
      bgemm, A = s[inverse].op.input_tensors
      kernel_pack, data_pack = s[bgemm].op.input_tensors
      input_tile, B = s[data_pack].op.input_tensors
      pad_data = s[input_tile].op.input_tensors[0]
  
      # data transform
      s[B].compute_inline()
  
      data_l = s.cache_write(data_pack, "local")
      omg, eps, nu, c, p = s[data_l].op.axis
      r_a, r_b, r_c = s[data_l].op.reduce_axis
      # TODO unrolling by omg, eps, nu may improve performance but
      # in some cases causes extremely long build times due to imperfect tiling.
      for axis in [r_a, r_b, r_c]:
          s[data_l].unroll(axis)
  
      omg, eps, nu, c, p = s[data_pack].op.axis
      p, pi = s[data_pack].split(p, 1)
      fused = s[data_pack].fuse(c, p)
      bb, tt = s[data_pack].split(fused, 128)
      s[data_pack].reorder(bb, tt, pi, omg, eps, nu)
      s[data_pack].bind(bb, te.thread_axis("blockIdx.x"))
      s[data_pack].bind(tt, te.thread_axis("threadIdx.x"))
  
      s[data_l].compute_at(s[data_pack], pi)
      s[input_tile].compute_at(s[data_pack], pi)
      s[pad_data].compute_inline()
  
      # transform kernel
      if not pre_computed and not autotvm.GLOBAL_SCOPE.in_tuning:
          kernel, G = s[kernel_pack].op.input_tensors
          omg, eps, nu, co, ci = s[kernel_pack].op.axis
          s[G].compute_inline()
          r_a, r_b, r_c = s[kernel_pack].op.reduce_axis
          # Could add additional unrolling by omg, eps, nu in the future.
          for axis in [r_a, r_b, r_c]:
              s[kernel_pack].unroll(axis)
  
          fused = s[kernel_pack].fuse(co, ci)
          bb, tt = s[kernel_pack].split(fused, 128)
          s[kernel_pack].reorder(bb, tt, omg, eps, nu, r_a, r_b, r_c)
          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 "dilate" in kernel.op.tag
      ):
          s[kernel].compute_inline()
  
      ##### space definition begin #####
      b1, b2, b3, y, x = s[bgemm].op.axis
      rc = s[bgemm].op.reduce_axis[0]
      alpha = get_const_int(b1.dom.extent)
  
      cfg.define_split(
          "tile_b",
          cfg.axis(alpha * alpha * alpha),
          num_outputs=4,
          filter=lambda x: x.size[-3:] == [1, 1, 1],
      )
      cfg.define_split("tile_y", y, num_outputs=4)
      cfg.define_split("tile_x", x, num_outputs=4)
      cfg.define_split("tile_rc", rc, num_outputs=2)
      cfg.define_knob("auto_unroll_max_step", [0, 128, 1500])
      target = tvm.target.Target.current()
      if target.kind.name in ["nvptx", "rocm"]:
          cfg.define_knob("unroll_explicit", [1])
      else:
          cfg.define_knob("unroll_explicit", [0, 1])
      ##### space definition end #####
  
      # batch gemm
      C = bgemm
      A0, B0 = kernel_pack, data_pack
  
      OL = s.cache_write(C, "local")
      AA = s.cache_read(A0, "shared", [OL])
      BB = s.cache_read(B0, "shared", [OL])
  
      b = s[bgemm].fuse(b1, b2, b3)
  
      # tile and bind spatial axes
      bgemm_scope, b = s[bgemm].split(b, nparts=1)
      bz, vz, tz, zi = cfg["tile_b"].apply(s, C, b)
      by, vy, ty, yi = cfg["tile_y"].apply(s, C, y)
      bx, vx, tx, xi = cfg["tile_x"].apply(s, C, x)
      s[C].bind(bz, te.thread_axis("blockIdx.z"))
      s[C].bind(by, te.thread_axis("blockIdx.y"))
      s[C].bind(bx, te.thread_axis("blockIdx.x"))
      s[C].bind(vz, te.thread_axis("vthread"))
      s[C].bind(vy, te.thread_axis("vthread"))
      s[C].bind(vx, te.thread_axis("vthread"))
      s[C].bind(tz, te.thread_axis("threadIdx.z"))
      s[C].bind(ty, te.thread_axis("threadIdx.y"))
      s[C].bind(tx, te.thread_axis("threadIdx.x"))
      s[C].reorder(
          bgemm_scope, bz, by, bx, vz, vy, vx, tz, ty, tx, zi, yi, xi
      )
  
      # tile reduction axes
      s[OL].compute_at(s[C], tx)
      b1, b2, b3, y, x = s[OL].op.axis
      b = s[OL].fuse(b1, b2, b3)
      (rc,) = s[OL].op.reduce_axis
      rco, rci = cfg["tile_rc"].apply(s, OL, rc)
      s[OL].reorder(rco, rci, b, y, x)
  
      s[AA].compute_at(s[OL], rco)
      s[BB].compute_at(s[OL], rco)
  
      # cooperative fetching
      for load in [AA, BB]:
          fused = s[load].fuse(*list(s[load].op.axis))
          fused, tx = s[load].split(fused, cfg["tile_x"].size[2])
          fused, ty = s[load].split(fused, cfg["tile_y"].size[2])
          fused, tz = s[load].split(fused, cfg["tile_b"].size[2])
          s[load].bind(tz, te.thread_axis("threadIdx.z"))
          s[load].bind(ty, te.thread_axis("threadIdx.y"))
          s[load].bind(tx, te.thread_axis("threadIdx.x"))
  
      s[C].pragma(
          bgemm_scope,
          "auto_unroll_max_step",
          cfg["auto_unroll_max_step"].val,
      )
      s[C].pragma(
          bgemm_scope, "unroll_explicit", cfg["unroll_explicit"].val
      )
  
      # 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]
  
      m = alpha - 3 + 1
      n, co, d, h, w = s[output].op.axis
      do, di = s[output].split(d, m)
      ho, hi = s[output].split(h, m)
      wo, wi = s[output].split(w, m)
      s[output].reorder(n, co, do, ho, wo, di, hi, wi)
      inverse_scope, n = s[output].split(n, nparts=1)
  
      fused = s[output].fuse(n, co, do, ho, wo)
      bb, tt = s[output].split(fused, 128)
  
      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)
  
      s[A].compute_inline()
      co, p, vd, vh, vw = s[inverse].op.axis
      r_a, r_b, r_c = s[inverse].op.reduce_axis
      # Could add additional unrolling of vd, vh, vw, in the future
      for axis in [r_a, r_b, r_c]:
          s[inverse].unroll(axis)
      s[inverse].compute_at(s[output], tt)
  
      return s
  
file_path: tvm/python/tvm/topi/cuda/conv3d_winograd.py
imports: import logging
  import tvm
  from tvm import te
  from tvm import autotvm
  from ..None import nn
  from ..utils import get_const_int, get_const_tuple, traverse_inline, simplify
  from ..nn.winograd_util import winograd_transform_matrices
label_desc: Use image node_id 4 for calling a global function with example usage: schedule_winograd_cuda(cfg, s, output, pre_computed) and returns: s