code: def avg_pool2d_NCHW(
      data,
      kernel,
      stride,
      padding,
      dilation,
      count_include_pad,
      oshape,
      odtype="float16",
  ):
      """avg_pool2d compute"""
      if odtype != "float16":
          raise RuntimeError(f"Unsupported output dtype '{odtype}'")
      kh, kw = kernel
      rh = te.reduce_axis((0, kh), name="rh")
      rw = te.reduce_axis((0, kw), name="rw")
      sh, sw = stride
      dh, dw = dilation
  
      dilated_kh = (kh - 1) * dh + 1
      dilated_kw = (kw - 1) * dw + 1
  
      pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
          get_const_tuple(padding), (dilated_kh, dilated_kw)
      )
  
      # DOPAD
  
      if (
          pad_top != 0
          or pad_down != 0
          or pad_left != 0
          or pad_right != 0
      ):
          pad_before = (0, 0, pad_top, pad_left)
          pad_after = (0, 0, pad_down, pad_right)
          data_pad = pad(data, pad_before, pad_after, name="data_pad")
      else:
          # By definition when True, zero-padding will be included in the averaging calculation
          # This is equivalent to PoolArea = (kh * kw)
          count_include_pad = True
          data_pad = data
  
      Sum = te.compute(
          oshape,
          lambda b, c, h, w: te.sum(
              data_pad[b, c, h * sh + dh * rh, w * sw + dw * rw].astype(
                  "float32"
              ),
              axis=[rh, rw],
          ),
          name="pool_sum",
      )
  
      if not count_include_pad:
          # Compute PoolArea using unpadded input tensor
          _, _, oh, ow = oshape
          _, _, ih, iw = data.shape
  
          PoolArea = te.compute(
              (oh, ow),
              lambda i, j: compute_PoolArea(
                  i,
                  j,
                  ih,
                  iw,
                  kh,
                  kw,
                  sh,
                  sw,
                  dh,
                  dw,
                  pad_top,
                  pad_left,
              ),
              name="pool_area",
          )
  
          InvArea = te.compute(
              (oh, ow),
              lambda i, j: tir.if_then_else(
                  tir.all(PoolArea[i, j] > 0),
                  (float(1) / PoolArea[i, j]),
                  0,
              ),
              name="inverse_area",
          )
  
          Avg = te.compute(
              oshape,
              lambda b, c, h, w: (
                  Sum[b, c, h, w] * InvArea[h, w]
              ).astype(odtype),
              name="pool_avg",
          )
      else:
          InvArea = float(1) / (kh * kw)
          Avg = te.compute(
              oshape,
              lambda b, c, h, w: (Sum[b, c, h, w] * InvArea).astype(
                  odtype
              ),
              name="pool_avg",
          )
  
      return Avg
  
file_path: tvm/python/tvm/topi/hexagon/slice_ops/avg_pool2d.py
imports: from tvm import te
  from tvm import tir
  from ..utils import get_layout_transform_fn
  from ...utils import get_const_tuple
  from ...nn.utils import get_pad_tuple
  from ...nn.pad import pad
  from ..compute_poolarea import compute_PoolArea
label_desc: Use image node_id 1 for calling a global function with example usage: avg_pool2d_NCHW(data, kernel, stride, padding, dilation, count_include_pad, oshape, odtype) and returns: Avg