code: def schedule_depthwise_conv2d_NHWC_HWOI(cfg, s, output):
      """
      schedule optimized for batch size = 1
  
      Algo:
      1. Split output axis to three parts: global work size, vthread, local worksize.
         The limitations for tuning includes heuristics from some tuned networks to limit
         search space and not pay much time for useles configurations.
      2. In case of 4d convolution schedule copying of the input (and filter) into
        5d tensors
      3. For depthwise convolution it's better to inline pad into the conv2d compute, the
         divergence in opencl kernel will not so significant as for regular conv2d.
      4. For 5d convolution we schedule the latest op with binding 5d axis and vectorize
         for textures
         For 4d tensor we are doing the same for the latest blocked stage, i.e. conversion
         of data type
      5. In case of 4d conv we need to schedule postops as well
      """
      latest = s.outputs[0].output(0)
      if len(latest.op.axis) == 4:
          latest_blocked = dummy = output.op.input_tensors[0]
          conv = dummy.op.input_tensors[0]
      else:
          conv = output.op.input_tensors[0]
          latest_blocked = latest
  
      ##### space definition begin #####
      n, y, x, fc, fb = s[conv].op.axis
      ry, rx = s[conv].op.reduce_axis
      cfg.define_split("tile_fc", fc, num_outputs=3)
      cfg.define_split("tile_y", y, num_outputs=3)
      cfg.define_split("tile_x", x, num_outputs=3)
      cfg.define_split("tile_ry", ry, num_outputs=2)
      cfg.define_split("tile_rx", rx, num_outputs=2)
      cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
      cfg.define_knob("unroll_explicit", [0, 1])
  
      cfg.multi_filter(
          filter=lambda entity: (  # pylint: disable=chained-comparison
              entity["tile_fc"].size[1]
              * entity["tile_y"].size[1]
              * entity["tile_x"].size[1]
          )
          <= 32
          and 32
          <= (
              entity["tile_fc"].size[2]
              * entity["tile_y"].size[2]
              * entity["tile_x"].size[2]
          )
          < 1024
      )
      if cfg.is_fallback:
          get_default_conv2d_config(
              cfg, conv.shape[3], conv.shape[1], conv.shape[2]
          )
      ##### space definition end #####
  
      pad_data, kernel = s[conv].op.input_tensors
      if (
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "filter_pack" in kernel.op.tag
      ):  # len(latest.op.axis) == 4:
          # manage scheduling of datacopy
          pad_data, kernel = s[conv].op.input_tensors
          if "pad_temp" in pad_data.op.name:
              pack_data = pad_data.op.input_tensors[0]
              bind_data_copy(s[pack_data])
          else:
              bind_data_copy(s[pad_data])
          bind_data_copy(s[kernel])
  
      pad_data, kernel = s[conv].op.input_tensors
  
      if "pad_temp" in pad_data.op.name:
          s[pad_data].compute_inline()
  
      s[conv].set_scope("local")
      if latest_blocked == latest and output != latest:
          s[output].compute_inline()
  
      if autotvm.GLOBAL_SCOPE.in_tuning or len(latest.op.axis) == 4:
          # create cache stage for tuning only or in case of 4d case
          AT = s.cache_read(
              pad_data, get_texture_storage(pad_data.shape), [conv]
          )
          bind_data_copy(s[AT])
          if kernel.shape[0] == 1 and kernel.shape[1] == 1:
              WT = s.cache_read(
                  kernel, get_texture_storage(kernel.shape), [conv]
              )
              bind_data_copy(s[WT])
  
      # tile and bind spatial axes
      n, y, x, fc, fb = s[latest_blocked].op.axis
      kernel_scope, n = s[latest_blocked].split(n, nparts=1)
  
      bf, vf, tf = cfg["tile_fc"].apply(s, latest_blocked, fc)
      by, vy, ty = cfg["tile_y"].apply(s, latest_blocked, y)
      bx, vx, tx = cfg["tile_x"].apply(s, latest_blocked, x)
  
      by = s[latest_blocked].fuse(n, by)
      s[latest_blocked].bind(bf, te.thread_axis("blockIdx.z"))
      s[latest_blocked].bind(by, te.thread_axis("blockIdx.y"))
      s[latest_blocked].bind(bx, te.thread_axis("blockIdx.x"))
      s[latest_blocked].bind(vf, te.thread_axis("vthread"))
      s[latest_blocked].bind(vy, te.thread_axis("vthread"))
      s[latest_blocked].bind(vx, te.thread_axis("vthread"))
      s[latest_blocked].bind(tf, te.thread_axis("threadIdx.z"))
      s[latest_blocked].bind(ty, te.thread_axis("threadIdx.y"))
      s[latest_blocked].bind(tx, te.thread_axis("threadIdx.x"))
      s[latest_blocked].reorder(bf, by, bx, vf, vy, vx, tf, ty, tx, fb)
      s[latest_blocked].vectorize(fb)
  
      s[conv].compute_at(s[latest_blocked], tx)
  
      # tile reduction axes
      n, y, x, fc, fb = s[conv].op.axis
  
      ry, rx = s[conv].op.reduce_axis
      ryo, ryi = cfg["tile_ry"].apply(s, conv, ry)
      rxo, rxi = cfg["tile_rx"].apply(s, conv, rx)
  
      s[conv].reorder(ryo, rxo, ryi, rxi, n, fc, y, x, fb)
      s[conv].vectorize(fb)
  
      # unroll
      s[latest_blocked].pragma(
          kernel_scope,
          "auto_unroll_max_step",
          cfg["auto_unroll_max_step"].val,
      )
      s[latest_blocked].pragma(
          kernel_scope, "unroll_explicit", cfg["unroll_explicit"].val
      )
      if latest_blocked != latest:
          s[latest].compute_root()
          bind_data_copy(s[latest], 1)
          if latest != output:
              s[output].compute_inline()
  
      N, OH, OW, OCC, OCB = get_const_tuple(latest_blocked.shape)
      KH, KW, _, _, _ = get_const_tuple(kernel.shape)
      KHKW = KH * KW
  
      if isinstance(N, int):
          cfg.add_flop(2 * N * OH * OW * OCC * OCB * KHKW)
  
file_path: tvm/python/tvm/topi/adreno/depthwise_conv2d_nhwc.py
imports: import tvm
  from tvm import te
  from tvm import autotvm
  from ..utils import get_const_tuple, traverse_inline
  from .utils import split_to_chunks, pack_input, pack_filter, expand_spatial_dimensions, add_pad, bind_data_copy, get_texture_storage, get_default_conv2d_config
label_desc: Use image node_id 3 for calling a global function with example usage: schedule_depthwise_conv2d_NHWC_HWOI(cfg, s, output) without return types