code: def _kernel(
      A,
      B,
      C,
      M,
      N,
      K,  #
      stride_am,
      stride_ak,  #
      stride_bk,
      stride_bn,  #
      stride_cm,
      stride_cn,  #
      acc_dtype: tl.constexpr,  #
      allow_tf32: tl.constexpr,  #
      fp8_fast_accum: tl.constexpr,  #
      BLOCK_M: tl.constexpr,
      BLOCK_N: tl.constexpr,
      BLOCK_K: tl.constexpr,  #
      GROUP_M: tl.constexpr,
      SPLIT_K: tl.constexpr,
      EVEN_K: tl.constexpr,
      AB_DTYPE: tl.constexpr,  #
  ):
      # matrix multiplication
      pid = tl.program_id(0)
      pid_z = tl.program_id(1)
      grid_m = tl.cdiv(M, BLOCK_M)
      grid_n = tl.cdiv(N, BLOCK_N)
      # re-order program ID for better L2 performance
      width = GROUP_M * grid_n
      group_id = pid // width
      group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
      pid_m = group_id * GROUP_M + (pid % group_size)
      pid_n = (pid % width) // (group_size)
      # do matrix multiplication
      rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
      rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
      ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
      rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
      rk = pid_z * BLOCK_K + tl.arange(0, BLOCK_K)
      # pointers
      A = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
      B = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
      acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=acc_dtype)
      for k in range(0, tl.cdiv(K, BLOCK_K * SPLIT_K)):
          if EVEN_K:
              a = tl.load(A)
              b = tl.load(B)
          else:
              k_remaining = K - k * (BLOCK_K * SPLIT_K)
              _0 = tl.zeros((1, 1), dtype=C.dtype.element_ty)
              a = tl.load(A, mask=rk[None, :] < k_remaining, other=_0)
              b = tl.load(B, mask=rk[:, None] < k_remaining, other=_0)
          if AB_DTYPE is not None:
              a = a.to(AB_DTYPE)
              b = b.to(AB_DTYPE)
          if fp8_fast_accum:
              acc = tl.dot(
                  a, b, acc, out_dtype=acc_dtype, allow_tf32=allow_tf32
              )
          else:
              acc += tl.dot(
                  a, b, out_dtype=acc_dtype, allow_tf32=allow_tf32
              )
          A += BLOCK_K * SPLIT_K * stride_ak
          B += BLOCK_K * SPLIT_K * stride_bk
      acc = acc.to(C.dtype.element_ty)
      # rematerialize rm and rn to save registers
      rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
      rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
      C = C + (rm[:, None] * stride_cm + rn[None, :] * stride_cn)
      mask = (rm < M)[:, None] & (rn < N)[None, :]
      # handles write-back with reduction-splitting
      if SPLIT_K == 1:
          tl.store(C, acc, mask=mask)
      else:
          tl.atomic_add(C, acc, mask=mask)
  
file_path: triton/python/triton/ops/matmul.py
imports: import torch
  from ..None import Config, autotune, cdiv, heuristics, jit
  from ..None import language
  from .matmul_perf_model import early_config_prune, estimate_matmul_time
label_desc: Use image node_id 5 for calling a global function with example usage: _kernel(A, B, C, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, acc_dtype, allow_tf32, fp8_fast_accum, BLOCK_M, BLOCK_N, BLOCK_K, GROUP_M, SPLIT_K, EVEN_K, AB_DTYPE) without return types