import torch
from task import input_t, output_t

import torch
import triton
import triton.language as tl

def is_cuda():
    return triton.runtime.driver.active.get_current_target().backend == "cuda"

def get_cuda_autotune_config():
    return [
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=3,
                      num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=5,
                      num_warps=2),
        triton.Config({'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}, num_stages=5,
                      num_warps=2),
        # Good config for fp8 inputs.
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=3,
                      num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=3,
                      num_warps=8),
        triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 128, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4),
        triton.Config({'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8}, num_stages=4,
                      num_warps=4)
    ]

def get_hip_autotune_config():
    return [
        # triton.Config(
        #     {'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 16, 'GROUP_SIZE_M': 1, 'waves_per_eu': 2},
        #     num_warps=4, num_stages=2),
        # triton.Config(
        #     {'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 16, 'GROUP_SIZE_M': 4, 'waves_per_eu': 2},
        #     num_warps=8, num_stages=2),
        # triton.Config(
        #     {'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 1, 'waves_per_eu': 2},
        #     num_warps=8, num_stages=2),
        # triton.Config(
        #     {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8, 'waves_per_eu': 3},
        #     num_warps=4, num_stages=2),
        # triton.Config(
        #     {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 1, 'waves_per_eu': 8},
        #     num_warps=4, num_stages=2),
        triton.Config(
            {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 16, 'GROUP_SIZE_M': 1, 'waves_per_eu': 8},
            num_warps=4, num_stages=2),
    ]

def get_autotune_config():
    if is_cuda():
        return get_cuda_autotune_config()
    else:
        return get_hip_autotune_config()

@triton.jit
def matmul_kernel(
    # Pointers to matrices
    A,
    B,
    C,
    As,
    Bs,
    # Matrix dimensions
    M, N, K,
    # The stride variables represent how much to increase the ptr by when moving by 1
    # element in a particular dimension. E.g. `stride_am` is how much to increase `a_ptr`
    # by to get the element one row down (A has M rows).
    # Block size for block-wise quantization
    group_n,
    group_k,
    # Stride for inputs and output
    stride_am, stride_ak,
    stride_bk, stride_bn,
    stride_cm, stride_cn,
    stride_As_m,
    stride_As_k,
    stride_Bs_k,
    stride_Bs_n,
    # Meta-parameters
    BLOCK_SIZE_M: tl.constexpr = 64,  # triton 内核的行块大小
    BLOCK_SIZE_N: tl.constexpr = 64,  # triton 内核的列块大小
    BLOCK_SIZE_K: tl.constexpr = 32,  # triton 内核的内积块大小
    GROUP_SIZE_M: tl.constexpr = 1,   # triton 内核的行分组大小
):
    """Kernel for computing the matmul C = A x B.
    A has shape (M, K), B has shape (K, N) and C has shape (M, N)
    """

    pid = tl.program_id(axis=0)
    num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
    num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
    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

    offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
    offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
    offs_k = tl.arange(0, BLOCK_SIZE_K)
    a_ptrs = A + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
    b_ptrs = B + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)

    As_ptrs = As + offs_am * stride_As_m
    offs_bsn = offs_bn // group_n
    Bs_ptrs = Bs + offs_bsn * stride_Bs_n

    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
        a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
        b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)

        k_start = k * BLOCK_SIZE_K
        offs_ks = k_start // group_k
        a_s = tl.load(As_ptrs + offs_ks * stride_As_k)
        b_s = tl.load(Bs_ptrs + offs_ks * stride_Bs_k)

        accumulator += tl.dot(a, b) * a_s[:, None] * b_s[None, :]
        a_ptrs += BLOCK_SIZE_K * stride_ak
        b_ptrs += BLOCK_SIZE_K * stride_bk

    c = accumulator.to(tl.bfloat16)

    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 + 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)

def dequantize_inputs(a, b, a_scale, b_scale, m, n, k, block_shape_n=128, block_shape_k=128):
    """
    将输入张量反量化为 float32。
    
    Args:
        a: [m, k] 形状的输入张量
        b: [n, k] 形状的输入张量
        a_scale: a 的缩放因子张量
        b_scale: b 的缩放因子张量
        m, n, k: 矩阵维度
        block_shape_n, block_shape_k: 缩放块大小
        
    Returns:
        a_dequantized: 反量化后的 a 张量
        b_dequantized: 反量化后的 b 张量
    """
    # 反量化 a
    a_scale_expanded = a_scale.unsqueeze(-1).repeat(1, 1, block_shape_k).reshape(m, a_scale.shape[1] * block_shape_k)[:, :k]
    a_dequantized = a.to(torch.float32) * a_scale_expanded
    
    # 反量化 b
    b_scale_expanded = (
        b_scale.view(-1, 1)
        .repeat(1, block_shape_n * block_shape_k)
        .view(b_scale.shape[0], b_scale.shape[1], block_shape_n, block_shape_k)
        .permute(0, 2, 1, 3)
        .reshape(b_scale.shape[0] * block_shape_n, b_scale.shape[1] * block_shape_k)
    )[:n, :k]
    b_dequantized = b.to(torch.float32) * b_scale_expanded
    
    return a_dequantized, b_dequantized

def compute_ref_output(a_fp8, b_fp8, a_scale, b_scale, m, n, k, block_shape_n=128, block_shape_k=128):
    a_dequantized, b_dequantized = dequantize_inputs(
        a_fp8, b_fp8, a_scale, b_scale, m, n, k, block_shape_n, block_shape_k
    )
    return (a_dequantized @ b_dequantized.T).to(torch.bfloat16)
    # return (a_fp8 @ b_fp8.T).to(torch.bfloat16)

def compute_native_output(a_dequantized, b_dequantized):
    return (a_dequantized @ b_dequantized.T).to(torch.bfloat16)

def custom_kernel(data: input_t) -> output_t:
    """
    Reference implementation of block-scale fp8 gemm using the matmul function.
    """
    a, b, a_scale, b_scale, c = data

    A = a.contiguous()
    B = b.contiguous()
    As = a_scale.contiguous()
    Bs = b_scale.contiguous()

    output_dtype = torch.bfloat16

    block_n = 128
    block_k = 128

    assert A.shape[-1] == B.shape[-1]
    assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous()
    assert triton.cdiv(A.shape[-1], block_k) == As.shape[-1]
    M = A.numel() // A.shape[-1]

    assert B.ndim == 2 and B.is_contiguous() and Bs.ndim == 2
    N, K = B.shape
    assert triton.cdiv(N, block_n) == Bs.shape[0]
    assert triton.cdiv(K, block_k) == Bs.shape[1]

    C_shape = A.shape[:-1] + (N,)
    C = A.new_empty(C_shape, dtype=output_dtype)

    # print("a:", A)
    # print("b:", B)
    # print("a_scale:", As)
    # print("b_scale:", Bs)

    grid = (triton.cdiv(M, 64) * triton.cdiv(N, 64), )
    matmul_kernel[grid](
        A, B, C,
        As, Bs,
        M, N, K,
        block_n, block_k,
        A.stride(-2),
        A.stride(-1),
        B.stride(1),
        B.stride(0),
        C.stride(-2),
        C.stride(-1),
        As.stride(-2),
        As.stride(-1),
        Bs.stride(1),
        Bs.stride(0),
    )
    
    print("finished custom kernel")
    return C

python eval.py benchmark /tmp/tmp7uontqcb
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
finished custom kernel
[...] 1204 lines omitted