code: def conv2d_NCHWc(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      layout,
      out_layout,
      out_dtype,
  ):
      """Compute conv2d with NCHWc layout."""
      # layout and out_layout are not used here,
      # we keep them for debug convenience when dumping autotvm workload
      if len(data.shape) == 5:
          n, ic_chunk, ih, iw, ic_bn = get_const_tuple(data.shape)
          (
              oc_chunk,
              ic_chunk_group,
              kernel_height,
              kernel_width,
              _,
              oc_bn,
          ) = get_const_tuple(kernel.shape)
          in_channel = ic_chunk * ic_bn
          num_filter = oc_chunk * oc_bn
      else:
          n, in_channel, ih, iw = get_const_tuple(data.shape)
          num_filter, _, kernel_height, kernel_width = get_const_tuple(
              kernel.shape
          )
  
      # Define autotvm tuning space
      is_kernel_1x1 = kernel_height == 1 and kernel_width == 1
      pt, pl, pb, pr = get_pad_tuple(
          padding, (kernel_height, kernel_width)
      )
      sh, sw = (
          strides
          if isinstance(strides, (tuple, list))
          else (strides, strides)
      )
      oh = (ih - kernel_height + pt + pb) // sh + 1
      ow = (iw - kernel_width + pl + pr) // sw + 1
  
      cfg.define_split("tile_ic", in_channel, num_outputs=2)
      cfg.define_split("tile_oc", num_filter, num_outputs=2)
      if isinstance(ow, (tvm.tir.IntImm, int)):
          cfg.define_split(
              "tile_ow",
              ow,
              num_outputs=2,
              filter=lambda y: y.size[-1] <= 64,
              policy="verbose",
          )
      if is_kernel_1x1:
          if isinstance(oh, (tvm.tir.IntImm, int)):
              cfg.define_knob("tile_oh", [1, 2] if oh > 1 else [1])
      else:
          cfg.define_knob("unroll_kw", [True, False])
  
      # If no config was set, we can fallback to default config.
      if cfg.is_fallback:
          _get_default_config(
              cfg,
              te.placeholder((n, in_channel, ih, iw), dtype=data.dtype),
              te.placeholder(
                  (num_filter, in_channel, kernel_height, kernel_width),
                  dtype=kernel.dtype,
              ),
              strides,
              padding,
              dilation,
              out_dtype,
          )
  
      # 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.
              dshape = (
                  n,
                  in_channel // cfg["tile_ic"].size[-1],
                  ih,
                  iw,
                  cfg["tile_ic"].size[-1],
              )
              data = tvm.te.placeholder(dshape, data.dtype, name="data")
              kshape = (
                  num_filter // cfg["tile_oc"].size[-1],
                  in_channel // cfg["tile_ic"].size[-1],
                  kernel_height,
                  kernel_width,
                  cfg["tile_ic"].size[-1],
                  cfg["tile_oc"].size[-1],
              )
              kernel = tvm.te.placeholder(
                  kshape, kernel.dtype, name="kernel"
              )
          else:
              data, kernel = _pack_data(cfg, data, kernel)
  
      return nn.conv2d_NCHWc(
          data,
          kernel,
          strides,
          padding,
          dilation,
          layout,
          out_layout,
          out_dtype,
      )
  
imports: import logging
  import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.contrib import dnnl
  from ..None import nn
  from ..generic import schedule_extern
  from ..nn.conv2d import conv2d_infer_layout, _get_workload
  from ..nn.conv2d import unpack_NCHWc_to_nchw
  from ..nn.depthwise_conv2d import _get_workload
  from ..nn.utils import get_pad_tuple
  from ..utils import get_const_tuple, traverse_inline
  from .None import conv2d_avx_1x1, conv2d_avx_common
label_desc: Use image node_id 7 for calling a global function with example usage: conv2d_NCHWc(cfg, data, kernel, strides, padding, dilation, layout, out_layout, out_dtype) and returns: nn