@triton.autotune(configs=[triton.Config({'BLOCK_SIZE_M': 256,
    'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4, 'grf_mode':
    'large'}, num_stages=2, num_warps=32)], key=['M', 'N', 'K'])
@triton.jit
def first_wave(a_ptr, b_ptr, c_ptr, M: tl.constexpr, N: tl.constexpr, K: tl
    .constexpr, stride_am: tl.constexpr, stride_ak: tl.constexpr, stride_bk:
    tl.constexpr, stride_bn: tl.constexpr, stride_cm: tl.constexpr,
    stride_cn: tl.constexpr, full_tiles, partial_tiles, iters_per_tile,
    BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K:
    tl.constexpr, GROUP_SIZE_M: tl.constexpr):
    pid = tl.program_id(axis=0)
    start_iter = pid * full_tiles + tl.minimum(pid, partial_tiles)
    last_iter = (pid + 1) * full_tiles + tl.minimum(pid + 1, partial_tiles)
    while start_iter < last_iter:
        end_iter = start_iter + (iters_per_tile - start_iter % iters_per_tile)
        end_iter = tl.minimum(end_iter, last_iter)
        mac_loop(a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak,
            stride_bk, stride_bn, stride_cm, stride_cn, iters_per_tile,
            start_iter, end_iter, BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K,
            GROUP_SIZE_M)
        start_iter = end_iter
