code: def test_op(
      BLOCK_M,
      BLOCK_N,
      BLOCK_K,
      SPLIT_K,
      NWARP,
      NSTAGE,
      M,
      N,
      K,
      AT,
      BT,
      ADTYPE,
      BDTYPE,
      ALLOW_TF32,
      F8_FASTACCUM,
      ACC_DTYPE,
      OUTPUT_DTYPE,
  ):
      capability = torch.cuda.get_device_capability()
      if capability[0] < 7:
          pytest.skip("Only test tl.dot() on devices with sm >= 70")
      if capability[0] < 8 and (
          ADTYPE == "bfloat16" or BDTYPE == "bfloat16"
      ):
          pytest.skip("Only test bfloat16 on devices with sm >= 80")
      if capability[0] < 9 and (
          ADTYPE == "float8e4nv" or BDTYPE == "float8e4nv"
      ):
          pytest.skip("Only test float8e4nv on devices with sm >= 90")
      if (
          ADTYPE == "bfloat16" or BDTYPE == "bfloat16"
      ) and SPLIT_K != 1:
          pytest.skip("bfloat16 matmuls don't allow split_k for now")
      torch.manual_seed(0)
      # nuke kernel decorators -- will set meta-parameters manually
      kwargs = {
          "BLOCK_M": BLOCK_M,
          "BLOCK_N": BLOCK_N,
          "BLOCK_K": BLOCK_K,
          "SPLIT_K": SPLIT_K,
      }
      pre_hook = (
          None if SPLIT_K == 1 else lambda nargs: nargs["C"].zero_()
      )
      configs = [
          triton.Config(
              kwargs=kwargs,
              num_warps=NWARP,
              num_stages=NSTAGE,
              pre_hook=pre_hook,
          )
      ]
      kernel = triton.ops._matmul.kernel
      kernel.configs = configs
      # kernel.run = kernel.run.run.run
  
      # get matrix shape
      M = BLOCK_M if M is None else M
      N = BLOCK_N if N is None else N
      K = BLOCK_K * SPLIT_K if K is None else K
  
      def is_fp8(dtype):
          return "float8" in dtype
  
      def f8_to_f16(x, dtype):
          @triton.jit
          def kernel(Y, X, N, BLOCK_SIZE: tl.constexpr):
              pid = tl.program_id(0)
              offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
              mask = offs < N
              x = tl.load(X + offs, mask=mask)
              tl.store(Y + offs, x, mask=mask)
  
          ret = torch.empty_strided(
              x.shape, x.stride(), dtype=torch.float16, device=x.device
          )
          grid = lambda META: (
              triton.cdiv(x.numel(), META["BLOCK_SIZE"]),
          )
          dtype = getattr(tl, dtype)
          kernel[grid](
              ret,
              triton.reinterpret(x, dtype),
              ret.numel(),
              BLOCK_SIZE=1024,
          )
          return ret
  
      def upcast_if_fp8(x, dtype):
          if is_fp8(dtype):
              return f8_to_f16(x, dtype)
          return x
  
      def init_input(m, n, dtype, acc_dtype):
          if "float8" in dtype:
              ewidth = {
                  "float8e4b15": 4,
                  "float8e4nv": 4,
                  "float8e5": 5,
              }[dtype]
              sign = (
                  torch.randint(
                      2, size=(m, n), device="cuda", dtype=torch.int8
                  )
                  * 128
              )
              val = (
                  torch.randint(
                      2**3 - 1,
                      size=(m, n),
                      device="cuda",
                      dtype=torch.int8,
                  )
                  << 7 - ewidth
              )
              return sign | val
          if dtype == "int8":
              return torch.randint(
                  -128, 127, (m, n), device="cuda", dtype=torch.int8
              )
          # Use small range of values to prevent numerical issues.
          min_exp = -4 if acc_dtype == "float16" else -10
          exponents = torch.randint(min_exp, 0, size=(m, n))
          ret = (2.0**exponents).to(getattr(torch, dtype)).to("cuda")
          return ret
  
      # allocate/transpose inputs
      a = init_input(M, K, ADTYPE, ACC_DTYPE)
      b = init_input(K, N, BDTYPE, ACC_DTYPE)
      a = a if not AT else a.T.contiguous().T
      b = b if not BT else b.T.contiguous().T
      # run test
      th_a = upcast_if_fp8(a, ADTYPE)
      th_b = upcast_if_fp8(b, BDTYPE)
      ab_dtype = triton.ops.get_higher_dtype(th_a.dtype, th_b.dtype)
      acc_dtype = getattr(torch, ACC_DTYPE) if ACC_DTYPE else ab_dtype
      output_dtype = (
          getattr(torch, OUTPUT_DTYPE) if OUTPUT_DTYPE else ab_dtype
      )
      th_c = torch.matmul(th_a.to(output_dtype), th_b.to(output_dtype))
      try:
          if is_fp8(ADTYPE):
              a = triton.reinterpret(a, getattr(tl, ADTYPE))
          if is_fp8(BDTYPE):
              b = triton.reinterpret(b, getattr(tl, BDTYPE))
          tt_c = triton.ops.matmul(
              a,
              b,
              acc_dtype if ACC_DTYPE else None,
              ALLOW_TF32,
              F8_FASTACCUM,
              output_dtype,
          )
          torch.testing.assert_close(th_c, tt_c)
      except triton.OutOfResources as e:
          pytest.skip(str(e))
  
imports: import itertools
  import pytest
  import torch
  import triton
  import triton.language
  import triton.ops
label_desc: Use image node_id 1 for calling a global function with example usage: test_op(BLOCK_M, BLOCK_N, BLOCK_K, SPLIT_K, NWARP, NSTAGE, M, N, K, AT, BT, ADTYPE, BDTYPE, ALLOW_TF32, F8_FASTACCUM, ACC_DTYPE, OUTPUT_DTYPE) and returns: unknown, ret, x, ret, f8_to_f16, unknown, torch