code: def build_and_run(sch):
      if tvm.testing.is_ampere_or_newer():
          with tvm.transform.PassContext(
              config={"tir.use_async_copy": 1}
          ):
              f = tvm.build(sch.mod["main"], target="cuda")
  
          dev = tvm.device("cuda", 0)
          a_np = np.random.uniform(size=(N, K)).astype("float16")
          b_np = np.random.uniform(size=(K, M)).astype("float16")
          c_np = np.dot(a_np.astype("float32"), b_np.astype("float32"))
          a = tvm.nd.array(a_np, dev)
          b = tvm.nd.array(b_np, dev)
          c = tvm.nd.array(np.zeros((N, M), dtype="float32"), dev)
          f(a, b, c)
          tvm.testing.assert_allclose(c.numpy(), c_np, rtol=1e-3)
  
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 43 for calling a global function with example usage: build_and_run(sch) without return types