code: def schedule_bgemm_direct(cfg, s, bgemm, data_pack, kernel_pack):
      """Schedule for bgemm direct"""
      b1, b2, y, x = s[bgemm].op.axis
      rc = s[bgemm].op.reduce_axis[0]
      alpha = get_const_int(b1.dom.extent)
  
      # Create tuning space
      cfg.define_split(
          "tile_b",
          cfg.axis(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("offset_bgemm", [0, 1, 2, 4, 8])
      cfg.define_knob("vector_bgemm", [1, 2, 4, 8])
      offset_bgemm = cfg["offset_bgemm"].val
      vector_bgemm = cfg["vector_bgemm"].val
  
      C = bgemm
      A0, B0 = kernel_pack, data_pack
  
      # Designate the memory hierarchy
      OL = s.cache_write(C, "local")
      AA = s.cache_read(A0, "shared", [OL])
      BB = s.cache_read(B0, "shared", [OL])
  
      # Tile and bind spatial axes
      b = s[bgemm].fuse(b1, b2)
      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, y, x = s[OL].op.axis
      b = s[OL].fuse(b1, b2)
      (rc,) = s[OL].op.reduce_axis
      rco, rci = cfg["tile_rc"].apply(s, OL, rc)
      s[OL].reorder(rco, b, y, x, rci)
  
      s[AA].compute_at(s[OL], rco)
      _, _, k, n = s[AA].op.axis
      AA_align = (
          offset_bgemm
          + cfg["tile_x"].size[1]
          * cfg["tile_x"].size[2]
          * cfg["tile_x"].size[3]
      )
      s[AA].storage_align(k, AA_align - 1, AA_align)
  
      s[BB].compute_at(s[OL], rco)
      _, _, m, k = s[BB].op.axis
      BB_align = offset_bgemm + cfg["tile_rc"].size[1]
      s[BB].storage_align(m, BB_align - 1, BB_align)
  
      # Schedule for A and B shared memory load
      for load in [AA, BB]:
          fused = s[load].fuse(*list(s[load].op.axis))
          fused, ti = s[load].split(fused, factor=vector_bgemm)
          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[load].vectorize(ti)
  
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 3 for calling a global function with example usage: schedule_bgemm_direct(cfg, s, bgemm, data_pack, kernel_pack) without return types