code: def test_1D_cascade_op_rolling_buffer():
      """The intermediate buffer must be allocated above rolling buffer's rolling loop,
      which is marked as opaque in consumer block's iter mappings."""
  
      @T.prim_func
      def before(
          A: T.Buffer((4, 16), "int32"), C: T.Buffer((4, 8), "int32")
      ):
          B = T.alloc_buffer((4, 6), "int32")
          for c in T.serial(4):
              for i in T.serial(0, 2):
                  for j in T.serial(0, 6):
                      for k in T.serial(3):
                          with T.block("P1"):
                              T.where(i < 1 or j >= 2)
                              cc, vi, vj, vk = T.axis.remap(
                                  "SSSR", [c, i, j, k]
                              )
                              if vk == 0:
                                  B[cc, T.floormod(vi * 4 + vj, 6)] = 0
                              B[cc, T.floormod(vi * 4 + vj, 6)] = (
                                  B[cc, T.floormod(vi * 4 + vj, 6)]
                                  + A[cc, vi * 4 + vj + vk]
                              )
                  for j in T.serial(0, 4):
                      for k in T.serial(3):
                          with T.block("P2"):
                              vi = T.axis.opaque(2, i)
                              cc, vj, vk = T.axis.remap(
                                  "SSR", [c, j, k]
                              )
                              if vk == 0:
                                  C[cc, vi * 4 + vj] = 0
                              C[cc, vi * 4 + vj] = (
                                  C[cc, vi * 4 + vj]
                                  + B[
                                      cc,
                                      T.floormod(vi * 4 + vj + vk, 6),
                                  ]
                              )
  
      @T.prim_func
      def after(
          A: T.Buffer((4, 16), "int32"), C: T.Buffer((4, 8), "int32")
      ):
          for c in T.serial(4):
              with T.block():
                  T.reads(A[c, 0:12], C[c, 0:8])
                  T.writes(C[c, 0:8])
                  B = T.alloc_buffer([4, 6], dtype="int32")
                  for i in T.serial(2):
                      for j, k in T.grid(6, 3):
                          with T.block("P1"):
                              T.where(i < 1 or j >= 2)
                              cc, vi, vj, vk = T.axis.remap(
                                  "SSSR", [c, i, j, k]
                              )
                              if vk == 0:
                                  B[cc, (vi * 4 + vj) % 6] = 0
                              B[cc, (vi * 4 + vj) % 6] = (
                                  B[cc, (vi * 4 + vj) % 6]
                                  + A[cc, vi * 4 + vj + vk]
                              )
                      for j, k in T.grid(4, 3):
                          with T.block("P2"):
                              vi = T.axis.opaque(2, i)
                              cc, vj, vk = T.axis.remap(
                                  "SSR", [c, j, k]
                              )
                              if vk == 0:
                                  C[cc, vi * 4 + vj] = 0
                              C[cc, vi * 4 + vj] = (
                                  C[cc, vi * 4 + vj]
                                  + B[cc, (vi * 4 + vj + vk) % 6]
                              )
  
      _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 16 for calling a global function with example usage: test_1D_cascade_op_rolling_buffer() without return types