code: def DivideConstants(const_dict):
      """This pass rewrites the IR and constant dict such that all constant
      accesses are at 0 offset and full length (i.e. they read the whole buffer).
  
      Where necessary, new constants are created in order to ensure the rewrite
      can take place. As an example, if a convolution is tiled along the channels
      axis, the accesses to the weights will need to be offset. This pass will
      create new constants consisting of 'slices' of the weights so each tile
      of the compute can access one of these 'slices'.
  
      The purpose of this pass is to transform the IR into a form we can apply
      constant encoding to (which will compress weights and encode biases).
      """
      buffer_to_const = {}  # type: ignore
      new_buffers = []
      new_consts = []
      keep_buffers = set()
      new_const_dict = {}
  
      def _visit(stmt):
          new_args = []
          # We don't want to divide the constant that will be executed on two cores in parallel
          is_u65_conv2d = (
              vela_api.get_accelerator_config()
              == vapi.NpuAccelerator.Ethos_U65_512
              and stmt.args[0] == "ethosu_conv2d"
          )
          for i, arg in enumerate(stmt.args):
              if isinstance(arg, tvm.tir.expr.BufferLoad):
                  # If we're trying to load a buffer that maps to a constant
                  if arg.buffer.data in buffer_to_const:
                      const = buffer_to_const[arg.buffer.data]
                      flattened_const_shape = np.prod(const.shape)
  
                      offset = int(arg.indices[0])
                      # Note by convention the arg after a constant read is the length of the read
                      length = int(stmt.args[i + 1])
                      # If it's anything other than a full read, create a new buffer
                      if (
                          offset != 0
                          or flattened_const_shape != length
                          and length > 0
                      ) and not is_u65_conv2d:
                          out_channels = const.shape[0]
                          offset_channels = int(
                              (offset * out_channels)
                              / flattened_const_shape
                          )
                          length_channels = int(
                              (length * out_channels)
                              / flattened_const_shape
                          )
                          # split the constant up across channels
                          split_const = np.split(
                              const, out_channels, axis=0
                          )
                          # create a new const out of the channels we want to keep
                          new_const = np.concatenate(
                              split_const[
                                  offset_channels : offset_channels
                                  + length_channels
                              ],
                              axis=0,
                          )
                          new_consts.append(new_const)
                          new_buffer = tvm.tir.decl_buffer(
                              (length,),
                              arg.dtype,
                              scope=arg.buffer.scope(),
                          )
                          new_buffers.append(new_buffer)
                          new_args.append(
                              tvm.tir.expr.BufferLoad(new_buffer, [0])
                          )
                          continue
                      keep_buffers.add(arg.buffer.data)
  
              new_args.append(arg)
  
          return tvm.tir.Call(stmt.dtype, stmt.op, new_args, stmt.span)
  
      def _ftransform(f, mod, ctx):
          for i, param in enumerate(f.params):
              if i in const_dict:
                  buffer_to_const[param] = const_dict[i]
                  buffer_to_const[
                      f.buffer_map[param].data
                  ] = const_dict[i]
  
          new_body = tvm.tir.stmt_functor.ir_transform(
              f.body, _visit, None, ["tir.Call"]
          )
          # Both the params and buffer map need updating for the newly introduced buffers
          new_params = []  # type: ignore
          new_buffer_map = {}
          for i, param in enumerate(f.params):
              buffer = f.buffer_map[param]
              pointer = buffer.data
              if pointer in buffer_to_const:
                  if pointer not in keep_buffers:
                      continue
                  new_const_dict[len(new_params)] = const_dict[i]
              new_params.append(param)
              new_buffer_map[param] = buffer
  
          for i, new_buffer in enumerate(new_buffers):
              handle = tvm.tir.Var("placeholder", "handle")
              new_params.append(handle)
              new_buffer_map[handle] = new_buffer
              new_const_dict[len(new_params) - 1] = new_consts[i]
  
          new_f = tvm.tir.PrimFunc(
              new_params,
              new_body,
              f.ret_type,
              new_buffer_map,
              f.attrs,
              f.span,
          )
          return new_f
  
      def _divide_constants(mod):
          transform_func = tvm.tir.transform.prim_func_pass(
              _ftransform,
              opt_level=0,
              name="tir.contrib.ethos-u.divide_constants",
          )
          new_func = transform_func(mod)
          return new_func, new_const_dict
  
      return _divide_constants
  
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 3 for calling a global function with example usage: DivideConstants(const_dict) and returns: _divide_constants, tvm, new_f, new_func,  new_const_dict