@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 full_tiles(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, streamk_tiles, BLOCK_SIZE_M: tl.constexpr,
    BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M:
    tl.constexpr):
    tile_id = tl.program_id(axis=0) + streamk_tiles
    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_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_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(0, tl.cdiv(K, BLOCK_SIZE_K)):
        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))
    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))
