code: def test_coproc_sync():
      @tvm.register_func("tvm.info.mem.global.cache")
      def meminfo_cache():
          return tvm.ir.make_node(
              "MemoryInfo",
              unit_bits=8,
              max_simd_bits=32,
              max_num_bits=128,
              head_address=tvm.tir.call_extern(
                  "handle", "global_cache"
              ),
          )
  
      ib = tvm.tir.ir_builder.create()
      n = te.size_var("n")
      cp = te.thread_axis((0, 1), "cop")
      A = ib.allocate("float32", 128, name="A", scope="global.cache")
      with ib.for_range(0, n, name="i") as i:
          A[i] = A[i] + 1
          with ib.for_range(0, 8, name="k") as k:
              with ib.for_range(0, 10, name="j") as j:
                  ib.scope_attr(cp, "coproc_scope", 1)
                  A[j] = A[j + k * 10] + 2
      stmt = ib.get()
  
      mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([n], stmt))
      stmt = tvm.tir.transform.CoProcSync()(mod)["main"].body
  
      body = stmt.body.body
      blist = tvm.tir.stmt_list(body)
  
      assert blist[1].value.op.same_as(
          tvm.ir.Op.get("tir.cop.coproc_read_barrier")
      )
      assert blist[1].value.args[3].value == 80
      assert blist[-2].value.op.same_as(
          tvm.ir.Op.get("tir.cop.coproc_sync")
      )
      assert blist[-1].value.op.same_as(
          tvm.ir.Op.get("tir.cop.coproc_write_barrier")
      )
      assert blist[-1].value.args[3].value == 10
  
imports: import tvm
  from tvm import te
label_desc: Use image node_id 1 for calling a global function with example usage: test_coproc_sync() and returns: tvm