## 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_078 ConvTranspose3d_Max_Max_Sum, 1.00x): 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, N, C, D, H, W, OD, OH, OW, total, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    idx = pid * BLOCK + tl.arange(0, BLOCK)
    idx64 = idx.to(tl.int64)
    mask = idx < total

    ow = idx64 % OW
    t = idx64 // OW
    oh = t % OH
    t = t // OH
    od = t % OD
    n = t // OD

    d0 = od * 6
    h0 = oh * 6
    w0 = ow * 6

    C64 = C.to(tl.int64)
    D64 = D.to(tl.int64)
    H64 = H.to(tl.int64)
    W64 = W.to(tl.int64)

    s = tl.zeros((BLOCK,), tl.float32)
    for c in range(0, C):
        base_nc = (n * C64 + c) * D64
        m = tl.full((BLOCK,), -1e30, tl.float32)
        for dz in tl.static_range(6):
            iz = d0 + dz
            for dy in tl.static_range(6):
                iy = h0 + dy
                for dx in tl.static_range(6):
                    ix = w0 + dx
                    off = ((base_nc + iz) * H64 + iy) * W64 + ix
                    v = tl.load(x_ptr + off, mask=mask, other=-1e30)
                    m = tl.maximum(m, v)
        s += m
    tl.store(out_ptr + idx64, s, mask=mask)


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

    def forward(self, x):
        c = self.conv_transpose(x).contiguous()
        N, C, D, H, W = c.shape
        P1D, P1H, P1W = D // 2, H // 2, W // 2
        OD, OH, OW = P1D // 3, P1H // 3, P1W // 3
        out = torch.empty((N, 1, OD, OH, OW), device=c.device, dtype=c.dtype)
        total = N * OD * OH * OW
        BLOCK = 128
        grid = (triton.cdiv(total, BLOCK),)
        fused_kernel[grid](c, out, N, C, D, H, W, OD, OH, OW, total, BLOCK=BLOCK)
        return out
```

### A verified kernel for a related conv_transpose problem (L2_074 ConvTranspose3d_LeakyReLU_Multiply_LeakyReLU_Max, 1.26x): 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, mult_ptr, N, C, D, H, W, OD, OH, OW,
                 slope, total, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    idx = pid * BLOCK + tl.arange(0, BLOCK)
    mask = idx < total

    ow = idx % OW
    t = idx // OW
    oh = t % OH
    t = t // OH
    od = t % OD
    nc = t // OD
    c = nc % C

    m = tl.load(mult_ptr + c, mask=mask, other=0.0)

    acc = tl.full((BLOCK,), -1e30, tl.float32)
    for dz in tl.static_range(2):
        for dy in tl.static_range(2):
            for dx in tl.static_range(2):
                iz = od * 2 + dz
                iy = oh * 2 + dy
                ix = ow * 2 + dx
                in_off = (((nc * D + iz) * H + iy) * W + ix)
                v = tl.load(x_ptr + in_off, mask=mask, other=0.0)
                v = tl.where(v >= 0, v, v * slope)
                v = v * m
                v = tl.where(v >= 0, v, v * slope)
                acc = tl.maximum(acc, v)
    tl.store(out_ptr + idx, acc, mask=mask)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, multiplier_shape):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding, output_padding=output_padding)
        self.multiplier = nn.Parameter(torch.randn(multiplier_shape))

    def forward(self, x):
        c = self.conv_transpose(x).contiguous()
        N, C, D, H, W = c.shape
        OD, OH, OW = D // 2, H // 2, W // 2
        out = torch.empty((N, C, OD, OH, OW), device=c.device, dtype=c.dtype)
        total = N * C * OD * OH * OW
        mult = self.multiplier.contiguous().view(-1)
        BLOCK = 256
        grid = (triton.cdiv(total, BLOCK),)
        fused_kernel[grid](c, out, mult, N, C, D, H, W, OD, OH, OW, 0.2, total, BLOCK=BLOCK)
        return out
```