code: def ReplaceOperators():
      """Replace operators represented as explicit loop nests with call_externs
      to NPU operators."""
      op_map = {
          "ethosu_conv2d": get_conv2d_params,
          "ethosu_copy": get_copy_params,
          "ethosu_depthwise_conv2d": get_depthwise_conv2d_params,
          "ethosu_pooling": get_pooling_params,
          "ethosu_binary_elementwise": get_binary_elementwise_params,
          "ethosu_identity": get_identity_params,
          "ethosu_unary_elementwise": get_unary_elementwise_params,
      }
      producers_consumers = ProducersConsumers()
      replace_output_pointer = {}
      pointer_to_extents = {}
      replaced_pointers = []
  
      ReplaceInfo = namedtuple("ReplaceInfo", ["pointer", "reallocate"])
  
      def _find_pointer_to_extent(stmt):
          if isinstance(stmt, tvm.tir.Allocate):
              pointer_to_extents[stmt.buffer_var] = stmt.extents
  
      def _resolve_pointers(stmt):
          """This pass determines information about the pointers present in the IR.
          In particular, it associates pointers with both the operations that
          produce them and the operations that consume them through the
          pointer_to_producer and pointer_to_consumer dicts.
  
          Additionally, it determines the extent (size/shape) of each pointer which
          is required for the _replace_pointers pass which runs later.
          """
          loads = []
  
          def _get_loads(stmt):
              if isinstance(stmt, tvm.tir.BufferLoad):
                  loads.append(stmt.buffer.data)
  
          buffer_var = None
  
          def _get_buffer_var(stmt):
              if isinstance(stmt, tvm.tir.BufferStore):
                  nonlocal buffer_var
                  buffer_var = stmt.buffer.data
  
          if isinstance(stmt, tvm.tir.AttrStmt):
              if stmt.attr_key == "pragma_op":
                  tvm.tir.stmt_functor.post_order_visit(
                      stmt, _get_buffer_var
                  )
                  producers_consumers.add_producer(buffer_var, stmt)
  
                  tvm.tir.stmt_functor.post_order_visit(
                      stmt, _get_loads
                  )
                  for load_pointer in loads:
                      if load_pointer != buffer_var:
                          producers_consumers.add_consumer(
                              load_pointer, stmt
                          )
  
      def _replace_operator(stmt):
          """Replace operators with call_externs, having derived the parameters
          from the relevant TIR expressions/statements.
  
          Note the complexity of this pass is mostly from the concept of 'replace
          pointers'. A call_extern may in principle require information from several
          loop nests in TIR (each corresponding to a different TE compute op). For
          example, a convolution operator will have other TE compute ops before and
          after corresponding to the input/output DMA functionality. Therefore, when
          the 'central' convolution op is replaced with a call_extern, the memory
          from the final DMA output op must be hoisted to the location/scope of
          the call_extern.
  
          The is done by replacing the pointer corresponding to the current operation
          with the 'true' output operator through the replace_output_pointer dict.
          Because of this, the param_func must provide a replace_pointer if the op
          isn't the true output but instead a no_compile op is."""
          if isinstance(stmt, tvm.tir.AttrStmt):
              op_name = stmt.value.value
              if stmt.attr_key == "pragma_op" and op_name in op_map:
                  # Get the parameters for the extern call
                  param_func = op_map[op_name]
                  (
                      info,
                      output_pointer,
                      replace_pointer,
                      is_allocator,
                  ) = param_func(stmt, producers_consumers)
                  if replace_pointer is not None:
                      # Allocate pointer only once
                      if replace_pointer in replaced_pointers:
                          is_allocator = False
                      replace_output_pointer[
                          output_pointer
                      ] = ReplaceInfo(replace_pointer, is_allocator)
                      replaced_pointers.append(replace_pointer)
                  # Make the extern call
                  irb = tvm.tir.ir_builder.create()
                  irb.emit(
                      tvm.tir.call_extern("handle", op_name, *info)
                  )
                  return irb.get()
          return None
  
      def _remove_no_compile(stmt):
          """Certain operators are marked as 'no compile' operators. This means they
          should be removed from the IR as they are compiled as part of other operators.
          The IFM DMA operations are an example of this, as they don't get compiled
          independently but instead get compiled into the operator they're associated with,
          e.g. a conv2d.
  
          There are potentially 2 parts to remove for an operator:
          the allocate for its output and the compute nest itself. For the
          allocate, we can check if the pointer they reference is produced by a 'no compile'
          operator. For the compute nest, we can just check the op pragma.
          """
          if isinstance(stmt, tvm.tir.AttrStmt):
              # Remove compute nests
              if (
                  stmt.attr_key == "pragma_op"
                  and stmt.value.value not in op_map
              ):
                  return tvm.tir.Evaluate(0)
  
          if isinstance(stmt, tvm.tir.Allocate):
              # Remove allocates
              producer = producers_consumers.get_last_producer(
                  stmt.buffer_var
              )
              if producer:
                  if (
                      producer.attr_key == "pragma_op"
                      and producer.value.value not in op_map
                  ):
                      return stmt.body
  
          return None
  
      def _replace_pointers(stmt):
          if isinstance(stmt, tvm.tir.Allocate):
              # If the allocate allocates a pointer that needs replacing
              if stmt.buffer_var in replace_output_pointer:
                  replace_pointer, reallocate = replace_output_pointer[
                      stmt.buffer_var
                  ]
                  if not reallocate:
                      return stmt.body
                  # Otherwise, rewrite the allocation statement with the new pointer
                  # and the new extent
                  replace_type = (
                      replace_pointer.type_annotation.element_type.dtype
                  )
                  replace_extents = pointer_to_extents[replace_pointer]
                  return tvm.tir.Allocate(
                      replace_pointer,
                      replace_type,
                      replace_extents,
                      stmt.condition,
                      stmt.body,
                  )
          return None
  
      def _remove_buffer_decl(stmt):
          if isinstance(stmt, tvm.tir.DeclBuffer):
              if stmt.buffer.data in replace_output_pointer:
                  return stmt.body
  
      def _post_transform(stmt):
          # Replace operators with call_externs
          result = _replace_operator(stmt)
          # Remove operators that don't need compiling
          result = result or _remove_no_compile(stmt)
          # Replace necessary pointers that were removed in the previous step
          result = result or _replace_pointers(stmt)
          # Replace BufferDecl, since only the tir.Var data pointer is
          # still used, and not the tir.Buffer
          result = result or _remove_buffer_decl(stmt)
  
          return result
  
      def _ftransform(f, mod, ctx):
          tvm.tir.stmt_functor.post_order_visit(
              f.body, _find_pointer_to_extent
          )
          tvm.tir.stmt_functor.post_order_visit(
              f.body, _resolve_pointers
          )
          producers_consumers.add_allocate_variables(
              pointer_to_extents.keys()
          )
          return f.with_body(
              tvm.tir.stmt_functor.ir_transform(
                  f.body,
                  None,
                  _post_transform,
                  ["tir.AttrStmt", "tir.Allocate"],
              )
          )
  
      return tvm.tir.transform.prim_func_pass(
          _ftransform,
          opt_level=0,
          name="tir.contrib.ethos-u.replace_operators",
      )
  
imports: from collections import namedtuple
  from typing import Optional
  import numpy
  from ethosu.vela import api
  import tvm
  from tvm.relay.backend.contrib.ethosu import vela_api
  from tvm.relay.backend.contrib.ethosu import tir_to_cs_translator
  from .convolution import get_conv2d_params
  from .depthwise import get_depthwise_conv2d_params
  from .pooling import get_pooling_params
  from .binary_elementwise import get_binary_elementwise_params
  from .identity import get_identity_params
  from .unary_elementwise import get_unary_elementwise_params
  from .transform import get_copy_params
  from .producers_consumers import ProducersConsumers
  from ..None import _ffi_api
label_desc: Use image node_id 2 for calling a global function with example usage: ReplaceOperators() and returns: tvm, None, None, None, result, f, irb, tvm, tvm, stmt, stmt, stmt