#!POPCORN leaderboard amd-fp8-mm
#!POPCORN gpus MI300
from task import input_t, output_t # Assuming input_t is defined elsewhere
import torch
import triton
import triton.language as tl
from typing import Optional
import os
import sys

# --- REMOVED global flag ---

# Set target architecture
# os.environ['TRITON_ROCM_ARCH'] = 'gfx942:xnack-'

# torch.float8_e4m3fnuz -> tl.float8e4b8 -> __hip_fp8_e4m3_fnuz

# --- Autotuner Configuration (Using the same list as before) ---
@triton.autotune(
    configs=[
        # Baseline config
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        # Vary K
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        # Vary M
        triton.Config({'BLOCK_SIZE_M': 64,  'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        # Vary N (Needs N >= 256 in test case)
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        # Vary num_warps
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=4),
        # Vary num_stages
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4}, num_stages=3, num_warps=8),
        # Vary GROUP_SIZE_M
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=2, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 1}, num_stages=2, num_warps=8),
        # Combinations
        triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 64,  'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=2, num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 4}, num_stages=2, num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 4}, num_stages=3, num_warps=4),
    ],
    key=['M', 'N', 'K'],
)
@triton.heuristics(
    {
        "EVEN_K": lambda args: args["K"] % args["SCALE_BLOCK_K"] == 0,
    }
)
@triton.jit
def _gemm_fp8_blockscale_kernel(
    # --- Kernel arguments (identical to previous version) ---
    a_ptr, b_ptr, c_ptr, a_scale_ptr, b_scale_ptr,
    M, N, K,
    stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,
    stride_ascale_m, stride_ascale_k, stride_bscale_k, stride_bscale_n,
    SCALE_BLOCK_K: tl.constexpr, SCALE_BLOCK_N: tl.constexpr,
    BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr,
    BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr,
    EVEN_K: tl.constexpr,
):
    """Autotuned FP8 GEMM kernel."""
    # --- Static asserts (identical to previous version) ---
    tl.static_assert(SCALE_BLOCK_K == 128, "SCALE_BLOCK_K must be 128")
    tl.static_assert(SCALE_BLOCK_N == 128, "SCALE_BLOCK_N must be 128")
    tl.static_assert(SCALE_BLOCK_K % BLOCK_SIZE_K == 0, "SCALE_BLOCK_K must be divisible by BLOCK_SIZE_K")
    tl.static_assert(BLOCK_SIZE_N % SCALE_BLOCK_N == 0, "Requires BLOCK_N multiple of SCALE_N")

    # --- Program ID and Block Mapping (identical to previous version) ---
    GRID_M = tl.cdiv(M, BLOCK_SIZE_M)
    GRID_N = tl.cdiv(N, BLOCK_SIZE_N)
    GRID_MN = GRID_M * GRID_N
    NUM_XCDS: tl.constexpr = 8
    pid = tl.program_id(axis=0)
    num_pid_m = GRID_M
    num_pid_n = GRID_N
    # PID Remapping logic (identical)
    pids_per_xcd = (GRID_MN + NUM_XCDS - 1) // NUM_XCDS
    tall_xcds = GRID_MN % NUM_XCDS
    tall_xcds = NUM_XCDS if tall_xcds == 0 else tall_xcds
    xcd = pid % NUM_XCDS
    local_pid = pid // NUM_XCDS
    if xcd < tall_xcds: pid = xcd * pids_per_xcd + local_pid
    else: pid = tall_xcds * pids_per_xcd + (xcd - tall_xcds) * (pids_per_xcd - 1) + local_pid
    if GROUP_SIZE_M == 1:
        pid_m = pid // num_pid_n
        pid_n = pid % num_pid_n
    else:
        num_pid_in_group = GROUP_SIZE_M * num_pid_n
        group_id = pid // num_pid_in_group
        first_pid_m = group_id * GROUP_SIZE_M
        group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
        pid_m = first_pid_m + (pid % group_size_m)
        pid_n = (pid % num_pid_in_group) // group_size_m

    # --- Rest of the kernel logic remains identical ---
    offs_m_block = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
    offs_n_block = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
    offs_k_inner = tl.arange(0, BLOCK_SIZE_K)
    a_start_ptr = a_ptr + offs_m_block[:, None] * stride_am
    b_start_ptr = b_ptr + offs_n_block[None, :] * stride_bn
    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    num_k_scale_blocks = tl.cdiv(K, SCALE_BLOCK_K)
    INNER_K_ITERATIONS = SCALE_BLOCK_K // BLOCK_SIZE_K

    for k_scale_idx in range(num_k_scale_blocks):
        block_accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
        for k_inner_idx in range(INNER_K_ITERATIONS):
            k_offset = k_scale_idx * SCALE_BLOCK_K + k_inner_idx * BLOCK_SIZE_K
            current_k = offs_k_inner + k_offset
            a_ptrs = a_start_ptr + (current_k[None, :] * stride_ak)
            b_ptrs = b_start_ptr + (current_k[:, None] * stride_bk)
            a_mask = current_k[None, :] < K
            b_mask = current_k[:, None] < K
            a = tl.load(a_ptrs, mask=a_mask, other=0.0)
            b = tl.load(b_ptrs, mask=b_mask, other=0.0)
            block_accumulator += tl.dot(a, b)

        # Apply Scaling
        a_scale_ptr_k = a_scale_ptr + offs_m_block * stride_ascale_m + k_scale_idx * stride_ascale_k
        offs_n_scale_block = offs_n_block // SCALE_BLOCK_N
        b_scale_ptr_k = b_scale_ptr + k_scale_idx * stride_bscale_k + offs_n_scale_block * stride_bscale_n
        m_mask_scale = offs_m_block < M
        n_mask_scale = offs_n_block < N
        a_scale = tl.load(a_scale_ptr_k, mask=m_mask_scale, other=1.0)
        b_scale_load_ptrs = b_scale_ptr + k_scale_idx * stride_bscale_k + (offs_n_block // SCALE_BLOCK_N) * stride_bscale_n
        b_scale = tl.load(b_scale_load_ptrs, mask=n_mask_scale, other=1.0)
        scaled_block_result = block_accumulator * a_scale[:, None] * b_scale[None, :]
        accumulator += scaled_block_result

    # Write results
    c = accumulator.to(c_ptr.type.element_ty)
    offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
    offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
    c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
    c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
    tl.store(c_ptrs, c, mask=c_mask)


# --- Updated Python Wrapper - REMOVED global flag ---
def custom_kernel(data: input_t) -> output_t:
    """
    Computes FP8 GEMM Y = X * W using Autotuned Triton kernel.
    Prints the selected best config to stderr every time tuning is run for a new M,N,K.
    """
    # --- Unpacking, Assertions, Transpose (identical to previous version) ---
    a, b_in, a_scale, b_scale_in, c = data
    assert a.device.type == 'cuda' # etc...
    assert a.dtype == torch.float8_e4m3fnuz # etc...
    assert b_in.dtype == torch.float8_e4m3fnuz
    assert a_scale.dtype == torch.float32
    assert b_scale_in.dtype == torch.float32
    assert c.dtype == torch.bfloat16

    b = b_in.T
    b_scale = b_scale_in.T

    M, K = a.shape
    K_b, N = b.shape
    M_c, N_c = c.shape

    assert K == K_b, f"K dimension mismatch between a ({K}) and transposed b ({K_b})"
    assert M == M_c, f"M dimension mismatch between a ({M}) and c ({M_c})"
    assert N == N_c, f"N dimension mismatch between transposed b ({N}) and c ({N_c})"

    SCALE_BLOCK_K = 128
    SCALE_BLOCK_N = 128
    expected_a_scale_shape = (M, triton.cdiv(K, SCALE_BLOCK_K))
    expected_b_scale_shape = (triton.cdiv(K, SCALE_BLOCK_K), triton.cdiv(N, SCALE_BLOCK_N))
    assert a_scale.shape == expected_a_scale_shape
    assert b_scale.shape == expected_b_scale_shape
    # --- End Setup ---

    # --- Grid Calculation for Autotuner ---
    def grid(meta):
        GRID_M = triton.cdiv(M, meta['BLOCK_SIZE_M'])
        GRID_N = triton.cdiv(N, meta['BLOCK_SIZE_N'])
        return (GRID_M * GRID_N,)

    # --- Launch the Autotuned Kernel ---
    _gemm_fp8_blockscale_kernel[grid](
        a, b, c, a_scale, b_scale,
        M, N, K,
        a.stride(0), a.stride(1),
        b.stride(0), b.stride(1),
        c.stride(0), c.stride(1),
        a_scale.stride(0), a_scale.stride(1),
        b_scale.stride(0), b_scale.stride(1),
        SCALE_BLOCK_K=SCALE_BLOCK_K,
        SCALE_BLOCK_N=SCALE_BLOCK_N
    )

    # --- Print best config (runs every time custom_kernel is called) ---
    # This will print after the kernel is launched. If tuning happened for this M,N,K,
    # it will print the newly selected or cached best config.
    try:
        # Access the best_config attribute
        # Note: This might print the same config multiple times if the benchmark runner
        # calls custom_kernel multiple times for timing the *same* M,N,K test case.
        # However, it *will* print the specific config used for each *different* M,N,K.
        best_config = _gemm_fp8_blockscale_kernel.best_config
        print(f"AUTOTUNER_BEST_CONFIG (for M={M}, N={N}, K={K}): {best_config}", file=sys.stderr)
    except AttributeError:
        print(f"AUTOTUNER_BEST_CONFIG (for M={M}, N={N}, K={K}): Not available yet.", file=sys.stderr)
    except Exception as e:
         print(f"AUTOTUNER_BEST_CONFIG (for M={M}, N={N}, K={K}): Error accessing best_config - {e}", file=sys.stderr)
    # -----------------------------------------------------

    return c

python eval.py benchmark /tmp/tmpb8nuo10u
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
AUTOTUNER_BEST_CONFIG (for M=1024, N=1536, K=7168): BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64, GROUP_SIZE_M: 4, num_warps: 8, num_ctas: 1, num_stages: 2, maxnreg: None
[...] 1830 lines omitted