## Technique library (KFORGE: techniques verified by the protected evaluator on training problems; use what fits this problem, ignore the rest)

### Verified technique (E-H1, admitted by the KFORGE engine): Dense Conv2d/Conv3d as a Triton implicit GEMM with an in-register bf16 3-product split and a fused epilogue
For dense convolutions with reduction length K = Ci*kh*kw (*kd) >= ~256 and Co >= 64, a Triton implicit-GEMM kernel should beat strict-fp32 cuDNN by 1.5-3x on the conv itself, and more once bias/activation/pool epilogues are fused. The kernel splits each loaded fp32 operand into bf16 hi+lo in registers and issues three bf16 tl.dot calls into one fp32 accumulator.
Physics: cuDNN under strict fp32 runs on the 19.5 TFLOP/s SIMT path. A100 bf16 mma peaks at 312 TFLOP/s, so even at 40% efficiency the three products deliver about 40 TFLOP/s of fp32-equivalent work. P1 already shows the split stays within about 0.34 of the tolerance for K <= 8192, and conv K is almost always below that. This is not the tf32x3 path that measured only 1.1-1.2x: bf16 mma has 2x the TF32 rate, and the hi/lo split is exact in fp32 adds. Implicit GEMM never materializes im2col, so it avoids the bandwidth loss of the AlexNet failure.
How: Map GEMM M = N*OH*OW (output pixels), GEMM N = Co, and K = (r, s, ci), with ci innermost. Use one cheap Triton prep kernel to convert x to channels_last and emit xh = bf16(x) and xl = bf16(x - float(xh)); these two bf16 tensors take the same bytes as fp32. Cache the weight reorganised as [kh*kw*Ci, Co] in hi/lo bf16 form in self.<x>, keyed on weight.data_ptr(). In the kernel, loop over (r, s), compute shifted input coordinates with padding masks, and load BK = 32 channel slabs of xh/xl and wh/wl. Accumulate acc += dot(xl, wh) + dot(xh, wl) first, then acc += dot(xh, wh), with fp32 accumulation. Use BM = 128, BN = 64 or 128, num_stages = 3-4, and a small autotune list. In the epilogue, add bias, apply the following elementwise ops (and pool if the window fits the tile), and store straight to NCHW with int64 offsets. Consecutive M indices are consecutive ow, so the stores coalesce. Validate against F.conv2d at atol = rtol = 1e-4 before timing, and fall back to cuDNN when K < 128 or Co < 32.
- Verified on training problems (protected evaluator): L1_080 conv_standard_2D_square_input_asymmetric_kernel___dilated____padded__: 2.29x vs 0.44x with the previous library; L2_032 Conv2d_Scaling_Min: 2.50x vs 1.11x with the previous library.
- Verified example implementing it (L2_032, 2.50x; adapt what applies):
```python
import torch
import torch.nn as nn
import triton
import triton.language as tl


@triton.autotune(
    configs=[
        triton.Config({"BM": 128, "BK": 32}, num_stages=3, num_warps=8),
        triton.Config({"BM": 64, "BK": 32}, num_stages=4, num_warps=4),
        triton.Config({"BM": 128, "BK": 64}, num_stages=2, num_warps=8),
    ],
    key=["M", "Ci", "Co"],
)
@triton.jit
def conv_min_kernel(x, wh, wl, b, out, M, Ci, H, W, OH, OW, Co, scale,
                    KH: tl.constexpr, KW: tl.constexpr, BN: tl.constexpr,
                    BM: tl.constexpr, BK: tl.constexpr):
    pid = tl.program_id(0)
    rm = pid * BM + tl.arange(0, BM)
    mm = rm < M
    rmc = tl.where(mm, rm, 0)
    ow = rmc % OW
    t = rmc // OW
    oh = t % OH
    n = t // OH
    base = n * (Ci * H * W) + oh * W + ow
    rk = tl.arange(0, BK)
    rn = tl.arange(0, BN)
    nm = rn < Co
    HW = H * W
    nci = tl.cdiv(Ci, BK)
    total = KH * KW * nci
    acc = tl.zeros((BM, BN), dtype=tl.float32)
    for kk in range(0, total):
        tap = kk // nci
        c0 = (kk - tap * nci) * BK
        r = tap // KW
        s = tap - r * KW
        ci = c0 + rk
        cm = ci < Ci
        a = tl.load(x + base[:, None] + (r * W + s) + ci[None, :] * HW,
                    mask=cm[None, :], other=0.0)
        krow = tap * Ci + ci
        boff = krow[:, None] * Co + rn[None, :]
        bmask = cm[:, None] & nm[None, :]
        bh = tl.load(wh + boff, mask=bmask, other=0.0)
        bl = tl.load(wl + boff, mask=bmask, other=0.0)
        ah = a.to(tl.bfloat16)
        al = (a - ah.to(tl.float32)).to(tl.bfloat16)
        acc = tl.dot(al, bh, acc)
        acc = tl.dot(ah, bl, acc)
        acc = tl.dot(ah, bh, acc)
    bias = tl.load(b + rn, mask=nm, other=0.0)
    v = (acc + bias[None, :]) * scale
    v = tl.where(nm[None, :], v, float("inf"))
    mn = tl.min(v, axis=1)
    tl.store(out + rm, mn, mask=mm)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, scale_factor):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.scale_factor = scale_factor

    def forward(self, x):
        x = x.contiguous()
        w = self.conv.weight
        Co, Ci, KH, KW = w.shape
        N, _, H, W = x.shape
        OH = H - KH + 1
        OW = W - KW + 1
        wr = w.permute(2, 3, 1, 0).reshape(KH * KW * Ci, Co).contiguous()
        whi = wr.to(torch.bfloat16)
        wlo = (wr - whi.float()).to(torch.bfloat16)
        out = torch.empty((N, 1, OH, OW), device=x.device, dtype=torch.float32)
        M = N * OH * OW
        BN = max(16, triton.next_power_of_2(Co))
        grid = lambda META: (triton.cdiv(M, META["BM"]),)
        conv_min_kernel[grid](x, whi, wlo, self.conv.bias, out, M, Ci, H, W, OH, OW, Co,
                              float(self.scale_factor), KH=KH, KW=KW, BN=BN)
        return out
```


## Kernel memory (items verified by the protected evaluator on training problems; may not all apply)

### Lessons about this environment (from official checker verdicts and verified outcomes)
- Pitfall: A reply without a ```python fenced code block was rejected by the checker. Always wrap the full ModelNew source in a ```python block.
- Checker rule: getattr needs a literal, non-private attribute name
- Pitfall: Output mismatch beyond 1e-4 occurred repeatedly. Fix by forcing exact fp32 tl.dot(allow_tf32=False) and exact masking: K bounds at block granularity, tril mask, and zero-initialized output for skipped tiles.
- Checker rule: forward modified its input tensors in place
- Checker rule: do not define a class named Model (the reference); define ModelNew only
- Process: never SUBMIT code that has not passed ACTION: CHECK in this episode. Untested submissions that fail to compile or mismatch the reference score 0; a CHECK costs one reply and shows the compiler output.
- Pitfall: cpp_sources is compiled by plain c++ without CUDA headers: never use cudaStream_t or at::cuda::getCurrentCUDAStream there. Put wrappers needing streams in cuda_sources, keeping only plain-type prototypes in cpp_sources.
- Pitfall: If cpp_sources must reference CUDA types or the current stream, explicitly #include <cuda_runtime.h> and <ATen/cuda/CUDAContext.h>; torch/extension.h does not provide them.
- Checker rule: assignment to attribute 'm.conv1' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.bn1' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.conv2' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.bn2' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.downsample' is not allowed (only self.<x>)
- Checker rule: At least one of your own kernels must actually launch during forward; a pure torch path (even with a compiled but unused kernel) is rejected as no_custom_kernel.

### Pitfalls seen on earlier conv problems (wrong results or slower than eager)
- Pitfall: Don't write direct convolution as one output channel per program with scalar weight loads; each input pixel gets re-read per output channel. Tile over output channels and use tl.dot (implicit GEMM).
- Pitfall: Avoid per-element div/mod index decoding plus per-tap weight gathers in stencil kernels. Use a 2D/3D grid (plane, row-tile, col-tile) and load per-channel weights once per program.
- Pitfall: For depthwise conv, don'...'s 2D output plane, reuse input rows via shared memory or registers, and make kernel dims constexpr.
- Pitfall: Fuse post-conv elementwise/reduction ops with following pooling into one kernel instead of separate launches that write and re-read full-size intermediates. Unfused memory-bound epilogues after cuDNN conv easily end up slower than eager.
- Pitfall: For softmax or reduction over the channel dim of NC... tensors, avoid one program per spatial position with channel-strided loads. Tile a 2D [channels × contiguous spatial block] so loads coalesce along the innermost dimension.
- Pitfall: Don't wrap the cuDNN/PyTorch op and then add a no-op copy kernel. The extra full-tensor read/write plus launch only adds time. Write a real fused kernel that replaces or eliminates work.

### A verified kernel for a related conv problem (L1_083 conv_depthwise_2D_square_input_asymmetric_kernel, 3.01x): adapt what applies
```python
import torch
import torch.nn as nn
import triton
import triton.language as tl


@triton.jit
def dwconv_kernel(x_ptr, w_ptr, b_ptr, out_ptr,
                  N, C, H, W, OH, OW,
                  stride, padding, dilation,
                  K: tl.constexpr, HAS_BIAS: tl.constexpr,
                  BLOCK: tl.constexpr, ROWS: tl.constexpr):
    pid_row = tl.program_id(0)
    pid_col = tl.program_id(1)

    oh_blk = pid_row % ((OH + ROWS - 1) // ROWS)
    tmp = pid_row // ((OH + ROWS - 1) // ROWS)
    c = tmp % C
    n = tmp // C

    ow = pid_col * BLOCK + tl.arange(0, BLOCK)
    mask = ow < OW

    iw = ow * stride - padding
    wvalid = mask & (iw >= 0) & (iw < W)

    base = (n.to(tl.int64) * C + c) * H * W
    wbase = c * K

    oh0 = oh_blk * ROWS

    for r in tl.static_range(ROWS):
        oh = oh0 + r
        if oh < OH:
            acc = tl.zeros((BLOCK,), tl.float32)
            ih0 = oh * stride - padding
            for k in tl.static_range(K):
                ih = ih0 + k * dilation
                if (ih >= 0) & (ih < H):
                    idx = base + ih.to(tl.int64) * W + iw
                    v = tl.load(x_ptr + idx, mask=wvalid, other=0.0)
                    wv = tl.load(w_ptr + wbase + k)
                    acc += v * wv
            if HAS_BIAS:
                bv = tl.load(b_ptr + c)
                acc += bv
            obase = ((n.to(tl.int64) * C + c) * OH + oh) * OW
            tl.store(out_ptr + obase + ow, acc, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels: int, kernel_size: int, stride: int = 1, padding: int = 0, dilation: int = 1, bias: bool = False):
        super(ModelNew, self).__init__()
        self.conv2d = nn.Conv2d(in_channels, in_channels, kernel_size=(kernel_size, 1),
                                stride=stride, padding=padding, dilation=dilation,
                                groups=in_channels, bias=bias)
        self.in_channels = in_channels
        self.kernel_size = kernel_size
        self.stride = stride
        self.padding = padding
        self.dilation = dilation
        self.has_bias = bias

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        x = x.contiguous()
        N, C, H, W = x.shape
        K = self.kernel_size
        stride = self.stride
        padding = self.padding
        dilation = self.dilation
        OH = (H + 2 * padding - dilation * (K - 1) - 1) // stride + 1
        OW = (W + 2 * padding - 1) // stride + 1
        out = torch.empty((N, C, OH, OW), device=x.device, dtype=x.dtype)
        w = self.conv2d.weight.contiguous().view(-1)
        if self.has_bias:
            b = self.conv2d.bias.contiguous()
        else:
            b = torch.empty(1, device=x.device, dtype=x.dtype)
        BLOCK = 256
        ROWS = 4
        row_blocks = (OH + ROWS - 1) // ROWS
        grid = (N * C * row_blocks, triton.cdiv(OW, BLOCK))
        dwconv_kernel[grid](x, w, b, out,
                            N, C, H, W, OH, OW,
                            stride, padding, dilation,
                            K=K, HAS_BIAS=self.has_bias, BLOCK=BLOCK, ROWS=ROWS,
                            num_warps=4)
        return out
```

### A verified kernel for a related conv problem (L2_092 Conv2d_GroupNorm_Tanh_HardSwish_ResidualAdd_LogSumExp, 2.37x): adapt what applies
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl
from triton.language.extra.cuda import libdevice


@triton.jit
def gn_stats_kernel(x_ptr, mean_ptr, invstd_ptr, CPG, S, GS, eps, BLOCK: tl.constexpr):
    pid = tl.program_id(0)  # over n*groups
    base = pid.to(tl.int64) * GS
    acc = 0.0
    acc2 = 0.0
    for off in range(0, GS, BLOCK):
        offs = off + tl.arange(0, BLOCK)
        mask = offs < GS
        v = tl.load(x_ptr + base + offs, mask=mask, other=0.0)
        acc += tl.sum(v, axis=0)
        acc2 += tl.sum(v * v, axis=0)
    mean = acc / GS
    var = acc2 / GS - mean * mean
    invstd = 1.0 / tl.sqrt(var + eps)
    tl.store(mean_ptr + pid, mean)
    tl.store(invstd_ptr + pid, invstd)


@triton.jit
def fused_kernel(x_ptr, mean_ptr, invstd_ptr, gamma_ptr, beta_ptr, out_ptr,
                 N, C, S, NS, G, CPG, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < NS
    n = offs // S
    s = offs % S
    base = n.to(tl.int64) * C * S + s
    maxv = tl.full((BLOCK,), -3.4e38, tl.float32)
    for c in range(C):
        g = c // CPG
        ng = n * G + g
        xc = tl.load(x_ptr + base + c * S, mask=mask, other=0.0)
        mean = tl.load(mean_ptr + ng, mask=mask, other=0.0)
        invstd = tl.load(invstd_ptr + ng, mask=mask, other=0.0)
        gamma = tl.load(gamma_ptr + c)
        beta = tl.load(beta_ptr + c)
        norm = (xc - mean) * invstd * gamma + beta
        t = libdevice.tanh(norm)
        hs = t * tl.minimum(tl.maximum(t + 3.0, 0.0), 6.0) / 6.0
        res = xc + hs
        maxv = tl.maximum(maxv, res)
    sumexp = tl.zeros((BLOCK,), tl.float32)
    for c in range(C):
        g = c // CPG
        ng = n * G + g
        xc = tl.load(x_ptr + base + c * S, mask=mask, other=0.0)
        mean = tl.load(mean_ptr + ng, mask=mask, other=0.0)
        invstd = tl.load(invstd_ptr + ng, mask=mask, other=0.0)
        gamma = tl.load(gamma_ptr + c)
        beta = tl.load(beta_ptr + c)
        norm = (xc - mean) * invstd * gamma + beta
        t = libdevice.tanh(norm)
        hs = t * tl.minimum(tl.maximum(t + 3.0, 0.0), 6.0) / 6.0
        res = xc + hs
        sumexp += tl.exp(res - maxv)
    out = maxv + tl.log(sumexp)
    tl.store(out_ptr + offs, out, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, groups, eps=1e-5):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.group_norm = nn.GroupNorm(groups, out_channels, eps=eps)
        self.tanh = nn.Tanh()
        self.hard_swish = nn.Hardswish()
        self.groups = groups
        self.eps = eps

    def forward(self, x):
        x_conv = self.conv(x).contiguous()
        N, C, H, W = x_conv.shape
        S = H * W
        G = self.groups
        CPG = C // G
        GS = CPG * S
        mean = torch.empty((N * G,), device=x_conv.device, dtype=torch.float32)
        invstd = torch.empty((N * G,), device=x_conv.device, dtype=torch.float32)
        gn_stats_kernel[(N * G,)](x_conv, mean, invstd, CPG, S, GS, self.eps, BLOCK=1024)
        out = torch.empty((N, 1, H, W), device=x_conv.device, dtype=x_conv.dtype)
        NS = N * S
        BLOCK = 256
        grid = (triton.cdiv(NS, BLOCK),)
        fused_kernel[grid](x_conv, mean, invstd,
                           self.group_norm.weight, self.group_norm.bias,
                           out, N, C, S, NS, G, CPG, BLOCK=BLOCK)
        return out
```