code: def match_ethosu_pooling(output_tensor, device_config):
      """Match a Tensor Expression corresponding to an NPU Pooling.
  
      If the Tensor Expression matches, an EthosuPart will be created that models the
      matched Tensor Expression. Otherwise, None will be returned.
  
      Parameters
      ----------
      output_tensor : tvm.te.Tensor
          The tensor to attempt to match with.
      device_config : EthosuDeviceConfig
          Target device configuration
  
      Returns
      -------
      Union[None, EthosuPart]
          The created EthosuPart if there was a match, otherwise None.
  
      """
      write = output_tensor
      if write.op.name != "ethosu_write":
          return None
      convert_to_nhcwb16 = write.op.input_tensors[0]
      if convert_to_nhcwb16.op.name != "ethosu_convert_to_nhcwb16":
          return None
      pool2d = convert_to_nhcwb16.op.input_tensors[0]
      if pool2d.op.name != "ethosu_pooling":
          return None
      pad = pool2d.op.input_tensors[0]
      if pad.op.name != "ethosu_pad":
          return None
      upscale = pad.op.input_tensors[0]
      if upscale.op.name != "ethosu_upscale":
          return None
      convert_to_nhwc = upscale.op.input_tensors[0]
      if convert_to_nhwc.op.name != "ethosu_convert_to_nhwc":
          return None
      read = convert_to_nhwc.op.input_tensors[0]
      if read.op.name != "ethosu_read":
          return None
  
      input_tensors = [
          read.op.input_tensors[0],
      ]
      subgraph = TESubgraph(input_tensors, output_tensor)
      propagators = [
          write.op.attrs["ifm_propagator"],
      ]
      ifm_dtype = input_tensors[0].dtype
      ofm_dtype = output_tensor.dtype
  
      # Use channels from a stage of TE graph where the IFM is always NHWC
      channels = int(pool2d.shape[3])
      pool_shape_h = int(pool2d.op.attrs["pool_shape_h"])
      pool_shape_w = int(pool2d.op.attrs["pool_shape_w"])
  
      subkernels = len(
          device_config.get_kernel_steps(
              pool2d.op.name, pool_shape_h, pool_shape_w, ifm_dtype
          )
      )
  
      output_layout = convert_to_nhcwb16.op.attrs["layout"]
      input_layout = convert_to_nhwc.op.attrs["layout"]
      output_quantum = device_config.get_output_quantum(output_layout)
  
      valid_block_configs = device_config.get_valid_block_configs(
          propagators[0],
          pool2d.op.attrs,
          output_tensor.shape,
          channels,
          channels,
          output_layout,
          input_layout,
          ifm_dtype,
          ofm_dtype,
          pool_shape_h,
          pool_shape_w,
      )
  
      return EthosuPart(
          subgraph,
          propagators,
          output_quantum,
          subkernels,
          valid_block_configs,
      )
  
file_path: tvm/python/tvm/relay/backend/contrib/ethosu/te/pooling.py
imports: from typing import Tuple
  import numpy
  from tvm import te
  from tvm.contrib.ethosu.cascader import TESubgraph, EthosuPart, Propagator, register_matcher
  from .dma import dma_ofm_compute, dma_ifm_compute
  from .common import get_layout_transform_matrices
label_desc: Use image node_id 2 for calling a global function with example usage: match_ethosu_pooling(output_tensor, device_config) and returns: EthosuPart, None, None, None, None, None, None, None