code: def test_async_pipelined_mma_gemm_simple():
      sch = get_mma_schedule()
  
      k0 = sch.get_loops(sch.get_block("C_o_update"))[3]
  
      sch.annotate(
          k0, ann_key="software_pipeline_stage", ann_val=[0, 0, 3]
      )
      sch.annotate(
          k0, ann_key="software_pipeline_order", ann_val=[0, 1, 2]
      )
      sch.annotate(
          k0, ann_key="software_pipeline_async_stages", ann_val=[0]
      )
  
      seq = tvm.transform.Sequential(
          [
              tvm.tir.transform.PlanAndUpdateBufferAllocationLocation(),
              tvm.tir.transform.ConvertBlocksToOpaque(),
              tvm.tir.transform.UnifyThreadBinding(),
              tvm.tir.transform.LowerMatchBuffer(),
              tvm.tir.transform.InjectSoftwarePipeline(),
          ]
      )
      mod = seq(sch.mod)
  
      pipeline = (
          mod["main"]
          .body.block.body.body.body.body.body.block.body[1]
          .block.body
      )
      prologue, body, epilogue = pipeline
  
      commit_queue_scope = prologue.block.body.body.block.body
      assert len(commit_queue_scope.body) == 2
      assert commit_queue_scope.value == 0
  
      commit_queue_scope = body.block.body.body[0].block.body
      assert len(commit_queue_scope.body) == 2
      assert commit_queue_scope.value == 0
  
      assert (
          body.block.body.body[1].block.body.body.attr_key
          == "async_wait_inflight_count"
      )
      assert body.block.body.body[1].block.body.body.value == 3
  
      assert (
          epilogue.block.body.body.block.body.body.attr_key
          == "async_wait_inflight_count"
      )
      assert (
          str(epilogue.block.body.body.block.body.body.value)
          == "2 - k_0_0"
      )
  
      build_and_run(sch)
  
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 44 for calling a global function with example usage: test_async_pipelined_mma_gemm_simple() without return types