## 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_transpose problems (wrong results or slower than eager)
- Pitfall: Don't rewrite a transposed conv as flipped-weight F.conv3d plus a weight-flip kernel. It still calls cuDNN, adds a launch and allocation per forward, and ran slower than eager ConvTranspose.
- Pitfall: Avoid launching a per-forward weight-transform kernel (flip/transpose) before a library call. Cache the transformed weight once, or fuse the transform, since extra launches and memory traffic make it slower than eager.
- Pitfall: Don'...'s native transposed conv is usually as fast or faster, so this measured slower.
- Pitfall: In Triton kernels over large tensors (e.g. 3D conv outputs), compute offsets in int64 (cast program_id/indices with .to(tl.int64)). int32 products like n*C*S can overflow and crash, leaving no output during timing.
- Pitfall: Never launch a pure copy/identity kernel into a fresh buffer just to have a custom kernel; that adds a full memory round-trip and guarantees slowdown versus eager.
- Pitfall: Don't rewrite conv_transpose as flipped-weight F.conv2d plus a separate elementwise kernel; the per-call weight flip and extra full-output pass make it slower. Fuse bias into a real kernel, or skip pointless copies.

### A verified kernel for a related conv_transpose problem (L2_015 ConvTranspose3d_BatchNorm_Subtract, 1.32x): adapt what applies
```python
import torch
import torch.nn as nn
import triton
import triton.language as tl


@triton.jit
def reduce_kernel(x_ptr, sum_ptr, sqsum_ptr, S, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = tl.arange(0, BLOCK)
    base = pid * S
    acc = tl.zeros((), dtype=tl.float32)
    accsq = tl.zeros((), dtype=tl.float32)
    for i in range(0, S, BLOCK):
        idx = i + offs
        m = idx < S
        v = tl.load(x_ptr + base + idx, mask=m, other=0.0)
        acc += tl.sum(v)
        accsq += tl.sum(v * v)
    tl.store(sum_ptr + pid, acc)
    tl.store(sqsum_ptr + pid, accsq)


@triton.jit
def out_kernel(x_ptr, out_ptr, scale_ptr, smean_ptr, S, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = tl.arange(0, BLOCK)
    base = pid * S
    sc = tl.load(scale_ptr + pid)
    sm = tl.load(smean_ptr + pid)
    for i in range(0, S, BLOCK):
        idx = i + offs
        m = idx < S
        v = tl.load(x_ptr + base + idx, mask=m, other=0.0)
        r = sc * (v - sm)
        tl.store(out_ptr + base + idx, r, mask=m)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, bias=True):
        super(ModelNew, self).__init__()
        self.conv_transpose = nn.ConvTranspose3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding, bias=bias)
        self.batch_norm = nn.BatchNorm3d(out_channels)

    def forward(self, x):
        c = self.conv_transpose(x)
        N, C, D, H, W = c.shape
        S = D * H * W
        c2 = c.contiguous().view(N * C, S)

        ssum = torch.empty(N * C, device=c.device, dtype=torch.float32)
        ssq = torch.empty(N * C, device=c.device, dtype=torch.float32)
        reduce_kernel[(N * C,)](c2, ssum, ssq, S, BLOCK=1024)

        ssum_nc = ssum.view(N, C)
        ssq_nc = ssq.view(N, C)
        total_sum = ssum_nc.sum(0)
        total_sq = ssq_nc.sum(0)
        M = N * S
        mean_c = total_sum / M
        var_c = total_sq / M - mean_c * mean_c
        scale_c = self.batch_norm.weight / torch.sqrt(var_c + self.batch_norm.eps)  # (C,)

        scale_nc = scale_c.unsqueeze(0).expand(N, C).contiguous().view(N * C)
        smean_nc = (ssum_nc / S).contiguous().view(N * C)

        out = torch.empty_like(c2)
        out_kernel[(N * C,)](c2, out, scale_nc, smean_nc, S, BLOCK=1024)
        return out.view(N, C, D, H, W)
```

### A verified kernel for a related conv_transpose problem (L2_089 ConvTranspose3d_MaxPool_Softmax_Subtract_Swish_Max, 1.06x): 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, sub_ptr, out_ptr, N, C, S, total, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    idx = pid * BLOCK + tl.arange(0, BLOCK)
    mask = idx < total

    # idx indexes (n, s) flat where s in [0,S)
    n = idx // S
    s = idx % S
    base = n * C * S + s

    # pass 1: max over channels
    mx = tl.full((BLOCK,), -1e30, tl.float32)
    for c in range(0, C):
        v = tl.load(x_ptr + base + c * S, mask=mask, other=-1e30)
        mx = tl.maximum(mx, v)

    # pass 2: sum of exp
    sm = tl.zeros((BLOCK,), tl.float32)
    for c in range(0, C):
        v = tl.load(x_ptr + base + c * S, mask=mask, other=0.0)
        sm += tl.exp(v - mx)

    # pass 3: softmax, subtract, swish, max over channels
    res = tl.full((BLOCK,), -1e30, tl.float32)
    for c in range(0, C):
        v = tl.load(x_ptr + base + c * S, mask=mask, other=0.0)
        sub = tl.load(sub_ptr + c)
        sm_val = tl.exp(v - mx) / sm
        y = sm_val - sub
        swish = (tl.sigmoid(y) * y)
        res = tl.maximum(res, swish)

    tl.store(out_ptr + idx, res, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, pool_kernel_size, pool_stride, pool_padding):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding, output_padding=output_padding)
        self.max_pool = nn.MaxPool3d(kernel_size=pool_kernel_size, stride=pool_stride, padding=pool_padding)
        self.subtract = nn.Parameter(torch.randn(out_channels))

    def forward(self, x):
        x = self.conv_transpose(x)
        x = self.max_pool(x).contiguous()
        N, C, D, H, W = x.shape
        S = D * H * W
        total = N * S

        out = torch.empty((N, D, H, W), device=x.device, dtype=x.dtype)
        sub = self.subtract.contiguous()

        BLOCK = 256
        grid = (triton.cdiv(total, BLOCK),)
        fused_kernel[grid](x, sub, out, N, C, S, total, BLOCK=BLOCK)
        return out
```