@triton.jit
def mac_loop(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, iters_per_tile, start_iter, end_iter,
    BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K:
    tl.constexpr, GROUP_SIZE_M: tl.constexpr):
    tile_id = start_iter // iters_per_tile
    remain_iters = start_iter % iters_per_tile
    if GROUP_SIZE_M > 0:
        pid_m, pid_n = swizzle_tile(tile_id, M, N, K, BLOCK_SIZE_M,
            BLOCK_SIZE_N, BLOCK_SIZE_K, GROUP_SIZE_M)
    else:
        pid_m, pid_n = linear_tile(tile_id, M, N, K, BLOCK_SIZE_M,
            BLOCK_SIZE_N, BLOCK_SIZE_K, GROUP_SIZE_M)
    a_ptr += BLOCK_SIZE_K * stride_ak * remain_iters
    a_block_ptr = tl.make_block_ptr(base=a_ptr, shape=(M, K), strides=(
        stride_am, stride_ak), offsets=(pid_m * BLOCK_SIZE_M, 0),
        block_shape=(BLOCK_SIZE_M, BLOCK_SIZE_K), order=(1, 0))
    b_ptr += BLOCK_SIZE_K * stride_bk * remain_iters
    b_block_ptr = tl.make_block_ptr(base=b_ptr, shape=(K, N), strides=(
        stride_bk, stride_bn), offsets=(0, pid_n * BLOCK_SIZE_N),
        block_shape=(BLOCK_SIZE_K, BLOCK_SIZE_N), order=(1, 0))
    acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for _ in range(start_iter, end_iter):
        a = tl.load(a_block_ptr, boundary_check=(0, 1))
        b = tl.load(b_block_ptr, boundary_check=(0, 1))
        acc += tl.dot(a, b)
        a_block_ptr = tl.advance(a_block_ptr, (0, BLOCK_SIZE_K))
        b_block_ptr = tl.advance(b_block_ptr, (BLOCK_SIZE_K, 0))
    if remain_iters == 0 and end_iter % iters_per_tile == 0:
        c_block_ptr = tl.make_block_ptr(base=c_ptr, shape=(M, N), strides=(
            stride_cm, stride_cn), offsets=(pid_m * BLOCK_SIZE_M, pid_n *
            BLOCK_SIZE_N), block_shape=(BLOCK_SIZE_M, BLOCK_SIZE_N), order=
            (1, 0))
        tl.store(c_block_ptr, acc, boundary_check=(0, 1))
    else:
        rm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
        rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
        c_ptr_ = c_ptr + rm[:, None] * stride_cm + rn[None, :] * stride_cn
        mask = (rm < M)[:, None] & (rn < N)[None, :]
        tl.atomic_add(c_ptr_, acc, mask=mask)
