code: def test_simple_compute_async():
      mod = tvm.IRModule.from_expr(
          gen_simple_compute(1).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_async_stages", ann_val=[0]
      )
      mod = tvm.tir.transform.InjectSoftwarePipeline()(sch.mod)
  
      @T.prim_func
      def ref(
          A: T.Buffer((16, 16), "float32"),
          C: T.Buffer((16, 16), "float32"),
      ):
          for tx in T.thread_binding(16, thread="threadIdx.x"):
              with T.block():
                  T.reads(A[tx, 0:16])
                  T.writes(C[tx, 0:16])
                  B = T.alloc_buffer(
                      [2, 16, 1], dtype="float32", scope="shared"
                  )
                  with T.block():
                      T.reads(A[tx, 0])
                      T.writes(B[T.FloorMod(0, 2), tx, 0])
                      with T.attr(0, "async_commit_queue_scope", 0):
                          with T.attr(0, "async_scope", 1):
                              B[T.FloorMod(0, 2), tx, 0] = A[
                                  tx, 0
                              ] * T.float32(2)
                  with T.block():
                      T.reads(A[tx, 1:16], B[0:2, tx, 0])
                      T.writes(B[0:2, tx, 0], C[tx, 0:15])
                      for i in T.serial(15):
                          with T.block():
                              T.where(i + 1 < 16)
                              T.reads(A[tx, i + 1])
                              T.writes(B[(i + 1) % 2, tx, 0])
                              with T.attr(
                                  0, "async_commit_queue_scope", 0
                              ):
                                  with T.attr(0, "async_scope", 1):
                                      B[(i + 1) % 2, tx, 0] = A[
                                          tx, i + 1
                                      ] * T.float32(2)
                          with T.block():
                              T.where(i + 1 - 1 < 16)
                              T.reads(B[(i - 1 + 1) % 2, tx, 0])
                              T.writes(C[tx, i - 1 + 1])
                              with T.attr(
                                  0, "async_wait_queue_scope", 0
                              ):
                                  with T.attr(
                                      0, "async_wait_inflight_count", 1
                                  ):
                                      C[tx, i - 1 + 1] = B[
                                          (i - 1 + 1) % 2, tx, 0
                                      ] + T.float32(1)
                  with T.block():
                      T.reads(B[T.FloorMod(15, 2), tx, 0])
                      T.writes(C[tx, 15])
                      with T.attr(0, "async_wait_queue_scope", 0):
                          with T.attr(
                              0, "async_wait_inflight_count", 0
                          ):
                              C[tx, 15] = B[
                                  T.FloorMod(15, 2), tx, 0
                              ] + T.float32(1)
  
      tvm.ir.assert_structural_equal(
          mod["main"], ref.with_attr("global_symbol", "main"), True
      )
  
      mod = tvm.IRModule.from_expr(
          gen_simple_compute(3).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_async_stages", ann_val=[0]
      )
      mod = tvm.tir.transform.InjectSoftwarePipeline()(sch.mod)
  
      @T.prim_func
      def ref(
          A: 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])
                  T.writes(C[tx, 0:16])
                  B = T.alloc_buffer(
                      [4, 16, 1], dtype="float32", scope="shared"
                  )
                  with T.block():
                      T.reads(A[tx, 0:3])
                      T.writes(B[0:3, tx, 0])
                      for i in T.unroll(3):
                          with T.block():
                              T.where(i < 16)
                              T.reads(A[tx, i])
                              T.writes(B[i % 4, tx, 0])
                              T.attr(0, "async_commit_queue_scope", 0)
                              T.attr(0, "async_scope", 1)
                              B[i % 4, tx, 0] = A[tx, i] * T.float32(2)
                  with T.block():
                      T.reads(A[tx, 3:16], B[0:4, tx, 0])
                      T.writes(B[0:4, tx, 0], C[tx, 0:13])
                      for i in T.serial(13):
                          with T.block():
                              T.where(i + 3 < 16)
                              T.reads(A[tx, i + 3])
                              T.writes(B[(i + 3) % 4, tx, 0])
                              T.attr(0, "async_commit_queue_scope", 0)
                              T.attr(0, "async_scope", 1)
                              B[(i + 3) % 4, tx, 0] = A[
                                  tx, i + 3
                              ] * T.float32(2)
                          with T.block():
                              T.where(i + 3 - 3 < 16)
                              T.reads(B[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", 3
                                  ):
                                      C[tx, i - 3 + 3] = B[
                                          (i - 3 + 3) % 4, tx, 0
                                      ] + T.float32(1)
                  with T.block():
                      T.reads(B[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(B[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] = B[
                                          (i - 3 + 16) % 4, tx, 0
                                      ] + T.float32(1)
  
      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 39 for calling a global function with example usage: test_simple_compute_async() without return types