code: def conv2d_NCHWc(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      layout,
      out_layout,
      out_dtype="float32",
  ):
      """Conv2D operator for Intel Graphics backend.
  
      Parameters
      ----------
      data : tvm.te.Tensor
          4-D with shape [batch, in_channel, in_height, in_width]
  
      kernel : tvm.te.Tensor
          5-D with shape [num_filter, in_channel, filter_height, filter_width, nnum_filter_vec]
  
      stride : int or a list/tuple of two ints
          stride size, or [stride_height, stride_width]
  
      padding : int or a list/tuple of two ints
          padding size, or [pad_height, pad_width]
  
      layout : str
          layout of data
  
      Returns
      -------
      output : tvm.te.Tensor
          4-D with shape [batch, out_channel, out_height, out_width]
      """
      if len(data.shape) == 5:
          batch, ic_chunk, ih, iw, ic_bn = get_const_tuple(data.shape)
          (
              oc_chunk,
              _,
              kernel_height,
              kernel_width,
              _,
              oc_bn,
          ) = get_const_tuple(kernel.shape)
          in_channel = ic_chunk * ic_bn
          num_filter = oc_chunk * oc_bn
      else:
          batch, in_channel, ih, iw = get_const_tuple(data.shape)
          num_filter, _, kernel_height, kernel_width = get_const_tuple(
              kernel.shape
          )
  
      dh, dw = (
          dilation
          if isinstance(dilation, (tuple, list))
          else (dilation, dilation)
      )
      pad_top, pad_left, pad_down, pad_right = nn.get_pad_tuple(
          padding, (kernel_height, kernel_width)
      )
      assert (dh, dw) == (1, 1), "Does not support dilation"
      if isinstance(strides, (tuple, list)):
          stride_h, stride_w = strides
      else:
          stride_h, stride_w = strides, strides
  
      data_shape = (batch, in_channel, ih, iw)
      kernel_shape = (
          num_filter,
          in_channel,
          kernel_height,
          kernel_width,
      )
      _create_schedule_template(
          cfg, data_shape, kernel_shape, strides, padding, dilation
      )
  
      if cfg.is_fallback:
          _get_default_config(
              cfg,
              te.placeholder(
                  (batch, in_channel, ih, iw), dtype=data.dtype
              ),
              te.placeholder(
                  (num_filter, in_channel, kernel_height, kernel_width),
                  dtype=kernel.dtype,
              ),
              strides,
              padding,
              out_dtype,
          )
  
      ic_bn = (
          cfg["tile_ic"].val
          if hasattr(cfg["tile_ic"], "val")
          else cfg["tile_ic"].size[-1]
      )
      oc_bn = (
          cfg["tile_oc"].val
          if hasattr(cfg["tile_oc"], "val")
          else cfg["tile_oc"].size[-1]
      )
  
      # Pack data if raw 4-D data is provided.
      if len(data.shape) == 4:
          data, kernel = _pack_data(data, kernel, ic_bn, oc_bn)
  
      out_channel = num_filter
      out_height = simplify(
          (ih - kernel_height + pad_top + pad_down) // stride_h + 1
      )
      out_width = simplify(
          (iw - kernel_width + pad_left + pad_right) // stride_w + 1
      )
      oshape = (
          batch,
          out_channel // oc_bn,
          out_height,
          out_width,
          oc_bn,
      )
  
      rc = te.reduce_axis((0, in_channel), name="rc")
      ry = te.reduce_axis((0, kernel_height), name="ry")
      rx = te.reduce_axis((0, kernel_width), name="rx")
  
      block_h = cfg["block_oh"].val
      block_w = cfg["block_ow"].val
  
      c_h = out_height
      c_w = out_width
  
      if out_height % block_h != 0:
          c_h = (out_height // block_h + 1) * block_h
  
      if out_width % block_w != 0:
          c_w = (out_width // block_w + 1) * block_w
  
      cshape = (batch, out_channel // oc_bn, c_h, c_w, oc_bn)
  
      pad_before = [0, 0, pad_top, pad_left, 0]
      pad_after = [
          0,
          0,
          pad_down + c_h - out_height,
          pad_right + c_w - out_width,
          0,
      ]
      DOPAD = (
          pad_top != 0
          or pad_left != 0
          or pad_down + c_h - out_height != 0
          or pad_right + c_w - out_width != 0
      )
      DOUNPACK = c_h - out_height != 0 or c_w - out_width != 0
      if DOPAD:
          temp = nn.pad(data, pad_before, pad_after, name="pad_temp")
      else:
          temp = data
  
      conv = te.compute(
          cshape,
          lambda nn, ff, yy, xx, ff_v: te.sum(
              temp[
                  nn,
                  rc // ic_bn,
                  yy * stride_h + ry,
                  xx * stride_w + rx,
                  rc % ic_bn,
              ].astype(out_dtype)
              * kernel[
                  ff, rc // ic_bn, ry, rx, rc % ic_bn, ff_v
              ].astype(out_dtype),
              axis=[rc, ry, rx],
          ),
          tag="conv2d_NCHWc",
          name="conv2d_NCHWc",
      )
  
      if DOUNPACK:
          output = te.compute(
              oshape,
              lambda nn, ff, yy, xx, ff_v: conv[nn][ff][yy][xx][ff_v],
              name="output_unpack",
              tag="conv2d_NCHWc_unpack",
          )
      else:
          output = conv
  
      return output
  
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 5 for calling a global function with example usage: conv2d_NCHWc(cfg, data, kernel, strides, padding, dilation, layout, out_layout, out_dtype) and returns: output