code: def depthwise_conv2d_nhwc(
      Input,
      Filter,
      stride,
      padding,
      dilation,
      kernel_layout="HWOI",
      out_dtype=None,
  ):
      """Depthwise convolution nhwc forward operator.
  
      Parameters
      ----------
      Input : tvm.te.Tensor
          4-D with shape [batch, in_height, in_width, in_channel]
  
      Filter : tvm.te.Tensor
          4-D with shape [filter_height, filter_width, in_channel, channel_multiplier]
  
      stride : tuple of two ints
          The spatial stride along height and width
  
      padding : int or str
          Padding size, or ['VALID', 'SAME']
  
      dilation: int or a list/tuple of two ints
          dilation size, or [dilation_height, dilation_width]
  
      out_dtype: str, optional
          Output data type
  
      Returns
      -------
      Output : tvm.te.Tensor
          4-D with shape [batch, out_height, out_width, out_channel]
      """
      out_dtype = Input.dtype if out_dtype is None else out_dtype
  
      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
  
      batch, in_height, in_width, in_channel = Input.shape
  
      # shape of dilated kernel
      if kernel_layout == "HWIO":
          (
              filter_height,
              filter_width,
              channel_multiplier,
              filter_channel,
          ) = Filter.shape
          kernel_permutation = [0, 1, 3, 2]
      else:
          (
              filter_height,
              filter_width,
              filter_channel,
              channel_multiplier,
          ) = Filter.shape
          kernel_permutation = [0, 1, 2, 3]
  
      dilated_kernel_h = (filter_height - 1) * dilation_h + 1
      dilated_kernel_w = (filter_width - 1) * dilation_w + 1
      pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
          padding, (dilated_kernel_h, dilated_kernel_w)
      )
      out_channel = simplify(in_channel * channel_multiplier)
      out_height = simplify(
          (in_height - dilated_kernel_h + pad_top + pad_down)
          // stride_h
          + 1
      )
      out_width = simplify(
          (in_width - dilated_kernel_w + pad_left + pad_right)
          // stride_w
          + 1
      )
  
      # padding stage
      pad_before = [0, pad_top, pad_left, 0]
      pad_after = [0, pad_down, pad_right, 0]
      PaddedInput = pad(
          Input, pad_before, pad_after, name="PaddedInput"
      )
      # depthconv stage
      idxdiv = tvm.tir.indexdiv
      idxmod = tvm.tir.indexmod
  
      di = te.reduce_axis((0, filter_height), name="di")
      dj = te.reduce_axis((0, filter_width), name="dj")
      Output = te.compute(
          (batch, out_height, out_width, out_channel),
          lambda b, i, j, c: te.sum(
              (
                  PaddedInput[
                      b,
                      i * stride_h + di * dilation_h,
                      j * stride_w + dj * dilation_w,
                      idxdiv(c, channel_multiplier),
                  ].astype(out_dtype)
                  * Filter[
                      tuple(
                          np.array(
                              [
                                  di,
                                  dj,
                                  idxdiv(c, channel_multiplier),
                                  idxmod(c, channel_multiplier),
                              ]
                          )[kernel_permutation]
                      )
                  ].astype(out_dtype)
              ),
              axis=[di, dj],
          ),
          name="DepthwiseConv2d",
          tag="depthwise_conv2d_nhwc",
      )
      return Output
  
file_path: tvm/python/tvm/topi/nn/depthwise_conv2d.py
imports: from __future__ import absolute_import
  from collections import namedtuple
  import numpy
  import tvm
  from tvm import te
  from .dilate import dilate
  from .pad import pad
  from .utils import get_pad_tuple
  from ..utils import simplify, get_const_tuple
label_desc: Use image node_id 3 for calling a global function with example usage: depthwise_conv2d_nhwc(Input, Filter, stride, padding, dilation, kernel_layout, out_dtype) and returns: Output