code: def test_matmul(M, N, K, dtype_str):
      stream = torch.cuda.Stream()
      torch.cuda.set_stream(stream)
      if dtype_str in ["float32", "int8"] and DEVICE_NAME != "a100":
          pytest.skip("Only test float32 & int8 on a100")
      if (M, N, K) in [
          (64, 4096, 4096),
          (64, 8192, 8192),
          (8192, 64, 8192),
      ] and dtype_str == "float32":
          pytest.skip("Out of shared memory in float32")
      dtype = {
          "float16": torch.float16,
          "float32": torch.float32,
          "int8": torch.int8,
      }[dtype_str]
      torch.manual_seed(0)
      ref_gpu_util = matmul_data[DEVICE_NAME][(M, N, K)][dtype_str]
      cur_sm_clock = nvsmi(["clocks.current.sm"])[0]
      max_gpu_perf = get_max_tensorcore_tflops(
          dtype, clock_rate=cur_sm_clock * 1e3
      )
      if dtype == torch.int8:
          a = torch.randint(
              -128, 127, (M, K), dtype=dtype, device="cuda"
          )
          b = torch.randint(
              -128, 127, (N, K), dtype=dtype, device="cuda"
          )
          b = b.t()  # only test row-col layout
      else:
          a = torch.randn((M, K), dtype=dtype, device="cuda")
          b = torch.randn((K, N), dtype=dtype, device="cuda")
      fn = lambda: triton.ops.matmul(a, b)
      ms = triton.testing.do_bench_cudagraph(fn)
      cur_gpu_perf = 2.0 * M * N * K / ms * 1e-9
      cur_gpu_util = cur_gpu_perf / max_gpu_perf
      print_perf(ms, cur_gpu_util, ref_gpu_util)
      triton.testing.assert_close(
          cur_gpu_util, ref_gpu_util, atol=0.02, rtol=0.01
      )
  
file_path: triton/python/test/regression/test_performance.py
imports: import pytest
  import torch
  import triton
  import triton.language
  import triton.ops
  from triton.testing import get_dram_gbps, get_max_tensorcore_tflops, nvsmi
label_desc: Use image node_id 2 for calling a global function with example usage: test_matmul(M, N, K, dtype_str) without return types