code: def vtcm_verify_pass(**kwargs):
      """Verify the validity of a hexagon kernel.
      This pass will check vtcm memory usage.
      """
  
      def verify_pass(f, *_):
          sizes = tvm.tir.analysis.calculate_allocated_bytes(f)
          vtcm_capacity = kwargs.get("vtcm_capacity", 0)
          vtcm_allocated = sizes.get("global.vtcm", 0)
          if 0 < vtcm_capacity < vtcm_allocated:
              raise InstantiationError(
                  "Skipped because of invalid vtcm memory usage limit"
              )
  
          return f
  
      return tvm.tir.transform.prim_func_pass(verify_pass, opt_level=0)
  
imports: import contextlib
  import logging
  import os
  import shutil
  import tempfile
  import threading
  import time
  import traceback
  import typing
  import warnings
  from collections import namedtuple
  from random import getrandbits
  import tvm._ffi
  import tvm.ir.transform
  from tvm import nd
  from tvm import rpc
  from tvm.autotvm.env import AutotvmGlobalScope, reset_global_scope
  from tvm.contrib import ndk, stackvm, tar
  from tvm.contrib.popen_pool import PopenPoolExecutor
  from tvm.driver import build
  from tvm.error import TVMError
  from tvm.target import Target
  from ..env import AutotvmGlobalScope
  from ..task.space import InstantiationError
  from ..utils import get_const_tuple
  from .measure import Builder, MeasureErrorNo, MeasureResult, Runner
label_desc: Use image node_id 8 for calling a global function with example usage: vtcm_verify_pass() and returns: tvm, f