submission 34598
Arseni Ivanov · python · License unknown
Use it
Vendorable · source mirrored · license unknownView source →
No package. Vendor the mirrored source: 416 lines, June 9 Researcher Reciprocity License v1.0.
triton_fully_fused_large_kernel.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-trimul-34598?include=source"interfacepython
Compatibility
measured onNVIDIA B200
declared hardwareNVIDIA B200
architecturessm_100
dtypesfp32
Benchmark evidence
1 measurement across 1 GPU, fastest first.
Operation / workload
Hardware
Latency
Rank
Observed
Reported · How evidence levels are derived →
Source and license
sourceavailable
revision digestsha256:c1404e1fe811c5cbb81b2dccaf422d21de489c7c5c2fcc1fc6c2b5db67680043
license declaredunknown
license concludedunknown
authorsArseni Ivanov
imported2026-08-15
Techniques
Extracted from the mirrored source by pattern, never inferred. Each row cites its line.
mma
accumulator_4way += tl.dot(x_norm_tile, w_tile)num-warps = 4
LN_EPS=1e-5, **config_k1, num_warps=4, num_stages=2stages = 2
LN_EPS=1e-5, **config_k1, num_warps=4, num_stages=2Kernel source
triton_fully_fused_large_kernel.py416 lines
#!POPCORN leaderboard trimul
import torch
import torch.nn.functional as F
import triton
import triton.language as tl
# Set PyTorch flags for performance
torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cuda.matmul.allow_fp16_reduced_precision_reduction = True
# Note: The @triton.autotune decorators have been removed from all kernels below.
@triton.jit
def fused_ln_dual_matmul_kernel(
# Pointers (9)
X_ptr, W_4way_ptr, W_og_ptr, Mask_ptr, Norm_Weight_ptr, Norm_Bias_ptr,
OutLeft_ptr, OutRight_ptr, OutOG_ptr,
# Metadata (5)
M, H, K, s1, s2,
# Strides (16)
stride_x_m, stride_x_k,
stride_w4_k, stride_w4_n,
stride_wog_k, stride_wog_n,
stride_ol_bs, stride_ol_h, stride_ol_s1, stride_ol_s2,
stride_or_t_bs, stride_or_t_h, stride_or_t_s2, stride_or_t_s1,
stride_og_m, stride_og_h,
stride_mask_m, stride_mask_h,
# Constexpr (now passed as arguments from the host)
LN_EPS: tl.constexpr,
BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
GROUP_SIZE_M: tl.constexpr, H_CHUNK_SIZE: tl.constexpr,
):
# --- PID Mapping: Based on the LARGER 4*H problem ---
pid = tl.program_id(axis=0)
N_4way = 4 * H
num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
num_pid_n = tl.cdiv(N_4way, 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
# --- SHARED LayerNorm calculation (done only ONCE) ---
offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
m_mask = offs_m < M
x_rows_base_ptr = X_ptr + offs_m[:, None] * stride_x_m
mean = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)
for k_offset in range(0, K, BLOCK_SIZE_K):
k_chunk_offs = tl.arange(0, BLOCK_SIZE_K)
x_ptrs = x_rows_base_ptr + (k_offset + k_chunk_offs)[None, :]
k_mask = (k_offset + k_chunk_offs) < K
x_chunk = tl.load(x_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)
mean += tl.sum(x_chunk, axis=1)
mean /= K
var = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)
for k_offset in range(0, K, BLOCK_SIZE_K):
k_chunk_offs = tl.arange(0, BLOCK_SIZE_K)
x_ptrs = x_rows_base_ptr + (k_offset + k_chunk_offs)[None, :]
k_mask = (k_offset + k_chunk_offs) < K
x_chunk = tl.load(x_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)
x_centered = x_chunk - mean[:, None]
var += tl.sum(x_centered * x_centered, axis=1)
var /= K
rstd = 1.0 / tl.sqrt(var + LN_EPS)
# --- Matmul Loop 1: For the 4-Way Projections ---
offs_n_4way = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
offs_k = tl.arange(0, BLOCK_SIZE_K)
w_4way_ptrs_base = W_4way_ptr + (offs_n_4way[None, :] * stride_w4_n)
accumulator_4way = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
accumulator_og = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
offs_n_og = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
k_block_start = k * BLOCK_SIZE_K;
x_ptrs = x_rows_base_ptr + (k_block_start + offs_k)[None, :] * stride_x_k
w_ptrs = w_4way_ptrs_base + (k_block_start + offs_k)[:, None] * stride_w4_k
x_mask = (offs_m[:, None] < M) & ((k_block_start + offs_k)[None, :] < K)
w_mask = ((k_block_start + offs_k)[:, None] < K) & (offs_n_4way[None, :] < N_4way)
x_tile = tl.load(x_ptrs, mask=x_mask, other=0.0).to(tl.float32)
norm_w_ptrs = Norm_Weight_ptr + k_block_start + offs_k
norm_b_ptrs = Norm_Bias_ptr + k_block_start + offs_k
nw = tl.load(norm_w_ptrs, mask=(k_block_start + offs_k) < K, other=0.0)
nb = tl.load(norm_b_ptrs, mask=(k_block_start + offs_k) < K, other=0.0)
x_norm_tile = (x_tile - mean[:, None]) * rstd[:, None]
x_norm_tile = (x_norm_tile * nw[None, :] + nb[None, :]).to(tl.float16)
w_tile = tl.load(w_ptrs, mask=w_mask, other=0.0)
accumulator_4way += tl.dot(x_norm_tile, w_tile)
#Some threads should calclate out_gate
if pid_n * BLOCK_SIZE_N < H:
w_og_ptrs_base = W_og_ptr + (offs_n_og[None, :] * stride_wog_n)
w_ptrs = w_og_ptrs_base + (k_block_start + offs_k)[:, None] * stride_wog_k
w_mask = ((k_block_start + offs_k)[:, None] < K) & (offs_n_og[None, :] < H);
w_tile = tl.load(w_ptrs, mask=w_mask, other=0.0)
accumulator_og += tl.dot(x_norm_tile, w_tile)
if pid_n * BLOCK_SIZE_N < H:
og_out = tl.sigmoid(accumulator_og)
outg_ptrs = OutOG_ptr + offs_m[:, None] * stride_og_m + offs_n_og[None, :] * stride_og_h
og_mask = m_mask[:, None] & (offs_n_og[None, :] < H)
tl.store(outg_ptrs, og_out, mask=og_mask)
# --- Fusion Logic for 4-Way Part ---
acc_reshaped = tl.reshape(accumulator_4way, (BLOCK_SIZE_M, H_CHUNK_SIZE, 4))
role_idx = tl.arange(0, 4)[None, None, :]
left_proj = tl.sum(tl.where(role_idx == 0, acc_reshaped, 0.0), axis=2)
left_gate = tl.sum(tl.where(role_idx == 1, acc_reshaped, 0.0), axis=2)
right_proj = tl.sum(tl.where(role_idx == 2, acc_reshaped, 0.0), axis=2)
right_gate = tl.sum(tl.where(role_idx == 3, acc_reshaped, 0.0), axis=2)
offs_h_chunk = (pid_n * H_CHUNK_SIZE) + tl.arange(0, H_CHUNK_SIZE)
mask_ptrs = Mask_ptr + offs_m[:, None] * stride_mask_m + offs_h_chunk[None, :] * stride_mask_h
m_mask_h = m_mask[:, None] & (offs_h_chunk[None, :] < H)
mask_tile = tl.load(mask_ptrs, mask=m_mask_h, other=0.0)
left_out = left_proj * tl.sigmoid(left_gate) * mask_tile
right_out = right_proj * tl.sigmoid(right_gate) * mask_tile
s1s2 = s1 * s2
offs_b = offs_m // s1s2
offs_s1 = (offs_m % s1s2) // s2
offs_s2 = offs_m % s2
offs_b_2d = tl.reshape(offs_b, (BLOCK_SIZE_M, 1))
offs_h_2d = tl.reshape(offs_h_chunk, (1, H_CHUNK_SIZE))
offs_s1_2d = tl.reshape(offs_s1, (BLOCK_SIZE_M, 1))
offs_s2_2d = tl.reshape(offs_s2, (BLOCK_SIZE_M, 1))
outl_ptrs = OutLeft_ptr + (offs_b_2d * stride_ol_bs + offs_h_2d * stride_ol_h +
offs_s1_2d * stride_ol_s1 + offs_s2_2d * stride_ol_s2)
outr_ptrs_t = OutRight_ptr + (offs_b_2d * stride_or_t_bs + offs_h_2d * stride_or_t_h +
offs_s2_2d * stride_or_t_s2 + offs_s1_2d * stride_or_t_s1)
tl.store(outl_ptrs, left_out, mask=m_mask_h)
tl.store(outr_ptrs_t, right_out, mask=m_mask_h)
@triton.jit
def bmm_coalesced_kernel(
# Pointers
Left_ptr, Right_ptr, Out_ptr,
# Dimensions
bs, s1, s2, H,
# Strides
stride_l_bs, stride_l_h, stride_l_s1, stride_l_s2,
stride_r_bs, stride_r_h, stride_r_s2, stride_r_s1,
stride_o_bs, stride_o_h, stride_o_s1, stride_o_s2,
# Kernel parameters
BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
GROUP_SIZE_M: tl.constexpr,
):
# Grid and program IDs
pid = tl.program_id(axis=0)
num_pid_m = tl.cdiv(s1, BLOCK_SIZE_M)
num_pid_n = tl.cdiv(s1, 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
pid_bh = tl.program_id(axis=1)
pid_b = pid_bh // H
pid_h = pid_bh % H
offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
offs_k = tl.arange(0, BLOCK_SIZE_K)
left_ptrs_base = Left_ptr + pid_b * stride_l_bs + pid_h * stride_l_h
right_ptrs_base = Right_ptr + pid_b * stride_r_bs + pid_h * stride_r_h
accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
for k in range(0, tl.cdiv(s2, BLOCK_SIZE_K)):
k_start = k * BLOCK_SIZE_K
a_ptrs = left_ptrs_base + (offs_m[:, None] * stride_l_s1 + (k_start + offs_k[None, :]) * stride_l_s2)
b_ptrs = right_ptrs_base + ((k_start + offs_k[:, None]) * stride_r_s2 + offs_n[None, :] * stride_r_s1)
a_mask = (offs_m[:, None] < s1) & ((k_start + offs_k[None, :]) < s2)
b_mask = ((k_start + offs_k[:, None]) < s2) & (offs_n[None, :] < s1)
a = tl.load(a_ptrs, mask=a_mask, other=0.0)
b = tl.load(b_ptrs, mask=b_mask, other=0.0)
accumulator += tl.dot(a, b)
out_ptrs = Out_ptr + pid_b * stride_o_bs + pid_h * stride_o_h + \
offs_m[:, None] * stride_o_s1 + offs_n[None, :] * stride_o_s2
c_mask = (offs_m[:, None] < s1) & (offs_n[None, :] < s1)
tl.store(out_ptrs, accumulator, mask=c_mask)
@triton.jit
def fused_final_kernel(
# Pointers
In_ptr, Gate_ptr, NormW_ptr, NormB_ptr, ProjW_ptr, Out_ptr,
# Metadata
M, H, D, s1,
# Strides
stride_in_bs, stride_in_h, stride_in_s1_row, stride_in_s1_col,
stride_gate_m, stride_gate_h,
stride_proj_d, stride_proj_h,
stride_out_bs, stride_out_s1_row, stride_out_s1_col, stride_out_d,
# Constants
LN_EPS: tl.constexpr,
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)
num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
num_pid_n = tl.cdiv(D, 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_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
m_mask = offs_m < M
s1s1 = s1 * s1
b = offs_m // s1s1
r = (offs_m % s1s1) // s1
c = offs_m % s1
sum_x = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)
sum_x2 = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)
in_ptr_base = In_ptr + b * stride_in_bs + r * stride_in_s1_row + c * stride_in_s1_col
for k_offset in range(0, H, BLOCK_SIZE_K):
offs_k = k_offset + tl.arange(0, BLOCK_SIZE_K)
k_mask = offs_k < H
in_ptrs = in_ptr_base[:, None] + offs_k[None, :] * stride_in_h
in_chunk = tl.load(in_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0).to(tl.float32)
sum_x += tl.sum(in_chunk, axis=1)
sum_x2 += tl.sum(in_chunk * in_chunk, axis=1)
mean = sum_x / H
var = (sum_x2 / H) - (mean * mean)
rstd = tl.math.rsqrt(var + LN_EPS)
acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
for k_offset in range(0, H, BLOCK_SIZE_K):
offs_k = k_offset + tl.arange(0, BLOCK_SIZE_K)
k_mask = offs_k < H
in_ptrs = in_ptr_base[:, None] + offs_k[None, :] * stride_in_h
a = tl.load(in_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)
a_norm = (a - mean[:, None]) * rstd[:, None]
norm_w = tl.load(NormW_ptr + offs_k, mask=k_mask, other=0.0)
norm_b = tl.load(NormB_ptr + offs_k, mask=k_mask, other=0.0)
a_norm = a_norm * norm_w[None, :] + norm_b[None, :]
proj_ptrs = ProjW_ptr + offs_n[None, :] * stride_proj_d + offs_k[:, None] * stride_proj_h
gate_ptrs = Gate_ptr + offs_m[:, None] * stride_gate_m + offs_k[None, :] * stride_gate_h
gate = tl.load(gate_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)
a_gated = a_norm * gate
b_w = tl.load(proj_ptrs, mask=k_mask[:, None] & (offs_n[None, :] < D), other=0.0)
acc += tl.dot(a_gated.to(b_w.dtype), b_w)
offs_d = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
out_ptr_base = Out_ptr + b*stride_out_bs + r*stride_out_s1_row + c*stride_out_s1_col
out_ptrs = out_ptr_base[:, None] + offs_d[None, :] * stride_out_d
tl.store(out_ptrs, acc, mask=m_mask[:, None] & (offs_d[None, :] < D))
def compiledtrimul_fused_interleaved_final(
x: torch.Tensor,
mask_mh: torch.Tensor,
norm_weight: torch.Tensor,
norm_bias: torch.Tensor,
W_4way: torch.Tensor,
W_og: torch.Tensor,
to_out_norm_weight: torch.Tensor,
to_out_norm_bias: torch.Tensor,
to_out_weight: torch.Tensor,
h: int,
):
bs, s1, s2, d = x.shape
M, K, H = bs * s1 * s2, x.shape[-1], h
x_flat = x.view(M, K)
left_final = torch.empty((bs, H, s1, s2), device=x.device, dtype=torch.float16)
right_final_t = torch.empty((bs, H, s2, s1), device=x.device, dtype=torch.float16)
og_mh = torch.empty((M, H), device=x.device, dtype=torch.float16)
# --- Kernel 1: Fused LN + Dual Matmul ---
# The grid is launched for the larger 4*H problem
N_4way = 4 * H
# Hardcoded best config from logs: M64-N128-K64-GM8-HC32-W4-S2
config_k1 = {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8, 'H_CHUNK_SIZE': 32}
grid = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE_M']) * triton.cdiv(N_4way, meta['BLOCK_SIZE_N']),)
fused_ln_dual_matmul_kernel[grid](
x_flat, W_4way, W_og, mask_mh, norm_weight, norm_bias,
left_final, right_final_t, og_mh,
M, H, K, s1, s2,
x_flat.stride(0), x_flat.stride(1), W_4way.stride(0), W_4way.stride(1),
W_og.stride(0), W_og.stride(1), left_final.stride(0), left_final.stride(1),
left_final.stride(2), left_final.stride(3), right_final_t.stride(0), right_final_t.stride(1),
right_final_t.stride(2), right_final_t.stride(3), og_mh.stride(0), og_mh.stride(1),
mask_mh.stride(0), mask_mh.stride(1),
LN_EPS=1e-5, **config_k1, num_warps=4, num_stages=2
)
# --- Kernel 2: Batched Matrix Multiplication ---
bmm_out_tmp = torch.empty((bs, H, s1, s1), device=x.device, dtype=torch.float16)
# Hardcoded best config from logs: M128-N128-K32-GM8-W8-S3
config_k2 = {'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}
grid_bmm = lambda meta: (triton.cdiv(s1, meta['BLOCK_SIZE_M']) * triton.cdiv(s1, meta['BLOCK_SIZE_N']), bs * H)
bmm_coalesced_kernel[grid_bmm](
left_final, right_final_t, bmm_out_tmp,
bs, s1, s2, H,
left_final.stride(0), left_final.stride(1), left_final.stride(2), left_final.stride(3),
right_final_t.stride(0), right_final_t.stride(1), right_final_t.stride(2), right_final_t.stride(3),
bmm_out_tmp.stride(0), bmm_out_tmp.stride(1), bmm_out_tmp.stride(2), bmm_out_tmp.stride(3),
**config_k2, num_warps=8, num_stages=3
)
# --- Kernel 3: Fully Fused Final Stage ---
final_out = torch.empty((bs, s1, s1, d), device=x.device, dtype=torch.float16)
# Hardcoded best config from logs: M32-N128-K32-GM8-W4-S3
config_k3 = {'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}
grid_final = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE_M']) * triton.cdiv(d, meta['BLOCK_SIZE_N']),)
fused_final_kernel[grid_final](
bmm_out_tmp, og_mh, to_out_norm_weight, to_out_norm_bias, to_out_weight, final_out,
M, H, d, s1,
bmm_out_tmp.stride(0), bmm_out_tmp.stride(1), bmm_out_tmp.stride(2), bmm_out_tmp.stride(3),
og_mh.stride(0), og_mh.stride(1), to_out_weight.stride(0), to_out_weight.stride(1),
final_out.stride(0), final_out.stride(1), final_out.stride(2), final_out.stride(3),
LN_EPS=1e-5, **config_k3, num_warps=4, num_stages=3
)
return final_out
def pack_w_4way_efficient(weights):
""" Packs L, LG, R, RG into a tight [K, 4*H] matrix. """
WL, WLG, WR, WRG = (weights[k] for k in ['left_proj.weight', 'left_gate.weight', 'right_proj.weight', 'right_gate.weight'])
H, K = WL.shape
ws = torch.stack([WL, WLG, WR, WRG], dim=0).permute(1, 0, 2).contiguous().view(4 * H, K)
return ws.t().to(torch.float16)
def get_w_og(weights):
""" Gets the transposed [K, H] out_gate weight matrix. """
return weights['out_gate.weight'].t().to(torch.float16)
@torch.compile()
def compiledtrimul(
x: torch.Tensor, mask: torch.Tensor, norm_weight: torch.Tensor, norm_bias: torch.Tensor,
w_concat: torch.Tensor, to_out_norm_weight: torch.Tensor, to_out_norm_bias: torch.Tensor,
to_out_weight: torch.Tensor, h: int
) -> torch.Tensor:
bs, s1, s2, d = x.shape
x_norm = F.layer_norm(x, (d,), norm_weight, norm_bias).view((bs * s1 * s2, d)).to(torch.float16)
all_projections = torch.mm(x_norm, w_concat)
left, right, lg, rg, og = all_projections.chunk(5, dim=1)
mask_expanded = mask.expand(-1, -1, -1, h).reshape(-1, h)
left = left * mask_expanded * torch.sigmoid(lg)
right = right * mask_expanded * torch.sigmoid(rg)
out_gate = torch.sigmoid(og)
left = left.view(bs, s1, s2, h).permute(0,3,1,2)
right = right.view(bs, s1, s2, h).permute(0,3,1,2)
out_p = torch.matmul(left.to(torch.float16), right.to(torch.float16).transpose(-1, -2))
out_einsum_flat = out_p.permute(0,2,3,1).reshape(bs * s1 * s1, h)
normed = F.layer_norm(out_einsum_flat, (h,), to_out_norm_weight, to_out_norm_bias).to(torch.float16)
gated = normed * out_gate
final_out_flat = gated @ to_out_weight.t()
return final_out_flat.view(bs, s1, s1, d)
def small_kernel_pt_path(data):
input_tensor, mask, weights, config = data
w_concat = torch.cat([
weights['left_proj.weight'], weights['right_proj.weight'], weights['left_gate.weight'],
weights['right_gate.weight'], weights['out_gate.weight']
], dim=0).t().contiguous().to(torch.float16)
return compiledtrimul(
x=input_tensor.to(torch.float32), mask=mask.unsqueeze(-1),
norm_weight=weights['norm.weight'].to(torch.float32),
norm_bias=weights['norm.bias'].to(torch.float32), w_concat=w_concat,
to_out_norm_weight=weights['to_out_norm.weight'].to(torch.float16),
to_out_norm_bias=weights['to_out_norm.bias'].to(torch.float16),
to_out_weight=weights['to_out.weight'].to(torch.float16),
h=config["hidden_dim"]
)
def custom_kernel(data):
input_tensor, mask, weights, config = data
bs, s1, s2, d = input_tensor.shape
if s1 < 800:
return small_kernel_pt_path(data)
H = config["hidden_dim"]
W_4way = pack_w_4way_efficient(weights)
W_og = get_w_og(weights)
M = bs * s1 * s2
mask_mh = mask.unsqueeze(-1).expand(-1, -1, -1, H).reshape(M, H).to(torch.float16)
return compiledtrimul_fused_interleaved_final(
x=input_tensor.to(torch.float32),
mask_mh=mask_mh,
norm_weight=weights['norm.weight'].to(torch.float32),
norm_bias=weights['norm.bias'].to(torch.float32),
W_4way=W_4way,
W_og=W_og,
to_out_norm_weight=weights['to_out_norm.weight'].to(torch.float16),
to_out_norm_bias=weights['to_out_norm.bias'].to(torch.float16),
to_out_weight=weights['to_out.weight'].to(torch.float16),
h=H,
)
scrolls · 416 lines total
Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0
Changes from previous submission
Against this author's previous submission submission 33982.
- \x2321504f50434f524e206c6561646572626f617264207472696d756c0a696d706f727420746f7263680a696d706f727420746f7263682e6e6e2e66756e6374696f6e616c20617320460a696d706f727420747269746f6e0a696d706f727420747269746f6e2e6c616e677561676520617320746c0a66726f6d207461736b20696d706f727420696e7075745f742c206f75747075745f740a746f7263682e6261636b656e64732e637564612e6d61746d756c2e616c6c6f775f74663332203d20547275650a746f7263682e6261636b656e64732e637564612e6d61746d756c2e616c6c6f775f667031365f726564756365645f707265636973696f6e5f726564756374696f6e203d20547275650a0a40747269746f6e2e6175746f74756e65280a20202020636f6e666967733d5b0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c202027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a2036342c202027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2031367d2c20206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a20202020202020200a20202020202020202320436f6e66696775726174696f6e732077697468206c617267657220626c6f636b2073697a657320666f722062657474657220646174612072657573650a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d382c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203235362c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2036347d2c206e756d5f77617270733d382c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203235362c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d382c206e756d5f7374616765733d34292c0a0a20202020202020202320436f6e66696775726174696f6e73207769746820646565706572204b2064696d656e73696f6e0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d342c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c202027424c4f434b5f53495a455f4e273a2036342c202027424c4f434b5f53495a455f4b273a203132382c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2031367d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a0a202020202020202023204d6f72652065787472656d6520636f6e66696775726174696f6e7320746f207465737420746865206c696d6974730a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203235362c2027424c4f434b5f53495a455f4e273a2036342c202027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2031367d2c206e756d5f77617270733d342c206e756d5f7374616765733d35292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c202027424c4f434b5f53495a455f4e273a203235362c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2036347d2c206e756d5f77617270733d342c206e756d5f7374616765733d35292c0a0a20202020202020202320436f6e66696775726174696f6e7320776974682066657765722077617270730a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c202027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20382c2027485f4348554e4b5f53495a45273a2033327d2c206e756d5f77617270733d322c206e756d5f7374616765733d34292c0a202020205d2c0a202020206b65793d5b274d272c20274e272c20274b275d2c0a290a40747269746f6e2e6a69740a6465662066757365645f6c6e5f6475616c5f6d61746d756c5f6b65726e656c280a202020202320506f696e74657273202839290a20202020585f7074722c20575f347761795f7074722c20575f6f675f7074722c204d61736b5f7074722c204e6f726d5f5765696768745f7074722c204e6f726d5f426961735f7074722c0a202020204f75744c6566745f7074722c204f757452696768745f7074722c204f75744f475f7074722c0a2020202023204d65746164617461202835290a202020204d2c20482c204b2c2073312c2073322c0a2020202023205374726964657320283136290a202020207374726964655f785f6d2c207374726964655f785f6b2c0a202020207374726964655f77345f6b2c207374726964655f77345f6e2c0a202020207374726964655f776f675f6b2c207374726964655f776f675f6e2c0a202020207374726964655f6f6c5f62732c207374726964655f6f6c5f682c207374726964655f6f6c5f73312c207374726964655f6f6c5f73322c0a202020207374726964655f6f725f745f62732c207374726964655f6f725f745f682c207374726964655f6f725f745f73322c207374726964655f6f725f745f73312c0a202020207374726964655f6f675f6d2c207374726964655f6f675f682c0a202020207374726964655f6d61736b5f6d2c207374726964655f6d61736b5f682c0a202020202320436f6e737465787072202866726f6d206465636f7261746f7220616e64206b7761726773290a202020204c4e5f4550533a20746c2e636f6e7374657870722c0a20202020424c4f434b5f53495a455f4d3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4e3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4b3a20746c2e636f6e7374657870722c0a2020202047524f55505f53495a455f4d3a20746c2e636f6e7374657870722c20485f4348554e4b5f53495a453a20746c2e636f6e7374657870722c0a293a0a2020202023202d2d2d20504944204d617070696e673a204261736564206f6e20746865204c415247455220342a482070726f626c656d202d2d2d0a20202020706964203d20746c2e70726f6772616d5f696428617869733d30290a202020204e5f34776179203d2034202a20480a202020206e756d5f7069645f6d203d20746c2e63646976284d2c20424c4f434b5f53495a455f4d290a202020206e756d5f7069645f6e203d20746c2e63646976284e5f347761792c20424c4f434b5f53495a455f4e290a202020206e756d5f7069645f696e5f67726f7570203d2047524f55505f53495a455f4d202a206e756d5f7069645f6e0a2020202067726f75705f6964203d20706964202f2f206e756d5f7069645f696e5f67726f75700a2020202066697273745f7069645f6d203d2067726f75705f6964202a2047524f55505f53495a455f4d0a2020202067726f75705f73697a655f6d203d206d696e286e756d5f7069645f6d202d2066697273745f7069645f6d2c2047524f55505f53495a455f4d290a202020207069645f6d203d2066697273745f7069645f6d202b202870696420252067726f75705f73697a655f6d290a202020207069645f6e203d20287069642025206e756d5f7069645f696e5f67726f757029202f2f2067726f75705f73697a655f6d0a202020200a2020202023202d2d2d20534841524544204c617965724e6f726d2063616c63756c6174696f6e2028646f6e65206f6e6c79204f4e434529202d2d2d0a202020206f6666735f6d203d207069645f6d202a20424c4f434b5f53495a455f4d202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4d290a202020206d5f6d61736b203d206f6666735f6d203c204d0a20202020785f726f77735f626173655f707472203d20585f707472202b206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f785f6d0a0a202020206d65616e203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c292c2064747970653d746c2e666c6f61743332290a20202020666f72206b5f6f666673657420696e2072616e676528302c204b2c20424c4f434b5f53495a455f4b293a0a20202020202020206b5f6368756e6b5f6f666673203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a2020202020202020785f70747273203d20785f726f77735f626173655f707472202b20286b5f6f6666736574202b206b5f6368756e6b5f6f666673295b4e6f6e652c203a5d0a20202020202020206b5f6d61736b203d20286b5f6f6666736574202b206b5f6368756e6b5f6f66667329203c204b0a2020202020202020785f6368756e6b203d20746c2e6c6f616428785f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a20202020202020206d65616e202b3d20746c2e73756d28785f6368756e6b2c20617869733d31290a202020206d65616e202f3d204b0a0a20202020766172203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c292c2064747970653d746c2e666c6f61743332290a20202020666f72206b5f6f666673657420696e2072616e676528302c204b2c20424c4f434b5f53495a455f4b293a0a20202020202020206b5f6368756e6b5f6f666673203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a2020202020202020785f70747273203d20785f726f77735f626173655f707472202b20286b5f6f6666736574202b206b5f6368756e6b5f6f666673295b4e6f6e652c203a5d0a20202020202020206b5f6d61736b203d20286b5f6f6666736574202b206b5f6368756e6b5f6f66667329203c204b0a2020202020202020785f6368756e6b203d20746c2e6c6f616428785f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a2020202020202020785f63656e7465726564203d20785f6368756e6b202d206d65616e5b3a2c204e6f6e655d0a2020202020202020766172202b3d20746c2e73756d28785f63656e7465726564202a20785f63656e74657265642c20617869733d31290a20202020766172202f3d204b0a2020202072737464203d20312e30202f20746c2e7371727428766172202b204c4e5f455053290a0a2020202023202d2d2d204d61746d756c204c6f6f7020313a20466f722074686520342d5761792050726f6a656374696f6e73202d2d2d0a202020206f6666735f6e5f34776179203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206f6666735f6b203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a20202020775f347761795f707472735f62617365203d20575f347761795f707472202b20286f6666735f6e5f347761795b4e6f6e652c203a5d202a207374726964655f77345f6e290a20202020616363756d756c61746f725f34776179203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a20202020616363756d756c61746f725f6f67203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a202020200a202020206f6666735f6e5f6f67203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a20202020666f72206b20696e2072616e676528302c20746c2e63646976284b2c20424c4f434b5f53495a455f4b29293a0a20202020202020206b5f626c6f636b5f7374617274203d206b202a20424c4f434b5f53495a455f4b3b0a2020202020202020785f70747273203d20785f726f77735f626173655f707472202b20286b5f626c6f636b5f7374617274202b206f6666735f6b295b4e6f6e652c203a5d202a207374726964655f785f6b0a2020202020202020775f70747273203d20775f347761795f707472735f62617365202b20286b5f626c6f636b5f7374617274202b206f6666735f6b295b3a2c204e6f6e655d202a207374726964655f77345f6b0a2020202020202020785f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c204d2920262028286b5f626c6f636b5f7374617274202b206f6666735f6b295b4e6f6e652c203a5d203c204b290a2020202020202020775f6d61736b203d2028286b5f626c6f636b5f7374617274202b206f6666735f6b295b3a2c204e6f6e655d203c204b29202620286f6666735f6e5f347761795b4e6f6e652c203a5d203c204e5f34776179290a2020202020202020785f74696c65203d20746c2e6c6f616428785f707472732c206d61736b3d785f6d61736b2c206f746865723d302e30292e746f28746c2e666c6f61743332290a20202020202020206e6f726d5f775f70747273203d204e6f726d5f5765696768745f707472202b206b5f626c6f636b5f7374617274202b206f6666735f6b0a20202020202020206e6f726d5f625f70747273203d204e6f726d5f426961735f707472202b206b5f626c6f636b5f7374617274202b206f6666735f6b0a20202020202020206e77203d20746c2e6c6f6164286e6f726d5f775f707472732c206d61736b3d286b5f626c6f636b5f7374617274202b206f6666735f6b29203c204b2c206f746865723d302e30290a20202020202020206e62203d20746c2e6c6f6164286e6f726d5f625f707472732c206d61736b3d286b5f626c6f636b5f7374617274202b206f6666735f6b29203c204b2c206f746865723d302e30290a2020202020202020785f6e6f726d5f74696c65203d2028785f74696c65202d206d65616e5b3a2c204e6f6e655d29202a20727374645b3a2c204e6f6e655d0a2020202020202020785f6e6f726d5f74696c65203d2028785f6e6f726d5f74696c65202a206e775b4e6f6e652c203a5d202b206e625b4e6f6e652c203a5d292e746f28746c2e666c6f61743136290a2020202020202020775f74696c65203d20746c2e6c6f616428775f707472732c206d61736b3d775f6d61736b2c206f746865723d302e30290a2020202020202020616363756d756c61746f725f34776179202b3d20746c2e646f7428785f6e6f726d5f74696c652c20775f74696c65290a0a202020202020202023536f6d6520746872656164732073686f756c642063616c636c617465206f75745f676174650a20202020202020206966207069645f6e202a20424c4f434b5f53495a455f4e203c20483a0a202020202020202020202020775f6f675f707472735f62617365203d20575f6f675f707472202b20286f6666735f6e5f6f675b4e6f6e652c203a5d202a207374726964655f776f675f6e290a202020202020202020202020775f70747273203d20775f6f675f707472735f62617365202b20286b5f626c6f636b5f7374617274202b206f6666735f6b295b3a2c204e6f6e655d202a207374726964655f776f675f6b0a202020202020202020202020775f6d61736b203d2028286b5f626c6f636b5f7374617274202b206f6666735f6b295b3a2c204e6f6e655d203c204b29202620286f6666735f6e5f6f675b4e6f6e652c203a5d203c2048293b0a202020202020202020202020775f74696c65203d20746c2e6c6f616428775f707472732c206d61736b3d775f6d61736b2c206f746865723d302e30290a202020202020202020202020616363756d756c61746f725f6f67202b3d20746c2e646f7428785f6e6f726d5f74696c652c20775f74696c65290a202020200a202020206966207069645f6e202a20424c4f434b5f53495a455f4e203c20483a0a20202020202020206f675f6f7574203d20746c2e7369676d6f696428616363756d756c61746f725f6f67290a20202020202020206f7574675f70747273203d204f75744f475f707472202b206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6f675f6d202b206f6666735f6e5f6f675b4e6f6e652c203a5d202a207374726964655f6f675f680a20202020202020206f675f6d61736b203d206d5f6d61736b5b3a2c204e6f6e655d202620286f6666735f6e5f6f675b4e6f6e652c203a5d203c2048290a2020202020202020746c2e73746f7265286f7574675f707472732c206f675f6f75742c206d61736b3d6f675f6d61736b290a0a2020202023202d2d2d20467573696f6e204c6f67696320666f7220342d5761792050617274202d2d2d0a202020206163635f7265736861706564203d20746c2e7265736861706528616363756d756c61746f725f347761792c2028424c4f434b5f53495a455f4d2c20485f4348554e4b5f53495a452c203429290a20202020726f6c655f696478203d20746c2e6172616e676528302c2034295b4e6f6e652c204e6f6e652c203a5d0a202020206c6566745f70726f6a20203d20746c2e73756d28746c2e776865726528726f6c655f696478203d3d20302c206163635f72657368617065642c20302e30292c20617869733d32290a202020206c6566745f6761746520203d20746c2e73756d28746c2e776865726528726f6c655f696478203d3d20312c206163635f72657368617065642c20302e30292c20617869733d32290a2020202072696768745f70726f6a203d20746c2e73756d28746c2e776865726528726f6c655f696478203d3d20322c206163635f72657368617065642c20302e30292c20617869733d32290a2020202072696768745f67617465203d20746c2e73756d28746c2e776865726528726f6c655f696478203d3d20332c206163635f72657368617065642c20302e30292c20617869733d32290a202020200a202020206f6666735f685f6368756e6b203d20287069645f6e202a20485f4348554e4b5f53495a4529202b20746c2e6172616e676528302c20485f4348554e4b5f53495a45290a202020206d61736b5f70747273203d204d61736b5f707472202b206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6d61736b5f6d202b206f6666735f685f6368756e6b5b4e6f6e652c203a5d202a207374726964655f6d61736b5f680a202020206d5f6d61736b5f68203d206d5f6d61736b5b3a2c204e6f6e655d202620286f6666735f685f6368756e6b5b4e6f6e652c203a5d203c2048290a202020206d61736b5f74696c65203d20746c2e6c6f6164286d61736b5f707472732c206d61736b3d6d5f6d61736b5f682c206f746865723d302e30290a0a202020206c6566745f6f7574203d206c6566745f70726f6a202a20746c2e7369676d6f6964286c6566745f6761746529202a206d61736b5f74696c650a2020202072696768745f6f7574203d2072696768745f70726f6a202a20746c2e7369676d6f69642872696768745f6761746529202a206d61736b5f74696c650a0a2020202073317332203d207331202a2073320a202020206f6666735f6220203d206f6666735f6d202f2f20733173320a202020206f6666735f7331203d20286f6666735f6d2025207331733229202f2f2073320a202020206f6666735f7332203d206f6666735f6d20252073320a202020206f6666735f625f326420203d20746c2e72657368617065286f6666735f622c202028424c4f434b5f53495a455f4d2c203129290a202020206f6666735f685f326420203d20746c2e72657368617065286f6666735f685f6368756e6b2c2028312c20485f4348554e4b5f53495a4529290a202020206f6666735f73315f3264203d20746c2e72657368617065286f6666735f73312c2028424c4f434b5f53495a455f4d2c203129290a202020206f6666735f73325f3264203d20746c2e72657368617065286f6666735f73322c2028424c4f434b5f53495a455f4d2c203129290a0a202020206f75746c5f70747273203d204f75744c6566745f707472202b20286f6666735f625f3264202a207374726964655f6f6c5f6273202b206f6666735f685f3264202a207374726964655f6f6c5f68202b0a202020202020202020202020202020202020202020202020202020202020202020202020206f6666735f73315f3264202a207374726964655f6f6c5f7331202b206f6666735f73325f3264202a207374726964655f6f6c5f7332290a202020206f7574725f707472735f74203d204f757452696768745f707472202b20286f6666735f625f3264202a207374726964655f6f725f745f6273202b206f6666735f685f3264202a207374726964655f6f725f745f68202b0a2020202020202020202020202020202020202020202020202020202020202020202020202020202020206f6666735f73325f3264202a207374726964655f6f725f745f7332202b206f6666735f73315f3264202a207374726964655f6f725f745f7331292023207332206f66667365742075736573207332207374726964652c207331206f66667365742075736573207331207374726964650a20202020746c2e73746f7265286f75746c5f707472732c206c6566745f6f75742c206d61736b3d6d5f6d61736b5f68290a20202020746c2e73746f7265286f7574725f707472735f742c2072696768745f6f75742c206d61736b3d6d5f6d61736b5f68290a0a40747269746f6e2e6175746f74756e65280a20202020636f6e666967733d5b0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d382c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d34292c0a202020205d2c0a202020206b65793d5b277331272c20277332272c202748275d2c0a290a40747269746f6e2e6a69740a64656620626d6d5f616e645f73746174735f6b65726e656c280a202020202320506f696e746572730a202020204c6566745f7074722c2052696768745f7074722c20426d6d5f4f75745f7074722c204d65616e5f4f75745f7074722c20527374645f4f75745f7074722c0a2020202023204d657461646174610a2020202062732c2073312c2073322c20482c0a202020202320537472696465730a202020207374726964655f6c5f62732c207374726964655f6c5f682c207374726964655f6c5f73312c207374726964655f6c5f73322c0a202020207374726964655f725f62732c207374726964655f725f682c207374726964655f725f73322c207374726964655f725f73312c0a202020207374726964655f626f5f62732c207374726964655f626f5f682c207374726964655f626f5f73312c207374726964655f626f5f73322c0a202020207374726964655f73746174735f62732c207374726964655f73746174735f73312c207374726964655f73746174735f73322c0a202020204c4e5f4550533a20746c2e636f6e7374657870722c0a2020202023204175746f74756e657220636f6e7374616e74730a20202020424c4f434b5f53495a455f4d3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4e3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4b3a20746c2e636f6e7374657870722c0a2020202047524f55505f53495a455f4d3a20746c2e636f6e7374657870722c0a293a0a2020202023202d2d2d204772696420616e6420504944204d617070696e67202d2d2d0a202020207069645f62617463685f7331203d20746c2e70726f6772616d5f696428617869733d30290a202020206e756d5f7069645f6d203d20746c2e636469762873312c20424c4f434b5f53495a455f4d290a202020206e756d5f7069645f6e203d20746c2e636469762873312c20424c4f434b5f53495a455f4e290a202020207069645f62203d207069645f62617463685f7331202f2f20286e756d5f7069645f6d202a206e756d5f7069645f6e290a202020207069645f73315f666c6174203d207069645f62617463685f7331202520286e756d5f7069645f6d202a206e756d5f7069645f6e290a202020206e756d5f7069645f696e5f67726f7570203d2047524f55505f53495a455f4d202a206e756d5f7069645f6e0a2020202067726f75705f6964203d207069645f73315f666c6174202f2f206e756d5f7069645f696e5f67726f75700a2020202066697273745f7069645f6d203d2067726f75705f6964202a2047524f55505f53495a455f4d0a2020202067726f75705f73697a655f6d203d206d696e286e756d5f7069645f6d202d2066697273745f7069645f6d2c2047524f55505f53495a455f4d290a202020207069645f6d203d2066697273745f7069645f6d202b20287069645f73315f666c617420252067726f75705f73697a655f6d290a202020207069645f6e203d20287069645f73315f666c61742025206e756d5f7069645f696e5f67726f757029202f2f2067726f75705f73697a655f6d0a202020206f6666735f6d203d207069645f6d202a20424c4f434b5f53495a455f4d202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4d290a202020206f6666735f6e203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206f6666735f6b203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a0a2020202023202d2d2d20424d4d20616e6420537461746973746963732043616c63756c6174696f6e202d2d2d0a202020206d65616e5f616363203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a202020207661725f616363203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a0a20202020666f72206820696e2072616e67652848293a0a2020202020202020626d6d5f736c696365203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a20202020202020206c6566745f707472735f62617365203d204c6566745f707472202b207069645f62202a207374726964655f6c5f6273202b2068202a207374726964655f6c5f680a202020202020202072696768745f707472735f62617365203d2052696768745f707472202b207069645f62202a207374726964655f725f6273202b2068202a207374726964655f725f680a2020202020202020666f72206b5f733220696e2072616e676528302c20746c2e636469762873322c20424c4f434b5f53495a455f4b29293a0a2020202020202020202020206b5f7374617274203d206b5f7332202a20424c4f434b5f53495a455f4b0a202020202020202020202020615f70747273203d206c6566745f707472735f62617365202b20286f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6c5f7331202b20286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29202a207374726964655f6c5f7332290a202020202020202020202020625f70747273203d2072696768745f707472735f62617365202b2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29202a207374726964655f725f7332202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f725f7331290a202020202020202020202020615f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c2073312920262028286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29203c207332290a202020202020202020202020625f6d61736b203d2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29203c20733229202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a20202020202020202020202061203d20746c2e6c6f616428615f707472732c206d61736b3d615f6d61736b2c206f746865723d302e30290a20202020202020202020202062203d20746c2e6c6f616428625f707472732c206d61736b3d625f6d61736b2c206f746865723d302e30290a202020202020202020202020626d6d5f736c696365202b3d20746c2e646f7428612c2062290a20202020202020200a20202020202020202320577269746520424d4d20726573756c7420666f72207468697320736c696365206f6620480a20202020202020206f75745f70747273203d20426d6d5f4f75745f707472202b207069645f62202a207374726964655f626f5f6273202b2068202a207374726964655f626f5f68202b205c0a202020202020202020202020202020202020206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f626f5f7331202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f626f5f73320a20202020202020206f75745f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c20733129202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a2020202020202020746c2e73746f7265286f75745f707472732c20626d6d5f736c6963652e746f28746c2e666c6f61743136292c206d61736b3d6f75745f6d61736b290a0a20202020202020202320416363756d756c6174652073746174730a20202020202020206d65616e5f616363202b3d20626d6d5f736c6963650a20202020202020207661725f616363202b3d20626d6d5f736c696365202a20626d6d5f736c6963650a0a2020202023202d2d2d2046696e616c697a6520616e642053746f72652053746174697374696373202d2d2d0a202020206d65616e203d206d65616e5f616363202f20480a20202020766172203d20287661725f616363202f204829202d20286d65616e202a206d65616e290a2020202072737464203d20746c2e6d6174682e727371727428766172202b204c4e5f455053290a0a2020202073746174735f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c20733129202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a202020206d65616e5f70747273203d204d65616e5f4f75745f707472202b207069645f62202a207374726964655f73746174735f6273202b206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f73746174735f7331202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f73746174735f73320a20202020727374645f70747273203d20527374645f4f75745f707472202b207069645f62202a207374726964655f73746174735f6273202b206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f73746174735f7331202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f73746174735f73320a20202020746c2e73746f7265286d65616e5f707472732c206d65616e2c206d61736b3d73746174735f6d61736b290a20202020746c2e73746f726528727374645f707472732c20727374642c206d61736b3d73746174735f6d61736b290a0a40747269746f6e2e6175746f74756e65280a20202020636f6e666967733d5b0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d382c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2033322c2027424c4f434b5f53495a455f4e273a2033322c2027424c4f434b5f53495a455f4b273a203132382c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a202020205d2c0a202020206b65793d5b277331272c20277332272c202748275d2c0a290a40747269746f6e2e6a69740a64656620626d6d5f6f75745f666c617474656e65645f6b65726e656c280a202020202320506f696e746572730a202020204c6566745f7074722c2052696768745f7074722c204f75745f7074722c0a20202020232044696d656e73696f6e730a2020202062732c2073312c2073322c20482c0a2020202023205374726964657320666f72204c6566742074656e736f72202862732c20682c2073312c207332290a202020207374726964655f6c5f62732c207374726964655f6c5f682c207374726964655f6c5f73312c207374726964655f6c5f73322c0a2020202023205374726964657320666f722052696768745f742074656e736f72202862732c20682c2073322c207331290a202020207374726964655f725f62732c207374726964655f725f682c207374726964655f725f73322c207374726964655f725f73312c0a2020202023205374726964657320666f72204f75747075742074656e736f7220286273202a207331202a2073312c2068290a202020207374726964655f6f5f666c61742c207374726964655f6f5f682c0a2020202023204b65726e656c20706172616d65746572730a20202020424c4f434b5f53495a455f4d3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4e3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4b3a20746c2e636f6e7374657870722c0a2020202047524f55505f53495a455f4d3a20746c2e636f6e7374657870722c0a293a0a2020202023204772696420616e642070726f6772616d2049442073657475700a20202020706964203d20746c2e70726f6772616d5f696428617869733d30290a202020206e756d5f7069645f6d203d20746c2e636469762873312c20424c4f434b5f53495a455f4d290a202020206e756d5f7069645f6e203d20746c2e636469762873312c20424c4f434b5f53495a455f4e290a202020206e756d5f7069645f696e5f67726f7570203d2047524f55505f53495a455f4d202a206e756d5f7069645f6e0a2020202067726f75705f6964203d20706964202f2f206e756d5f7069645f696e5f67726f75700a2020202066697273745f7069645f6d203d2067726f75705f6964202a2047524f55505f53495a455f4d0a2020202067726f75705f73697a655f6d203d206d696e286e756d5f7069645f6d202d2066697273745f7069645f6d2c2047524f55505f53495a455f4d290a202020207069645f6d203d2066697273745f7069645f6d202b202870696420252067726f75705f73697a655f6d290a202020207069645f6e203d20287069642025206e756d5f7069645f696e5f67726f757029202f2f2067726f75705f73697a655f6d0a0a202020202320426174636820616e642048656164204944732066726f6d20617869733d310a202020207069645f6268203d20746c2e70726f6772616d5f696428617869733d31290a202020207069645f62203d207069645f6268202f2f20480a202020207069645f68203d207069645f6268202520480a0a2020202023204f66667365747320666f7220746865204d2c204e2c20616e64204b2064696d656e73696f6e730a202020206f6666735f6d203d207069645f6d202a20424c4f434b5f53495a455f4d202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4d290a202020206f6666735f6e203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206f6666735f6b203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a0a202020202320506f696e7465727320746f207468652063757272656e7420626174636820616e6420686561640a202020206c6566745f707472735f62617365203d204c6566745f707472202b207069645f62202a207374726964655f6c5f6273202b207069645f68202a207374726964655f6c5f680a2020202072696768745f707472735f62617365203d2052696768745f707472202b207069645f62202a207374726964655f725f6273202b207069645f68202a207374726964655f725f680a202020200a202020202320416363756d756c61746f7220666f7220746865206d6174726978206d756c7469706c69636174696f6e20726573756c740a20202020616363756d756c61746f72203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a202020200a2020202023204c6f6f70206f76657220746865204b2064696d656e73696f6e20287332290a20202020666f72206b20696e2072616e676528302c20746c2e636469762873322c20424c4f434b5f53495a455f4b29293a0a20202020202020206b5f7374617274203d206b202a20424c4f434b5f53495a455f4b0a20202020202020200a20202020202020202320506f696e7465727320746f2074686520626c6f636b73206f66204120616e642042206d617472696365730a2020202020202020615f70747273203d206c6566745f707472735f62617365202b20286f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6c5f7331202b20286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29202a207374726964655f6c5f7332290a2020202020202020625f70747273203d2072696768745f707472735f62617365202b2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29202a207374726964655f725f7332202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f725f7331290a20202020202020200a202020202020202023204d61736b7320746f2068616e646c65206e6f6e2d66756c6c20626c6f636b730a2020202020202020615f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c2073312920262028286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29203c207332290a2020202020202020625f6d61736b203d2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29203c20733229202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a20202020202020200a202020202020202023204c6f6164204120616e64204220626c6f636b730a202020202020202061203d20746c2e6c6f616428615f707472732c206d61736b3d615f6d61736b2c206f746865723d302e30290a202020202020202062203d20746c2e6c6f616428625f707472732c206d61736b3d625f6d61736b2c206f746865723d302e30290a20202020202020200a20202020202020202320506572666f726d206d6174726978206d756c7469706c69636174696f6e0a2020202020202020616363756d756c61746f72202b3d20746c2e646f7428612c2062290a0a2020202023202d2d2d2044697265637420777269746520746f20666c617474656e6564206f7574707574202d2d2d0a20202020232043616c63756c61746520746865206261736520726f7720696e64657820696e2074686520666c617474656e6564206f75747075742074656e736f720a20202020666c61745f726f775f62617365203d207069645f62202a207331202a2073310a202020200a20202020232043616c63756c6174652074686520726f77206f66667365747320666f72207468652063757272656e7420626c6f636b0a20202020666c61745f726f775f6f666673203d206f6666735f6d5b3a2c204e6f6e655d202a207331202b206f6666735f6e5b4e6f6e652c203a5d0a202020200a202020202320436f6d62696e6520746f20676574207468652066696e616c20726f7720696e64696365730a2020202066696e616c5f666c61745f726f7773203d20666c61745f726f775f62617365202b20666c61745f726f775f6f6666730a202020200a20202020232043616c63756c61746520746865206f757470757420706f696e746572730a202020206f75745f70747273203d204f75745f707472202b2066696e616c5f666c61745f726f7773202a207374726964655f6f5f666c6174202b207069645f68202a207374726964655f6f5f680a0a2020202023204372656174652061206d61736b20666f722073746f72696e6720746865206f75747075740a20202020635f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c20733129202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a202020200a20202020232053746f72652074686520726573756c740a20202020746c2e73746f7265286f75745f707472732c20616363756d756c61746f722c206d61736b3d635f6d61736b290a0a40747269746f6e2e6175746f74756e65280a20202020636f6e666967733d5b0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d382c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2033322c2027424c4f434b5f53495a455f4e273a2033322c2027424c4f434b5f53495a455f4b273a203132382c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a202020205d2c0a202020206b65793d5b277331272c20277332272c202748275d2c0a290a40747269746f6e2e6a69740a64656620626d6d5f636f616c65736365645f6b65726e656c280a202020202320506f696e746572730a202020204c6566745f7074722c2052696768745f7074722c204f75745f7074722c0a20202020232044696d656e73696f6e730a2020202062732c2073312c2073322c20482c0a202020202320537472696465730a202020207374726964655f6c5f62732c207374726964655f6c5f682c207374726964655f6c5f73312c207374726964655f6c5f73322c0a202020207374726964655f725f62732c207374726964655f725f682c207374726964655f725f73322c207374726964655f725f73312c0a202020207374726964655f6f5f62732c207374726964655f6f5f682c207374726964655f6f5f73312c207374726964655f6f5f73322c0a2020202023204b65726e656c20706172616d65746572730a20202020424c4f434b5f53495a455f4d3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4e3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4b3a20746c2e636f6e7374657870722c0a2020202047524f55505f53495a455f4d3a20746c2e636f6e7374657870722c0a293a0a2020202023204772696420616e642070726f6772616d204944730a20202020706964203d20746c2e70726f6772616d5f696428617869733d30290a202020206e756d5f7069645f6d203d20746c2e636469762873312c20424c4f434b5f53495a455f4d290a202020206e756d5f7069645f6e203d20746c2e636469762873312c20424c4f434b5f53495a455f4e290a202020206e756d5f7069645f696e5f67726f7570203d2047524f55505f53495a455f4d202a206e756d5f7069645f6e0a2020202067726f75705f6964203d20706964202f2f206e756d5f7069645f696e5f67726f75700a2020202066697273745f7069645f6d203d2067726f75705f6964202a2047524f55505f53495a455f4d0a2020202067726f75705f73697a655f6d203d206d696e286e756d5f7069645f6d202d2066697273745f7069645f6d2c2047524f55505f53495a455f4d290a202020207069645f6d203d2066697273745f7069645f6d202b202870696420252067726f75705f73697a655f6d290a202020207069645f6e203d20287069642025206e756d5f7069645f696e5f67726f757029202f2f2067726f75705f73697a655f6d0a0a202020207069645f6268203d20746c2e70726f6772616d5f696428617869733d31290a202020207069645f62203d207069645f6268202f2f20480a202020207069645f68203d207069645f6268202520480a0a202020206f6666735f6d203d207069645f6d202a20424c4f434b5f53495a455f4d202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4d290a202020206f6666735f6e203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206f6666735f6b203d20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a0a202020206c6566745f707472735f62617365203d204c6566745f707472202b207069645f62202a207374726964655f6c5f6273202b207069645f68202a207374726964655f6c5f680a2020202072696768745f707472735f62617365203d2052696768745f707472202b207069645f62202a207374726964655f725f6273202b207069645f68202a207374726964655f725f680a202020200a20202020616363756d756c61746f72203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a202020200a20202020666f72206b20696e2072616e676528302c20746c2e636469762873322c20424c4f434b5f53495a455f4b29293a0a20202020202020206b5f7374617274203d206b202a20424c4f434b5f53495a455f4b0a2020202020202020615f70747273203d206c6566745f707472735f62617365202b20286f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6c5f7331202b20286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29202a207374726964655f6c5f7332290a2020202020202020625f70747273203d2072696768745f707472735f62617365202b2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29202a207374726964655f725f7332202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f725f7331290a20202020202020200a2020202020202020615f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c2073312920262028286b5f7374617274202b206f6666735f6b5b4e6f6e652c203a5d29203c207332290a2020202020202020625f6d61736b203d2028286b5f7374617274202b206f6666735f6b5b3a2c204e6f6e655d29203c20733229202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a20202020202020200a202020202020202061203d20746c2e6c6f616428615f707472732c206d61736b3d615f6d61736b2c206f746865723d302e30290a202020202020202062203d20746c2e6c6f616428625f707472732c206d61736b3d625f6d61736b2c206f746865723d302e30290a20202020202020200a2020202020202020616363756d756c61746f72202b3d20746c2e646f7428612c2062290a0a2020202023202d2d2d20436f616c6573636564205772697465202d2d2d0a202020202320577269746520746f2061207374616e64617264202862732c20482c2073312c20733129206c61796f75740a202020206f75745f70747273203d204f75745f707472202b207069645f62202a207374726964655f6f5f6273202b207069645f68202a207374726964655f6f5f68202b205c0a2020202020202020202020202020206f6666735f6d5b3a2c204e6f6e655d202a207374726964655f6f5f7331202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f6f5f73320a0a20202020635f6d61736b203d20286f6666735f6d5b3a2c204e6f6e655d203c20733129202620286f6666735f6e5b4e6f6e652c203a5d203c207331290a20202020746c2e73746f7265286f75745f707472732c20616363756d756c61746f722c206d61736b3d635f6d61736b290a200a0a40747269746f6e2e6a69740a6465662072656f726465725f746f5f666c61745f6b65726e656c280a20202020496e5f7074722c204f75745f7074722c0a2020202062732c2073312c20482c0a2020202023205374726964657320666f7220496e7075743a202862732c20482c2073312c207331290a202020207374726964655f696e5f62732c207374726964655f696e5f682c207374726964655f696e5f73315f726f772c207374726964655f696e5f73315f636f6c2c0a2020202023205374726964657320666f72204f75747075743a20286273202a207331202a2073312c2048290a202020207374726964655f6f75745f666c61742c207374726964655f6f75745f682c0a20202020424c4f434b5f53495a455f483a20746c2e636f6e7374657870722c0a293a0a202020202320456163682070726f6772616d20636f6d7075746573206f6e65206f757470757420726f772c207768696368206973206120766563746f72206f662073697a6520482e0a2020202023205468652070726f6772616d20494420636f72726573706f6e647320746f2074686520666c617474656e65642028622c20722c20632920696e6465782e0a20202020706964203d20746c2e70726f6772616d5f696428617869733d30290a0a2020202023204465636f6d706f7365207468652031442070726f6772616d204944206261636b20696e746f2033442062617463682c20726f772c20616e6420636f6c756d6e20696e64696365730a2020202073317331203d207331202a2073310a202020206220203d20706964202f2f20733173310a202020207220203d20287069642025207331733129202f2f2073310a202020206320203d2070696420252073310a0a202020202320477561726420616761696e7374206f75742d6f662d626f756e64732070726f6772616d732c2077686963682063616e2068617070656e206966207468652067726964206973206e6f7420706572666563746c7920646976697369626c650a2020202069662062203e3d2062733a0a202020202020202072657475726e0a0a20202020232043616c63756c61746520706f696e7465727320666f7220746869732073706563696669632066696265720a2020202023204261736520706f696e74657220746f20746865207374617274206f6620746865206f757470757420726f770a202020206f75745f7074725f62617365203d204f75745f707472202b20706964202a207374726964655f6f75745f666c61740a2020202023204261736520706f696e74657220746f20746865207374617274206f662074686520696e70757420646174610a20202020696e5f7074725f62617365203d20496e5f707472202b2062202a207374726964655f696e5f6273202b2072202a207374726964655f696e5f73315f726f77202b2063202a207374726964655f696e5f73315f636f6c0a0a2020202023204c6f6f70206f7665722074686520482064696d656e73696f6e20696e20626c6f636b730a20202020666f7220685f6f666673657420696e2072616e676528302c20746c2e6364697628482c20424c4f434b5f53495a455f4829293a0a202020202020202023204f66667365747320666f72207468652063757272656e7420626c6f636b206f6620480a20202020202020206f6666735f68203d20685f6f6666736574202a20424c4f434b5f53495a455f48202b20746c2e6172616e676528302c20424c4f434b5f53495a455f48290a2020202020202020685f6d61736b203d206f6666735f68203c20480a0a202020202020202023202d2d2d20436f616c6573636564205772697465205061747465726e202d2d2d0a2020202020202020232043616c63756c617465206f757470757420706f696e746572733a2074686573652061726520636f6e746967756f75730a20202020202020206f75745f70747273203d206f75745f7074725f62617365202b206f6666735f68202a207374726964655f6f75745f680a0a202020202020202023202d2d2d20537472696465642052656164205061747465726e202d2d2d0a2020202020202020232043616c63756c61746520696e70757420706f696e746572733a207468657365206172652073747269646564206279207374726964655f696e5f680a2020202020202020696e5f70747273203d20696e5f7074725f62617365202b206f6666735f68202a207374726964655f696e5f680a0a202020202020202023204c6f61642066726f6d20696e70757420616e642073746f726520746f206f75747075740a202020202020202064617461203d20746c2e6c6f616428696e5f707472732c206d61736b3d685f6d61736b2c206f746865723d302e30290a2020202020202020746c2e73746f7265286f75745f707472732c20646174612c206d61736b3d685f6d61736b290a0a64656620746f7263685f707432286f75745f65696e73756d5f666c61742c2062732c2073312c2073322c20642c20682c20746f5f6f75745f6e6f726d5f7765696768742c20746f5f6f75745f6e6f726d5f626961732c206f675f6d682c20746f5f6f75745f776569676874293a0a2020202023204170706c79206c61796572206e6f726d20616e642066696e616c20676174696e670a202020206e6f726d6564203d20462e6c617965725f6e6f726d286f75745f65696e73756d5f666c61742c2028682c292c20746f5f6f75745f6e6f726d5f7765696768742c20746f5f6f75745f6e6f726d5f62696173292e746f28746f7263682e666c6f61743136290a202020206761746564203d206e6f726d6564202a206f675f6d680a0a20202020232046696e616c2070726f6a656374696f6e0a2020202066696e616c5f6f75745f666c6174203d206761746564204020746f5f6f75745f7765696768742e7428290a2020202066696e616c5f6f7574203d2066696e616c5f6f75745f666c61742e766965772862732c2073312c2073322c2064290a2020202072657475726e2066696e616c5f6f75740a0a40747269746f6e2e6175746f74756e65280a20202020636f6e666967733d5b0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a2033322c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2033322c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d33292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a203132382c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d382c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a203132382c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2033322c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d382c206e756d5f7374616765733d34292c0a2020202020202020747269746f6e2e436f6e666967287b27424c4f434b5f53495a455f4d273a2036342c2027424c4f434b5f53495a455f4e273a2036342c2027424c4f434b5f53495a455f4b273a2036342c202747524f55505f53495a455f4d273a20387d2c206e756d5f77617270733d342c206e756d5f7374616765733d34292c0a202020205d2c0a202020206b65793d5b2748272c202744275d2c0a290a40747269746f6e2e6a69740a6465662066757365645f66696e616c5f6b65726e656c280a202020202320506f696e746572730a20202020496e5f7074722c20476174655f7074722c204e6f726d575f7074722c204e6f726d425f7074722c2050726f6a575f7074722c204f75745f7074722c0a2020202023204d657461646174610a202020204d2c20482c20442c2073312c204d5f676174652c2023204d5f67617465203d2062732a73312a73320a202020202320537472696465730a202020207374726964655f696e5f62732c207374726964655f696e5f682c207374726964655f696e5f73315f726f772c207374726964655f696e5f73315f636f6c2c0a202020207374726964655f676174655f6d2c207374726964655f676174655f682c0a202020207374726964655f70726f6a5f6b2c207374726964655f70726f6a5f6e2c0a202020207374726964655f6f75745f62732c207374726964655f6f75745f73315f726f772c207374726964655f6f75745f73315f636f6c2c207374726964655f6f75745f642c0a202020202320436f6e7374616e74730a202020204c4e5f4550533a20746c2e636f6e7374657870722c0a20202020424c4f434b5f53495a455f4d3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4e3a20746c2e636f6e7374657870722c20424c4f434b5f53495a455f4b3a20746c2e636f6e7374657870722c0a2020202047524f55505f53495a455f4d3a20746c2e636f6e7374657870722c0a293a0a2020202023202d2d2d204772696420616e642050494420536574757020666f72204d61746d756c202d2d2d0a20202020706964203d20746c2e70726f6772616d5f696428617869733d30290a202020206e756d5f7069645f6d203d20746c2e63646976284d2c20424c4f434b5f53495a455f4d290a202020206e756d5f7069645f6e203d20746c2e6364697628442c20424c4f434b5f53495a455f4e290a202020200a202020206e756d5f7069645f696e5f67726f7570203d2047524f55505f53495a455f4d202a206e756d5f7069645f6e0a2020202067726f75705f6964203d20706964202f2f206e756d5f7069645f696e5f67726f75700a2020202066697273745f7069645f6d203d2067726f75705f6964202a2047524f55505f53495a455f4d0a2020202067726f75705f73697a655f6d203d206d696e286e756d5f7069645f6d202d2066697273745f7069645f6d2c2047524f55505f53495a455f4d290a202020207069645f6d203d2066697273745f7069645f6d202b202870696420252067726f75705f73697a655f6d290a202020207069645f6e203d20287069642025206e756d5f7069645f696e5f67726f757029202f2f2067726f75705f73697a655f6d0a0a202020206f6666735f6d203d207069645f6d202a20424c4f434b5f53495a455f4d202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4d290a202020206f6666735f6e203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206d5f6d61736b203d206f6666735f6d203c204d0a0a2020202023204465636f6d706f7365204d206261636b20746f2028622c20722c20632920666f722072656f72646572696e67206c6f6f6b7570730a2020202073317331203d207331202a2073310a2020202062203d206f6666735f6d202f2f20733173310a2020202072203d20286f6666735f6d2025207331733129202f2f2073310a2020202063203d206f6666735f6d20252073310a0a2020202023202d2d2d205061737320313a20526f772d77697365204d65616e2043616c63756c6174696f6e202d2d2d0a202020206d65616e203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c292c2064747970653d746c2e666c6f61743332290a20202020696e5f7074725f62617365203d20496e5f707472202b2062202a207374726964655f696e5f6273202b2072202a207374726964655f696e5f73315f726f77202b2063202a207374726964655f696e5f73315f636f6c0a20202020666f72206b5f6f666673657420696e2072616e676528302c20482c20424c4f434b5f53495a455f4b293a0a20202020202020206f6666735f6b203d206b5f6f6666736574202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a20202020202020206b5f6d61736b203d206f6666735f6b203c20480a20202020202020200a2020202020202020696e5f70747273203d20696e5f7074725f626173655b3a2c204e6f6e655d202b206f6666735f6b5b4e6f6e652c203a5d202a207374726964655f696e5f680a2020202020202020696e5f6368756e6b203d20746c2e6c6f616428696e5f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a20202020202020206d65616e202b3d20746c2e73756d28696e5f6368756e6b2c20617869733d31290a202020206d65616e202f3d20480a0a2020202023202d2d2d205061737320323a204e756d65726963616c6c7920537461626c652056617269616e63652043616c63756c6174696f6e202d2d2d0a20202020766172203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c292c2064747970653d746c2e666c6f61743332290a20202020666f72206b5f6f666673657420696e2072616e676528302c20482c20424c4f434b5f53495a455f4b293a0a20202020202020206f6666735f6b203d206b5f6f6666736574202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a20202020202020206b5f6d61736b203d206f6666735f6b203c20480a20202020202020200a2020202020202020696e5f70747273203d20696e5f7074725f626173655b3a2c204e6f6e655d202b206f6666735f6b5b4e6f6e652c203a5d202a207374726964655f696e5f680a2020202020202020696e5f6368756e6b203d20746c2e6c6f616428696e5f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a20202020202020200a202020202020202023205375627472616374206d65616e206265666f7265207371756172696e670a2020202020202020696e5f63656e7465726564203d20696e5f6368756e6b202d206d65616e5b3a2c204e6f6e655d0a2020202020202020766172202b3d20746c2e73756d28696e5f63656e7465726564202a20696e5f63656e74657265642c20617869733d31290a20202020766172202f3d20480a2020202072737464203d20746c2e6d6174682e727371727428766172202b204c4e5f455053290a0a0a2020202023202d2d2d205061737320333a20467573656420476174696e6720616e64204d61746d756c202d2d2d0a20202020616363203d20746c2e7a65726f732828424c4f434b5f53495a455f4d2c20424c4f434b5f53495a455f4e292c2064747970653d746c2e666c6f61743332290a20202020666f72206b5f6f666673657420696e2072616e676528302c20482c20424c4f434b5f53495a455f4b293a0a20202020202020206f6666735f6b203d206b5f6f6666736574202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4b290a20202020202020206b5f6d61736b203d206f6666735f6b203c20480a20202020202020200a2020202020202020696e5f70747273203d20696e5f7074725f626173655b3a2c204e6f6e655d202b206f6666735f6b5b4e6f6e652c203a5d202a207374726964655f696e5f680a202020202020202061203d20746c2e6c6f616428696e5f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a0a202020202020202070726f6a5f70747273203d2050726f6a575f707472202b206f6666735f6b5b3a2c204e6f6e655d202a207374726964655f70726f6a5f6b202b206f6666735f6e5b4e6f6e652c203a5d202a207374726964655f70726f6a5f6e0a2020202020202020625f77203d20746c2e6c6f61642870726f6a5f707472732c206d61736b3d6b5f6d61736b5b3a2c204e6f6e655d202620286f6666735f6e5b4e6f6e652c203a5d203c2044292c206f746865723d302e30290a20202020202020200a2020202020202020232053696e63652073313d73322c204d3d4d5f6761746520616e6420746865206d6f64756c75732069732061206e6f2d6f702c20627574206973206861726d6c6573732e0a2020202020202020676174655f6d5f696478203d206f6666735f6d2025204d5f676174650a2020202020202020676174655f70747273203d20476174655f707472202b20676174655f6d5f6964785b3a2c204e6f6e655d202a207374726964655f676174655f6d202b206f6666735f6b5b4e6f6e652c203a5d202a207374726964655f676174655f680a202020202020202067617465203d20746c2e6c6f616428676174655f707472732c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d2026206b5f6d61736b5b4e6f6e652c203a5d2c206f746865723d302e30290a20202020202020200a20202020202020206e6f726d5f77203d20746c2e6c6f6164284e6f726d575f707472202b206f6666735f6b2c206d61736b3d6b5f6d61736b2c206f746865723d302e30290a20202020202020206e6f726d5f62203d20746c2e6c6f6164284e6f726d425f707472202b206f6666735f6b2c206d61736b3d6b5f6d61736b2c206f746865723d302e30290a0a2020202020202020615f6e6f726d203d202861202d206d65616e5b3a2c204e6f6e655d29202a20727374645b3a2c204e6f6e655d0a2020202020202020615f6e6f726d203d20615f6e6f726d202a206e6f726d5f775b4e6f6e652c203a5d202b206e6f726d5f625b4e6f6e652c203a5d0a2020202020202020615f6761746564203d20615f6e6f726d202a20676174650a0a2020202020202020616363202b3d20746c2e646f7428615f67617465642e746f28625f772e6474797065292c20625f77290a20202020202020200a2020202023202d2d2d2053746f72652046696e616c204f7574707574202d2d2d0a202020206f6666735f64203d207069645f6e202a20424c4f434b5f53495a455f4e202b20746c2e6172616e676528302c20424c4f434b5f53495a455f4e290a202020206f75745f7074725f62617365203d204f75745f707472202b20622a7374726964655f6f75745f6273202b20722a7374726964655f6f75745f73315f726f77202b20632a7374726964655f6f75745f73315f636f6c0a202020206f75745f70747273203d206f75745f7074725f626173655b3a2c204e6f6e655d202b206f6666735f645b4e6f6e652c203a5d202a207374726964655f6f75745f640a202020200a20202020746c2e73746f7265286f75745f707472732c206163632c206d61736b3d6d5f6d61736b5b3a2c204e6f6e655d202620286f6666735f645b4e6f6e652c203a5d203c204429290a0a20200a0a64656620636f6d70696c65647472696d756c5f66757365645f696e7465726c6561766564280a20202020783a20746f7263682e54656e736f722c0a202020206d61736b5f6d683a20746f7263682e54656e736f722c0a202020206e6f726d5f7765696768743a20746f7263682e54656e736f722c0a202020206e6f726d5f626961733a20746f7263682e54656e736f722c0a20202020575f347761793a20746f7263682e54656e736f722c20232055736520746865206e657720776569676874206d617472696365730a20202020575f6f673a20746f7263682e54656e736f722c0a20202020746f5f6f75745f6e6f726d5f7765696768743a20746f7263682e54656e736f722c0a20202020746f5f6f75745f6e6f726d5f626961733a20746f7263682e54656e736f722c0a20202020746f5f6f75745f7765696768743a20746f7263682e54656e736f722c0a20202020683a20696e742c0a293a0a2020202062732c2073312c2073322c2064203d20782e73686170650a202020204d2c204b2c2048203d206273202a207331202a2073322c20782e73686170655b2d315d2c20680a20202020785f666c6174203d20782e76696577284d2c204b290a0a202020206c6566745f66696e616c20203d20746f7263682e656d707479282862732c20482c2073312c207332292c206465766963653d782e6465766963652c2064747970653d746f7263682e666c6f61743136290a2020202072696768745f66696e616c5f74203d20746f7263682e656d707479282862732c20482c2073322c207331292c206465766963653d782e6465766963652c2064747970653d746f7263682e666c6f61743136290a202020206f675f6d68203d20746f7263682e656d70747928284d2c2048292c206465766963653d782e6465766963652c2064747970653d746f7263682e666c6f61743136290a0a2020202023205468652067726964206973206c61756e6368656420666f7220746865206c617267657220342a482070726f626c656d0a202020204e5f34776179203d2034202a20480a2020202067726964203d206c616d626461206d6574613a2028747269746f6e2e63646976284d2c206d6574615b27424c4f434b5f53495a455f4d275d29202a20747269746f6e2e63646976284e5f347761792c206d6574615b27424c4f434b5f53495a455f4e275d292c290a202020200a2020202066757365645f6c6e5f6475616c5f6d61746d756c5f6b65726e656c5b677269645d280a20202020202020202320506f696e74657273202839290a2020202020202020785f666c61742c20575f347761792c20575f6f672c206d61736b5f6d682c206e6f726d5f7765696768742c206e6f726d5f626961732c0a20202020202020206c6566745f66696e616c2c2072696768745f66696e616c5f742c206f675f6d682c0a202020202020202023204d6574616461746120283529202d204d2c20482c204b2c2073312c2073320a20202020202020204d2c20482c204b2c2073312c2073322c0a202020202020202023205374726964657320283136290a2020202020202020785f666c61742e7374726964652830292c20785f666c61742e7374726964652831292c0a2020202020202020575f347761792e7374726964652830292c20575f347761792e7374726964652831292c0a2020202020202020575f6f672e7374726964652830292c20575f6f672e7374726964652831292c0a20202020202020206c6566745f66696e616c2e7374726964652830292c206c6566745f66696e616c2e7374726964652831292c206c6566745f66696e616c2e7374726964652832292c206c6566745f66696e616c2e7374726964652833292c0a202020202020202072696768745f66696e616c5f742e7374726964652830292c2072696768745f66696e616c5f742e7374726964652831292c2072696768745f66696e616c5f742e7374726964652832292c2072696768745f66696e616c5f742e7374726964652833292c0a20202020202020206f675f6d682e7374726964652830292c206f675f6d682e7374726964652831292c0a20202020202020206d61736b5f6d682e7374726964652830292c206d61736b5f6d682e7374726964652831292c0a20202020202020202320436f6e737465787072202831290a20202020202020204c4e5f4550533d31652d350a20202020290a202020200a20202020626d6d5f6f75745f746d70203d20746f7263682e656d707479282862732c20482c2073312c207331292c206465766963653d782e6465766963652c2064747970653d746f7263682e666c6f61743136290a0a20202020677269645f626d6d203d206c616d626461206d6574613a2028747269746f6e2e636469762873312c206d6574615b27424c4f434b5f53495a455f4d275d29202a20747269746f6e2e636469762873312c206d6574615b27424c4f434b5f53495a455f4e275d292c206273202a2048290a20202020626d6d5f636f616c65736365645f6b65726e656c5b677269645f626d6d5d280a20202020202020206c6566745f66696e616c2c2072696768745f66696e616c5f742c20626d6d5f6f75745f746d702c0a202020202020202062732c2073312c2073322c20482c0a20202020202020206c6566745f66696e616c2e7374726964652830292c206c6566745f66696e616c2e7374726964652831292c206c6566745f66696e616c2e7374726964652832292c206c6566745f66696e616c2e7374726964652833292c0a202020202020202072696768745f66696e616c5f742e7374726964652830292c2072696768745f66696e616c5f742e7374726964652831292c2072696768745f66696e616c5f742e7374726964652832292c2072696768745f66696e616c5f742e7374726964652833292c0a2020202020202020626d6d5f6f75745f746d702e7374726964652830292c20626d6d5f6f75745f746d702e7374726964652831292c20626d6d5f6f75745f746d702e7374726964652832292c20626d6d5f6f75745f746d702e7374726964652833292c0a20202020290a0a2020202023202d2d2d204b65726e656c20333a2046756c6c792046757365642046696e616c205374616765202d2d2d0a2020202066696e616c5f6f7574203d20746f7263682e656d707479282862732c2073312c2073312c2064292c206465766963653d782e6465766963652c2064747970653d746f7263682e666c6f61743136290a202020204d5f66696e616c203d206273202a207331202a2073310a202020200a20202020677269645f66696e616c203d206c616d626461206d6574613a2028747269746f6e2e63646976284d5f66696e616c2c206d6574615b27424c4f434b5f53495a455f4d275d29202a20747269746f6e2e6364697628642c206d6574615b27424c4f434b5f53495a455f4e275d292c290a0a2020202023202a2a2a2a2a2054484520464958202a2a2a2a2a0a202020202320546865206b65726e656c2065787065637473206050726f6a575f7074726020746f20686176652061207368617065206f662028482c2044292e0a20202020232060746f5f6f75745f776569676874602069732028442c2048292c20736f207765204d555354207472616e73706f736520616e64206d616b6520697420636f6e746967756f75732e0a2020202070726f6a5f775f7472616e73706f736564203d20746f5f6f75745f7765696768742e7428292e636f6e746967756f757328290a0a2020202066757365645f66696e616c5f6b65726e656c5b677269645f66696e616c5d280a20202020202020202320506f696e746572730a2020202020202020626d6d5f6f75745f746d702c206f675f6d682c20746f5f6f75745f6e6f726d5f7765696768742c20746f5f6f75745f6e6f726d5f626961732c2070726f6a5f775f7472616e73706f7365642c2066696e616c5f6f75742c0a202020202020202023204d657461646174610a20202020202020204d5f66696e616c2c20482c20642c2073312c204d2c0a20202020202020202320537472696465730a2020202020202020626d6d5f6f75745f746d702e7374726964652830292c20626d6d5f6f75745f746d702e7374726964652831292c20626d6d5f6f75745f746d702e7374726964652832292c20626d6d5f6f75745f746d702e7374726964652833292c0a20202020202020206f675f6d682e7374726964652830292c206f675f6d682e7374726964652831292c0a202020202020202070726f6a5f775f7472616e73706f7365642e7374726964652830292c2070726f6a5f775f7472616e73706f7365642e7374726964652831292c2023205573652073747269646573206f662074686520636f727265637465642074656e736f720a202020202020202066696e616c5f6f75742e7374726964652830292c2066696e616c5f6f75742e7374726964652831292c2066696e616c5f6f75742e7374726964652832292c2066696e616c5f6f75742e7374726964652833292c0a20202020202020202320436f6e7374616e74730a20202020202020204c4e5f4550533d31652d352c0a20202020290a0a2020202072657475726e2066696e616c5f6f75742e766965772862732c2073312c2073322c2064290a0a646566207061636b5f775f347761795f656666696369656e742877656967687473293a0a20202020222222205061636b73204c2c204c472c20522c20524720696e746f2061207469676874205b4b2c20342a485d206d61747269782e202222220a20202020574c203d20776569676874735b276c6566745f70726f6a2e776569676874275d0a20202020574c47203d20776569676874735b276c6566745f676174652e776569676874275d0a202020205752203d20776569676874735b2772696768745f70726f6a2e776569676874275d0a20202020575247203d20776569676874735b2772696768745f676174652e776569676874275d0a20202020482c204b203d20574c2e73686170650a202020207773203d20746f7263682e737461636b285b574c2c20574c472c2057522c205752475d2c2064696d3d30292e7065726d75746528312c20302c2032290a202020207773203d2077732e636f6e746967756f757328292e766965772834202a20482c204b290a2020202072657475726e2077732e7428292e636f6e746967756f757328292e746f28746f7263682e666c6f61743136290a0a646566206765745f775f6f672877656967687473293a0a20202020222222204765747320746865207472616e73706f736564205b4b2c20485d206f75745f6761746520776569676874206d61747269782e202222220a20202020574f47203d20776569676874735b276f75745f676174652e776569676874275d0a2020202072657475726e20574f472e7428292e636f6e746967756f757328292e746f28746f7263682e666c6f61743136290a0a64656620637573746f6d5f6b65726e656c2864617461293a0a20202020696e7075745f74656e736f722c206d61736b2c20776569676874732c20636f6e666967203d20646174610a2020202048203d20636f6e6669675b2268696464656e5f64696d225d0a0a20202020575f34776179203d207061636b5f775f347761795f656666696369656e742877656967687473290a20202020575f6f67203d206765745f775f6f672877656967687473290a0a2020202062732c2073312c2073322c2064203d20696e7075745f74656e736f722e73686170650a202020204d203d206273202a207331202a2073320a202020206d61736b5f6d68203d206d61736b2e756e73717565657a65282d31292e657870616e64282d312c202d312c202d312c2048292e72657368617065284d2c2048292e746f28746f7263682e666c6f617431362920236d6f766520696e746f206b65726e656c20706f737369626c790a0a2020202072657475726e20636f6d70696c65647472696d756c5f66757365645f696e7465726c6561766564280a2020202020202020783d696e7075745f74656e736f722e746f28746f7263682e666c6f61743332292c0a20202020202020206d61736b5f6d683d6d61736b5f6d682c0a20202020202020206e6f726d5f7765696768743d776569676874735b276e6f726d2e776569676874275d2e746f28746f7263682e666c6f61743332292c0a20202020202020206e6f726d5f626961733d776569676874735b276e6f726d2e62696173275d2e746f28746f7263682e666c6f61743332292c0a2020202020202020575f347761793d575f347761792c2023205061737320746865206e657720342d776179206d61747269780a2020202020202020575f6f673d575f6f672c202020202023205061737320746865206e6577206f75745f67617465206d61747269780a2020202020202020746f5f6f75745f6e6f726d5f7765696768743d776569676874735b27746f5f6f75745f6e6f726d2e776569676874275d2e746f28746f7263682e666c6f61743136292c0a2020202020202020746f5f6f75745f6e6f726d5f626961733d776569676874735b27746f5f6f75745f6e6f726d2e62696173275d2e746f28746f7263682e666c6f61743136292c0a2020202020202020746f5f6f75745f7765696768743d776569676874735b27746f5f6f75742e776569676874275d2e746f28746f7263682e666c6f61743136292c0a2020202020202020683d482c0a20202020290aNo newline at end of file+ #!POPCORN leaderboard trimul+ import torch+ import torch.nn.functional as F+ import triton+ import triton.language as tl++ # Set PyTorch flags for performance+ torch.backends.cuda.matmul.allow_tf32 = True+ torch.backends.cuda.matmul.allow_fp16_reduced_precision_reduction = True++ # Note: The @triton.autotune decorators have been removed from all kernels below.++ @triton.jit+ def fused_ln_dual_matmul_kernel(+ # Pointers (9)+ X_ptr, W_4way_ptr, W_og_ptr, Mask_ptr, Norm_Weight_ptr, Norm_Bias_ptr,+ OutLeft_ptr, OutRight_ptr, OutOG_ptr,+ # Metadata (5)+ M, H, K, s1, s2,+ # Strides (16)+ stride_x_m, stride_x_k,+ stride_w4_k, stride_w4_n,+ stride_wog_k, stride_wog_n,+ stride_ol_bs, stride_ol_h, stride_ol_s1, stride_ol_s2,+ stride_or_t_bs, stride_or_t_h, stride_or_t_s2, stride_or_t_s1,+ stride_og_m, stride_og_h,+ stride_mask_m, stride_mask_h,+ # Constexpr (now passed as arguments from the host)+ LN_EPS: tl.constexpr,+ BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,+ GROUP_SIZE_M: tl.constexpr, H_CHUNK_SIZE: tl.constexpr,+ ):+ # --- PID Mapping: Based on the LARGER 4*H problem ---+ pid = tl.program_id(axis=0)+ N_4way = 4 * H+ num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)+ num_pid_n = tl.cdiv(N_4way, 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++ # --- SHARED LayerNorm calculation (done only ONCE) ---+ offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)+ m_mask = offs_m < M+ x_rows_base_ptr = X_ptr + offs_m[:, None] * stride_x_m++ mean = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)+ for k_offset in range(0, K, BLOCK_SIZE_K):+ k_chunk_offs = tl.arange(0, BLOCK_SIZE_K)+ x_ptrs = x_rows_base_ptr + (k_offset + k_chunk_offs)[None, :]+ k_mask = (k_offset + k_chunk_offs) < K+ x_chunk = tl.load(x_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)+ mean += tl.sum(x_chunk, axis=1)+ mean /= K++ var = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)+ for k_offset in range(0, K, BLOCK_SIZE_K):+ k_chunk_offs = tl.arange(0, BLOCK_SIZE_K)+ x_ptrs = x_rows_base_ptr + (k_offset + k_chunk_offs)[None, :]+ k_mask = (k_offset + k_chunk_offs) < K+ x_chunk = tl.load(x_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)+ x_centered = x_chunk - mean[:, None]+ var += tl.sum(x_centered * x_centered, axis=1)+ var /= K+ rstd = 1.0 / tl.sqrt(var + LN_EPS)++ # --- Matmul Loop 1: For the 4-Way Projections ---+ offs_n_4way = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)+ offs_k = tl.arange(0, BLOCK_SIZE_K)+ w_4way_ptrs_base = W_4way_ptr + (offs_n_4way[None, :] * stride_w4_n)+ accumulator_4way = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)+ accumulator_og = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)++ offs_n_og = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)+ for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):+ k_block_start = k * BLOCK_SIZE_K;+ x_ptrs = x_rows_base_ptr + (k_block_start + offs_k)[None, :] * stride_x_k+ w_ptrs = w_4way_ptrs_base + (k_block_start + offs_k)[:, None] * stride_w4_k+ x_mask = (offs_m[:, None] < M) & ((k_block_start + offs_k)[None, :] < K)+ w_mask = ((k_block_start + offs_k)[:, None] < K) & (offs_n_4way[None, :] < N_4way)+ x_tile = tl.load(x_ptrs, mask=x_mask, other=0.0).to(tl.float32)+ norm_w_ptrs = Norm_Weight_ptr + k_block_start + offs_k+ norm_b_ptrs = Norm_Bias_ptr + k_block_start + offs_k+ nw = tl.load(norm_w_ptrs, mask=(k_block_start + offs_k) < K, other=0.0)+ nb = tl.load(norm_b_ptrs, mask=(k_block_start + offs_k) < K, other=0.0)+ x_norm_tile = (x_tile - mean[:, None]) * rstd[:, None]+ x_norm_tile = (x_norm_tile * nw[None, :] + nb[None, :]).to(tl.float16)+ w_tile = tl.load(w_ptrs, mask=w_mask, other=0.0)+ accumulator_4way += tl.dot(x_norm_tile, w_tile)++ #Some threads should calclate out_gate+ if pid_n * BLOCK_SIZE_N < H:+ w_og_ptrs_base = W_og_ptr + (offs_n_og[None, :] * stride_wog_n)+ w_ptrs = w_og_ptrs_base + (k_block_start + offs_k)[:, None] * stride_wog_k+ w_mask = ((k_block_start + offs_k)[:, None] < K) & (offs_n_og[None, :] < H);+ w_tile = tl.load(w_ptrs, mask=w_mask, other=0.0)+ accumulator_og += tl.dot(x_norm_tile, w_tile)++ if pid_n * BLOCK_SIZE_N < H:+ og_out = tl.sigmoid(accumulator_og)+ outg_ptrs = OutOG_ptr + offs_m[:, None] * stride_og_m + offs_n_og[None, :] * stride_og_h+ og_mask = m_mask[:, None] & (offs_n_og[None, :] < H)+ tl.store(outg_ptrs, og_out, mask=og_mask)++ # --- Fusion Logic for 4-Way Part ---+ acc_reshaped = tl.reshape(accumulator_4way, (BLOCK_SIZE_M, H_CHUNK_SIZE, 4))+ role_idx = tl.arange(0, 4)[None, None, :]+ left_proj = tl.sum(tl.where(role_idx == 0, acc_reshaped, 0.0), axis=2)+ left_gate = tl.sum(tl.where(role_idx == 1, acc_reshaped, 0.0), axis=2)+ right_proj = tl.sum(tl.where(role_idx == 2, acc_reshaped, 0.0), axis=2)+ right_gate = tl.sum(tl.where(role_idx == 3, acc_reshaped, 0.0), axis=2)++ offs_h_chunk = (pid_n * H_CHUNK_SIZE) + tl.arange(0, H_CHUNK_SIZE)+ mask_ptrs = Mask_ptr + offs_m[:, None] * stride_mask_m + offs_h_chunk[None, :] * stride_mask_h+ m_mask_h = m_mask[:, None] & (offs_h_chunk[None, :] < H)+ mask_tile = tl.load(mask_ptrs, mask=m_mask_h, other=0.0)++ left_out = left_proj * tl.sigmoid(left_gate) * mask_tile+ right_out = right_proj * tl.sigmoid(right_gate) * mask_tile++ s1s2 = s1 * s2+ offs_b = offs_m // s1s2+ offs_s1 = (offs_m % s1s2) // s2+ offs_s2 = offs_m % s2+ offs_b_2d = tl.reshape(offs_b, (BLOCK_SIZE_M, 1))+ offs_h_2d = tl.reshape(offs_h_chunk, (1, H_CHUNK_SIZE))+ offs_s1_2d = tl.reshape(offs_s1, (BLOCK_SIZE_M, 1))+ offs_s2_2d = tl.reshape(offs_s2, (BLOCK_SIZE_M, 1))++ outl_ptrs = OutLeft_ptr + (offs_b_2d * stride_ol_bs + offs_h_2d * stride_ol_h ++ offs_s1_2d * stride_ol_s1 + offs_s2_2d * stride_ol_s2)+ outr_ptrs_t = OutRight_ptr + (offs_b_2d * stride_or_t_bs + offs_h_2d * stride_or_t_h ++ offs_s2_2d * stride_or_t_s2 + offs_s1_2d * stride_or_t_s1)+ tl.store(outl_ptrs, left_out, mask=m_mask_h)+ tl.store(outr_ptrs_t, right_out, mask=m_mask_h)++ @triton.jit+ def bmm_coalesced_kernel(+ # Pointers+ Left_ptr, Right_ptr, Out_ptr,+ # Dimensions+ bs, s1, s2, H,+ # Strides+ stride_l_bs, stride_l_h, stride_l_s1, stride_l_s2,+ stride_r_bs, stride_r_h, stride_r_s2, stride_r_s1,+ stride_o_bs, stride_o_h, stride_o_s1, stride_o_s2,+ # Kernel parameters+ BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,+ GROUP_SIZE_M: tl.constexpr,+ ):+ # Grid and program IDs+ pid = tl.program_id(axis=0)+ num_pid_m = tl.cdiv(s1, BLOCK_SIZE_M)+ num_pid_n = tl.cdiv(s1, 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++ pid_bh = tl.program_id(axis=1)+ pid_b = pid_bh // H+ pid_h = pid_bh % H++ offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)+ offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)+ offs_k = tl.arange(0, BLOCK_SIZE_K)++ left_ptrs_base = Left_ptr + pid_b * stride_l_bs + pid_h * stride_l_h+ right_ptrs_base = Right_ptr + pid_b * stride_r_bs + pid_h * stride_r_h++ accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)++ for k in range(0, tl.cdiv(s2, BLOCK_SIZE_K)):+ k_start = k * BLOCK_SIZE_K+ a_ptrs = left_ptrs_base + (offs_m[:, None] * stride_l_s1 + (k_start + offs_k[None, :]) * stride_l_s2)+ b_ptrs = right_ptrs_base + ((k_start + offs_k[:, None]) * stride_r_s2 + offs_n[None, :] * stride_r_s1)++ a_mask = (offs_m[:, None] < s1) & ((k_start + offs_k[None, :]) < s2)+ b_mask = ((k_start + offs_k[:, None]) < s2) & (offs_n[None, :] < s1)++ a = tl.load(a_ptrs, mask=a_mask, other=0.0)+ b = tl.load(b_ptrs, mask=b_mask, other=0.0)++ accumulator += tl.dot(a, b)++ out_ptrs = Out_ptr + pid_b * stride_o_bs + pid_h * stride_o_h + \+ offs_m[:, None] * stride_o_s1 + offs_n[None, :] * stride_o_s2++ c_mask = (offs_m[:, None] < s1) & (offs_n[None, :] < s1)+ tl.store(out_ptrs, accumulator, mask=c_mask)++ @triton.jit+ def fused_final_kernel(+ # Pointers+ In_ptr, Gate_ptr, NormW_ptr, NormB_ptr, ProjW_ptr, Out_ptr,+ # Metadata+ M, H, D, s1,+ # Strides+ stride_in_bs, stride_in_h, stride_in_s1_row, stride_in_s1_col,+ stride_gate_m, stride_gate_h,+ stride_proj_d, stride_proj_h,+ stride_out_bs, stride_out_s1_row, stride_out_s1_col, stride_out_d,+ # Constants+ LN_EPS: tl.constexpr,+ 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)+ num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)+ num_pid_n = tl.cdiv(D, 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_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)+ offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)+ m_mask = offs_m < M++ s1s1 = s1 * s1+ b = offs_m // s1s1+ r = (offs_m % s1s1) // s1+ c = offs_m % s1++ sum_x = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)+ sum_x2 = tl.zeros((BLOCK_SIZE_M,), dtype=tl.float32)+ in_ptr_base = In_ptr + b * stride_in_bs + r * stride_in_s1_row + c * stride_in_s1_col++ for k_offset in range(0, H, BLOCK_SIZE_K):+ offs_k = k_offset + tl.arange(0, BLOCK_SIZE_K)+ k_mask = offs_k < H+ in_ptrs = in_ptr_base[:, None] + offs_k[None, :] * stride_in_h+ in_chunk = tl.load(in_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0).to(tl.float32)+ sum_x += tl.sum(in_chunk, axis=1)+ sum_x2 += tl.sum(in_chunk * in_chunk, axis=1)++ mean = sum_x / H+ var = (sum_x2 / H) - (mean * mean)+ rstd = tl.math.rsqrt(var + LN_EPS)++ acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)+ for k_offset in range(0, H, BLOCK_SIZE_K):+ offs_k = k_offset + tl.arange(0, BLOCK_SIZE_K)+ k_mask = offs_k < H+ in_ptrs = in_ptr_base[:, None] + offs_k[None, :] * stride_in_h+ a = tl.load(in_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)+ a_norm = (a - mean[:, None]) * rstd[:, None]+ norm_w = tl.load(NormW_ptr + offs_k, mask=k_mask, other=0.0)+ norm_b = tl.load(NormB_ptr + offs_k, mask=k_mask, other=0.0)+ a_norm = a_norm * norm_w[None, :] + norm_b[None, :]+ proj_ptrs = ProjW_ptr + offs_n[None, :] * stride_proj_d + offs_k[:, None] * stride_proj_h+ gate_ptrs = Gate_ptr + offs_m[:, None] * stride_gate_m + offs_k[None, :] * stride_gate_h+ gate = tl.load(gate_ptrs, mask=m_mask[:, None] & k_mask[None, :], other=0.0)+ a_gated = a_norm * gate+ b_w = tl.load(proj_ptrs, mask=k_mask[:, None] & (offs_n[None, :] < D), other=0.0)+ acc += tl.dot(a_gated.to(b_w.dtype), b_w)++ offs_d = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)+ out_ptr_base = Out_ptr + b*stride_out_bs + r*stride_out_s1_row + c*stride_out_s1_col+ out_ptrs = out_ptr_base[:, None] + offs_d[None, :] * stride_out_d+ tl.store(out_ptrs, acc, mask=m_mask[:, None] & (offs_d[None, :] < D))++ def compiledtrimul_fused_interleaved_final(+ x: torch.Tensor,+ mask_mh: torch.Tensor,+ norm_weight: torch.Tensor,+ norm_bias: torch.Tensor,+ W_4way: torch.Tensor,+ W_og: torch.Tensor,+ to_out_norm_weight: torch.Tensor,+ to_out_norm_bias: torch.Tensor,+ to_out_weight: torch.Tensor,+ h: int,+ ):+ bs, s1, s2, d = x.shape+ M, K, H = bs * s1 * s2, x.shape[-1], h+ x_flat = x.view(M, K)++ left_final = torch.empty((bs, H, s1, s2), device=x.device, dtype=torch.float16)+ right_final_t = torch.empty((bs, H, s2, s1), device=x.device, dtype=torch.float16)+ og_mh = torch.empty((M, H), device=x.device, dtype=torch.float16)++ # --- Kernel 1: Fused LN + Dual Matmul ---+ # The grid is launched for the larger 4*H problem+ N_4way = 4 * H+ # Hardcoded best config from logs: M64-N128-K64-GM8-HC32-W4-S2+ config_k1 = {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8, 'H_CHUNK_SIZE': 32}+ grid = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE_M']) * triton.cdiv(N_4way, meta['BLOCK_SIZE_N']),)++ fused_ln_dual_matmul_kernel[grid](+ x_flat, W_4way, W_og, mask_mh, norm_weight, norm_bias,+ left_final, right_final_t, og_mh,+ M, H, K, s1, s2,+ x_flat.stride(0), x_flat.stride(1), W_4way.stride(0), W_4way.stride(1),+ W_og.stride(0), W_og.stride(1), left_final.stride(0), left_final.stride(1),+ left_final.stride(2), left_final.stride(3), right_final_t.stride(0), right_final_t.stride(1),+ right_final_t.stride(2), right_final_t.stride(3), og_mh.stride(0), og_mh.stride(1),+ mask_mh.stride(0), mask_mh.stride(1),+ LN_EPS=1e-5, **config_k1, num_warps=4, num_stages=2+ )++ # --- Kernel 2: Batched Matrix Multiplication ---+ bmm_out_tmp = torch.empty((bs, H, s1, s1), device=x.device, dtype=torch.float16)+ # Hardcoded best config from logs: M128-N128-K32-GM8-W8-S3+ config_k2 = {'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}+ grid_bmm = lambda meta: (triton.cdiv(s1, meta['BLOCK_SIZE_M']) * triton.cdiv(s1, meta['BLOCK_SIZE_N']), bs * H)++ bmm_coalesced_kernel[grid_bmm](+ left_final, right_final_t, bmm_out_tmp,+ bs, s1, s2, H,+ left_final.stride(0), left_final.stride(1), left_final.stride(2), left_final.stride(3),+ right_final_t.stride(0), right_final_t.stride(1), right_final_t.stride(2), right_final_t.stride(3),+ bmm_out_tmp.stride(0), bmm_out_tmp.stride(1), bmm_out_tmp.stride(2), bmm_out_tmp.stride(3),+ **config_k2, num_warps=8, num_stages=3+ )++ # --- Kernel 3: Fully Fused Final Stage ---+ final_out = torch.empty((bs, s1, s1, d), device=x.device, dtype=torch.float16)+ # Hardcoded best config from logs: M32-N128-K32-GM8-W4-S3+ config_k3 = {'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 32, 'GROUP_SIZE_M': 8}+ grid_final = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE_M']) * triton.cdiv(d, meta['BLOCK_SIZE_N']),)++ fused_final_kernel[grid_final](+ bmm_out_tmp, og_mh, to_out_norm_weight, to_out_norm_bias, to_out_weight, final_out,+ M, H, d, s1,+ bmm_out_tmp.stride(0), bmm_out_tmp.stride(1), bmm_out_tmp.stride(2), bmm_out_tmp.stride(3),+ og_mh.stride(0), og_mh.stride(1), to_out_weight.stride(0), to_out_weight.stride(1),+ final_out.stride(0), final_out.stride(1), final_out.stride(2), final_out.stride(3),+ LN_EPS=1e-5, **config_k3, num_warps=4, num_stages=3+ )+ return final_out++ def pack_w_4way_efficient(weights):+ """ Packs L, LG, R, RG into a tight [K, 4*H] matrix. """+ WL, WLG, WR, WRG = (weights[k] for k in ['left_proj.weight', 'left_gate.weight', 'right_proj.weight', 'right_gate.weight'])+ H, K = WL.shape+ ws = torch.stack([WL, WLG, WR, WRG], dim=0).permute(1, 0, 2).contiguous().view(4 * H, K)+ return ws.t().to(torch.float16)++ def get_w_og(weights):+ """ Gets the transposed [K, H] out_gate weight matrix. """+ return weights['out_gate.weight'].t().to(torch.float16)++ @torch.compile()+ def compiledtrimul(+ x: torch.Tensor, mask: torch.Tensor, norm_weight: torch.Tensor, norm_bias: torch.Tensor,+ w_concat: torch.Tensor, to_out_norm_weight: torch.Tensor, to_out_norm_bias: torch.Tensor,+ to_out_weight: torch.Tensor, h: int+ ) -> torch.Tensor:+ bs, s1, s2, d = x.shape+ x_norm = F.layer_norm(x, (d,), norm_weight, norm_bias).view((bs * s1 * s2, d)).to(torch.float16)+ all_projections = torch.mm(x_norm, w_concat)+ left, right, lg, rg, og = all_projections.chunk(5, dim=1)+ mask_expanded = mask.expand(-1, -1, -1, h).reshape(-1, h)+ left = left * mask_expanded * torch.sigmoid(lg)+ right = right * mask_expanded * torch.sigmoid(rg)+ out_gate = torch.sigmoid(og)+ left = left.view(bs, s1, s2, h).permute(0,3,1,2)+ right = right.view(bs, s1, s2, h).permute(0,3,1,2)+ out_p = torch.matmul(left.to(torch.float16), right.to(torch.float16).transpose(-1, -2))+ out_einsum_flat = out_p.permute(0,2,3,1).reshape(bs * s1 * s1, h)+ normed = F.layer_norm(out_einsum_flat, (h,), to_out_norm_weight, to_out_norm_bias).to(torch.float16)+ gated = normed * out_gate+ final_out_flat = gated @ to_out_weight.t()+ return final_out_flat.view(bs, s1, s1, d)++ def small_kernel_pt_path(data):+ input_tensor, mask, weights, config = data+ w_concat = torch.cat([+ weights['left_proj.weight'], weights['right_proj.weight'], weights['left_gate.weight'],+ weights['right_gate.weight'], weights['out_gate.weight']+ ], dim=0).t().contiguous().to(torch.float16)+ return compiledtrimul(+ x=input_tensor.to(torch.float32), mask=mask.unsqueeze(-1),+ norm_weight=weights['norm.weight'].to(torch.float32),+ norm_bias=weights['norm.bias'].to(torch.float32), w_concat=w_concat,+ to_out_norm_weight=weights['to_out_norm.weight'].to(torch.float16),+ to_out_norm_bias=weights['to_out_norm.bias'].to(torch.float16),+ to_out_weight=weights['to_out.weight'].to(torch.float16),+ h=config["hidden_dim"]+ )++ def custom_kernel(data):+ input_tensor, mask, weights, config = data+ bs, s1, s2, d = input_tensor.shape++ if s1 < 800:+ return small_kernel_pt_path(data)++ H = config["hidden_dim"]+ W_4way = pack_w_4way_efficient(weights)+ W_og = get_w_og(weights)+ M = bs * s1 * s2+ mask_mh = mask.unsqueeze(-1).expand(-1, -1, -1, H).reshape(M, H).to(torch.float16)++ return compiledtrimul_fused_interleaved_final(+ x=input_tensor.to(torch.float32),+ mask_mh=mask_mh,+ norm_weight=weights['norm.weight'].to(torch.float32),+ norm_bias=weights['norm.bias'].to(torch.float32),+ W_4way=W_4way,+ W_og=W_og,+ to_out_norm_weight=weights['to_out_norm.weight'].to(torch.float16),+ to_out_norm_bias=weights['to_out_norm.bias'].to(torch.float16),+ to_out_weight=weights['to_out.weight'].to(torch.float16),+ h=H,+ )
scrolls · 417 diff lines total
Best evidence level for this revision: reported
JSON