code: def conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding):
      assert N == 1, "Only consider batch_size = 1 in this template"
  
      data = te.placeholder((N, CI, H, W), name="data")
      kernel = te.placeholder((CO, CI, KH, KW), name="kernel")
      conv = topi.nn.conv2d_nchw(
          data, kernel, stride, padding, dilation=1, out_dtype="float32"
      )
      s = te.create_schedule([conv.op])
  
      ##### space definition begin #####
      n, f, y, x = s[conv].op.axis
      rc, ry, rx = s[conv].op.reduce_axis
  
      cfg = autotvm.get_config()
      cfg.define_split("tile_f", f, num_outputs=4)
      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=3)
      cfg.define_split("tile_ry", ry, num_outputs=3)
      cfg.define_split("tile_rx", rx, num_outputs=3)
      cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
      cfg.define_knob("unroll_explicit", [0, 1])
      ##### space definition end #####
  
      # inline padding
      pad_data = s[conv].op.input_tensors[0]
      s[pad_data].compute_inline()
      data, raw_data = pad_data, data
  
      output = conv
      OL = s.cache_write(conv, "local")
  
      # create cache stage
      AA = s.cache_read(data, "shared", [OL])
      WW = s.cache_read(kernel, "shared", [OL])
      AL = s.cache_read(AA, "local", [OL])
      WL = s.cache_read(WW, "local", [OL])
  
      # tile and bind spatial axes
      n, f, y, x = s[output].op.axis
      bf, vf, tf, fi = cfg["tile_f"].apply(s, output, f)
      by, vy, ty, yi = cfg["tile_y"].apply(s, output, y)
      bx, vx, tx, xi = cfg["tile_x"].apply(s, output, x)
      kernel_scope = n  # this is the scope to attach global config inside this kernel
  
      s[output].bind(bf, te.thread_axis("blockIdx.z"))
      s[output].bind(by, te.thread_axis("blockIdx.y"))
      s[output].bind(bx, te.thread_axis("blockIdx.x"))
      s[output].bind(vf, te.thread_axis("vthread"))
      s[output].bind(vy, te.thread_axis("vthread"))
      s[output].bind(vx, te.thread_axis("vthread"))
      s[output].bind(tf, te.thread_axis("threadIdx.z"))
      s[output].bind(ty, te.thread_axis("threadIdx.y"))
      s[output].bind(tx, te.thread_axis("threadIdx.x"))
      s[output].reorder(
          n, bf, by, bx, vf, vy, vx, tf, ty, tx, fi, yi, xi
      )
      s[OL].compute_at(s[output], tx)
  
      # tile reduction axes
      n, f, y, x = s[OL].op.axis
      rc, ry, rx = s[OL].op.reduce_axis
      rco, rcm, rci = cfg["tile_rc"].apply(s, OL, rc)
      ryo, rym, ryi = cfg["tile_rx"].apply(s, OL, ry)
      rxo, rxm, rxi = cfg["tile_ry"].apply(s, OL, rx)
      s[OL].reorder(
          rco, ryo, rxo, rcm, rym, rxm, rci, ryi, rxi, n, f, y, x
      )
  
      s[AA].compute_at(s[OL], rxo)
      s[WW].compute_at(s[OL], rxo)
      s[AL].compute_at(s[OL], rxm)
      s[WL].compute_at(s[OL], rxm)
  
      # cooperative fetching
      for load in [AA, WW]:
          n, f, y, x = s[load].op.axis
          fused = s[load].fuse(n, f, y, x)
          tz, fused = s[load].split(fused, nparts=cfg["tile_f"].size[2])
          ty, fused = s[load].split(fused, nparts=cfg["tile_y"].size[2])
          tx, fused = s[load].split(fused, nparts=cfg["tile_x"].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"))
  
      # tune unroll
      s[output].pragma(
          kernel_scope,
          "auto_unroll_max_step",
          cfg["auto_unroll_max_step"].val,
      )
      s[output].pragma(
          kernel_scope, "unroll_explicit", cfg["unroll_explicit"].val
      )
  
      return s, [raw_data, kernel, conv]
  
file_path: tvm/gallery/how_to/tune_with_autotvm/tune_conv2d_cuda.py
imports: import logging
  import sys
  import numpy
  import tvm
  from tvm import te, topi, testing
  from tvm.topi.testing import conv2d_nchw_python
  import tvm.testing
  from tvm import autotvm
label_desc: Use image node_id 1 for calling a global function with example usage: conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding) and returns: s,  list