code: def _schedule_cl_spatialpack(s, op):
      output = op.output(0)
      _, _, out_height, out_width = [
          utils.get_const_int(x) for x in output.shape
      ]
  
      conv = op.input_tensors[0]
      temp = s[conv].op.input_tensors[0]
      kernel_vec = s[conv].op.input_tensors[1]
      kernel = s[kernel_vec].op.input_tensors[0]
      temp_W = s.cache_read(temp, "shared", [conv])
      conv_L = s.cache_write(conv, "local")
  
      kernel_L = s.cache_read(kernel_vec, "local", [conv_L])
      _, in_channel, temp_h, temp_w = [
          utils.get_const_int(x) for x in temp.shape
      ]
  
      attrs = s[conv].op.attrs
      OUTPUT_BLOCK_HEIGHT = attrs["block_h"]
      OUTPUT_BLOCK_WIDTH = attrs["block_w"]
  
      # schedule conv
      y_factor = z_factor = 1
      x_factor = 16
      thread_z = te.thread_axis((0, z_factor), "threadIdx.z")
      thread_y = te.thread_axis((0, y_factor), "threadIdx.y")
      thread_x = te.thread_axis((0, x_factor), "threadIdx.x")
      _, co, oh, ow, vc = s[conv].op.axis
      ooh, ioh = s[conv].split(oh, factor=OUTPUT_BLOCK_HEIGHT)
      oow, iow = s[conv].split(ow, factor=OUTPUT_BLOCK_WIDTH)
      s[conv].reorder(_, co, ooh, oow, vc, ioh, iow)
      coo, coi = s[conv].split(co, nparts=1)
      ooho, oohi = s[conv].split(ooh, factor=z_factor)
      oowo, oowi = s[conv].split(oow, factor=y_factor)
      vco, vci = s[conv].split(vc, factor=x_factor)
      s[conv].reorder(
          _, coo, vco, ooho, oowo, coi, oohi, oowi, vci, ioh, iow
      )
      s[conv].bind(oohi, thread_z)
      s[conv].bind(oowi, thread_y)
      s[conv].bind(vci, thread_x)
      s[conv].bind(ooho, te.thread_axis("blockIdx.z"))
      s[conv].bind(oowo, te.thread_axis("blockIdx.y"))
      s[conv].bind(coi, te.thread_axis("blockIdx.x"))
  
      # schedule conv_L
      s[conv_L].compute_at(s[conv], vci)
      i, oc, h, w, vc = s[conv_L].op.axis
      rc, ry, rx = s[conv_L].op.reduce_axis
      s[conv_L].reorder(i, oc, rc, ry, rx, vc, h, w)
      s[temp_W].compute_at(s[conv_L], rc)
      if kernel.shape[3].value != 7:
          s[conv_L].unroll(ry)
          s[conv_L].unroll(rx)
  
      # schedule temp
      _, ci, h, w = s[temp].op.axis
      tile_and_bind3d(s, temp, ci, h, w, 1, 16, 16)
  
      # schedule temp_W
      _, ci, h, w = s[temp_W].op.axis
      zo, zi = s[temp_W].split(ci, 1)
      yo, yi = s[temp_W].split(h, 1)
      xo, xi = s[temp_W].split(w, 16)
      s[temp_W].reorder(zo, yo, xo, zi, yi, xi)
      s[temp_W].bind(zi, thread_z)
      s[temp_W].bind(yi, thread_y)
      s[temp_W].bind(xi, thread_x)
      s[temp_W].storage_align(s[temp_W].op.axis[2], 16, 0)
  
      s[kernel_vec].compute_inline()
  
      # schedule kernel_L
      if OUTPUT_BLOCK_HEIGHT == 2 and OUTPUT_BLOCK_WIDTH == 14:
          s[kernel_L].compute_at(s[conv_L], ry)
      else:
          s[kernel_L].compute_at(s[conv_L], rx)
  
      # schedule output
      if output.op in s.outputs:
          out = output
      else:
          s[output].compute_inline()
          out = s.outputs[0]
  
      _, co, h, w = s[out].op.axis
      tile_and_bind3d(s, out, w, h, co, 4, 8, 8)
  
file_path: tvm/python/tvm/topi/intel_graphics/conv2d.py
imports: from __future__ import absolute_import
  import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.autotvm.task.space import SplitEntity, OtherOptionEntity
  from ..None import nn
  from ..None import utils
  from ..utils import simplify, get_const_tuple, traverse_inline
label_desc: Use image node_id 11 for calling a global function with example usage: _schedule_cl_spatialpack(s, op) without return types