code: def schedule_nhwc_winograd_cuda(
      cfg, s, output, use_tensorcore, pre_computed
  ):
      """Schedule winograd template"""
      # Get stages
      inverse = s[output].op.input_tensors[0]
      bgemm, A = s[inverse].op.input_tensors
      Transdata, TransFilter = s[bgemm].op.input_tensors
      data_pack = s[Transdata].op.input_tensors[0]
      kernel_pack = s[TransFilter].op.input_tensors[0]
      s[Transdata].compute_inline()
      s[TransFilter].compute_inline()
  
      input_tile, B = s[data_pack].op.input_tensors
      pad_data = s[input_tile].op.input_tensors[0]
  
      # Define the stride of intrin functions
      cfg.define_knob("thread_num_inverse", [1, 32, 64, 128, 256])
      cfg.define_knob("thread_num_data", [1, 32, 64, 128, 256])
      cfg.define_knob("thread_num_kernel", [1, 32, 64, 128, 256])
      cfg.define_knob("offset_inverse", [0, 2, 4])
      cfg.define_knob("offset_data", [0, 1, 2, 4])
      cfg.define_knob("offset_kernel", [0, 1, 2, 4])
      cfg.define_knob("inverse_in_vector", [1, 2, 4])
  
      thread_num_data = cfg["thread_num_data"].val
      thread_num_kernel = cfg["thread_num_kernel"].val
      thread_num_inverse = cfg["thread_num_inverse"].val
      offset_data = cfg["offset_data"].val
      offset_kernel = cfg["offset_kernel"].val
      offset_inverse = cfg["offset_inverse"].val
      inverse_in_vector = cfg["inverse_in_vector"].val
  
      # Data transform
      s[B].compute_inline()
      data_weight_transform(
          s,
          data_pack,
          input_tile,
          thread_num_data,
          offset_data,
          trans_tag="data",
      )
      s[input_tile].compute_inline()
      s[pad_data].compute_inline()
  
      # Kernel transform
      if not pre_computed and not autotvm.GLOBAL_SCOPE.in_tuning:
          kernel, G = s[kernel_pack].op.input_tensors
          s[G].compute_inline()
          data_weight_transform(
              s,
              kernel_pack,
              kernel,
              thread_num_kernel,
              offset_kernel,
              trans_tag="kernel",
          )
      else:
          kernel = kernel_pack
  
      if (
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "dilate" in kernel.op.tag
      ):
          s[kernel].compute_inline()
  
      b1, b2, y, x = s[bgemm].op.axis
      alpha = get_const_int(b1.dom.extent)
      _, _, P, CI = get_const_tuple(Transdata.shape)
      _, _, _, CO = get_const_tuple(TransFilter.shape)
  
      # Determine whether the shape is available with tensorcore
      shape_judge = (
          (P % 16 == 0 and CI % 16 == 0 and CO % 16 == 0)
          or (P % 8 == 0 and CI % 16 == 0 and CO % 32 == 0)
          or (P % 32 == 0 and CI % 16 == 0 and CO % 8 == 0)
      )
  
      if shape_judge and use_tensorcore:
          schedule_bgemm_tensorcore(
              cfg, s, bgemm, Transdata, TransFilter
          )
      else:
          schedule_bgemm_direct(cfg, s, bgemm, Transdata, TransFilter)
  
      # 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]
  
      s[A].compute_inline()
      inverse_s = s.cache_read(bgemm, "shared", [inverse])
  
      m = alpha - 3 + 1
      offset_inverse_in = offset_inverse
      vector_width_inverse_in = inverse_in_vector
  
      # Schedule for output
      n, h, w, co = s[output].op.axis
      ho, wo, hi, wi = s[output].tile(h, w, m, m)
      s[output].reorder(n, ho, wo, co, hi, wi)
      fused = s[output].fuse(n, ho, wo)
  
      block_x_s, thread_x_s = s[output].split(co, thread_num_inverse)
      block_x_s = s[output].fuse(fused, block_x_s)
      s[output].reorder(block_x_s, thread_x_s, hi, wi)
  
      if OL is not None:
          s[OL].compute_inline()
  
      # Schedule for inverse
      s[inverse].compute_at(s[output], thread_x_s)
      p_inv, co_inv, eps_inv, nu_inv = s[inverse].op.axis
      block_x_inv, thread_x_inv = s[inverse].split(
          co_inv, thread_num_inverse
      )
      r_a, r_b = s[inverse].op.reduce_axis
      for axis in [eps_inv, nu_inv, r_a, r_b]:
          s[inverse].unroll(axis)
  
      # Schedule for share memory load
      s[inverse_s].compute_at(s[output], block_x_s)
      eps_inv_s, nu_inv_s, p_inv_s, co_inv_s = s[inverse_s].op.axis
      inverse_in_align = offset_inverse_in + thread_num_inverse
      s[inverse_s].storage_align(
          p_inv_s, inverse_in_align - 1, inverse_in_align
      )
      block_x_inv_s, thread_x_inv_s = s[inverse_s].split(
          co_inv_s, thread_num_inverse
      )
      block_x_inv_s = s[inverse_s].fuse(p_inv_s, block_x_inv_s)
      s[inverse_s].reorder(
          block_x_inv_s, eps_inv_s, nu_inv_s, thread_x_inv_s
      )
      t = s[inverse_s].fuse(eps_inv_s, nu_inv_s, thread_x_inv_s)
      t, ti = s[inverse_s].split(t, factor=vector_width_inverse_in)
      t, tx = s[inverse_s].split(t, factor=thread_num_inverse)
      s[inverse_s].bind(tx, te.thread_axis("threadIdx.x"))
      s[inverse_s].vectorize(ti)
  
      s[output].bind(thread_x_s, te.thread_axis("threadIdx.x"))
      s[output].bind(block_x_s, te.thread_axis("blockIdx.x"))
      return s
  
file_path: tvm/python/tvm/topi/cuda/conv2d_nhwc_winograd.py
imports: import tvm
  from tvm import autotvm, te
  from ..None import nn
  from ..nn.winograd_util import winograd_transform_matrices
  from ..utils import get_const_int, get_const_tuple, traverse_inline
  from .tensor_intrin import intrin_wmma_gemm, intrin_wmma_load_matrix_A, intrin_wmma_load_matrix_W, intrin_wmma_store_matrix
label_desc: Use image node_id 6 for calling a global function with example usage: schedule_nhwc_winograd_cuda(cfg, s, output, use_tensorcore, pre_computed) and returns: s