code: def ir_lower_vtcm():
      """Create a VTCM lowering pass.
  
      VTCM memory has to be allocated using special functions.
      """
  
      def get_alloc(stmt, align):
          assert isinstance(stmt, tvm.tir.Allocate)
          return tvm.tir.call_extern(
              "handle",
              "HexagonBackendAllocateVTCM",
              ft.reduce(lambda x, y: x * y, stmt.extents, 1),
              align,
          )
  
      def get_free(var):
          return tvm.tir.call_extern(
              "handle", "HexagonBackendFreeVTCM", var
          )
  
      # pylint: disable=bad-whitespace
      @tvm.tir.transform.prim_func_pass(
          opt_level=0, name="Lower VTCM pass"
      )
      def transform(func, mod, ctx):
          return lower_vtcm_(get_alloc, get_free, 2048, func, mod, ctx)
  
      return transform
  
imports: import functools
  import numpy
  import tvm
  from tvm import relay
  from tvm.relay.dataflow_pattern import DFPatternCallback, is_constant, is_op, is_tuple, rewrite, wildcard
  from tvm.topi.utils import get_const_tuple
  from tvm.relay.expr import Call
  from tvm.runtime import ndarray
  from ..._ffi.registry import register_func
label_desc: Use image node_id 3 for calling a global function with example usage: ir_lower_vtcm() and returns: transform, tvm, tvm, lower_vtcm_