code: def _schedule_dense_int8(cfg, s, output):
      data, weight = s[output].op.input_tensors
      if (
          len(weight.op.input_tensors) == 1
          and weight.op.input_tensors[0] == data
      ):
          s[weight].compute_inline()
  
      batch, in_dim = get_const_tuple(data.shape)
      out_dim, _ = get_const_tuple(weight.shape)
  
      in_dim_factor = 4
      assert (
          in_dim % in_dim_factor == 0
      ), f"Input dimension must divide {in_dim_factor}"
      if in_dim % 16 == 0:
          in_dim_factor = 16
  
      # create tuning space
      cfg.define_split("tile_y", batch, num_outputs=4)
      cfg.define_split("tile_x", out_dim, num_outputs=4)
      cfg.define_split("tile_k", in_dim // in_dim_factor, num_outputs=2)
      cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
  
      # create cache stage
      AA = s.cache_read(data, "shared", [output])
      WW = s.cache_read(weight, "shared", [output])
      CC = s.cache_write(output, "local")
  
      # handle bias
      if output.op not in s.outputs:
          s[output].compute_inline()
          output = s.outputs[0].output(0)
  
      n, x = s[output].op.axis
  
      # this is the scope to attach global config inside this kernel
      kernel_scope, n = s[output].split(n, nparts=1)
  
      ko = CC.op.reduce_axis[0]
      ko, ki = s[CC].split(ko, factor=4)
      ko, kt = cfg["tile_k"].apply(s, CC, ko)
      target = tvm.target.Target.current(allow_none=False)
      do_tensorize = (
          "+dotprod" in target.mattr
          or target.supports_integer_dot_product
      )
  
      if do_tensorize:
          dtypes = (data.dtype, weight.dtype)
          s[CC].tensorize(ki, dp4a("shared", "shared", "local", dtypes))
      by, vy, ty, yi = cfg["tile_y"].apply(s, output, n)
      bx, vx, tx, xi = cfg["tile_x"].apply(s, output, x)
  
      s[output].reorder(by, bx, vy, vx, ty, tx, yi, xi)
      s[output].bind(by, te.thread_axis("blockIdx.y"))
      s[output].bind(bx, te.thread_axis("blockIdx.x"))
      s[output].bind(vy, te.thread_axis("vthread"))
      s[output].bind(vx, te.thread_axis("vthread"))
      s[output].bind(ty, te.thread_axis("threadIdx.y"))
      s[output].bind(tx, te.thread_axis("threadIdx.x"))
      n_ty = cfg["tile_y"].size[2]
      n_tx = cfg["tile_x"].size[2]
  
      s[CC].compute_at(s[output], tx)
      yo, xo = CC.op.axis[:2]
      s[CC].reorder(ko, kt, yo, xo, ki)
  
      for load in [AA, WW]:
          s[load].compute_at(s[CC], ko)
  
          outer, inner = s[load].split(
              s[load].op.axis[-1], factor=in_dim_factor
          )
          s[load].vectorize(inner)
          fused = s[load].op.axis[:-1] + [outer]
          fused = s[load].fuse(*fused)
  
          fused, tx = s[load].split(fused, factor=n_tx)
          fused, ty = s[load].split(fused, factor=n_ty)
          s[load].bind(tx, te.thread_axis("threadIdx.x"))
          s[load].bind(ty, te.thread_axis("threadIdx.y"))
  
      s[output].pragma(
          kernel_scope,
          "auto_unroll_max_step",
          cfg["auto_unroll_max_step"].val,
      )
      s[output].pragma(kernel_scope, "unroll_explicit", False)
      return s
  
file_path: tvm/python/tvm/topi/cuda/dense.py
imports: import logging
  import tvm
  from tvm import te, autotvm
  from tvm.contrib import cublas
  from .tensor_intrin import dp4a
  from ..None import tag
  from ..None import generic
  from ..utils import traverse_inline, get_const_tuple
label_desc: Use image node_id 8 for calling a global function with example usage: _schedule_dense_int8(cfg, s, output) and returns: s