code: def _benchmark_hexagon_elementwise_add_kernel(
      hexagon_session: Session, shape: list, dtype: str, mem_scope: str
  ):
      """
      Generate and benchmark a single elementwise-add kernel for Hexagon.
  
      Produce these outputs:
        - Printed status updates / results to stdout and/or stderr.
  
        - Create a new subdirectory under _HOST_OUTPUT_DIR, and populate it with
          various logs and intermediate files.
  
        - Add to _BT a row describing this benchmark run.
      """
      # Represent the benchmark details in a form required by the benchmark table
      # and for other logging...
      keys_dict = {
          "basic_kernel": "ewise-add",
          "dtype": dtype,
          "shape": shape,
          "mem_scope": mem_scope,
      }
  
      desc = bu.get_benchmark_decription(keys_dict)
  
      # Create the host-side directory for this benchmark run's files / logs...
      host_files_dir_name = bu.get_benchmark_id(keys_dict)
      host_files_dir_path = os.path.join(
          _HOST_OUTPUT_DIR, host_files_dir_name
      )
      os.mkdir(host_files_dir_path)
  
      keys_dict["host_files_dir_path"] = host_files_dir_path
  
      log_file_path = os.path.join(host_files_dir_path, "out.txt")
      with open(log_file_path, "w", encoding="UTF-8") as log_file:
          print(f"CONFIGURATION: {desc}")
          log_file.write(f"CONFIGURATION: {desc}\n")
  
          try:
              ns_tir_module = _get_irmod_elemwise_add(
                  shape, dtype, mem_scope
              )
  
              # Dump the primfunc NS-TIR (as text) to the log file...
              lowered_mod = tvm.lower(ns_tir_module, _PRIMFUNC_NAME)
              log_file.write("LOWERED IR MODULE:\n")
              log_file.write(str(lowered_mod))
              log_file.write("\n")
  
              # Lower the primfunc's IRModule to Hexagon object code...
              input1 = tvm.te.placeholder(shape, dtype=dtype)
              input2 = tvm.te.placeholder(shape, dtype=dtype)
              output = tvm.te.placeholder(shape, dtype=dtype)
  
              built_module: tvm.driver.build_module.OperatorModule = (
                  tvm.build(
                      ns_tir_module,
                      [
                          input1,
                          input2,
                          output,
                      ],
                      get_hexagon_target("v69"),
                      name=_PRIMFUNC_NAME,
                  )
              )
  
              # Create an actual Hexagon-native shared object file, initially stored on the
              # host's file system...
              host_dso_binary_path = os.path.join(
                  host_files_dir_path, "test_binary.so"
              )
              built_module.save(host_dso_binary_path)
              print(
                  f"SAVED BINARY TO HOST PATH: {host_dso_binary_path}"
              )
  
              # Upload the .so to the Android device's file system (or wherever is appropriate
              # when using the Hexagon simulator)...
              target_dso_binary_filename = "test_binary.so"
              target_dso_binary_pathname = hexagon_session.upload(
                  host_dso_binary_path, target_dso_binary_filename
              )
  
              # Generate our testing / validation data...
              (
                  host_numpy_input1_data,
                  host_numpy_input2_data,
                  host_numpy_output_data_expected,
              ) = _get_elemwise_add_reference_value_tensors(
                  shape, dtype
              )
  
              # On the target device / simulator, make our Hexagon-native shared object
              # available for use...
              loaded_hexagon_module: tvm.runtime.module.Module = (
                  hexagon_session.load_module(
                      target_dso_binary_pathname
                  )
              )
  
              # Create the target-side tensors to hold the primfunc's inputs and outputs...
              input1_data = tvm.nd.empty(
                  shape, dtype, hexagon_session.device, mem_scope
              )
              input2_data = tvm.nd.empty(
                  shape, dtype, hexagon_session.device, mem_scope
              )
              output_data = tvm.nd.empty(
                  shape, dtype, hexagon_session.device, mem_scope
              )
  
              # Populate the primfunc's input tensors...
              input1_data.copyfrom(host_numpy_input1_data)
              input2_data.copyfrom(host_numpy_input2_data)
  
              # Actually benchmark the primfunc...
              timer = loaded_hexagon_module.time_evaluator(
                  "main", hexagon_session.device, number=10, repeat=1
              )
              timing_result = timer(
                  input1_data, input2_data, output_data
              )
  
              print(f"TIMING RESULT: {timing_result}")
              log_file.write(f"TIMING RESULT: {timing_result}\n")
  
              # Verify that the computation actually happened, and produced the correct result.
              result = output_data.numpy()
  
              if dtype == "float16":
                  # These are the closest tolerance we currently expect / require for these
                  # kernels.  They may be changed in the future.
                  rel_tolerance = 0.005
                  abs_tolerance = 2.0
              elif dtype == "int8":
                  rel_tolerance = 0
                  abs_tolerance = 0
              else:
                  raise Exception(f"Unexpected dtype: {dtype}")
  
              # TODO: We're assuming that *any* assertion thrown by 'assert_allclose' is because
              # the numerical differences were too large.  But ideally this code would
              # differentiate between (a) numerical difference errors, which should simply be
              # recorded as a failed benchmark run, vs. (b) more serious errors that should
              # kill the overall script.
              try:
                  tvm.testing.assert_allclose(
                      result,
                      host_numpy_output_data_expected,
                      rel_tolerance,
                      abs_tolerance,
                  )
              except AssertionError as err:
                  raise bu.NumericalAccuracyException(str(err))
  
              _BT.record_success(timing_result, **keys_dict)
  
          except bu.NumericalAccuracyException as err:
              print()
              print("FAIL: Numerical accuracy error. See log file.")
  
              log_file.write("\n")
              log_file.write(f"FAIL: {err}\n")
  
              _BT.record_fail(
                  **keys_dict,
                  comments="Numerical accuracy error. See log file.",
              )
  
          except bu.UnsupportedException as err:
              print()
              print(f"SKIP: {err}")
  
              log_file.write("\n")
              log_file.write(f"SKIP: {err}\n")
  
              _BT.record_skip(
                  **keys_dict,
                  comments=f"Unsupported configuration: {err}",
              )
  
file_path: tvm/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py
imports: import os
  import os.path
  import sys
  import tempfile
  import numpy
  import pytest
  import tvm.script
  import tvm.testing
  from tvm.contrib.hexagon.session import Session
  from tvm.script import tir
  from .None import benchmark_util
  from .infrastructure import get_hexagon_target
label_desc: Use image node_id 2 for calling a global function with example usage: _benchmark_hexagon_elementwise_add_kernel(hexagon_session, shape, dtype, mem_scope) without return types