code: def test_double_buffer():
      dtype = "int64"
      n = 100
      m = 4
      tx = te.thread_axis("threadIdx.x")
      ib = tvm.tir.ir_builder.create()
      A = ib.pointer("float32", name="A")
      C = ib.pointer("float32", name="C")
      ib.scope_attr(tx, "thread_extent", 1)
      with ib.for_range(0, n) as i:
          B = ib.allocate("float32", m, name="B", scope="shared")
          with ib.new_scope():
              ib.scope_attr(B.asobject().data, "double_buffer_scope", 1)
              with ib.for_range(0, m) as j:
                  B[j] = A[i * 4 + j]
          with ib.for_range(0, m) as j:
              C[j] = B[j] + 1
  
      stmt = ib.get()
      mod = tvm.IRModule(
          {"db": tvm.tir.PrimFunc([A.asobject(), C.asobject()], stmt)}
      )
  
      opt = tvm.transform.Sequential(
          [
              tvm.tir.transform.InjectDoubleBuffer(),
              tvm.tir.transform.Simplify(),
          ]
      )
  
      with tvm.transform.PassContext(
          config={"tir.InjectDoubleBuffer": {"split_loop": 2}}
      ):
          mod = opt(mod)
      stmt = mod["db"].body
  
      assert isinstance(stmt.body, tvm.tir.Allocate)
      assert list(stmt.body.extents) == [m * 2]
  
      f = tvm.tir.transform.ThreadSync("shared")(mod)["db"]
      count = [0]
  
      def count_sync(op):
          if isinstance(op, tvm.tir.Call) and op.op.same_as(
              tvm.ir.Op.get("tir.tvm_storage_sync")
          ):
              count[0] += 1
  
      tvm.tir.stmt_functor.post_order_visit(f.body, count_sync)
      assert count[0] == 4
  
file_path: tvm/tests/python/tir-transform/test_tir_transform_inject_double_buffer.py
imports: import tvm
  import tvm.testing
  from tvm.script import tir, ir
  from tvm import te
label_desc: Use image node_id 1 for calling a global function with example usage: test_double_buffer() without return types