code: def test_async_producer_interleaving():
      @T.prim_func
      def simple_compute(
          A: T.Buffer((16, 16), "float32"),
          B: T.Buffer((16, 16), "float32"),
          C: T.Buffer((16, 16), "float32"),
      ):
          for tx in T.thread_binding(0, 16, thread="threadIdx.x"):
              for i in range(16):
                  with T.block("compute"):
                      T.reads(A[tx, i])
                      T.writes(C[tx, i])
                      A_shared = T.alloc_buffer(
                          (16, 1), dtype="float32", scope="shared"
                      )
                      B_shared = T.alloc_buffer(
                          (16, 1), dtype="float32", scope="shared"
                      )
                      with T.block():
                          T.reads(A[tx, i])
                          T.writes(A_shared[tx, 0])
                          A_shared[tx, 0] = A[tx, i]
                      with T.block():
                          T.reads(B[tx, i])
                          T.writes(B_shared[tx, 0])
                          B_shared[tx, 0] = B[tx, i]
                      with T.block():
                          T.reads(A_shared[tx, 0], B_shared[tx, 0])
                          T.writes(C[tx, i])
                          C[tx, i] = A_shared[tx, 0] + B_shared[tx, 0]
  
      mod = tvm.IRModule.from_expr(
          simple_compute.with_attr("global_symbol", "main")
      )
      sch = tvm.tir.Schedule(mod)
  
      _, loop = sch.get_loops(sch.get_block("compute"))
      sch.annotate(
          loop, ann_key="software_pipeline_stage", ann_val=[0, 0, 3]
      )
      sch.annotate(
          loop, ann_key="software_pipeline_order", ann_val=[0, 2, 1]
      )
      sch.annotate(
          loop, ann_key="software_pipeline_async_stages", ann_val=[0]
      )
      mod = tvm.tir.transform.InjectSoftwarePipeline()(sch.mod)
  
      @T.prim_func
      def ref(
          A: T.Buffer((16, 16), "float32"),
          B: T.Buffer((16, 16), "float32"),
          C: T.Buffer((16, 16), "float32"),
      ) -> None:
          for tx in T.thread_binding(16, thread="threadIdx.x"):
              with T.block():
                  T.reads(A[tx, 0:16], B[tx, 0:16])
                  T.writes(C[tx, 0:16])
                  A_shared = T.alloc_buffer(
                      [4, 16, 1], dtype="float32", scope="shared"
                  )
                  B_shared = T.alloc_buffer(
                      [4, 16, 1], dtype="float32", scope="shared"
                  )
                  with T.block():
                      T.reads(A[tx, 0:3], B[tx, 0:3])
                      T.writes(
                          A_shared[0:3, tx, 0], B_shared[0:3, tx, 0]
                      )
                      for i in T.unroll(3):
                          with T.block():
                              T.where(i < 16)
                              T.reads(A[tx, i], B[tx, i])
                              T.writes(
                                  A_shared[i % 4, tx, 0],
                                  B_shared[i % 4, tx, 0],
                              )
                              with T.attr(
                                  0, "async_commit_queue_scope", 0
                              ):
                                  with T.attr(0, "async_scope", 1):
                                      A_shared[i % 4, tx, 0] = A[tx, i]
                                  with T.attr(0, "async_scope", 1):
                                      B_shared[i % 4, tx, 0] = B[tx, i]
                  with T.block():
                      T.reads(
                          A[tx, 3:16],
                          A_shared[0:4, tx, 0],
                          B_shared[0:4, tx, 0],
                          B[tx, 3:16],
                      )
                      T.writes(
                          A_shared[0:4, tx, 0],
                          C[tx, 0:13],
                          B_shared[0:4, tx, 0],
                      )
                      for i in T.serial(13):
                          with T.block():
                              T.where(i + 3 < 16)
                              T.reads(A[tx, i + 3])
                              T.writes(A_shared[(i + 3) % 4, tx, 0])
                              with T.attr(
                                  0, "async_commit_queue_scope", 0
                              ):
                                  with T.attr(0, "async_scope", 1):
                                      A_shared[(i + 3) % 4, tx, 0] = A[
                                          tx, i + 3
                                      ]
                          with T.block():
                              T.where(i + 3 - 3 < 16)
                              T.reads(
                                  A_shared[0:4, tx, 0],
                                  B_shared[0:4, tx, 0],
                              )
                              T.writes(C[tx, i - 3 + 3])
                              with T.attr(
                                  0, "async_wait_queue_scope", 0
                              ):
                                  with T.attr(
                                      0, "async_wait_inflight_count", 5
                                  ):
                                      C[tx, i - 3 + 3] = (
                                          A_shared[
                                              (i - 3 + 3) % 4, tx, 0
                                          ]
                                          + B_shared[
                                              (i - 3 + 3) % 4, tx, 0
                                          ]
                                      )
                          with T.block():
                              T.where(i + 3 < 16)
                              T.reads(B[tx, i + 3])
                              T.writes(B_shared[(i + 3) % 4, tx, 0])
                              with T.attr(
                                  0, "async_commit_queue_scope", 0
                              ):
                                  with T.attr(0, "async_scope", 1):
                                      B_shared[(i + 3) % 4, tx, 0] = B[
                                          tx, i + 3
                                      ]
                  with T.block():
                      T.reads(
                          A_shared[0:4, tx, 0], B_shared[0:4, tx, 0]
                      )
                      T.writes(C[tx, 13:16])
                      for i in T.unroll(3):
                          with T.block():
                              T.where(i + 16 - 3 < 16)
                              T.reads(
                                  A_shared[0:4, tx, 0],
                                  B_shared[0:4, tx, 0],
                              )
                              T.writes(C[tx, i - 3 + 16])
                              with T.attr(
                                  0, "async_wait_queue_scope", 0
                              ):
                                  with T.attr(
                                      0,
                                      "async_wait_inflight_count",
                                      2 - i,
                                  ):
                                      C[tx, i - 3 + 16] = (
                                          A_shared[
                                              (i - 3 + 16) % 4, tx, 0
                                          ]
                                          + B_shared[
                                              (i - 3 + 16) % 4, tx, 0
                                          ]
                                      )
  
      tvm.ir.assert_structural_equal(
          mod["main"], ref.with_attr("global_symbol", "main"), True
      )
  
imports: import sys
  import numpy
  import pytest
  import tvm
  import tvm.testing
  import tvm.tir.tensor_intrin.cuda
  from tvm import TVMError, te, tir
  from tvm.meta_schedule.testing import te_workload
  from tvm.script import tir
  from tvm.testing.tir import mma_schedule
  from tvm.tir.tensor_intrin.cuda import LDMATRIX_f16_A_DYN_INTRIN, LDMATRIX_f16_B_DYN_INTRIN, MMA_f16f16f32_INTRIN, MMA_fill_16x16_f32_INTRIN, MMA_store_16x16_f32_global_INTRIN, shared_16x16_to_ldmatrix_32x8_layout
label_desc: Use image node_id 40 for calling a global function with example usage: test_async_producer_interleaving() without return types