code: def test_simplex_data_transferring():
      r"""
      Test the heterogeneous execution of a simple network where data
      transferring is from the target device to the host processor at runtime.
      The host processor is always assumed to be cpu, and the device varies.
      """
      host = "cpu"
      target_host = "llvm"
      host_dev = tvm.device(host)
      if not tvm.runtime.enabled(target_host):
          print("Skip test because llvm is not enabled.")
          return
  
      def check_device(device, target_device):
          if not tvm.runtime.enabled(target_device):
              print(
                  "Skip test because {} is not enabled.".format(
                      target_device
                  )
              )
              return
  
          device_dev = tvm.device(device)
          graph = get_simplex_graph(
              host_dev.device_type, device_dev.device_type
          )
          shape = (4,)
  
          # Create module for add whose target is the device.
          tensor_a = te.placeholder(shape, name="A")
          tensor_b = te.placeholder(shape, name="B")
          elemwise_add = te.compute(
              shape,
              lambda *i: tensor_a(*i) + tensor_b(*i),
              name="elemwise_add",
          )
          target = topi.cpp.TEST_create_target(device)
          schedule_add = topi.cpp.cuda.schedule_injective(
              target, [elemwise_add]
          )
          lower_add = tvm.lower(
              schedule_add,
              [tensor_a, tensor_b, elemwise_add],
              name="elemwise_add",
          )
  
          # Insert copy. Neither compute nor schedule is required for the copy
          # node. The compute will be performed at runtime which is just data
          # copy from the input to the output.
          tensor_copy = te.placeholder(shape, name="__copy")
  
          # Create module for sub whose target is the host.
          tensor_c = te.placeholder(shape, name="C")
          elemwise_sub = te.compute(
              shape,
              lambda *i: tensor_copy(*i) - tensor_c(*i),
              name="elemwise_sub",
          )
          schedule_sub = te.create_schedule(elemwise_sub.op)
          lower_sub = tvm.lower(
              schedule_sub,
              [tensor_copy, tensor_c, elemwise_sub],
              name="elemwise_sub",
          )
  
          target_flist = {
              target_device: lower_add,
              target_host: lower_sub,
          }
          target = tvm.target.Target(target, target_host)
          mhost = tvm.build(target_flist, target=target)
          dev = [host_dev, device_dev]
          mod = graph_executor.create(graph, mhost, dev)
          params = {}
          params["A"] = tensor_a = np.random.uniform(size=shape).astype(
              tensor_a.dtype
          )
          params["B"] = tensor_b = np.random.uniform(size=shape).astype(
              tensor_b.dtype
          )
          params["C"] = tensor_c = np.random.uniform(size=shape).astype(
              tensor_c.dtype
          )
          mod.set_input(**params)
          mod.run()
          out = mod.get_output(0, tvm.nd.empty(shape))
          np.testing.assert_equal(
              out.numpy(), (tensor_a + tensor_b) - tensor_c
          )
  
      dev_tar = {"cuda": "cuda", "opencl": "opencl"}
      for device, target in dev_tar.items():
          with tvm.target.Target(device):
              check_device(device, target)
  
file_path: tvm/tests/python/runtime/test_runtime_heterogeneous.py
imports: import json
  import numpy
  import tvm
  from tvm import te
  from tvm.contrib import graph_executor, utils
  from tvm import topi
label_desc: Use image node_id 2 for calling a global function with example usage: test_simplex_data_transferring() without return types