code: def test_loop_carried_dependency():
      """The buffer allocation should be above opaque iter var's loop scopes
      such that buffer accesses with loop carried dependencies are covered,
      and the allocate buffer should keep the order."""
  
      @T.prim_func
      def before(
          A: T.Buffer((8, 8, 8), "int32"),
          B: T.Buffer((8, 8, 8), "int32"),
      ):
          C = T.alloc_buffer([8, 8, 8], dtype="int32")
          D = T.alloc_buffer([8, 8, 8], dtype="int32")
          for i in T.serial(8):
              for j in T.serial(8):
                  for k in T.serial(8):
                      with T.block("b0"):
                          vi, vj, vk = T.axis.remap("SSS", [i, j, k])
                          C[vi, vj, vk] = A[vi, vj, vk] + 1
                  for k in T.serial(8):
                      with T.block("b1"):
                          vi, vj, vk = T.axis.remap("SSS", [i, j, k])
                          D[vi, vj, vk] = A[vi, vj, vk] + 2
                  for k in T.serial(8):
                      with T.block("b2"):
                          vi, vk = T.axis.remap("SS", [i, k])
                          vj = T.axis.opaque(8, j)
                          B[vi, vj, vk] = (
                              C[vi, vj, vk]
                              + T.if_then_else(
                                  0 < vj,
                                  C[vi, vj - 1, vk],
                                  0,
                                  dtype="int32",
                              )
                              + D[vi, vj, vk]
                          )
  
      @T.prim_func
      def after(
          A: T.Buffer((8, 8, 8), "int32"),
          B: T.Buffer((8, 8, 8), "int32"),
      ) -> None:
          for i in T.serial(8):
              with T.block():
                  T.reads(A[i, 0:8, 0:8])
                  T.writes(B[i, 0:8, 0:8])
                  C = T.alloc_buffer([8, 8, 8], dtype="int32")
                  D = T.alloc_buffer([8, 8, 8], dtype="int32")
                  for j in T.serial(8):
                      for k in T.serial(8):
                          with T.block("b0"):
                              vi, vj, vk = T.axis.remap(
                                  "SSS", [i, j, k]
                              )
                              C[vi, vj, vk] = A[vi, vj, vk] + 1
                      for k in T.serial(8):
                          with T.block("b1"):
                              vi, vj, vk = T.axis.remap(
                                  "SSS", [i, j, k]
                              )
                              D[vi, vj, vk] = A[vi, vj, vk] + 2
                      for k in T.serial(8):
                          with T.block("b2"):
                              vi, vk = T.axis.remap("SS", [i, k])
                              vj = T.axis.opaque(8, j)
                              B[vi, vj, vk] = (
                                  C[vi, vj, vk]
                                  + T.if_then_else(
                                      0 < vj,
                                      C[vi, vj - 1, vk],
                                      0,
                                      dtype="int32",
                                  )
                                  + D[vi, vj, vk]
                              )
  
      _check(before, after)
  
imports: import numpy
  import tvm
  import tvm.testing
  from tvm import te
  from tvm.script import tir
  from tvm import relay, tir
  from tvm.relay.backend.te_compiler import lower_to_primfunc
  from tvm.tir.tensor_intrin.hexagon import VRMPY_u8u8i32_INTRIN
label_desc: Use image node_id 15 for calling a global function with example usage: test_loop_carried_dependency() without return types