code: def test_add_pipeline():
      nn = 64
      max_threads = 4
      n = tvm.runtime.convert(nn)
      A = te.placeholder((n,), name="A")
  
      def extern_generator(ins, outs):
          """Manually write the IR for the extern function, add pipeline"""
          ib = tvm.tir.ir_builder.create()
          with ib.for_range(0, (n + 1) // 2) as i:
              ib.emit(
                  outs[0].vstore(
                      i * 2,
                      ins[0].vload(i * 2, "float32x2")
                      + tvm.tir.const(1, "float32x2"),
                  )
              )
          return ib.get()
  
      def extern_generator_gpu(ins, outs):
          """Manually write the IR for the extern function, add pipeline"""
          ib = tvm.tir.ir_builder.create()
          bx = te.thread_axis("blockIdx.x")
          tx = te.thread_axis("threadIdx.x")
          ib.scope_attr(
              bx, "thread_extent", (nn + max_threads - 1) // max_threads
          )
          ib.scope_attr(tx, "thread_extent", max_threads)
          idx = bx.var * max_threads + tx.var
          with ib.if_scope(ib.likely(idx < n)):
              ib.emit(
                  outs[0].vstore(
                      idx * 2,
                      ins[0].vload(idx * 2, "float32x2")
                      + tvm.tir.const(1, "float32x2"),
                  )
              )
          return ib.get()
  
      C_cpu = te.extern(A.shape, [A], extern_generator, name="C")
      C_gpu = te.extern(A.shape, [A], extern_generator_gpu, name="C")
      s_cpu = te.create_schedule(C_cpu.op)
      s_gpu = te.create_schedule(C_gpu.op)
      print(tvm.lower(s_cpu, [A, C_cpu], simple_mode=True))
      print(tvm.lower(s_gpu, [A, C_gpu], simple_mode=True))
  
      def check_target(target):
          if not tvm.testing.device_enabled(target):
              return
          s = s_gpu if target in ["opencl", "cuda"] else s_cpu
          C = C_gpu if target in ["opencl", "cuda"] else C_cpu
          # build and invoke the kernel.
          f = tvm.build(s, [A, C], target)
          dev = tvm.device(target, 0)
          # launch the kernel.
          n = nn
          a = tvm.nd.array(
              np.random.uniform(size=n).astype(A.dtype), dev
          )
          c = tvm.nd.array(np.zeros(n, dtype=C.dtype), dev)
          f(a, c)
          tvm.testing.assert_allclose(c.numpy(), a.numpy() + 1)
  
      check_target("llvm")
      check_target("opencl")
      check_target("cuda")
  
imports: import tvm
  from tvm import te
  import numpy
  import tvm.testing
label_desc: Use image node_id 1 for calling a global function with example usage: test_add_pipeline() and returns: ib, ib