code: def match_ethosu_conv2d(output_tensor, device_config):
      """Match a Tensor Expression corresponding to an NPU Conv2D.
  
      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
      conv2d = convert_to_nhcwb16.op.input_tensors[0]
      if conv2d.op.name != "ethosu_conv2d":
          return None
      pad = conv2d.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],
          conv2d.op.input_tensors[1],
          conv2d.op.input_tensors[2],
      ]
  
      subgraph = TESubgraph(input_tensors, output_tensor)
      propagators = [
          write.op.attrs["ifm_propagator"],
          write.op.attrs["weights_propagator"],
          write.op.attrs["bias_propagator"],
      ]
      ifm_dtype = input_tensors[0].dtype
      ofm_dtype = output_tensor.dtype
  
      # Use channels from the weights tensor since that its shape doesn't change during layout
      # conversion
      ifm_channels = int(input_tensors[1].shape[3])
      ofm_channels, kernel_height, kernel_width = (
          int(axis) for axis in input_tensors[1].shape[0:3]
      )
      kernel_elements = kernel_height * kernel_width
  
      is_part_kernel = device_config.is_partkernel(
          conv2d.op.name, ifm_channels, ifm_dtype, kernel_elements
      )
      subkernels = len(
          device_config.get_kernel_steps(
              conv2d.op.name,
              kernel_height,
              kernel_width,
              ifm_dtype,
              is_part_kernel,
          )
      )
  
      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],
          conv2d.op.attrs,
          output_tensor.shape,
          ofm_channels,
          ifm_channels,
          output_layout,
          input_layout,
          ifm_dtype,
          ofm_dtype,
          kernel_height,
          kernel_width,
      )
  
      return EthosuPart(
          subgraph,
          propagators,
          output_quantum,
          subkernels,
          valid_block_configs,
          1,
      )
  
file_path: tvm/python/tvm/relay/backend/contrib/ethosu/te/convolution.py
imports: from typing import Tuple, Union, List
  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_conv2d(output_tensor, device_config) and returns: EthosuPart, None, None, None, None, None, None, None