code: def schedule_conv2d_NHWC(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
      4. pad should be scheduled separately to create independent opencl kernel. If pad is
         inlined into convolution, this gives 1.5x performance drop
      5. We are using cache_read for intermediate tensors to produce texture and guarantee
         the best performance on the next stage.
         The weights are managed through static texture planning mechanism and guarantied come
         in texture memory scope.
         Thus way we are calling cache_read only for data tensor
      6. 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
      7. 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
  
      pad_data, kernel = s[conv].op.input_tensors
      filter_pack_rt = bool(
          isinstance(kernel.op, tvm.te.ComputeOp)
          and "filter_pack" in kernel.op.tag
      )
  
      if "pad_temp" in pad_data.op.name:
          input_pad_temp = pad_data.op.input_tensors[0]
      else:
          input_pad_temp = pad_data
  
      input_pack_rt = bool(
          isinstance(input_pad_temp.op, tvm.te.ComputeOp)
          and "input_pack" in input_pad_temp.op.tag
      )
      ##### space definition begin #####
      n, y, x, fc, fb = s[conv].op.axis
      ry, rx, rcc, rcb = s[conv].op.reduce_axis
  
      if conv.shape[3] % 2 == 0:
          min_threads_div = 2
      else:
          min_threads_div = 1
  
      cfg.define_split(
          "tile_fc",
          fc,
          num_outputs=3,
          filter=lambda entity: entity.size[1] <= 8
          and entity.size[2] >= min_threads_div
          and entity.size[2] < 256,
      )
      cfg.define_split(
          "tile_y",
          y,
          num_outputs=3,
          filter=lambda entity: entity.size[1] <= 8
          and entity.size[2] <= 16,
      )
      cfg.define_split(
          "tile_x",
          x,
          num_outputs=3,
          filter=lambda entity: entity.size[1] <= 8
          and entity.size[2] <= 16,
      )
  
      cfg.define_split("tile_rcc", rcc, num_outputs=2)
      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]
          )
          <= 24
          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
      # There are several conditions that have to be handled:
      # 1. If we are in the tuning, we always add cache read for data to main conv kernel
      #    to get texture in tuning opencl kernel
      # 2. If we are repacking input in runtime, we should always explicit schedule this one more
      #    stage of data copy from 4d to 5d (referred as pack_data).
      # 3. If we have pad (independently if we have runtime repack or not) we should inline it in the
      #    cache_read("texture")
      if autotvm.GLOBAL_SCOPE.in_tuning or input_pack_rt:
          if autotvm.GLOBAL_SCOPE.in_tuning:
              if "pad_temp" in pad_data.op.name:
                  s[pad_data].compute_inline()
          else:
              if "pad_temp" in pad_data.op.name:
                  s[pad_data].compute_inline()
                  pack_data = pad_data.op.input_tensors[0]
                  bind_data_copy(s[pack_data])
              else:
                  pack_data = pad_data
                  bind_data_copy(s[pack_data])
  
          AT = s.cache_read(
              pad_data, get_texture_storage(pad_data.shape), [conv]
          )
          bind_data_copy(s[AT])
      elif "pad_temp" in pad_data.op.name:
          s[pad_data].compute_inline()
          # create cache stage
          AT = s.cache_read(
              pad_data, get_texture_storage(pad_data.shape), [conv]
          )
          bind_data_copy(s[AT])
  
      if autotvm.GLOBAL_SCOPE.in_tuning or filter_pack_rt:
          if not autotvm.GLOBAL_SCOPE.in_tuning:
              bind_data_copy(s[kernel])
          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])
  
      s[conv].set_scope("local")
      if latest_blocked == latest and output != latest:
          s[output].compute_inline()
  
      # 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, rcc, rcb = s[conv].op.reduce_axis
      rco, rci = cfg["tile_rcc"].apply(s, conv, rcc)
      ryo, ryi = cfg["tile_ry"].apply(s, conv, ry)
      rxo, rxi = cfg["tile_rx"].apply(s, conv, rx)
  
      s[conv].reorder(
          rco, ryo, rxo, rci, ryi, rxi, rcb, n, fc, y, x, fb
      )
      s[conv].vectorize(fb)
      s[conv].unroll(rcb)
  
      # 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, IC, _, _ = get_const_tuple(kernel.shape)
      ICKHKW = IC * KH * KW
  
      if isinstance(N, int):
          cfg.add_flop(2 * N * OH * OW * OCC * OCB * ICKHKW)
  
file_path: tvm/python/tvm/topi/adreno/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_conv2d_NHWC(cfg, s, output) without return types