code: def depthwise_conv2d_nchwc(
      cfg, Input, Filter, stride, padding, dilation, out_dtype
  ):
      """
      Depthwise convolution operator in NCHWc layout.
      Algo:
        1. Convert into blocked format if we have 4d original tensor.
           In case of AutoTVM we override the convert by just tensors since such conversion
           will be absent for real blocked convolution, no sense to include into tuning
        2. Expand spatial dimensions to have width and height be dividable by factor 4
           This leads to slightly bigger amount of compute but allow utilize GPU much better
        3. Add paddings. This happens even if we do not need pad originaly. This is useful
           due to work arounding of the gaps of texture annotation between Primary Functions
           and limited support of textures in schedules. Later on this pad will be executed
           separately and will produce texture
        4. 5d Convolution compute with accumulating into out_dtype
        5. Cast to the origin output data type
        6. For case of 4d convolution: convert of output from 5d to 4d
      """
      if out_dtype is None:
          out_dtype = Input.dtype
      assert isinstance(stride, int) or len(stride) == 2
      assert isinstance(dilation, int) or len(dilation) == 2
  
      if isinstance(stride, int):
          stride_h = stride_w = stride
      else:
          stride_h, stride_w = stride
  
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
  
      convert_from4d = False
      if len(Input.shape) == 4:
          batch, in_channels, in_height, in_width = Input.shape
          (
              out_channles,
              in_filter_channels,
              kernel_h,
              kernel_w,
          ) = Filter.shape
  
          (
              in_channel_chunks,
              in_channel_block,
              in_channel_tail,
          ) = split_to_chunks(in_channels, 4)
          (
              out_channel_chunks,
              out_channel_block,
              out_channel_tail,
          ) = split_to_chunks(out_channles, 4)
  
          if autotvm.GLOBAL_SCOPE.in_tuning:
              dshape = (
                  batch,
                  in_channel_chunks,
                  in_height,
                  in_width,
                  in_channel_block,
              )
              Input = tvm.te.placeholder(
                  dshape, Input.dtype, name="data_placeholder"
              )
              kshape = (
                  out_channel_chunks,
                  in_filter_channels,
                  kernel_h,
                  kernel_w,
                  out_channel_block,
              )
              Filter = tvm.te.placeholder(
                  kshape, Filter.dtype, name="kernel_placeholder"
              )
          else:
              convert_from4d = True
              Input = pack_input(
                  Input,
                  "NCHW",
                  batch,
                  in_channel_chunks,
                  in_channel_block,
                  in_channel_tail,
                  in_height,
                  in_width,
              )
              Filter = pack_filter(
                  Filter,
                  "OIHW",
                  out_channel_chunks,
                  out_channel_block,
                  out_channel_tail,
                  in_filter_channels,
                  in_channel_chunks,
                  in_channel_block,
                  in_channel_tail,
                  kernel_h,
                  kernel_w,
              )
  
      else:
          (
              batch,
              in_channel_chunks,
              in_height,
              in_width,
              in_channel_block,
          ) = Input.shape
          (
              out_channel_chunks,
              in_filter_channels,
              kernel_h,
              kernel_w,
              out_channel_block,
          ) = Filter.shape
  
      (
          out_height_orig,
          out_height,
          out_width_orig,
          out_width,
      ) = expand_spatial_dimensions(
          in_height,
          in_width,
          kernel_h,
          kernel_w,
          dilation_h,
          dilation_w,
          padding,
          stride_h,
          stride_w,
      )
  
      temp = add_pad(
          Input,
          "NCHW",
          out_height_orig,
          out_width_orig,
          kernel_h,
          kernel_w,
          dilation_h,
          dilation_w,
          padding,
          stride_h,
          stride_w,
      )
  
      ry = te.reduce_axis((0, kernel_h), name="ry")
      rx = te.reduce_axis((0, kernel_w), name="rx")
      conv = te.compute(
          (
              batch,
              out_channel_chunks,
              out_height,
              out_width,
              out_channel_block,
          ),
          lambda nn, ffc, yy, xx, ffb: te.sum(
              (
                  temp[
                      nn,
                      ffc // in_filter_channels,
                      yy * stride_h + ry * dilation_h,
                      xx * stride_w + rx * dilation_w,
                      ffb,
                  ]
                  * Filter[
                      ffc // in_filter_channels,
                      ffc % in_filter_channels,
                      ry,
                      rx,
                      ffb,
                  ]
              ).astype(out_dtype),
              axis=[ry, rx],
          ),
          tag="depthwise_conv2d_nchwc",
      )
  
      if convert_from4d and not autotvm.GLOBAL_SCOPE.in_tuning:
          dummy_cast = te.compute(
              (
                  batch,
                  out_channel_chunks,
                  out_height_orig,
                  out_width_orig,
                  out_channel_block,
              ),
              lambda n, fc, y, x, fb: conv[n, fc, y, x, fb].astype(
                  out_dtype
              ),
              tag="dummy_cast",
          )
          return te.compute(
              (batch, out_channles, out_height_orig, out_width_orig),
              lambda n, c, y, x: dummy_cast[
                  n, c // out_channel_block, y, x, c % out_channel_block
              ],
              tag="adreno_dw_conv2d_latest_op",
          )
      else:
          return te.compute(
              (
                  batch,
                  out_channel_chunks,
                  out_height_orig,
                  out_width_orig,
                  out_channel_block,
              ),
              lambda n, ffc, y, x, ffb: conv[n, ffc, y, x, ffb].astype(
                  out_dtype
              ),
              tag="adreno_dw_conv2d_latest_op",
          )
  
file_path: tvm/python/tvm/topi/adreno/depthwise_conv2d_nchw.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 2 for calling a global function with example usage: depthwise_conv2d_nchwc(cfg, Input, Filter, stride, padding, dilation, out_dtype) and returns: te, te