## 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: Before hand-writing a dense conv, benchmark against cuDNN; a naive Triton or CUDA kernel rarely wins. Reserve custom kernels for fusing cheap elementwise ops around the conv.
- 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.

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


@triton.jit
def fused_kernel(x_ptr, out_ptr, sub, N, C, H, W, OH, OW, P, TOTAL, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < TOTAL
    ow = offs % OW
    tmp = offs // OW
    oh = tmp % OH
    tmp = tmp // OH
    c = tmp % C
    n = tmp // C

    base = (n.to(tl.int64) * C + c) * H * W
    h0 = oh * P
    w0 = ow * P
    acc = tl.full((BLOCK,), -3.4e38, tl.float32)
    for ph in range(0, P):
        for pw in range(0, P):
            ih = h0 + ph
            iw = w0 + pw
            vmask = mask & (ih < H) & (iw < W)
            v = tl.load(x_ptr + base + ih.to(tl.int64) * W + iw, mask=vmask, other=-3.4e38)
            v = v - sub
            hs = v * tl.minimum(tl.maximum(v + 3.0, 0.0), 6.0) / 6.0
            acc = tl.maximum(acc, hs)
    # mish: x * tanh(softplus(x))
    sp = tl.log(1.0 + tl.exp(acc))
    th = (tl.exp(sp) - tl.exp(-sp)) / (tl.exp(sp) + tl.exp(-sp))
    res = acc * th
    tl.store(out_ptr + offs, res, mask=mask)


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

    def forward(self, x):
        x = self.conv(x).contiguous()
        N, C, H, W = x.shape
        P = self.pool_k
        OH = H // P
        OW = W // P
        out = torch.empty((N, C, OH, OW), device=x.device, dtype=torch.float32)
        TOTAL = out.numel()
        BLOCK = 256
        grid = (triton.cdiv(TOTAL, BLOCK),)
        fused_kernel[grid](x, out, float(self.subtract_value),
                           N, C, H, W, OH, OW, P, TOTAL, BLOCK=BLOCK)
        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
```