## 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 rnn problems (wrong results or slower than eager)
- Pitfall: Don'...'t beat eager. Replace or fuse real compute instead.
- Pitfall: Don't wrap the reference library op (e.g., cuDNN LSTM) and then add a trivial copy kernel. The extra allocation and launch only add overhead. Write a real fused kernel for the heavy compute, or return results directly.
- Pitfall: Avoid one-program-per-output-element dot-product kernels for matmuls; they ignore tiling and data reuse and lose to cuBLAS. Use tiled tl.dot with blocked M/N/K.
- Pitfall: Don't wrap the library op (e.g., cuDNN LSTM) and replace only a tiny trailing linear with a custom kernel; the extra launch cost makes it slower. Fuse substantial work or leave cuBLAS in place.
- Pitfall: Don't wrap the cuDNN-backed PyTorch op and add a trivial Triton copy kernel. The extra launch and memory traffic only make it slower. Either fuse real work or write a genuinely faster kernel.
- Pitfall: Before writing a kernel, check whether the eager baseline is already one fused vendor-library call (cuDNN RNN, cuBLAS). Decomposing it into many smaller launches is almost always slower, even when correct.

### A verified kernel for a related rnn problem (L3_034 VanillaRNNHidden, 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 bias_tanh_kernel(pre_ptr, b_ptr, out_ptr, n_rows, hidden_size, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = tl.arange(0, BLOCK)
    mask = offs < hidden_size
    v = tl.load(pre_ptr + pid * hidden_size + offs, mask=mask, other=0.0)
    b = tl.load(b_ptr + offs, mask=mask, other=0.0)
    r = v + b
    e = tl.exp(2.0 * r)
    r = (e - 1.0) / (e + 1.0)
    tl.store(out_ptr + pid * hidden_size + offs, r, mask=mask)


def bias_tanh(pre, b, out):
    n_rows, hidden_size = pre.shape
    BLOCK = triton.next_power_of_2(hidden_size)
    bias_tanh_kernel[(n_rows,)](pre, b, out, n_rows, hidden_size, BLOCK=BLOCK)
    return out


class ModelNew(nn.Module):
    def __init__(self, input_size: int, hidden_size: int, output_size: int):
        super(ModelNew, self).__init__()
        self.input_size = input_size
        self.hidden_size = hidden_size
        self.output_size = output_size

        self.i2h = nn.Linear(input_size + hidden_size, hidden_size)
        self.h2o = nn.Linear(hidden_size, output_size)
        self.tanh = nn.Tanh()

    def forward(self, x: torch.Tensor, h0: torch.Tensor) -> torch.Tensor:
        seq_len, batch_size, _ = x.size()
        hidden = h0.to(x.device)

        W = self.i2h.weight
        W_x = W[:, :self.input_size]
        W_h = W[:, self.input_size:]
        b = self.i2h.bias

        x2d = x.reshape(seq_len * batch_size, self.input_size)
        pre_all = (x2d @ W_x.t()).reshape(seq_len, batch_size, self.hidden_size)

        Wh_t = W_h.t().contiguous()

        hiddens = torch.empty(seq_len, batch_size, self.hidden_size,
                              device=x.device, dtype=x.dtype)

        for t in range(seq_len):
            pre = pre_all[t].addmm(hidden, Wh_t)
            hidden = bias_tanh(pre, b, hiddens[t])

        h2d = hiddens.reshape(seq_len * batch_size, self.hidden_size)
        out = self.h2o(h2d).reshape(seq_len, batch_size, self.output_size)
        return out
```