code: def depthwise_conv2d_NCHWc(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      layout,
      out_layout,
      out_dtype=None,
  ):
      """Compute depthwise conv2d with NCHWc layout"""
      out_dtype = data.dtype if out_dtype is None else out_dtype
  
      if len(data.shape) == 5:
          (
              batch,
              in_channel_chunk,
              in_height,
              in_width,
              in_channel_block,
          ) = get_const_tuple(data.shape)
          (
              out_channel_chunk,
              cm_chunk,
              filter_height,
              filter_width,
              cm_block,
              out_channel_block,
          ) = get_const_tuple(kernel.shape)
          in_channel = in_channel_chunk * in_channel_block
          out_channel = out_channel_chunk * out_channel_block
          channel_multiplier = cm_chunk * cm_block
          assert channel_multiplier * in_channel == out_channel
      else:
          batch, in_channel, in_height, in_width = get_const_tuple(
              data.shape
          )
          (
              out_channel,
              channel_multiplier,
              filter_height,
              filter_width,
          ) = get_const_tuple(kernel.shape)
      assert channel_multiplier == 1
  
      strides = (
          strides
          if isinstance(strides, (tuple, list))
          else (strides, strides)
      )
      HSTR, WSTR = strides
  
      dh, dw = (
          dilation
          if isinstance(dilation, (tuple, list))
          else (dilation, dilation)
      )
  
      dilated_kernel_h = (filter_height - 1) * dh + 1
      dilated_kernel_w = (filter_width - 1) * dw + 1
      pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
          padding, (dilated_kernel_h, dilated_kernel_w)
      )
      HPAD = pad_top + pad_down
      WPAD = pad_left + pad_right
  
      out_height = (in_height + HPAD - dilated_kernel_h) // HSTR + 1
      out_width = (in_width + WPAD - dilated_kernel_w) // WSTR + 1
  
      cfg.define_split("tile_ic", in_channel, num_outputs=2)
      cfg.define_split("tile_oc", out_channel, num_outputs=2)
      cfg.define_split(
          "tile_ow",
          out_width,
          num_outputs=2,
          filter=lambda y: y.size[-1] <= 64,
      )
      cfg.define_knob("unroll_kw", [True, False])
  
      # get workload and related schedule config
      wkl = _get_workload(
          te.placeholder(
              (batch, in_channel, in_height, in_width), dtype=data.dtype
          ),
          te.placeholder(
              (
                  out_channel,
                  channel_multiplier,
                  filter_height,
                  filter_width,
              ),
              dtype=kernel.dtype,
          ),
          strides,
          (pad_top, pad_down),
          dilation,
          out_dtype,
      )
      if cfg.is_fallback:
          _fallback_schedule(cfg, wkl)
  
      # Pack data if raw 4-D data is provided.
      # This can only happen when autotuning.
      if len(data.shape) == 4:
          if autotvm.GLOBAL_SCOPE.in_tuning:
              # Directly use modified data layout placeholder.
              in_channel_block = cfg["tile_ic"].size[-1]
              in_channel_chunk = in_channel // in_channel_block
              out_channel_block = cfg["tile_oc"].size[-1]
              out_channel_chunk = out_channel // out_channel_block
              dshape = (
                  batch,
                  in_channel_chunk,
                  in_height,
                  in_width,
                  in_channel_block,
              )
              data = tvm.te.placeholder(dshape, data.dtype, name="data")
              kshape = (
                  out_channel_chunk,
                  1,
                  filter_height,
                  filter_width,
                  1,
                  out_channel_block,
              )
              kernel = tvm.te.placeholder(
                  kshape, kernel.dtype, name="kernel"
              )
          else:
              data, kernel = _pack_data(cfg, data, kernel)
              _, _, _, _, in_channel_block = get_const_tuple(data.shape)
              (
                  out_channel_chunk,
                  _,
                  _,
                  _,
                  _,
                  out_channel_block,
              ) = get_const_tuple(kernel.shape)
  
      # padding stage
      DOPAD = (
          pad_top != 0
          or pad_left != 0
          or pad_down != 0
          or pad_right != 0
      )
      if DOPAD:
          pad_before = [0, 0, pad_top, pad_left, 0]
          pad_after = [0, 0, pad_down, pad_right, 0]
          data_pad = pad(
              data, pad_before, pad_after, name="PaddedInput"
          )
      else:
          data_pad = data
  
      # depthconv stage
      idxdiv = tvm.tir.indexdiv
      idxmod = tvm.tir.indexmod
  
      kh = te.reduce_axis((0, filter_height), name="kh")
      kw = te.reduce_axis((0, filter_width), name="kw")
      Output = te.compute(
          (
              batch,
              out_channel_chunk,
              out_height,
              out_width,
              out_channel_block,
          ),
          lambda b, oco, oh, ow, oci: te.sum(
              (
                  data_pad[
                      b,
                      idxdiv(
                          idxdiv(
                              oco * out_channel_block + oci,
                              channel_multiplier,
                          ),
                          in_channel_block,
                      ),
                      oh * HSTR + kh * dh,
                      ow * WSTR + kw * dw,
                      idxmod(
                          idxdiv(
                              oco * out_channel_block + oci,
                              channel_multiplier,
                          ),
                          in_channel_block,
                      ),
                  ].astype(out_dtype)
                  * kernel[oco, 0, kh, kw, 0, oci].astype(out_dtype)
              ),
              axis=[kh, kw],
          ),
          name="DepthwiseConv2d",
          tag="depthwise_conv2d_NCHWc",
      )
      return Output
  
file_path: tvm/python/tvm/topi/x86/depthwise_conv2d.py
imports: import tvm
  from tvm import autotvm, te
  from tvm.autotvm.task.space import OtherOptionEntity, SplitEntity
  from tvm.target.x86 import get_simd_32bit_lanes
  from ..nn.conv2d import unpack_NCHWc_to_nchw
  from ..nn.depthwise_conv2d import _get_workload, depthwise_conv2d_infer_layout
  from ..nn.pad import pad
  from ..nn.utils import get_pad_tuple
  from ..utils import get_const_tuple, traverse_inline
label_desc: Use image node_id 5 for calling a global function with example usage: depthwise_conv2d_NCHWc(cfg, data, kernel, strides, padding, dilation, layout, out_layout, out_dtype) and returns: Output