## 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 attention problems (wrong results or slower than eager)
- Pitfall: Avoid per-row Triton LayerNorm with one program per row and huge num_warps for short rows; launch overhead plus extra .contiguous() copies make it slower than PyTorch's fused layer_norm.
- Pitfall: In large transformer models, replacing only LayerNorm (or fused LN+residual) with custom kernels while matmuls, attention, and window permutes stay eager yields slowdown; fuse the attention core or the dominant ops instead.
- Pitfall: Before investing in a full custom kernel for a matmul-dominated op under strict fp32, estimate whether it can beat cuBLAS SGEMM. If not, fuse only the memory-bound elementwise and reduction parts around the GEMMs.
- Pitfall: Avoid hand-written Triton flash attention with tl.dot(allow_tf32=False). Strict fp32 dots skip tensor cores and lose badly to cuBLAS. Keep the QK^T and PV matmuls in torch/cuBLAS and fuse only scale+mask+softmax.
- Pitfall: Avoid extra .contiguous()/reshape copies and per-call Python wrapper overhead around small Triton launches; on small tensors launch latency outweighs the gain over cuDNN/ATen fused kernels.
- Pitfall: Don'...'t beat eager.