## 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 (L1_065 conv_transposed_2D__square_input__asymmetric_kernel, 2.21x): 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 flip_kernel(w_ptr, o_ptr, IC, OC, KH, KW, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = IC * OC * KH * KW
    mask = offs < total
    o = offs
    kw = o % KW
    t = o // KW
    kh = t % KH
    t = t // KH
    oc = t % OC
    ic = t // OC
    src = ((ic * OC + oc) * KH + kh) * KW + kw
    fkh = KH - 1 - kh
    fkw = KW - 1 - kw
    dst = ((oc * IC + ic) * KH + fkh) * KW + fkw
    v = tl.load(w_ptr + src, mask=mask, other=0.0)
    tl.store(o_ptr + dst, v, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels: int, out_channels: int, kernel_size: tuple, stride: int = 1, padding: int = 0, output_padding: int = 0, groups: int = 1, bias: bool = False):
        super(ModelNew, self).__init__()
        self.conv_transpose2d = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride=stride, padding=padding, output_padding=output_padding, groups=groups, bias=bias)
        self.stride = stride
        self.padding = padding
        self.output_padding = output_padding
        self.groups = groups
        self.has_bias = bias

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        pad = self.padding
        if isinstance(pad, tuple):
            ph, pw = pad
        else:
            ph = pw = pad
        simple = (self.stride == 1 and self.output_padding == 0 and self.groups == 1 and x.is_cuda)
        if simple:
            w = self.conv_transpose2d.weight  # (IC, OC, KH, KW)
            IC, OC, KH, KW = w.shape
            wf = torch.empty((OC, IC, KH, KW), device=w.device, dtype=w.dtype)
            total = IC * OC * KH * KW
            BLOCK = 256
            flip_kernel[(triton.cdiv(total, BLOCK),)](
                w.contiguous(), wf, IC, OC, KH, KW, BLOCK=BLOCK)
            pad_h = KH - 1 - ph
            pad_w = KW - 1 - pw
            b = self.conv_transpose2d.bias if self.has_bias else None
            return F.conv2d(x, wf, bias=b, stride=1, padding=(pad_h, pad_w))
        return self.conv_transpose2d(x)
```

### A verified kernel for a related conv_transpose problem (L1_061 conv_transposed_3D__square_input__square_kernel, 1.01x): 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 flip_kernel(w_ptr, o_ptr, IC, OC, KD, KH, KW, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = IC * OC * KD * KH * KW
    mask = offs < total
    o = offs
    kw = o % KW
    t = o // KW
    kh = t % KH
    t = t // KH
    kd = t % KD
    t = t // KD
    oc = t % OC
    ic = t // OC
    src = ((((ic * OC + oc) * KD + kd) * KH + kh) * KW + kw)
    fkd = KD - 1 - kd
    fkh = KH - 1 - kh
    fkw = KW - 1 - kw
    dst = ((((oc * IC + ic) * KD + fkd) * KH + fkh) * KW + fkw)
    v = tl.load(w_ptr + src, mask=mask, other=0.0)
    tl.store(o_ptr + dst, v, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels: int, out_channels: int, kernel_size: int, stride: int = 1, padding: int = 0, output_padding: int = 0, groups: int = 1, bias: bool = False):
        super(ModelNew, self).__init__()
        self.conv_transpose3d = nn.ConvTranspose3d(in_channels, out_channels, kernel_size=(kernel_size, kernel_size, kernel_size), stride=stride, padding=padding, output_padding=output_padding, groups=groups, bias=bias)
        self.stride = stride
        self.padding = padding
        self.output_padding = output_padding
        self.groups = groups
        self.has_bias = bias

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        simple = (self.stride == 1 and self.output_padding == 0 and self.groups == 1 and x.is_cuda)
        if simple:
            w = self.conv_transpose3d.weight  # (IC, OC, KD, KH, KW)
            IC, OC, KD, KH, KW = w.shape
            wf = torch.empty((OC, IC, KD, KH, KW), device=w.device, dtype=w.dtype)
            total = IC * OC * KD * KH * KW
            BLOCK = 256
            flip_kernel[(triton.cdiv(total, BLOCK),)](
                w.contiguous(), wf, IC, OC, KD, KH, KW, BLOCK=BLOCK)
            p = KD - 1 - self.padding
            b = self.conv_transpose3d.bias if self.has_bias else None
            return F.conv3d(x, wf, bias=b, stride=1, padding=(p, p, p))
        return self.conv_transpose3d(x)
```