code: def data_weight_transform(
      s,
      data_trans,
      input_tile,
      thread_num_trans,
      offset_trans,
      trans_tag,
  ):
      """Schedule for data or kernel transform"""
      kernel_align = thread_num_trans + offset_trans
      indata_s = s.cache_read(input_tile, "shared", [data_trans])
      data_l = s.cache_write(data_trans, "local")
      # Schedule for data or kernel transform
      eps, nu, p, c = s[data_trans].op.axis
  
      block_x, thread_x = s[data_trans].split(c, thread_num_trans)
      block_x = s[data_trans].fuse(p, block_x)
      s[data_trans].reorder(block_x, thread_x, eps, nu)
      s[data_trans].bind(thread_x, te.thread_axis("threadIdx.x"))
      s[data_trans].bind(block_x, te.thread_axis("blockIdx.x"))
  
      s[data_l].compute_at(s[data_trans], thread_x)
      eps_l, nu_l, p_l, c_l = s[data_l].op.axis
      r_a, r_b = s[data_l].op.reduce_axis
      block_x_l, thread_x_l = s[data_l].split(c_l, thread_num_trans)
      block_x_l = s[data_l].fuse(p_l, block_x_l)
  
      s[data_l].reorder(block_x_l, thread_x_l, eps_l, nu_l, r_a, r_b)
  
      for axis in [eps_l, nu_l, r_a, r_b]:
          s[data_l].unroll(axis)
  
      # Schedule for share memory load
      s[indata_s].compute_at(s[data_l], block_x_l)
      if trans_tag == "data":
          p_is, c_is, eps_is, nu_is = s[indata_s].op.axis
          data_align = (
              get_const_int(eps_is.dom.extent)
              * get_const_int(nu_is.dom.extent)
              + offset_trans
          )
          s[indata_s].storage_align(c_is, data_align - 1, data_align)
          block_x_is, thread_x_is = s[indata_s].split(
              c_is, thread_num_trans
          )
          s[indata_s].bind(thread_x_is, te.thread_axis("threadIdx.x"))
      else:
          eps_is, nu_is, ci_is, co_is = s[indata_s].op.axis
          s[indata_s].storage_align(
              nu_is, kernel_align - 1, kernel_align
          )
          block_x_is, thread_x_is = s[indata_s].split(
              co_is, thread_num_trans
          )
          s[indata_s].reorder(
              ci_is, block_x_is, eps_is, nu_is, thread_x_is
          )
          s[indata_s].bind(thread_x_is, te.thread_axis("threadIdx.x"))
  
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 5 for calling a global function with example usage: data_weight_transform(s, data_trans, input_tile, thread_num_trans, offset_trans, trans_tag) without return types