code: def test_elementwise(N, dtype_str):
      stream = torch.cuda.Stream()
      torch.cuda.set_stream(stream)
      torch.manual_seed(0)
      if dtype_str in ["bfloat16"] and DEVICE_NAME != "a100":
          pytest.skip("Only test bfloat16 on a100")
      dtype = {
          "float16": torch.float16,
          "bfloat16": torch.bfloat16,
          "float32": torch.float32,
      }[dtype_str]
      ref_dtype_str = (
          "float16" if dtype_str == "bfloat16" else dtype_str
      )
      ref_gpu_util = elementwise_data[DEVICE_NAME][N][ref_dtype_str]
      max_gpu_perf = get_dram_gbps()
      z = torch.empty((N,), dtype=dtype, device="cuda")
      x = torch.randn_like(z)
      y = torch.randn_like(z)
      grid = lambda args: (triton.cdiv(N, args["BLOCK_SIZE"]),)
      fn = lambda: _add[grid](x, y, z, N, BLOCK_SIZE=1024)
      ms = triton.testing.do_bench_cudagraph(fn)
      cur_gpu_perf = 3.0 * N * z.element_size() / ms * 1e-6
      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 4 for calling a global function with example usage: test_elementwise(N, dtype_str) without return types