code: def benchmark(M, N, K, provider):
      a = torch.randn((M, K), device="cuda", dtype=torch.float16)
      b = torch.randn((K, N), device="cuda", dtype=torch.float16)
      if provider == "cublas":
          ms, min_ms, max_ms = triton.testing.do_bench(
              lambda: torch.matmul(a, b)
          )
      if provider == "triton":
          ms, min_ms, max_ms = triton.testing.do_bench(
              lambda: matmul(a, b)
          )
      if provider == "cublas + relu":
          torch_relu = torch.nn.ReLU(inplace=True)
          ms, min_ms, max_ms = triton.testing.do_bench(
              lambda: torch_relu(torch.matmul(a, b))
          )
      if provider == "triton + relu":
          ms, min_ms, max_ms = triton.testing.do_bench(
              lambda: matmul(a, b, activation="leaky_relu")
          )
      perf = lambda ms: 2 * M * N * K * 1e-12 / (ms * 1e-3)
      return perf(ms), perf(max_ms), perf(min_ms)
  
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 4 for calling a global function with example usage: benchmark(M, N, K, provider) and returns: perf,  perf,  perf