## 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_073 Conv2d_BatchNorm_Scaling, 1.68x): adapt what applies
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl


@triton.jit
def stats_kernel2(x_ptr, sum_ptr, sumsq_ptr, N, C, S, SBLK: tl.constexpr):
    c = tl.program_id(0)
    sblk = tl.program_id(1)
    offs = sblk * SBLK + tl.arange(0, SBLK)
    smask = offs < S
    acc_sum = tl.zeros((SBLK,), tl.float32)
    acc_sq = tl.zeros((SBLK,), tl.float32)
    cS = c.to(tl.int64) * S
    NCS = C * S
    for n in range(N):
        base = n.to(tl.int64) * NCS + cS + offs
        v = tl.load(x_ptr + base, mask=smask, other=0.0)
        acc_sum += v
        acc_sq += v * v
    s = tl.sum(acc_sum, axis=0)
    sq = tl.sum(acc_sq, axis=0)
    tl.atomic_add(sum_ptr + c, s)
    tl.atomic_add(sumsq_ptr + c, sq)


@triton.jit
def bn_apply_kernel(x_ptr, out_ptr, mean_ptr, invstd_ptr, w_ptr, b_ptr,
                    scaling, C, S, TOTAL, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < TOTAL
    c = (offs // S) % C
    v = tl.load(x_ptr + offs, mask=mask, other=0.0)
    mean = tl.load(mean_ptr + c, mask=mask, other=0.0)
    invstd = tl.load(invstd_ptr + c, mask=mask, other=0.0)
    w = tl.load(w_ptr + c, mask=mask, other=0.0)
    b = tl.load(b_ptr + c, mask=mask, other=0.0)
    r = ((v - mean) * invstd * w + b) * scaling
    tl.store(out_ptr + offs, r, mask=mask)


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

    def forward(self, x):
        x = self.conv(x).contiguous()
        N, C, H, W = x.shape
        S = H * W
        count = N * S
        device = x.device

        sums = torch.zeros(C, device=device, dtype=torch.float32)
        sumsq = torch.zeros(C, device=device, dtype=torch.float32)
        SBLK = 1024
        nsblk = triton.cdiv(S, SBLK)
        stats_kernel2[(C, nsblk)](x, sums, sumsq, N, C, S, SBLK=SBLK)

        if self.training:
            mean = sums / count
            var = sumsq / count - mean * mean
            with torch.no_grad():
                self.bn.running_mean.mul_(1 - self.bn.momentum).add_(self.bn.momentum * mean.detach())
                unbiased = var * count / (count - 1)
                self.bn.running_var.mul_(1 - self.bn.momentum).add_(self.bn.momentum * unbiased.detach())
        else:
            mean = self.bn.running_mean
            var = self.bn.running_var

        invstd = torch.rsqrt(var + self.bn.eps)
        w = self.bn.weight
        b = self.bn.bias

        out = torch.empty_like(x)
        TOTAL = x.numel()
        BLOCK = 1024
        grid = (triton.cdiv(TOTAL, BLOCK),)
        bn_apply_kernel[grid](x, out, mean.contiguous(), invstd.contiguous(),
                              w.contiguous(), b.contiguous(),
                              float(self.scaling_factor), C, S, TOTAL, BLOCK=BLOCK)
        return out
```

### A verified kernel for a related conv problem (L2_071 Conv2d_Divide_LeakyReLU, 1.26x): adapt what applies
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl


@triton.jit
def div_lrelu_kernel(x_ptr, out_ptr, inv_div, TOTAL, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < TOTAL
    v = tl.load(x_ptr + offs, mask=mask, other=0.0) * inv_div
    r = tl.where(v >= 0.0, v, v * 0.01)
    tl.store(out_ptr + offs, r, mask=mask)


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

    def forward(self, x):
        x = self.conv(x).contiguous()
        out = torch.empty_like(x)
        TOTAL = x.numel()
        BLOCK = 1024
        grid = (triton.cdiv(TOTAL, BLOCK),)
        div_lrelu_kernel[grid](x, out, 1.0 / float(self.divisor), TOTAL, BLOCK=BLOCK)
        return out
```