ACE playbook: reusable strategies from prior completed tasks. Apply only when relevant. When reporting your attempt, cite any useful or misleading bullet IDs if possible.

## strategies
[str-00001] helpful=150 harmful=0 :: load_inline compile pitfall: cpp_sources is compiled by plain c++ (main.cpp) with only <torch/extension.h>, so CUDA runtime types such as cudaStream_t and helpers such as at::cuda::getCurrentCUDAStream() are undeclared there. Errors look like 'cudaStream_t has not been declared' or 'getCurrentCUDAStream is not a member of at::cuda'. Safest layout: put the torch::Tensor wrapper function, the stream retrieval and the kernel launch all in cuda_sources. The .cu file should include <torch/extension.h>, <cuda_runtime.h> and <ATen/cuda/CUDAContext.h>. Leave cpp_sources with only a forward declaration that uses torch types, e.g. `torch::Tensor my_op(torch::Tensor x);`, and list that name in functions=[...]. If CUDA types must appear in cpp_sources, explicitly include <cuda_runtime.h> and <ATen/cuda/CUDAContext.h> there.
[str-00002] helpful=19 harmful=1 :: Always compile and run a local correctness check of the extension or kernel before submitting. A first-turn SUBMIT without any check risks a trivial compile_error, which scores 0. Also time the candidate locally against the reference op (e.g., torch.matmul) with proper warmup and cuda synchronize before submitting. Without a timing check, a slow result (e.g., 0.56x) is only discovered after submission, with no chance to switch strategy. Keep the local timing bounded: use a few warmup runs and a fixed small iteration count. This also reveals the runtime scale, so a candidate that might hit the evaluator's timing limit (status 'timeout', 'timing exceeded 240s', scored 0) is caught before SUBMIT. Passing CHECK alone does not mean a candidate is ready to submit. A further failure mode: CHECK passed, but the official run reported status 'incorrect' with 'candidate produced no output', even though the kernel did run (own_kernel_time_frac > 0). The likely causes are a crash, OOM or timeout at full scale. To catch this, run the forward locally at the FULL problem size (actual get_inputs shapes), with memory and time printed, before SUBMIT.
[str-00003] helpful=14 harmful=0 :: For conv+BN+activation networks (e.g., MBConv/inverted-residual blocks): folding BN into conv weights and bias is a reasonable first step, but it is only valid when BN is in eval mode. A freshly constructed nn.Module defaults to training mode, where BN uses batch statistics and cannot be folded; check which mode the reference runs in. Keeping cuDNN convs and adding a standalone elementwise activation kernel (ReLU6, etc.) after each one gives no speedup when it does not reduce the number of memory passes compared with the reference. In that case it just replaces cuDNN/PyTorch's own activation with an extra full pass over large hidden tensors. A standalone activation kernel CAN win when the reference has multiple separate elementwise ops after the conv, which folding plus fusion collapses into a single pass. Examples are scale/div plus activation, or an activation composed of several eager ops such as Mish = x*tanh(softplus(x)). Otherwise, the activation must be fused into a conv epilogue or a custom kernel. Better still, fuse chains such as 1x1-expand→depthwise→1x1-project into fewer kernels, so the large expanded intermediate is not round-tripped through global memory.
[str-00004] helpful=29 harmful=11 :: Large DENSE fp32 GEMM (KernelBench matmul family, including plain A @ B, A @ B.T, A.T @ B and A.T @ B.T variants and Linear stacks in MLP networks): a hand-written Triton kernel using tl.dot(..., allow_tf32=False) runs pure-FP32 FMA without tensor cores. It loses heavily to cuBLAS, even with grouped tiling and autotune.

Measured failures (strict FP32):
- About 0.56x on an irregular ~8k x 3k x 6k GEMM.
- 0.71x (11.06 vs 7.81ms, custom kernel 100% of runtime) on A @ B.T at 2048x8192x4096.
- 0.60x (17.88 vs 10.68ms, custom kernel 100% of runtime) on a 3-layer MLP (M=128, K=N=16384) with fused bias+ReLU, BM=128, BN=64, BK=32, 4 warps.

Measured WINS (tensor cores at near-FP32 accuracy), all with own kernel at 100% of runtime and passing correctness:
- A.T @ B with A stored [K,M], M=2048, K=8192, N=4096: 1.59x (8.08 -> 5.09ms).
- Plain square A @ B, N=4096: 1.43x (7.28 -> 5.11ms).
- A.T @ B.T (A stored [K,M], B stored [N,K], M=2048, K=8192, N=4096) with the OPERAND-SWAP fix below: 1.46x (7.53 -> 5.15ms).
Recipe for all: tl.dot(a, b, acc, input_precision='tf32x3'), @triton.autotune (key M,N,K) over BM/BN in {128,256} and BK in {32,64}, num_warps=8, num_stages 3-4, and grouped pid ordering (GM=8). All followed CHECK-then-SUBMIT with no local timing, so the tf32 vs tf32x3 and unmasked-loop headroom below remains untested.

LAYOUT-SENSITIVE: the key rule is that BOTH tl.dot operands must be loaded in native orientation (A tile [BM,BK], B tile [BK,BN] with the contiguous dim matching storage where possible). On A.T @ B.T, reading both operands via swapped strides (B tile [BK,BN] with K contiguous) scored only 0.90x (8.36 vs 7.55ms). MEASURED FIX (1.46x, same autotune set): compute C^T = B_stored @ A_stored. Load the B tile as [BN,BK] from B[N,K] (K contiguous) and the A tile as [BK,BM] from A[K,M] (M contiguous), acc = tl.dot(b_tile, a_tile, acc, input_precision='tf32x3') of shape [BN,BM], and store through swapped C strides (c_ptrs = C + offs_m[None,:]*sc_m + offs_n[:,None]*sc_n), with no transpose copy. Then TT performs like NN/TN. Do NOT assume the recipe is layout-agnostic: for any transposed layout, first look for an operand swap / output-transpose identity that makes both loads native, and time it. Also try plain tf32 if the tolerance allows, and an unmasked K loop.

Scope limit: the tf32x3 default is proven only for dense GEMMs with contiguous operand loads. It did NOT transfer to a gather-bound conv implicit GEMM, which scored 0.29x vs cuDNN with tf32x3 (see the dense conv bullet, failure mode C). For convs and other gather-heavy kernels, prefer plain TF32 when tolerance allows, and time before submitting.

Decide the precision path BEFORE writing the kernel, and read the task's actual tolerance first (see the tolerance pitfall bullet). Tolerance is NOT always ~1e-2: the MLP task checked at atol=rtol=1e-4, and a TF32 tl.dot failed CHECK there.
- Unknown tolerance (dense GEMM): default to input_precision='tf32x3'. It keeps near-FP32 accuracy and beat cuBLAS strict FP32 in the NN, TN and (operand-swapped) TT layouts.
- If tolerance is ~1e-2: plain allow_tf32=True (or input_precision='tf32') is likely faster still (the 5.5x Linear+GELU recipe); time tf32 vs tf32x3 if turns allow. The reference cuBLAS path may itself use TF32 if torch.backends.cuda.matmul.allow_tf32 is set. The same applies to every tl.dot inside fused kernels such as flash attention (QK^T and PV). A strict-FP32 causal attention kernel scored 0.11x; see the causal attention bullet.
- If tolerance is 1e-4: follow the 1e-4 options in the tolerance pitfall bullet (tf32x3, or a cuBLAS fallback). Do not submit a strict-FP32 Triton GEMM.

For transposed operands, read through strides so no transpose copy is made (never use a .T.contiguous() copy per forward). For A.T with A stored [K,M], a_ptrs = A + offs_k[None,:]*stride_k + offs_m[:,None]*stride_m works (TN win). For B.T with B stored [N,K] at compute-bound M (nn.Linear NT), the swapped-stride [BK,BN] read loses with tf32x3; see str-00075/str-00076 (cached W.t() NN conversion), and for A.T@B.T use the operand swap above. If strict FP32 is needed for a dense problem, a custom GEMM is unlikely to beat cuBLAS, so aim for parity with it.

Scope exception: when the operands or output have structural zeros (triangular, banded, block-sparse), strict-FP32 Triton can still win big by skipping that work (4.1x on tril@tril; see the structured-sparsity GEMM bullet).

Headroom: when K % BK == 0, run the main K loop unmasked; for non-multiple-of-BK K dimensions, peel the K-tail iteration. Drop the %M/%N wrap when dims divide the tiles. Try a fixed 128x128x64 config vs the 256-wide autotune set (str-00032 found autotune ~20% slower), and larger BK/stages configs. Use the turns left after CHECK to time locally against torch.matmul (and, if tolerance permits, tf32 vs tf32x3). A sub-1x result should trigger a precision, operand-swap or strategy switch rather than a SUBMIT.
[str-00005] helpful=5 harmful=0 :: Conv followed by a scalar multiply or divide and then an elementwise activation (e.g., Conv2d → x/d → LeakyReLU): fold the scalar into the conv parameters (w/d, b/d; this is exact for a linear conv). Then call cuDNN via F.conv2d and run one fused activation kernel. This removes a full memory pass compared with the eager conv → div → act chain (3 passes become 2). It gave about 1.24x over eager and about 1.55x over torch.compile, even though the custom kernel was only about 25% of runtime. CONFIRMED a second time on Conv2d 8->64, 3x3, N=128, 128px -> x/2 -> LeakyReLU(0.01): 1.2271x (3.68 -> 3.00ms, lo95 1.2247), own_kernel_time_frac 0.26. That run cached w/d and b/d in self.<x> attributes keyed on conv.weight.data_ptr() (str-00014), then F.conv2d plus ONE in-place grid-stride LeakyReLU kernel (256 threads, 4096-block cap, long n per str-00058, str-00001 layout). It was a turn-0 SUBMIT with no CHECK or timing, correct only because the recipe was copied. Treat ~1.23x as the SOLVED BASELINE for this layout: resubmitting it cannot raise the score. After the first passing CHECK, spend the spare turns timing levers at full size against the reference and keep the fastest: conv-side (cuDNN ~74% of runtime) cuDNN TF32 within tolerance (str-00029), cudnn.benchmark, channels_last (non-dilated only, str-00051); kernel-side (~26% of runtime) float4 loads/stores with a scalar tail and a grid cap near SMs x 8. Refinements: (1) precompute or cache the scaled weight and bias rather than recomputing w/d on every forward. For small conv weights this is cheap but nonzero. When the weight is as large as or larger than the output, per-forward rescaling erases the gain: Linear 8192x8192 gave 0.96x (see the GEMM+epilogue bullet). In that case, cache the scaled weight, or skip the fold and apply the scale in the activation pass instead. (2) For more speed, fuse the activation into the conv epilogue so the conv output is never re-read from global memory. Also apply the activation in place on the conv output to avoid a new allocation (done in the confirmed run).
[str-00006] helpful=24 harmful=0 :: Conv followed by a long chain of eager elementwise ops (e.g., relu, then x * clamp((x+3)/6, 0, 1) as hand-written HardSwish; or several add/div/clamp/mul steps): keep the conv in cuDNN (nn.Conv2d / F.conv2d) and collapse the whole chain into ONE fused elementwise CUDA kernel. Each eager op is its own kernel doing a full read+write pass over the large conv output, so a chain of about 5-6 ops becomes 1 pass. This gave about 2.17x on a 128x64x126x126 output, even though the custom kernel was only about 26% of runtime. The win grows with the number of separate elementwise ops in the reference; contrast this with a single activation, where a standalone kernel gives no gain (see the conv+BN+activation bullet). Even a chain of just TWO built-in activations pays off. Example: Conv2d → F.hardswish → torch.relu (128x64x126x126 output) was replaced by one in-place grid-stride kernel computing hs = x*clamp((x+3)/6,0,1), then max(hs,0). It gave 1.25x (3.64 → 2.92ms), with the kernel at about 39% of runtime. CONFIRMED a second time on the same task: 1.2354x (4.04 → 3.27ms, lo95 1.19), kernel ~31% of runtime, using scalar loads, a 1024-block grid cap and long n (str-00058). That run was a turn-0 SUBMIT with no CHECK or timing, correct only by luck (str-00002), and left every headroom lever untried. Treat ~1.24x as the BASELINE to beat on this chain, not the target: use spare turns to CHECK and time float4 loads/stores, a grid cap near SMs×8, and conv-side cuDNN TF32 / channels_last within tolerance. Algebraic simplification: hardswish followed by relu equals x>0 ? x*min(x+3,6)/6 : 0 (hardswish is <= 0 for all x <= 0, so relu zeroes it). A simple 1-thread-per-element kernel with __restrict__ pointers, launched on at::cuda::getCurrentCUDAStream(), is enough. Per-channel variant: Conv3d → x*scale[c] → tanh → *bias[c] → sigmoid (4 eager passes over a ~1GB 128x16x14x62x62 output) as one grid-stride kernel gave 1.287x vs eager (7.83 → 6.09ms, kernel ~14% of runtime) but only 0.88x vs torch.compile, which finds the same fusion. When the conv is ~85%+ of runtime, beating compile requires changing the conv itself (cuDNN TF32 / cudnn.benchmark within tolerance, timed channels_last_3d, or an epilogue-fused conv; see str-00053), not just kernel cleanups. Extension: if the chain ends in a pooling step (e.g., MaxPool3d(2)), fold the pool into the same kernel. Use one thread per POOLED output. Each thread reads its 2x2x2 window, applies the activation chain to each element in registers, and writes only the max. Apply the chain before the max, because a per-channel multiplier can be negative and LeakyReLU*mult does not commute with max. The conv output is read once, and the output is only 1/8 the size. Example, CONFIRMED TWICE: ConvTranspose3d(stride 2) → LeakyReLU → *per-channel mult → LeakyReLU → MaxPool3d(2) gave 1.25x (5.95 vs 4.75ms) and again 1.2518x (5.87 → 4.69ms, lo95 1.251, own_kernel_time_frac 0.05) with a grid-stride 4096-block-cap kernel, long indices and the str-00001 layout (compiled first try, CHECK then SUBMIT on turn 1, no timing, no conv lever tried). The cuDNN ConvTranspose3d is ~95% of runtime, so ~1.25x is the SOLVED baseline for this task: resubmitting it cannot improve the score, and kernel micro-tweaks (float2/float4, 2D grid, in-place) move at most ~5% of runtime. After the first passing CHECK, spend the remaining turns timing conv-side levers at full size against the reference before SUBMIT: cuDNN TF32 within tolerance, cudnn.benchmark, channels_last_3d (str-00051), and the stride-2 phase decomposition (str-00030). Scope limit for the one-thread-per-pooled-output pattern: it assumes the pooled output is large enough to fill the GPU (small windows such as 2x2x2, with no extra channel reduction). When the window is large and/or a channel sum is folded in, too few threads launch. Example: a 6x6x6 window × 64 channels gave only ~16k threads for ~1GB of reads, and scored 0.97x (see the double-MaxPool → sum-over-C bullet). In that case, spread the window and channels across threads and reduce with shuffles. Further headroom: write in place on the conv output (no empty_like allocation), use vectorized float4 loads/stores, multiply by a precomputed 1/6 instead of dividing, use a 2D grid (blockIdx.y = n*C + c) so per-channel params load once per block and there is no per-element 64-bit div/mod, cap a grid-stride grid near SMs×8 rather than 65535 or a fixed 1024, use __expf only if tolerance allows, try channels_last or TF32 for the cuDNN conv (within tolerance), or fuse the chain (and pool) into the conv epilogue.
[str-00007] helpful=3 harmful=0 :: Conv -> multi-op activation -> BatchNorm2d in its default training mode (e.g., Mish written as torch.multiply(torch.tanh(F.softplus(x)), x)): keep cuDNN conv and nn.BatchNorm2d, and replace the softplus->tanh->multiply chain with ONE fused elementwise kernel. This collapses about 3-4 eager memory passes and temporary allocations into 1.

Measured, CONFIRMED TWICE on a 64x128x126x126 output:
- Triton kernel: about 1.20x, custom kernel about 7% of runtime.
- load_inline CUDA kernel (float4 grid-stride, 4096-block cap, scalar tail, x>20 guard with log1pf(expf)/tanhf, empty_like output, str-00001 layout; compiled first try): 1.2067x (10.96 -> 9.08ms), own_kernel_time_frac 0.074, but only 0.9998x vs torch.compile. It went CHECK, CHECK, SUBMIT with no local timing and no headroom lever tried.
Treat ~1.2x as the BASELINE, not the target: it equals torch.compile. Do NOT submit this layout on the turn after CHECK without timing the levers below.

Numerics: guard softplus for large x (x > 20 ? x : log1p(exp(x))). PyTorch-exact formulas (str-00029) passed. To cut exp count, tanh(softplus(x)) = (e^{2x} + 2e^x) / (e^{2x} + 2e^x + 2), with the same large-x guard (verify tolerance first). Training-mode BN can NOT be folded into the conv, because it needs batch statistics.

Next headroom, in order (all untested; time each locally):
(1) Fuse the per-channel BN sum/sumsq into the activation kernel: 2D grid with blockIdx.y = n*C+c, apply Mish in place on the conv output, accumulate fp32 per thread, shared-memory tree, then fp32 atomicAdd into zeroed per-channel buffers. Then a finalize kernel (one thread per channel: mean, biased var, rsqrt, fold gamma/beta into scale/shift, update running_mean/running_var in-kernel with unbiased var*n/(n-1) and momentum, buffers passed as tensors per str-00014), then ONE in-place normalize pass. That is 3 passes instead of eager's activation pass plus BN's 2-3 passes (same layout as the str-00043 win). Provide an eval branch using running stats so a custom kernel runs in both modes (str-00031).
(2) Apply the activation in place on the conv output to avoid empty_like.
(3) Conv-side levers within tolerance: cuDNN TF32, cudnn.benchmark, channels_last (non-dilated, timed per str-00051).
[str-00008] helpful=67 harmful=0 :: Submission format pitfall: EVERY reply, whether a CHECK turn or a SUBMIT, must contain the full solution inside a ```python fenced code block. Otherwise it is rejected as no_code and the turn is wasted, even if the code was already validated by CHECK on an earlier turn.

This includes the fix-up turn right after a failed CHECK, where one MobileNetV1 attempt lost a turn. It also includes the SUBMIT turn right after successful CHECKs. A ConvTranspose3d attempt lost its turn-2 SUBMIT this way while this bullet was shown, and so did a Conv2d+BN+scale attempt, a Conv2d+subtract+HardSwish+MaxPool+Mish attempt, and a Linear+training BN+GELU+ReLU attempt (that lost turn could have been the missing local timing that would have caught its 0.95x tail).

Re-paste the complete final source in each reply. This has recurred across multiple attempts even with this bullet present. Treat it as a mandatory final self-check before sending any reply: confirm that the reply contains a ```python block holding the whole module.
[str-00009] helpful=9 harmful=4 :: Stride-1 ConvTranspose2d/3d (KernelBench conv_transpose family): rewriting it as a regular forward conv is a CANDIDATE, not a guaranteed win. Whether it wins depends on the shape, because cuDNN runs the transposed path as conv backward-data, and that algorithm can be faster or slower than its forward-conv algorithm.

STOP CHECK FIRST: compare the task's EXACT kernel shape and padding against the list below BEFORE applying the recipe. Matching only the input dims is not enough. The 2.30x win and the 0.74x loss share the same 8x32x512x1024 input and 32->32 channels. They differ only in the kernel (3x7 with padding (1,3) wins; 3x3 with padding 0 loses). A square 3x3, padding 0 kernel on 8x32x512x1024, 32->32 is a SOLVED LOSS (confirmed twice): do NOT submit the flip+F.conv2d rewrite for it.

Measured:
- WIN (2D), CONFIRMED TWICE: 8x32x512x1024, 32->32 channels, asymmetric 3x7 kernel, padding (1,3). The latest run gave 2.2987x (35.57 -> 15.47ms, lo95 2.293), with own_kernel_time_frac 0.0001, and reproduced the earlier ~2.3x.
- WIN (2D): 2.22x (44.5 -> 20.0ms) at 8x64x512x512, 64->64 channels, asymmetric 3x7 kernel, padding 0. The flip kernel was 0.02% of runtime (own_kernel_time_frac 0.0002), so the whole gain was cuDNN's forward algorithm beating backward-data.
- WIN (2D): a 5x5 kernel, padding 1, 512x64x32x32 -> 128 channels, as part of a BN/tanh/pool/GN chain. The rewrite supplied most of a 3.81x total (see the stride-1 ConvTranspose -> training-BN -> activation -> MaxPool bullet).
- LOSS (2D), CONFIRMED TWICE: 8x32x512x1024 with a square 3x3 kernel, padding 0, 32->32 channels. Run 1: 0.75x (4.45 vs 3.32ms). Run 2: 0.7377x (4.45 vs 3.28ms, lo95 0.737). The flip kernel was 0.07-0.08% of runtime both times, so the whole slowdown came from cuDNN's forward algorithm being slower than its backward-data path for this shape. Both went CHECK then immediate SUBMIT on turn 1 with no timing; run 2 did so while this bullet listed the shape as a loss, apparently pattern-matching the 3x7 win on the same input dims. A trivial-kernel wrapper around the direct transposed conv only reaches parity, which also fails. For this shape, after the correctness CHECK, time these at full size against nn.ConvTranspose2d and submit only a variant above 1x: (1) direct F.conv_transpose2d with cudnn.benchmark and cuDNN TF32 (if tolerance allows), plus a custom kernel doing real work; (2) the rewrite with TF32/benchmark enabled; (3) a TF32 implicit-GEMM or register-tiled direct kernel spanning all 32 OC per tile.
- PARITY (3D), CONFIRMED TWICE: ConvTranspose3d, square 3x3x3 kernel, padding 0, 48->48 channels, 8x48x64^3: run 1 1.016x (42.75 -> 42.08ms); run 2 1.0138x (42.71 -> 42.13ms, lo95 1.012, 1.009x vs compile). The flip kernel was 0.01% of runtime both times. Both went CHECK then immediate SUBMIT on turn 1 with no timing, the second time while this bullet explicitly warned against it. The rewrite is a SOLVED ~1.01x baseline on this shape: resubmitting it cannot improve the score. On this exact shape, after the correctness CHECK, spend every remaining turn timing conv levers at full size against nn.ConvTranspose3d (cuDNN TF32 if tolerance allows, cudnn.benchmark, channels_last_3d, and a direct F.conv_transpose3d with those levers), keeping the fastest; beyond cuDNN, a TF32 implicit-GEMM kernel spanning all OC per tile.
- LOSS (3D, ASYMMETRIC): ConvTranspose3d 32->16, k=(3,5,7), padding 0, input 16x32x16x32x64: 0.849x (16.86 vs 14.31ms). The flip kernel was 0.02% of runtime, so the whole loss was cuDNN's 3D forward algorithm being slower than backward-data. It went CHECK then immediate SUBMIT on turn 1 with no timing, nudged by the earlier 'asymmetric wins' pattern, and no conv lever was tried.

Pattern so far:
- 2D: asymmetric (3x7) and larger (5x5) kernels have won; a small square 3x3 lost (twice).
- 3D: the rewrite has NOT won yet. Square 3x3x3 gave parity (twice) and asymmetric 3x5x7 gave a loss. The 2D 'asymmetric wins' signal does NOT carry over to 3D.
In 3D, and for small square kernels in 2D, treat the rewrite as a hypothesis with a baseline of ~1.0x or below. Spend the remaining turns on the conv levers below. Even in 2D the pattern is a signal, not a rule.

MANDATORY for any shape NOT listed above as a measured win: time the rewrite locally against nn.ConvTranspose2d/3d at FULL size, with bounded warmup and iterations and cuda synchronize, before SUBMIT. If it is below 1x (or near parity with turns left), do not submit it as is. Never SUBMIT on turn 1 right after a correctness-only CHECK for such shapes. The 0.75x, 0.7377x, 1.016x, 1.014x and 0.849x attempts all went CHECK-then-SUBMIT with no timing while a turn remained.
Narrow exception: when the task's EXACT shape (dims, channels, kernel, padding) matches a measured WIN above, CHECK-then-SUBMIT has reproduced the result (2.30x confirmed on the identical 3x7/(1,3) shape). Even then, spare turns are better spent timing the untried headroom levers (TF32, cudnn.benchmark, channels_last) on top of the rewrite.

The weight-prep kernel is negligible, so 'flip kernel + cuDNN F.conv2d/F.conv3d' is effectively a cuDNN wrapper. It wins only if the algorithm swap itself wins. When it does not (or for extra headroom when it does), the real levers are on the conv itself. Time each locally and keep whichever variant (rewrite or direct transposed) is fastest:
- cuDNN TF32 (torch.backends.cudnn.allow_tf32), within tolerance;
- channels_last / channels_last_3d (non-dilated only, see str-00051);
- cudnn.benchmark-style algorithm selection;
- a TF32 implicit-GEMM kernel with all OC per tile (see the dense conv bullet).
If the rewrite loses, also time a direct F.conv_transpose2d/3d call with these levers applied. Make the custom kernel do real work, e.g. an epilogue or a layout conversion, not an unused weight prep. A trivial-kernel wrapper alone lands at parity (dense conv failure mode B).

Steps (only for shapes where the rewrite is a measured win or has been timed above 1x):
(1) Flip the weight spatially and swap the in/out channel dims: [IC,OC,kH,kW] -> [OC,IC,kH,kW], where out[o,i,kH-1-h,kW-1-w] = in[i,o,h,w]. In 3D: [IC,OC,kD,kH,kW] -> [OC,IC,kD,kH,kW], flipping all three spatial dims. A tiny custom flip/permute CUDA kernel is enough (one thread per weight element, launched on the current stream from cuda_sources per str-00001). In torch this is w.flip(2,3).permute(1,0,2,3).contiguous() (2D) or w.flip(2,3,4).permute(1,0,2,3,4).contiguous() (3D).
(2) Call F.conv2d/F.conv3d(x, w_eq, bias, stride=1, padding=k-1-p per dim) with the same bias.

Assumptions and fallback: this assumes groups=1, dilation=1 and output_padding=0, and the padding must be <= k-1. Fall back to nn.ConvTranspose when stride != 1 (or groups/output_padding non-default), or else derive the zero-dilated-input form.

[str-00010] helpful=10 harmful=0 :: BatchNorm (training or eval) followed directly by global average pooling (AdaptiveAvgPool to 1x1(x1)): skip the normalize pass entirely. The output depends only on per-(N,C) spatial means and per-channel stats. Use ONE reduction kernel, with one block per (n,c), to produce the per-(N,C) sum and sum-of-squares. Then compute: mean_c = sum over N of psum / (N*S); var_c = E[x^2] - mean_c^2 (biased). The output is (psum/S - mean_c) * gamma * rsqrt(var_c + eps) + beta, shaped [N,C,1,1,1]. In training mode, also update running_mean and running_var (the latter with unbiased var*total/(total-1), using momentum). num_batches_tracked can be skipped (str-00014). In eval mode, use the running stats. A preceding scalar multiply (x*s) should NOT be an eager full-tensor pass. Fold it into the conv weight and bias, or apply it to the reduced stats: sum scales by s, sumsq scales by s^2 (a small per-channel finalize kernel accumulating in double works).

Measured, CONFIRMED TWICE at ~1.009x on L2_077 (ConvTranspose3d 64->128, k=5, stride 1, padding 0, 16x64x16x32x32 -> scale -> BN3d -> global avgpool):
- Run 1: ~1.0085x, reduction 0.27% of runtime.
- Run 2: 1.0088x (64.15 -> 63.59ms, lo95 1.0084), own_kernel_time_frac 0.0038. It used the str-00001 layout. It went CHECK on turn 0 then SUBMIT on turn 1 with no timing and no conv lever tried, even though turns remained.
This epilogue is a SOLVED ~1.01x baseline: resubmitting it cannot raise the score. Keep it as the custom kernel, but do NOT go CHECK -> immediate SUBMIT on this task. The cuDNN ConvTranspose3d is ~99.6% of runtime, so real gains must come from the conv itself.

After the first passing CHECK, time these at full size against the reference (bounded warmup/iterations, cuda synchronize) and keep the fastest:
(a) The stride-1 forward-conv rewrite (flip/swap weight + F.conv3d, padding k-1-p; see str-00009). In 3D it has only reached parity (square 3x3x3, 48->48) or a loss (asymmetric 3x5x7). The square 5x5x5 3D case is UNTESTED; 5x5 won in 2D.
(b) torch.backends.cudnn.benchmark=True.
(c) cuDNN TF32 (torch.backends.cudnn.allow_tf32), within tolerance (str-00029). Global avgpool after BN may shrink absolute error, but print the max error first.
(d) channels_last_3d (non-dilated, str-00051).
(e) Longer term: fuse the stats reduction into the conv.
[str-00011] helpful=19 harmful=0 :: GEMM + elementwise epilogue (e.g., Linear → relu → /d, Linear → +bias → relu, Linear → +add_value → swish/tanh/gelu/hardtanh chain, Linear → *scale → hardtanh → erf-GELU, Linear → sigmoid → *scale → +residual (the GEMM output itself), or Linear → Mish → Mish, KernelBench matmul family) at large sizes, e.g., 1024x8192 @ 8192x8192 or 2048x8192 @ 8192x8192 fp32: the GEMM is about 99% of runtime, so a standalone elementwise kernel saves almost nothing.

Evidence:
- Attempt 1 recomputed w/d and b/d every forward, then called F.linear and an in-place ReLU kernel. It scored only 0.96x, with the custom kernel at 0.64% of runtime. The w/d step is a full read+write pass over the 256MB weight, which is larger than the 32MB output, so it costs more than the saved output pass.
- Attempt 2 used cuBLAS F.linear plus an in-place grid-stride bias+ReLU CUDA kernel, with no weight rescaling. It passed CHECK but was submitted on turn 1 with no local timing. It scored 0 because of an evaluator timeout ('timing exceeded 240s'). The custom kernel was only 0.86% of runtime.
- Attempt 3 (sigmoid*s + residual) used cuBLAS F.linear + one in-place grid-stride __expf kernel computing y = sig(y)*s + y. It was correct but scored only 1.0126x (8.07 vs 7.97ms), with the custom kernel at 0.77% of runtime. It collapsed 3 eager passes (sigmoid, mul, add) over the 32MB output into 1, and that is all this layout can bank.
- Attempt 4 (Mish → Mish) used cuBLAS F.linear + one in-place double-Mish kernel (log1pf(__expf), tanhf, x>20 guard). It was submitted on turn 0 with no CHECK or timing. It was correct but scored 1.0012x (7.97 vs 7.98ms), with the custom kernel at 1.0% of runtime. Collapsing the eager multi-op Mish passes did not help here, because the GEMM dominates.
- Attempt 5 (scale → hardtanh → erf-GELU, batch 2048) used cuBLAS F.linear + one in-place scale/clamp/erff kernel. It was correct but scored 1.014x (16.11 vs 15.89ms), with the custom kernel at 0.84% of runtime. It again ignored rule (0). Two turns were also lost: one to a static-check failure on a w._version cache key (see the static-check pitfall bullet), and one to a SUBMIT with no code block.
- Attempt 6 (same scale → hardtanh → erf-GELU task, batch 2048) did change the GEMM, but used a swapped-stride tf32x3 Triton kernel with the fused epilogue. It was correct but scored 0.944x (16.84 vs 15.90ms) after CHECK then immediate SUBMIT with no timing. tf32x3 on the nn.Linear NT layout loses; use the plain-TF32 str-00032 recipe instead (see the nn.Linear NT layout bullet).

The 'cuBLAS + standalone epilogue kernel' layout reliably gives ~1.0x, whatever the epilogue chain.

Rules:
(0) Decide on turn 0 to change the GEMM itself. Do not spend the attempt on the 'cuBLAS + standalone epilogue kernel' layout, and do not SUBMIT it on turn 0. Use the remaining turns to CHECK and time the GEMM rewrite.
(1) Never rescale a large weight per forward. Either cache the scaled weight and bias, or leave the weights alone and do one fused epilogue pass. Cache keys must not use private attributes like ._version; use data_ptr() or compute once. ReLU commutes with positive scaling, so out = relu(y) * (1/d) for d > 0. By contrast, folding an extra per-output-feature additive vector into the bias is cheap and sound, because it touches only an out_features-sized vector: bias = linear.bias + add_value, then F.linear(x, W, bias). Follow this with ONE fused activation pass for the whole chain. Writing that pass in place on the fresh F.linear output is safe: nothing else holds the intermediate, even when the chain reuses it as a residual. It also avoids an allocation. Expect only a small gain from this alone.
(2) For a real win, change the GEMM itself and fold bias+act+scale(+residual) into it.
  - Option a: cuBLASLt with CUBLASLT_EPILOGUE_RELU_BIAS and alpha = 1/d.
  - Option b (preferred): a Triton tl.dot kernel with allow_tf32=True (plain TF32, NOT tf32x3) and a fused bias/activation/scale epilogue (str-00032 recipe, 5.5x). TF32 fits within the ~1e-2 tolerance and gives tensor-core throughput. tf32x3 on this NT layout measured 0.944x and 0.85x.
  - Minimum fallback if staying on cuBLAS: run the matmul in TF32, e.g., by enabling torch.backends.cuda.matmul.allow_tf32 around F.linear. Check correctness and timing locally.
(3) Time locally against the reference before submitting, using bounded warmup and iteration counts. Near-parity and sub-1x results are typical here, and slow or hanging timing runs must be caught while turns remain.
[str-00012] helpful=13 harmful=0 :: Triton API portability pitfall: math function names differ across Triton versions. `tl.math.tanh` does not exist in the evaluator's Triton, and `tl.math.erf` may also be missing. The failure is a CompilationError/AttributeError that appears only at JIT compile time, on the first kernel launch, so it scores 0 as runtime_error. Portable fixes: (1) Build tanh from tl.exp, e.g. tanh(x) = 2/(1+tl.exp(-2x)) - 1. This is overflow-safe, because exp going to inf yields -1, not NaN. (2) For GELU, use top-level `tl.erf`. It is confirmed to compile and run in the evaluator's Triton, including inside a fused tl.dot GEMM epilogue. Otherwise use the tanh-approximation GELU built from tl.exp. (3) Reach libdevice via `from triton.language.extra.cuda import libdevice` (libdevice.tanh/erf) inside a try/except import fallback to the exp-based form. Always actually launch the Triton kernel locally (CHECK) before SUBMIT, since importing the module does not trigger compilation.
[str-00013] helpful=5 harmful=0 :: Linear followed by a linear reduction over output features, KernelBench matmul family. Examples: AvgPool1d(k) over output features (e.g., Linear 8192->8192 -> AvgPool1d(k) -> GELU -> scale -> row max), torch.sum(dim=1, keepdim=True) over all output features, or torch.mean(dim=1, keepdim=True) after a per-feature subtract. Exploit linearity before writing any kernels.

Case 1, pooling: pool(x W^T + b) = x W'^T + b'. Build W' by reshaping W [M, K] to [M/k, k, K] and taking the mean over the k axis. Build b' by reshaping b to [M/k, k] and averaging the same way. Build both in fp64, then cast to fp32. If M is not divisible by k, drop the tail rows first, because AvgPool1d floors. This shrinks the GEMM k times (16x for 8192->512 outputs). Then run F.linear(x, W', b') followed by one small fused per-row kernel for GELU -> scale -> max. A working row kernel is one Triton program per row that loads the whole row (BLOCK_N = triton.next_power_of_2(N)), masks with other=-inf and tl.where(mask, g, -inf), computes erf-GELU via top-level tl.erf (str-00012), multiplies by the scale, then tl.max and a single store.

Case 2, full sum or mean over output features: sum_j (x W^T + b)_j = x . (sum_j W_j) + sum_j b_j; for mean, use mean_j(W_j) and mean_j(b_j). Any per-output-feature additive vector (e.g., a subtract parameter of shape [out_features]) folds into the constant exactly like the bias: c = mean(b - s). The large GEMM collapses to a GEMV against one cached K-length vector. Steps:
- Build w_sum/w_mean = W.double().sum(0)/mean(0) and the constant in fp64, then cast to fp32 (or keep fp64).
- Run a per-row dot-product kernel (one program per row, looping K in BLOCK_K=2048 chunks, num_warps=8), with fp64 or pairwise accumulation if tolerances are tight (see the tolerance bullet). Fuse the following per-row elementwise op (e.g., erf-GELU via top-level tl.erf, str-00012) into the same kernel.
- Also trace the ops that follow. max, mean and logsumexp over a size-1 dim each return their single element, so a chain like sum(keepdim) -> max -> mean -> logsumexp after the reduction is an exact identity. Drop it.
- A trailing residual add of the original input (x + g[row]) can be done either as a separate broadcast-add kernel (2D grid = (cdiv(K,BLOCK), B), loading g[row] once per program) or fused into the GEMV kernel (after the dot product + GELU, a second loop re-reads the row, L2-hot, and stores x + g). Both measured the same (see evidence).

Cache W', b', w_sum/w_mean and the constant once, in __init__ or lazily keyed on weight.data_ptr() (and the subtract param's data_ptr; no private attrs, str-00014), and recompute only if the weight changes. Never rebuild them every forward: that is a full pass over the large weight and erases the gain.

Evidence:
- CONFIRMED end to end (pooling variant): Linear 1024x8192x8192 -> AvgPool1d(16) -> erf-GELU -> *2.0 -> row max. It used fp64-built W'/b' cached on data_ptr, F.linear down to 512 outputs, and the Triton per-row GELU/scale/max kernel described above. Result: correct, 9.37x (8.00 -> 0.853ms), via CHECK then SUBMIT.
- CONFIRMED TWICE end to end (mean variant): Linear 2048x8192x8192 -> subtract[out] -> mean(dim=1, keepdim) -> logsumexp(size-1) -> erf-GELU -> + original x. Run 1: fp64-built cached w_mean and c, a Triton GEMV with fused GELU, plus a separate Triton broadcast-add kernel: correct, 62x (15.97 -> 0.257ms), own_kernel_time_frac 1.0, via CHECK, CHECK, SUBMIT. Run 2: ONE fused Triton kernel per row (dot product, tl.erf GELU, second loop re-reading the row and writing x + g; BLOCK_K=2048, num_warps=8, grid=B): correct, 60.8x (15.94 -> 0.262ms), via CHECK then SUBMIT. Fusing the residual gave NO measurable gain over the two-kernel version; at grid=B=2048 the kernel is already near memory-bound. Treat ~61-62x as the SOLVED baseline for this task; resubmitting either layout cannot meaningfully improve the score. The 'excessive_speedup' audit flag is expected for a legitimate algebraic fold.
- Pooling case without the fold: keeping cuBLAS on the full 8k GEMM and fusing pool+erf-GELU+scale+max into a per-row reduction kernel was correct but gave only 1.004x (GEMM 99.6% of runtime). This layout is NOT acceptable; always fold.
- Sum case without the fold: keeping the full cuBLAS GEMM plus a Triton row-sum kernel gave 1.0099x. Two earlier CHECKs in that attempt failed at atol=rtol=1e-4 and were abandoned rather than debugged.

Remaining headroom after the fold (untested):
- Pooling case: own_kernel_time_frac was only 0.0067, so the reduced cuBLAS GEMM (e.g., 1024x8192x512) dominates. Next levers: a TF32 tensor-core GEMM within tolerance (plain TF32 per str-00032 if tolerance ~1e-2, tf32x3 otherwise; mind the NT-layout warnings in str-00075), or a Triton GEMM on W' whose epilogue does bias+GELU+scale and a per-tile row max combined across N-tiles with tl.atomic_max (or a small partials buffer). Time locally against the folded baseline before SUBMIT.
- GEMV case: residual fusion is already measured as neutral. Only small levers remain (e.g., float4-friendly larger BLOCK_K, splitting each row across a few programs, or a 2D-grid residual pass with more parallelism); expect at most a few percent.

The same idea applies to any linear op after a GEMM: sum or mean pooling, fixed linear combinations of output features, per-feature add/subtract, or a scalar scale. It does NOT apply if a nonlinear op sits between the GEMM and the reduction (see str-00067). Alternative if folding is impossible: a Triton tl.dot kernel whose epilogue does the reduction, so the large GEMM output is never written to global memory. Use allow_tf32=True only if the tolerance allows it.
[str-00014] helpful=42 harmful=0 :: Static-check pitfall: the evaluator's static checker rejects some Python constructs outright, and the turn is wasted.
(a) Any attribute assignment whose target is not self.<x> is rejected. For example, `bn.num_batches_tracked += 1` or `layer.foo = ...` on a submodule fails with 'assignment to attribute ... is not allowed (only self.<x>)'.
(b) Reading private or underscore attributes is rejected. For example, `w._version` or `b._version` fails with 'private attribute ... is not allowed'. This often appears in weight-cache invalidation code, such as caching a scaled, folded or flipped weight.
(c) getattr with a computed or non-literal name is rejected ('getattr needs a literal, non-private attribute name'). For example, getattr(self.gru, f'weight_hh_l{l}') fails. For RNN modules, use self.gru.all_weights[2*l + dir] (a list of [w_ih, w_hh, b_ih, b_hh]) instead. Otherwise, use literal attribute names or an explicit list.
(d) Any non-ASCII character anywhere in the source, including comments and docstrings (e.g., an arrow U+2192), causes a syntax-error rejection. Keep the whole file ASCII, writing '->' for arrows.
(e) ANY reference to 'torch.backends' is rejected ('torch.backends' is not allowed (streams/graphs/events/compile/flags/file IO)), and assigning torch.backends.cudnn.benchmark also trips rule (a). MEASURED on L3 RegNet: a module-level `torch.backends.cudnn.benchmark = True` failed the turn-0 CHECK. So the cuDNN flag levers listed in many bullets (cudnn.benchmark, cudnn.allow_tf32, cuda.matmul.allow_tf32) are NOT available through torch.backends in submitted code. The same message lists streams, CUDA graphs, events and torch.compile as forbidden, so treat CUDA-graph capture as likely rejected too (verify with CHECK before planning around it). Conv-side levers that remain: channels_last via your own transpose kernel (str-00102, non-dilated only per str-00051), TF32 inside your OWN kernels (Triton tl.dot input_precision), or custom fused conv epilogues.

Always CHECK before SUBMIT so these static failures are caught early.
[str-00015] helpful=1 harmful=3 :: DenseNet-style BN->ReLU->conv networks (KernelBench network family, e.g., DenseNet121/DenseNet201 at batch 10, 224px, or a single DenseNet121 dense block). cuDNN convs and the repeated torch.cat dominate, so swapping only the BN+ReLU pass gives ~1.0x or worse. On FULL networks every per-layer BN+ReLU replacement tried so far has LOST; treat it as a measured dead end unless a local full-size timing in BOTH train() and eval() says otherwise.

Measured, training mode: a custom fused BN+ReLU replacement was correct but gave only about 1.04x. It used three steps: a per-channel sum/sumsq stats kernel, a tiny finalize kernel that also updates running_mean and running_var (unbiased var*n/(n-1), with momentum), and an apply+ReLU pass. The custom kernels were about 18% of runtime.

Measured FAILURE 1, eval mode (DenseNet201): 0.986x (39.70 vs 39.13ms), 0.936x vs torch.compile, own_kernel_time_frac 0.037. The candidate replaced eager BN+ReLU with one fused bn_relu kernel per layer, but kept torch.cat. It also recomputed scale = gamma/sqrt(var+eps) and shift = beta - mean*scale on EVERY forward. Those are ~4 tiny torch launches per BN across ~100+ BN layers, plus an empty_like per layer. The eager reference is already one cuDNN BN pass plus an in-place ReLU, so the single saved elementwise pass per layer cost less than the added launch and allocation overhead. It was submitted via CHECK-then-SUBMIT with no local timing.

Measured FAILURE 2, eval mode (DenseNet201, full network), WORSE than failure 1: 0.9523x (40.48 vs 38.55ms), 0.917x vs torch.compile, own_kernel_time_frac 0.027. It applied the block-buffer recipe below items (0), (1) and (1b): preallocated block buffer, BN fold cached in a dict keyed on bn.weight.data_ptr(), a strided bn_relu_slice kernel reading the buffer prefix directly but writing a FRESH compact pre-activation tensor, F.conv2d into a temp, then copy_ into the buffer slice. The transition still ran BN+ReLU -> 1x1 conv -> pool (item (4) skipped), and there was no CUDA graph (item (5) skipped). Each layer still pays a full read+write of the growing prefix, an extra conv-output copy and two allocations, which is comparable traffic to the reference torch.cat, plus extra Python/launch overhead per layer across 98 layers at 7x7-56x56. CHECK then SUBMIT with no timing.

Measured FAILURE 3, CATASTROPHIC (DenseNet121, full network, batch 10, 224px): 0.4898x (31.74 vs 15.55ms, lo95 0.488), own_kernel_time_frac 0.058. It applied (1), (1b) AND lever (4): block buffer, strided bn_relu into a fresh prefix tensor, conv-to-temp + copy_ into the slice, transition BN+ReLU+avgpool fused BEFORE the 1x1 conv, eval fold cached once (keyed on a data_ptr). The training branch folded BN with eager x.mean/x.var over the non-contiguous growing prefix view plus running-stat mul_/add_ (~10 launches per BN x ~120 BNs, a full O(L^2) re-reduction of the prefix every layer). The ~16ms of extra time is NOT in the custom kernels (only 5.8%); it is per-layer torch/Python work: most likely that eager training-mode stats path, otherwise per-layer allocations, copies and dispatch. Lever (4) did NOT rescue it. Turn 0 failed no_custom_kernel; turns 1-2 correctness CHECKs; turn 3 SUBMIT with no timing, while this bullet already recorded the DenseNet201 losses.
CONCLUSION: the single-block win below does NOT extrapolate to full networks; block buffer + BN+ReLU swap (with or without the pool-first transition) is a measured LOSS on DenseNet121 and DenseNet201. Do not submit it there.

Measured WIN, eval mode, SINGLE BLOCK only (L3 DenseNet121DenseBlock: batch 10, 224px, 6 layers, C_in 32, growth 32): 1.098x (15.50 -> 14.12ms), own kernels ~20% of runtime. Recipe: (i) ONE preallocated block-wide buffer of C_in + L*growth channels, with the input copied into the first slice and each layer's conv output copied into its channel slice (no torch.cat); (ii) BN fold scale/shift cached per layer in a self.<x> dict keyed on bn.weight.data_ptr(); (iii) one fused bn_relu kernel per layer, then cuDNN F.conv2d; (iv) a training branch with its own bn_relu_train kernel so kernels run in both modes (str-00031). Here cutting the O(L^2) concat traffic turned the ~1.0x swap into a win. It went CHECK-then-SUBMIT with no timing.

What to do on full networks:
(A) MANDATORY: never submit any per-layer BN+ReLU replacement without a local full-size timing against the reference in BOTH train() and eval() modes (bounded warmup/iterations, cuda synchronize). A sub-1x timing blocks SUBMIT.
(B) Never compute training-mode BN stats with torch reductions over the whole prefix (x.mean/x.var on the strided view) per layer. Make stats incremental: compute sum/sumsq in-kernel only for each newly written channel slice (e.g., in the kernel that writes it), cache them in a per-block stats buffer, and let a tiny finalize kernel per BN fold gamma/beta and update that module's running stats in-kernel (buffers passed as tensors, str-00014).
(C) The only remaining untested route with real upside: CUDA-graph the whole forward (static input/output and intermediate buffers) with the custom kernels inside the graph (str-00026), plus cuDNN TF32 within tolerance (str-00029).
(D) Fallback if local timing stays below 1x: keep the eager torch.cat and nn.BatchNorm/ReLU per layer, and add only the fused transition BN+ReLU+avgpool kernel (pool before the 1x1 conv) as the custom piece, routed through both branches.

Detail levers (necessary, not sufficient):
(0) In eval mode, fold each BN's scale/shift ONCE and cache it as self.<x> attributes (see str-00038). Never recompute the fold per forward.
(1) Attack the concat: preallocate one block-wide feature buffer per dense block (C_in + L*growth channels), copy the input into the first slice, have each layer's conv write into its channel slice, and read buffer.narrow(1, 0, C_i) with no cat. Wins on a single short block; lost on full networks (failures 2, 3).
(1b) buf[:, :cur].contiguous() copies the growing prefix every layer; have the bn_relu kernel read the strided view directly (batch stride C_total*HW). The bn_relu output is still a fresh prefix-sized tensor per layer and the conv output goes to a temp + copy_, so traffic and allocations stay comparable to torch.cat; reuse a preallocated pre-activation buffer and write the conv output in place where possible.
(2) Eval mode: each layer's BN has its own gamma/beta/running stats per channel, so verify reuse is exact before caching BN-applied slices.
(3) Vectorize the normalize+ReLU apply with float4 loads/stores and use 2D grid indexing to avoid per-element 64-bit div/mod (minor here).
(4) Fuse the transition's BN+ReLU with the AvgPool BEFORE the 1x1 conv (avgpool commutes with the 1x1 conv; 4x fewer conv FLOPs): BN+ReLU per element, average the 2x2 window, write the quarter-size tensor, then F.conv2d 1x1. Correct, but measured NOT to rescue the block-buffer layout (failure 3); keep it as the fallback custom piece in (D).
(5) Route custom kernels through the eval branch and preferably both branches (str-00031); a kernel only in an unchecked path fails no_custom_kernel.
[str-00016] helpful=3 harmful=2 :: Linear (large GEMM, e.g., 1024x8192 @ 8192x8192 fp32, or 16384x4096 @ 4096x4096 fp32) followed by training-mode BatchNorm1d and an elementwise chain (bias, divide, swish, GELU, ReLU, scale -> softmax, etc.), KernelBench matmul family. Training-mode BN uses batch statistics, so it can NOT be folded into the weight.

The low-gain layout keeps cuBLAS F.linear and collapses the eager BN/elementwise passes into 2-3 custom kernels:
(1) a per-column sum/sumsq stats kernel that also updates running_mean and running_var in-kernel (unbiased var*N/(N-1), with momentum), or a separate finalize kernel for that;
(2) one in-place normalize+affine+activation pass.
Results:
- Linear 8192 + swish chain: about 1.01x vs eager and 0.997x vs torch.compile. The GEMM was about 98% of runtime and the custom kernels about 1.9%.
- 16384x4096x4096 + BN + erf-GELU + ReLU (L2_041, atol=rtol=1e-4): 1.014x (30.05 vs 30.47ms). The custom kernels were about 3.4% of runtime.
- SAME L2_041 task with a SLOW tail: 0.9517x (32.04 vs 30.49ms), own_kernel_time_frac 0.088. The stats kernel used ONE THREAD PER COLUMN (N/256 = 16 blocks, 4096 threads total), each thread serially looping over all 16384 rows with double FMAs: latency-bound with most SMs idle. The apply kernels also did a 64-bit idx % N per element, and the eval kernel recomputed rsqrtf per element. A slow custom tail does not merely fail to gain; it drops the score BELOW 1x (~1.5ms worse than the 1.014x tail).
~1.01x is the CEILING of this layout. Do NOT submit it expecting more, and never submit its tail untimed. Decide on the first turn whether to rewrite the GEMM, and do the rewrite while turns remain.

Measured FAILURES of a tf32x3 GEMM rewrite on Linear 1024x8192x8192 -> training BN -> *scale -> softmax(dim=1). Both used an autotuned tf32x3 Triton GEMM with fused bias, reading the nn.Linear weight as W.t() through swapped strides ([BK,BN] tile, K contiguous) with a masked K loop. Both were correct, and both went CHECK then SUBMIT with no timing.
- Slow tail: a one-block-per-column stride-N stats kernel with double FMAs, a finalize kernel, a 3-read-pass softmax, and a .item() sync on the scale every forward. It scored 0.84x (9.62 vs 8.10ms), own kernels 99.9%.
- Lean tail: column-spanning coalesced stats kernel, a finalize kernel reading the scale by pointer and updating running stats in-kernel, and a single-read register-resident row softmax (whole 8192 row loaded once). It scored 0.9228x (8.77 vs 8.09ms), own kernels 100%. The tail is only ~0.5-1ms, so the GEMM itself ran ~7.5-8ms. That is far from the ~5.1ms tf32x3 reaches on NN/TN layouts.
Conclusion: the NT-layout swapped-stride tf32x3 GEMM is the bottleneck (see str-00075). Fixing the surrounding passes alone cannot produce a win, and a lean tail is necessary but not sufficient. At M=16384 the fixed tf32x3 NN config failed catastrophically (0.0336x, str-00090), and plain TF32 likely fails at 1e-4, so any GEMM rewrite on L2_041-class shapes must be timed ALONE against F.linear before building the pipeline around it.

Unscored attempt (L2_097, Linear 1024x8192x8192 -> BN -> +bias -> /d -> swish): cached W.t() NN tf32x3 fixed 128x128x64/8 warps/3 stages with per-column sum/sumsq tl.atomic_add in the epilogue, Linear bias folded into the stats algebraically, CUDA finalize (running stats in-kernel), in-place apply+swish. CHECK passed, then an immediate untimed SUBMIT crashed on an evaluator Triton-cache error (str-00077 (c)), scoring 0. Its GEMM design matched the str-00090 catastrophe signature anyway.

Correctness details that worked:
- Compute biased var as E[x^2]-m^2 for normalization.
- Use the running stats when not self.training, or fall back to the eager modules.
- Skip num_batches_tracked.
- Keep the stream retrieval and kernel launch in cuda_sources.
- The Linear bias can be folded into raw-GEMM stats: sum' = sum + M*b, sumsq' = sumsq + 2*b*sum + M*b^2, then shift += b*scale.

For a real gain, change the GEMM AND make the tail lean. Follow a STAGED process: GEMM variants, in order (time the GEMM alone vs F.linear after the first CHECK):
(1) Plain TF32 (str-00032 recipe), after printing the max error against the reference. Softmax after BN tends to shrink absolute errors, so ~1e-2 tolerance is likely to pass. BN followed by GELU/ReLU does not hide relative error.
(2) If tf32x3 is required, convert to NN via a cached W.t() (str-00075 lever 4) with a LIGHT (bias-only) epilogue, subject to the str-00090 guard.
(3) Restrict to BM=BN=128 configs and run the K loop unmasked when K % BK == 0.
(4) DO NOT put per-column stats (tl.sum + tl.atomic_add or a partials write) into the epilogue of the fixed 128x128x64 tf32x3 GEMM by default: that is exactly the str-00090 ~30x-slowdown signature (measured 0.0336x on L2_039). Compute the stats in a separate 2D-grid stats kernel (layout below). Only try epilogue stats after the light-epilogue GEMM has been timed, and re-time; a ~10x jump means revert.
Tail:
- A tiny finalize kernel reduces the partials (double math is fine here, it is only N elements), updates running stats, and folds gamma, beta and any scalar scale into per-column sc/sh. Read the scale parameter via its pointer inside this kernel (no .item()).
- Then ONE normalize+activation pass. For a softmax tail, use a row kernel that loads the row once into registers (or an online max+sum), then one write pass. For an elementwise tail, use a 2D grid (column index from blockIdx.x/threadIdx.x, rows from blockIdx.y with a grid-stride over rows), sc/sh held in registers, float4 loads/stores, and no per-element 64-bit idx % N.
- At minimum, run the GEMM in TF32, if within tolerance.
- Time the GEMM alone against F.linear, and the full pipeline (or at least the tail against eager BN+activation) against the reference, before SUBMIT, so a sub-1x result can be traced to a stage. Never go CHECK -> immediate SUBMIT.

Inefficiencies to avoid in the stats pass:
- A one-block-per-column stats kernel that reads with column stride (stride C) is uncoalesced: each warp load touches 32 sectors (~8x traffic).
- One thread per column looping over ALL rows is coalesced but badly UNDERFILLED when N is a few thousand (N=4096 gives only 16 blocks; measured 0.95x above). Threads spanning columns is necessary but not sufficient.
- Correct layout: a 2D grid of (column tiles of 256 threads, coalesced across columns) x (row chunks of ~64-128 rows each), giving thousands of blocks; fp32 accumulation per thread; then atomicAdd into zeroed per-column sum/sumsq (or write a [row_chunks, N] partials buffer) for the finalize kernel. Do not use double FMAs per element.
- Reading a scalar parameter (bias or scale) via .item() forces a host sync every forward. Pass the tensor pointer to the kernel instead.
[str-00017] helpful=2 harmful=0 :: Large transformer networks (Swin/ViT/SwinMLP, KernelBench attention/network family, e.g., SwinV2 or SwinMLP at batch 10, 224px). Swapping a single already-fused eager op for a standalone Triton kernel makes the model SLOWER. Two measured cases: (a) swapping only nn.GELU (act_layer) for a standalone erf-GELU kernel in SwinV2 scored 0.84x (ref 18.9ms, cand 22.4ms), with the custom kernel at about 2.3% of runtime; (b) swapping only nn.LayerNorm for a one-row-per-program Triton LN kernel (num_warps=4) in SwinMLP scored 0.70x (ref 8.54ms, cand 12.21ms, 0.50x vs torch.compile), with the custom kernel at about 4.6% of runtime. Eager nn.GELU and F.layer_norm are each already a single fused pass, so the replacement removes no memory traffic. It adds launch overhead (~25 LN calls per forward), Python module dispatch, empty_like allocations and reshape(...).contiguous() copies. A non-contiguous input, such as the PatchEmbed flatten+transpose output, forces a full copy; read such inputs via strides instead. A one-program-per-row grid on tiny rows (C=96..768) also underutilizes the GPU; process several rows per program. This is the same failure as a standalone single activation after a conv (see the conv+BN+activation bullet). A real win must remove memory passes or launches by targeting multi-op eager chains. (1) Cosine window attention (SwinV2): F.normalize(q), F.normalize(k) → q@k^T → * clamp(logit_scale).exp() → + relative-position bias → + shift mask → softmax → @v. These are about 8+ kernels over B_×heads×N×N, with small N (49 for a 7x7 window). Replace them with ONE fused per-window Triton kernel: one program per (window, head), with N padded to a power of 2 (49→64) and masked loads (-inf on padded columns). Keep the qkv and proj linears in cuBLAS. (2) The relative-position bias, 16*sigmoid(cpb_mlp(coords_table)) gathered by relative_position_index, and clamp(logit_scale).exp() depend only on parameters. Compute them once and cache them, recomputing only if the parameters change, instead of rerunning the MLP and gather every forward. Precombine bias+mask into one [nW, heads, N, N] additive tensor per block. (3) Fold torch.roll / F.pad + window_partition / head transpose / window_reverse / crop permute+contiguous copies into gather/scatter index math inside the attention or qkv-load path. In SwinMLP, the grouped Conv1d spatial MLP is a per-head [49x49] matrix applied over the window-token axis. Run it as one batched matmul, or as a gather-matmul-scatter kernel that reads directly from the [B,H,W,C] normed tensor and writes back in place of window_reverse+crop. This removes about 4 permute+contiguous copies per block. (4) Further fusions: the fc1 bias with GELU; the residual add with LayerNorm in one kernel that writes both the sum and the normed output (pre-norm: x = shortcut + spatial; y = LN(x); post-norm: x = shortcut + norm(y)); and PatchMerging's strided 2x2 gather + torch.cat fused with its LN. Profile which ops dominate and time locally against the reference before submitting. A CHECK that only verifies correctness does not catch a slowdown; both swaps above were found slow only after a CHECK-then-SUBMIT flow.
[str-00018] helpful=7 harmful=2 :: Multi-layer LSTM/GRU/RNN, unidirectional or bidirectional (KernelBench rnn family, e.g., LSTM returning h_n or c_n, bidirectional LSTM -> fc(out[:, -1, :]), bidirectional GRU returning only `output`, bidirectional GRU returning only h_n, or unidirectional GRU returning `output`, with B=10, T=512, H=256, L=6).

Do NOT reimplement the recurrence as a Python loop over layers x timesteps, where each step does an addmm with W_hh, empty_like allocations for h/c, a tiny custom gate/cell kernel and a slice copy. That is about 3k steps with several launches each. Host launch and dispatch overhead dwarfs the tiny per-step work, while cuDNN runs the whole recurrence as fused persistent kernels. One attempt was correct but scored 0.10x (315ms vs 32ms); its custom kernels were only about 17% of runtime. Precomputing the input projection as one big GEMM per layer does not rescue it.

Also do NOT use the 'cuDNN wrapper' layout. This means keeping nn.LSTM/nn.GRU, dropping dead work (fc, out[:, -1] slice, unused outputs), and adding a trivial custom kernel (e.g., a copy of c_n, h_n or `output`, or replacing only the tiny final fc Linear). It is GUARANTEED to land at about 0.77-0.80x, which scores as a failure. Measured FIVE times:
- LSTM returning c_n: 0.79x (40.3 vs 31.8ms), own_kernel_time_frac=0.0003.
- LSTM returning c_n again (grid-stride c_n copy kernel, submitted via CHECK, CHECK, SUBMIT while this bullet was shown): 0.785x (40.3 vs 31.6ms), own_kernel_time_frac=0.0004. An exact reproduction of the first result.
- Bidirectional GRU, passthrough copy of output: 0.80x (75.8 vs 60.4ms), own_kernel_time_frac=0.0003.
- Bidirectional GRU returning h_n, with self.gru.flatten_parameters() at forward start plus a copy kernel: 0.771x (80.2 vs 61.9ms), own_kernel_time_frac=0.0001.
- Bidirectional LSTM -> fc(out[:, -1, :]), nn.LSTM kept and only the Linear(512->10) replaced by a naive custom fc kernel: 0.7948x (101.5 vs 80.6ms), own_kernel_time_frac=0.0003. The fc is dead-cheap and is NOT a lever.
In all five the extra time is inside the cuDNN RNN call when it is wrapped in ModelNew, not the custom kernel. Calling flatten_parameters() does NOT fix the gap, so the re-flatten hypothesis is unconfirmed, and no known tweak rescues this layout. An unused output (h_n or output) is not a lever either: cuDNN computes it anyway at negligible cost (e.g., c_n does not depend on fc or on the last layer's `output` writes, but cuDNN still does them). Never submit this layout, even as a 'safe' fallback after a custom recurrence fails a CHECK: it is not safe, it is a guaranteed failing score. If a custom-recurrence CHECK crashes (e.g., 'candidate process died'), debug it instead of retreating: suspect a grid-sync deadlock from non-co-resident blocks (grid > SMs x occupancy), an illegal smem size (needs cudaFuncSetAttribute for >48KB dynamic smem), or out-of-bounds indexing. Test on a tiny shape (small T, L) first, and fall back to option (b) below rather than to the wrapper.

Also do NOT write a persistent kernel parallelized over BATCH with weights streamed. One GRU attempt used grid=B (10 blocks), block=H threads, one thread per hidden unit, with each thread re-reading its 3 W_hh rows (3x256 floats) from global/L2 every timestep in a serial 256-long dot product. It was correct but scored 0.07x (457 vs 32ms, own_kernel_time_frac 0.998). The causes: only 10 SMs busy; about 786KB of weights per block per step; reads uncoalesced (each lane reads a different row at the same k); the whole weight set re-read once per batch element; and 2 __syncthreads per step.

Options that can actually win; commit to one on turn 0:
(a) Write one persistent kernel per layer, or a wavefront across all layers, that loops over time on the GPU. Decompose over HIDDEN UNITS, not batch:
- Partition the 3H (GRU) / 4H (LSTM) rows of W_hh across many blocks (e.g., 64-128 blocks, each owning a slice of hidden units j with all its gate rows, about 6-12KB). Load that slice into shared memory or registers ONCE and keep it there for all T steps.
- Each block computes its slice for ALL B batch elements at once, so every weight load is reused B times.
- Keep h_t in a double-buffered global/L2 array, read it into shared memory each step, and write the block's new h slice to the other buffer.
- Synchronize per step with a cooperative launch + grid.sync(), or a global atomic-counter barrier with all blocks guaranteed co-resident (grid <= SM count x occupancy).
- Optionally run a wavefront across layers (layer l step t concurrently with layer l+1 step t-1) to multiply the active blocks.
Feed it the precomputed input-projection GEMM, one big GEMM per layer (e.g., torch.addmm), with biases b_ih (+ b_hh for LSTM) folded in. For GRU, keep b_hn separate, since it sits inside r*(W_hn h + b_hn). Fuse the gate nonlinearities and the h (and c) update. When only c_n/h_n is returned, write the final cell/hidden state straight into the c_n/h_n slot per layer; the last layer's per-step output need not be stored. For bidirectional models, compute the input projection for BOTH directions in one GEMM. Run the forward and backward recurrences concurrently in the same persistent launch (disjoint block sets), writing directly into the two halves of the [T,B,2H] output, which becomes the next layer's input. Write each direction's final h straight into its h_n slot.
(b) At minimum, preallocate the h/c and output buffers and capture the timestep loop (or the whole forward) in a CUDA graph with static buffers, to remove per-step launch and allocation overhead.

Dead-work pruning lever when the model reads only out[:, -1, :] (and not h_n/c_n): the LAST layer's backward direction contributes only its output at t=T-1, which is its FIRST step (from h0/c0 for that layer/direction, using the input projection at t=T-1). Compute just that one step and skip its full T-step backward recurrence; the last layer's forward direction and all lower layers (both directions) must still run fully. This saves about 1/(2L) of the recurrence work (~1/12 for L=6). It is only available in a custom recurrence, not via nn.LSTM.

Always time locally against a reference nn.LSTM/nn.GRU before submitting, ideally after load_state_dict from a reference instance, to mimic the evaluator's weight sync. Never SUBMIT on turn 1 straight after a correctness-only CHECK. All five ~0.8x attempts and the 0.07x batch-parallel attempt were submitted right after CHECK with no timing; a single timing run would have exposed each slowdown.
[str-00019] helpful=4 harmful=0 :: Structured-sparsity GEMM (triangular, banded or block-sparse operands or outputs, e.g., KernelBench tril(A@B) where A and B are lower triangular, or triu(A@B) where A and B are upper triangular). A custom Triton tile kernel can beat dense cuBLAS by a wide margin, even with strict FP32 tl.dot(allow_tf32=False). The gain comes from skipping structurally zero work, not from computing faster.

Lower-triangular A and B (tril): two kinds of work can be skipped.
(1) Output tiles fully above the diagonal: return early when row0+BM-1 < col0.
(2) Out-of-range K iterations: for tile (row0, col0), loop K only from floor(col0/BK)*BK to min(row0+BM, N). This works because A[i,k]=0 for k>i and B[k,j]=0 for k<j. It cuts FLOPs to about 1/6 of the dense count.
Epilogue: mask the store with offs_m>=offs_n (zero the upper part of diagonal tiles via tl.where). Preallocate C with torch.zeros so skipped tiles remain zero.
Results at N=4096 fp32, plain 2D grid, 1D kmask broadcast per operand (str-00042):
- Strict FP32: 4.10x (7.40 vs 1.81ms) with BM=BN=128, BK=32, num_warps=8, num_stages=3.
- input_precision='tf32x3': 6.37x (7.38 -> 1.16ms, lo95 6.33, own kernel ~96%), passed correctness. Autotune of 4 configs centred on BM=BN=128, BK 32/64, 8 warps, 3-4 stages (plus one 128x64/4-warp config). Treat ~6.4x as the SOLVED BASELINE for tril: resubmitting it cannot raise the score. It went CHECK on turn 0 then SUBMIT on turn 1 with no timing.

Upper-triangular A and B (triu): mirror the roles, since A[i,k]=0 for k<i and B[k,j]=0 for k>j.
- Return early when row0 > col0+BN-1 (tile fully below the diagonal).
- Loop K only from floor(row0/BK)*BK to min(col0+BN, N).
- Mask the store with offs_m <= offs_n, and preallocate C with torch.zeros.
Results at N=4096 fp32 (plain 2D grid, 1D kmask broadcast per operand per str-00042, masked loads):
- Strict FP32 (allow_tf32=False), fixed BM=BN=128, BK=32, 8 warps, 3 stages: 3.33x (7.36 vs 2.21ms, own kernel 98%).
- tf32x3, same fixed config: 4.39x (7.41 -> 1.69ms, own kernel 97.7%).
- tf32x3 with the tril-style 4-config autotune set (128x128 BK 32/64, stages 3-4, plus 128x64/4 warps): 4.4582x (7.41 -> 1.66ms, lo95 4.442, own kernel 97.7%). CONFIRMED: autotune is roughly NEUTRAL and does NOT close the gap to tril's 6.37x.
Treat ~4.45x (~1.66ms) as the SOLVED BASELINE for triu with this layout: resubmitting it (fixed or autotuned) cannot raise the score. All three triu runs went CHECK then SUBMIT with no local timing, so every structural lever below is still untested.
Prefer tf32x3 over strict FP32 for both orientations (confirmed on both). Likely cause of the triu-vs-tril gap: 2D-grid scheduling order. In triu the long-K tiles (small row, large col) are interleaved with programs that exit early / short near-diagonal tiles, giving wave imbalance and a long tail; plus programs launched only to exit.

Untried headroom (applies to both orientations; after the first passing CHECK, time each locally against the tf32x3 baseline, ~1.16ms tril / ~1.66ms triu, instead of submitting; on triu start with (b)):
(a') Plain tf32 (allow_tf32=True) if tolerance allows; K can be as long as N, so print the max error first.
(b) A 1D grid enumerating only the valid-triangle tiles, so no programs are launched just to exit. Order the tiles longest-K-range first for load balance.
(c) Drop the torch.zeros pass (~64MB write) by having the skipped/below-diagonal tiles write zeros instead (torch.empty output).
(d) Unmasked loads when N % BK == 0 (4096 divides BK 32/64).

General rule: before writing any matmul kernel, check whether the inputs or the output have known zero structure that the reference ignores.
[str-00020] helpful=3 harmful=0 :: Training-mode BatchNorm followed by subtracting the per-(n,c) spatial mean (x - x.mean(dim=(2,3,4), keepdim=True)), KernelBench conv_transpose family (L2_015). Simplify the algebra before writing kernels: the BN shift (beta) and the per-channel batch mean both cancel. So out = gamma_c * rsqrt(var_c + eps) * (x - spatial_mean_nc), where var_c is the biased batch variance.

Kernel layout (2 full-tensor passes instead of about 5 for eager BN stats + normalize + torch.mean + subtract):
(1) ONE reduction kernel, one block per (n,c), producing sum and sumsq (same as the BN+global-avgpool bullet).
(2) A tiny per-channel finalize kernel: sum over N; mean_c = csum/(N*S); var_c = E[x^2] - mean_c^2 (accumulate in double); scale_c = gamma_c*rsqrt(var_c+eps).
(3) ONE apply pass: scale_c * (x - sum_nc/S), 2D grid (blockIdx.y = n*C+c).
Use the str-00001 layout and long S (str-00058).

Measured, CONFIRMED TWICE on ConvTranspose3d 16->32, k=3, stride 2, p=1 (16x16x16x32x32 input) -> BN3d -> subtract:
- Run 1: 1.30x over eager and 1.06x over torch.compile; custom kernels ~13% of runtime, cuDNN conv ~87%.
- Run 2: 1.3278x (7.57 -> 5.70ms, lo95 1.3265), own_kernel_time_frac 0.10. Double-accumulating stats kernel (256 threads), one-thread-per-channel finalize, 2D-grid apply into empty_like. It went CHECK on turn 0 then SUBMIT on turn 1 with no timing and no conv lever tried.
~1.33x is the SOLVED BASELINE: resubmitting it cannot raise the score. Do NOT go CHECK -> immediate SUBMIT on this task.

After the first passing CHECK, spend the spare turns timing at full size against the reference (bounded warmup/iterations, cuda synchronize) and keep the fastest:
(a) Conv side (~90% of runtime, the main payoff): cuDNN TF32 within tolerance (print max error first, str-00029), torch.backends.cudnn.benchmark=True, channels_last_3d (timed per str-00051). Do NOT use the per-phase F.conv3d decomposition: k=3, s=2, p=1, dilation 1 is the non-shortcut case, a measured loss (str-00030). Stride 2 also rules out the stride-1 forward-conv rewrite.
(b) Kernel side (~10% of runtime, small gains): float4 fp32 stats with per-block double combine instead of double FMAs per element (str-00043); write the apply in place on the conv output instead of empty_like; atomics into per-channel buffers instead of the finalize's serial loop over N.
(c) Update running_mean/running_var in the finalize kernel only if module state must match; pass the buffers to the kernel (str-00014).

General rule: when a normalization is followed by a mean-subtraction or a pooling step, look for terms that cancel or reduce to per-(n,c) statistics before writing any normalize pass.
[str-00021] helpful=14 harmful=0 :: Conv → reduction over one dim → softmax over channels (e.g., Conv3d → min over D → softmax over C, KernelBench conv family). Keeping cuDNN conv and adding a badly laid out reduction kernel is SLOWER than eager. One attempt used one block per (n,h,w) with only C=24 threads, each thread striding through D by H*W, with thread-0 serial max/sum loops in shared memory and syncthreads. It was correct but scored 0.84x (2.84 vs 2.38ms), with the custom kernel at about 26% of runtime. The problems: warp loads are uncoalesced (lanes are Dp*H*W floats apart), 75% of each warp is idle, and the serial reductions add latency.

Better layout (avoids the slowdown, but only reaches PARITY): one thread per output pixel (n,h,w), with w the fastest-varying thread index so warps read contiguous W. Each thread loops over c and d, keeping the per-c reduction results in a register array (C is small and fixed, so template or #define it). It then does the max/exp/sum/normalize softmax over C entirely in registers and writes the C outputs, which are also coalesced across w. This reads the conv output once, fully coalesced, with no shared memory or __syncthreads. MEASURED on IC=3, OC=24, k=3, input 128x3x24x32x32: correct but 0.9887x (2.404 vs 2.377ms, lo95 0.988), 0.90x vs torch.compile, own_kernel_time_frac 0.115. It was a turn-0 SUBMIT with no CHECK or timing. The eager min+softmax tail is already cheap (two small passes), and the cuDNN conv dominates, so fixing the reduction layout cannot produce a win. Do NOT submit this layout expecting >1x. Minor kernel weaknesses: a 64-bit multiply per D iteration and serial dependent loads across 24x22 values per thread (low memory-level parallelism).

To beat 1x, change the conv (decide on turn 0, time each at full size before SUBMIT):
(a) A fused direct conv. This is feasible when IC and k are tiny: here the weight is only 24x81 floats, which fits in shared/constant memory. Each thread handles one (n,h,w) pixel for all OC, computes every D output, keeps a running min per channel in registers, and applies the softmax in its epilogue. The full 5D conv output (~0.6GB here) is never written or re-read. Verify indexing against nn.Conv3d on a small shape first.
(b) cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv) within tolerance (str-00029), and/or cudnn.benchmark, paired with the register-resident tail kernel.

General rule: in any reduction kernel, map consecutive threads to consecutive memory addresses (the innermost dim), and reduce along the strided dims inside each thread's loop. Time locally against the reference before SUBMIT; a correctness-only CHECK does not catch a slowdown or a parity result, and a turn-0 SUBMIT forfeits the chance to switch to (a)/(b).
[str-00022] helpful=2 harmful=0 :: Conv -> per-channel multiply -> InstanceNorm (no affine) -> clamp -> multiply again -> max over channels (e.g., Conv3d -> *mult[C,1,1,1] -> InstanceNorm3d -> clamp -> *mult -> torch.max(dim=1), KernelBench conv family, L2_079).

Recipe: keep the cuDNN conv and replace the whole tail with TWO kernels.
(1) Stats kernel: one block per (n,c) computes sum and sumsq of the RAW conv output; mean = sum/S, biased var = E[x^2]-mean^2.
(2) Apply + channel-max kernel: one thread per output position (n,s), with s the fastest-varying thread index so the loads across the strided c loop are coalesced. Each thread loops over c and computes the norm directly.
- Fold the pre-norm multiplier into the stats algebraically, so x*m is never materialized: mean(x*m) = m*mean and var(x*m) = m^2*var, giving norm = (v*m - m*mean) * rsqrt(m^2*var + eps).
- Then clamp, multiply by m again, and keep a running max in a register. Apply the clamp and the second multiply BEFORE the max, because a negative multiplier does not commute with max.
- Only the [N,D,H,W] output is written.
About 6 eager passes become 2 reads of the conv output plus a small write.

Measured, CONFIRMED THREE TIMES:
- Run 1: 1.41x (2.17 -> 1.53 ms), custom kernels ~14% of runtime.
- Run 2: 1.3923x (2.17 -> 1.56ms, lo95 1.391), only 0.919x vs torch.compile, own_kernel_time_frac 0.17. Used 256-thread stats blocks with a shared tree, long sizes (str-00058) and the str-00001 layout. It compiled first try, then went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no lever tried.
- Run 3: 1.3828x (2.18 -> 1.57ms, lo95 1.381), own_kernel_time_frac 0.16. Recipe copied nearly verbatim (grid-stride apply kernel, 4096-block cap, per-element recompute of mean/var/rsqrt per channel, 64-bit div per thread); CHECK on turn 0, SUBMIT on turn 1, no timing, no lever. It gained NO new information.
~1.38-1.41x is the SOLVED BASELINE for this task: resubmitting it cannot raise the score, and it is still below torch.compile. Do NOT go CHECK -> immediate SUBMIT again.

After the first passing CHECK, spend the spare turns timing these levers at full size against the reference (~1.57ms baseline) before SUBMIT, and keep the fastest. NOTE: torch.backends flags (cudnn.benchmark, cudnn.allow_tf32) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Realistic levers:
(a) Tail (~16-17% of runtime), cheap and safe:
- Precompute per-(n,c) scale = m*rsqrt(m^2*var+eps) and shift = -scale*mean ONCE (in the stats kernel's thread 0 or a tiny finalize step) so the apply loop is just fmaf(v,scale,shift) -> clamp -> *m -> max, instead of recomputing mean/var/rsqrt per element per channel.
- A 2D grid (blockIdx.y = n) in the apply kernel to avoid per-thread 64-bit div/mod.
- float4 loads in the stats kernel, guarded on S%4==0 plus 16-byte base alignment with a scalar fallback (str-00035 misaligned-address pitfall).
(b) Conv side (~83% of runtime), the large lever: a fused direct conv. IC=3, k=3, so the weight is only 16x81 floats and fits in smem/constant memory. Compute the conv output tile for all 16 OC, write it AND accumulate per-(n,c) sum/sumsq in the epilogue (block reduce + atomicAdd), so the separate stats read disappears; or go further and avoid writing/re-reading the conv output entirely. Verify indexing vs nn.Conv3d on a small shape first. Alternatively, channels_last_3d via your own transpose kernel (str-00102 pattern, timed per str-00051), with the tail reading NDHWC.
(c) Accumulate the variance with Welford or in double for robustness when the mean is large (E[x^2]-mean^2 in fp32 passed here).

General rule: a per-channel scale before an instance or batch norm can be folded into the norm statistics instead of being run as a separate pass.
[str-00023] helpful=1 harmful=0 :: Mamba2/SSD chunked-scan problems (KernelBench attention family, e.g., L3 Mamba2ReturnFinalState returning only new_states[:, -1], checked at tight atol=rtol=1e-4 with randn A, B, C).

(1) Prune dead work first. Trace which tensors reach the returned output. In the final-state variant, Y_diag, the L = exp(segsum) matrix and the C einsum never reach the output, so drop them. Only two pieces are needed: the intra-chunk states, sum_l B[l,n]*exp(Acum[-1]-Acum[l])*X[l,p], and the inter-chunk decay combination. This is a large, safe win.

(2) Reproduce the reference's numerics; do not 'improve' them. Exponents built from cumsums of randn reach about e^+-20, and B is mixed-sign. The cancellation makes the fp32 reference's own rounding error exceed 1e-4, so a more accurate kernel still reports 'output mismatch'. One attempt used fp64 accumulation, tl.sum chunk totals, an inter-chunk factor of total-consumed-cs_j, and a 3D broadcast-multiply reduction. Its algebra was correct, yet it failed 3 CHECKs and scored 0.

Match the reference instead:
- Use the same fp32 cumsum-then-difference order: decay = exp(Acum[-1] - Acum).
- Take the inter-chunk factor from the padded cumsum of chunk totals, as segsum does: ycum = cumsum(pad(Acum[..., -1], (1,0))), then factor_c = exp(ycum[-1] - ycum[1:]).
- Use matmul/dot-style accumulation like einsum/bmm, not separately summed totals.

(3) Build up safely. First write a pure-torch version of the pruned states-only path and confirm it passes CHECK. Then replace pieces with kernels one at a time, rechecking after each.

MEASURED WIN (final-state variant, b=2048, T=128, h=8, p=64, n=16, block_len=64): 3.58x (11.97 -> 3.35ms), correct at 1e-4, own_kernel_time_frac 0.039. Layout: torch cumsum + decay + the reference's own states einsum ('bclhn,bhcl,bclhp->bchpn'), torch padded-cumsum factors, then a tiny Triton kernel computing final = sum_c factor_c * states_c (grid = (b*h, cdiv(p*n, 256))). The turn-0 CHECK failed at 1e-4; turn 1 passed with this reference-ordered path. Nearly the whole gain is pruning. The attempt submitted right after the first passing CHECK with no timing, so treat ~3.6x as the BASELINE. Next lever (untested): fuse the states einsum (the dominant cost; B*decay*X reduced over l per chunk) into a custom kernel, and fold the inter-chunk weighted combine into it, since final = sum_c factor_c * states_c. Only change the accumulation order as far as 1e-4 tolerates: print the max abs/rel error against the reference (str-00024) and time it at full size before adopting.
[str-00024] helpful=1 harmful=0 :: Debugging output mismatches: when CHECK reports 'output mismatch beyond atol/rtol', do not resubmit a blind variant. Locally, print the max abs and max rel error against the reference, plus the reference output's magnitude range. Then compare intermediate stages (e.g., per-chunk states, decay factors, final combine) to find where they diverge. A small error relative to large output values points to numerics or op-order differences, so reproduce the reference's computation order. A large or structured error points to a logic, indexing or layout bug. Never SUBMIT a kernel that just failed CHECK when a simpler passing version, even pure torch with dead work removed, is available. A correct modest speedup beats a guaranteed 0.
[str-00025] helpful=0 harmful=1 :: Causal self-attention in fp32 (KernelBench attention family, e.g., minGPT L3_043 causal attention with B=128, nh=8, T=512, hs=96, eager ref ~30.9ms). This task checks at atol=rtol=1e-4, NOT ~1e-2. Read the tolerance first (str-00029).

Two hand-written Triton flash kernels have both lost heavily to eager cuBLAS:
- Strict FP32 (allow_tf32=False): correct but 0.11x (277 vs 31ms), custom kernel at 94% of runtime.
- tf32x3: correct but 0.2257x (136.8 vs 30.9ms), own_kernel_time_frac 0.88. It read q/k/v from c_attn via strides and wrote [B,T,C] directly. It still kept BD=128 padding for hs=96, fixed BLOCK 64x64 with 4 warps and 2 stages, and causal plus bounds masks on every tile. tf32x3 triples the MMAs in both QK^T and PV. It went CHECK then SUBMIT with a turn still left and no timing.
- Plain TF32 dots FAILED CHECK at 1e-4 on this task. An earlier version of this bullet claimed '~1e-2 tolerance accepts TF32', and that claim cost a CHECK. Do not use plain TF32 here unless a max-error probe passes.

Causes common to the losing flash kernels:
(1) No tensor cores (strict FP32), or 3x tensor-core work (tf32x3).
(2) hs=96 padded to BD=128, which wastes 25% of the work.
(3) Fixed BLOCK_M=BLOCK_N=64, num_warps=4, num_stages=2, with no autotune.
(4) Causal masking applied on every K tile, not just the diagonal tile.
(5) .contiguous() transposes in and out. This was fixed in the tf32x3 run but was not the bottleneck.

Preferred routes at 1e-4, in order. Time each at full size against the eager reference before SUBMIT.
(a) Keep strict-FP32 cuBLAS for both GEMMs: torch.baddbmm/bmm on strided views of the c_attn output for QK^T (scale folded into alpha) and for att@v. Replace scale + masked_fill + softmax with ONE in-place row kernel. That kernel reads only the causal prefix of each row, writes zeros above the diagonal, and uses accurate expf (no __expf). This saves about 2-3 full passes over the ~1GB att tensor (B*nh*T*T fp32). Untested.
(b) F.scaled_dot_product_attention(q, k, v, is_causal=True), if the task's custom-kernel rules allow it. Pair it with real custom work.
(c) If you stay with a flash kernel:
- Split hs=96 into 64+32 dot chunks instead of padding to 128.
- Run unmasked full tiles below the diagonal plus one masked diagonal tile.
- Autotune BLOCK_M/N up to 128 with 8 warps and 3-4 stages.
- Use tf32x3 at 1e-4, or plain TF32 only if a max-error probe passes.
- Expect to need all of these together. Submit only if it times above 1x.

If writing a custom kernel, also:
- Read q/k/v directly from the [B,T,3C] c_attn output via strides (head offset h*hs, q/k/v offsets 0/C/2C).
- Write the output directly in [B,T,C] layout.
- Keep c_attn and c_proj in cuBLAS.

Always time locally after the first passing CHECK. Both flash attempts used CHECK-then-SUBMIT flows that hid 4-9x slowdowns while turns remained.
[str-00026] helpful=5 harmful=0 :: Small-batch CNN networks (KernelBench network family, e.g., MobileNetV2/EfficientNet at batch 2-10, 224px; ref about 3.4-7ms spread over about 150 small kernels) are launch- and dispatch-bound. Anything that adds kernel launches or Python work per layer makes them SLOWER.

Measured failure 1: 0.65x (10.9 vs 7.1ms), 0.59x vs torch.compile, with the custom kernel only about 8% of runtime. The candidate had three overheads:
(a) a standalone Triton ReLU6 launch plus .contiguous() after every conv, about 35 sites, whereas the reference's in-place ReLU6 is already one cheap pass;
(b) the BN fold (sqrt/mul/sub plus weight scaling) recomputed on every forward for about 52 conv+BN pairs, adding several tiny kernels each;
(c) the module tree re-flattened in Python on every forward.
It was submitted right after CHECK with no local timing, so the slowdown went unnoticed.

Measured failure 2 (EfficientNetB0): 0.80x (7.01 vs 5.63ms), 0.76x vs torch.compile, with own_kernel_time_frac about 1%. It followed rules (1)/(2): BN folded once and cached, F.conv2d with folded bias, in-place hardtanh_, and only one custom global-avg-pool kernel. It still lost because it kept one cuDNN launch per conv plus a separate activation launch, so the launch count barely dropped. It also added Python per-block dispatch: cache-key data_ptr checks, dict lookups, and a non-inplace residual `x = x + identity`. Caching the BN fold alone is NOT a speedup lever here; it only avoids a slowdown.

Measured failure 3 (EfficientNetB2, batch 2, atol=rtol=1e-4): 0.80x (4.28 vs 3.44ms), 0.67x vs torch.compile, own_kernel_time_frac 0.0012. This is the network-family version of the 'cuDNN wrapper + trivial kernel' layout (see str-00018 and dense-conv failure mode B). It reliably lands at about 0.8x, because the extra module and launch overhead hurts a launch-bound reference. Never submit it. Earlier turns of that attempt had BN-fold+bias kernels that either never ran in eval (str-00031) or failed CHECK at 1e-4. They were abandoned without printing the max abs/rel error (str-00024) instead of being fixed with the exact fp32 fold from str-00038.

Measured failure 4 (MobileNetV2, batch 10, a repeat of failure 2): 0.7921x (8.78 vs 6.96ms), own_kernel_time_frac 0.0008. BN folded once lazily and cached as a per-block plan list, F.conv2d with the folded bias, then a standalone in-place grid-stride relu6 CUDA kernel after each of ~35 activation sites, and non-inplace `out = out + identity` residuals. The training branch also ran a meaningless relu6 on the pooled tensor. Turns 0-1 failed no_custom_kernel (kernel only reachable outside the checked eval path, str-00031); turn 2 passed; turn 3 submitted with no timing. Routing the kernel through the eval path only satisfies the check; it gives no speed. The str-00038 MobileNetV1 recipe (and its 'pass b' to F.conv2d + custom in-place activation' variant) does NOT transfer to MobileNetV2.

Read the reference module structure literally before fusing. Example: KernelBench's EfficientNetB2 MBConv puts the SE layers (AvgPool → 1x1 → ReLU → 1x1 → Sigmoid) sequentially in nn.Sequential. The SE output REPLACES the activation tensor rather than gating it, so after SE the block continues from an [N,C,1,1] tensor. Fuse to match what the code computes, not the textbook architecture.

Rules:
(1) Fold BN into conv weight and bias ONCE and cache the result. Do this lazily on the first eval-mode forward, recompute only if the parameters change, and guard on not self.training, since training-mode BN cannot be folded. Store folded tensors as plain attributes or buffers, not dict entries, and keep per-forward Python checks minimal (or fold once at first call with no per-call key check).
(2) Never add a separate activation kernel. Fuse it into the conv epilogue, or keep the reference's in-place clamp_. Keep residual adds in place (add_).
(3) Precompute the flat layer and fused-op plan in __init__, not in forward.
(4) For a real win, cut the launch count. These are the only levers that beat eager here:
  (a) Capture the whole eval forward in a CUDA graph with static input and output buffers, with the custom kernel inside the graphed path.
  (b) Write custom fused kernels per inverted-residual block, e.g. depthwise+bias+ReLU6 (str-00059 depthwise layout), ideally fused with expand or project, so hidden tensors are not round-tripped. Where SE squeezes to 1x1, fuse the squeeze reduction and the tiny SE matmuls.
  (c) Fuse the tail: the final 1x1 conv + bias + ReLU(6) + global avg pool in one reduction kernel, so the large 1280/1408xHxW tensor is never written.
(5) Time locally against the reference right after the first passing CHECK, before SUBMIT. A sub-1x result should trigger a switch to (4), not a submit.
[str-00027] helpful=3 harmful=1 :: Dense standard conv with large batch and/or many output channels (KernelBench conv family, e.g., AlexNet conv1: Conv2d 3->96, 11x11, stride 4, pad 2, N=256, 224px; or Conv2d 64->128, 5x7 asymmetric kernel, N=8, 512x256 input, bias=False; or Conv2d 32->64, 5x9 kernel, N=8, 512x512 input, possibly with dilation e.g. (2,3) and padding (2,4)). The same failure mode (B) also applies to pure transposed convs (conv_transpose family); see the strided/grouped ConvTranspose bullet. Four failure modes are known.

(A) A naive direct-conv CUDA kernel is hopeless against cuDNN. One attempt was correct but scored 0.17x (31.0 vs 5.33ms), with the custom kernel at 100% of runtime. Its flaws: (1) one block per (n, oc), giving N*OC = 24576 blocks with gridDim.x capped at 8; (2) only that oc's 363 weights cached in shared memory; (3) each thread re-read its full IC*KS*KS input window from global memory for every oc, so input reuse across the 96 output channels was zero; (4) bounds-check branches inside the unrolled inner loop; (5) plain FP32 FMA with no tensor cores.

(B) Wrapping cuDNN F.conv2d (or F.conv_transpose3d) with a trivial custom kernel (a no-op launch, a 1-thread 'touch' adding 0 to out[0], or a bias-add when bias=False makes it unused) gives parity at best. Measured: 0.998x (31.50 vs 31.43ms, own_kernel_time_frac=0.0) on Conv2d, and 0.997x (4.40 vs 4.39ms, own_kernel_time_frac=0.0006) on a grouped stride-2 ConvTranspose3d. Both attempts counted as failed. It is also suspect as a 'custom kernel' solution. There is no epilogue to fuse in a pure conv, so the only lever is the conv itself. Decide on turn 0 to change it; this warning has been ignored twice.

(C) A Triton implicit GEMM at near-FP32 precision over an NCHW gather is much slower than cuDNN, CONFIRMED TWICE on Conv2d 32->64, 5x9, N=8, 512x512 (K=IC*KH*KW=1440, M~2M):
- Run 1: 0.29x (35.0 vs 10.0ms, own kernel 100%). Design: tl.dot(input_precision='tf32x3'), NCHW input, K ordered (ic,kh,kw) and decoded with per-element div/mod every K step, runtime (non-constexpr) KH/KW, scattered A-tile gathers with K not the contiguous dim, and the weight tile reloaded per M tile.
- Run 2: 0.3727x (26.8 vs 9.99ms, own kernel 100%). It FIXED only the div/mod part: constexpr KH/KW with tl.static_range kh/kw loops and an inner ic block (BK=32, autotuned BM 64-256, BN 64-128). It kept tf32x3, the NCHW gather (ic strided by H*W, so the [BM,BK] A tile is uncoalesced) and the weight read in place from [OC,IC,KH,KW] (B tile strided by KH*KW). With C=32 the ic loop is a single BK step, so each of the 45 taps issues one tiny dot fed by poorly coalesced loads.
Lesson: removing div/mod is NOT enough; the gather layout and tf32x3 are the dominant costs. tf32x3 triples the tensor-core work; it only paid off on dense GEMMs with contiguous loads, not on this gather-bound conv. Both runs were submitted after correctness-only CHECKs with no timing (run 2: CHECK, CHECK, SUBMIT).

(D) CATASTROPHIC: cuDNN in channels_last for a DILATED fp32 conv. Conv2d 32->64, 5x9, dilation (2,3), padding (2,4), N=8, 512x512: a custom NCHW->NHWC transpose kernel plus F.conv2d with a channels_last weight was correct but scored 0.0157x (1285 vs 20.1ms). The transpose kernel was only 0.4% of runtime, so cuDNN itself fell back to a very slow dilated NHWC algorithm (~64x slower). It went CHECK then SUBMIT with no timing; one timing run would have exposed it.

Options that can compete:
(a) Implicit GEMM / im2col with M = N*OH*OW, N_gemm = OC, K = IC*KH*KW (e.g., 363; pad K to a BK multiple, or peel and mask the tail). Gather the patch indices on the fly rather than materializing im2col. Each output tile should span many or all OC (e.g., all 64 or 128, BM 128-256), so every gathered input patch is reused across all output channels. For groups>1, run one GEMM per group (program axis over groups). Make the gather cheap, per failure (C); ALL of these are needed together, not just the first one you try:
  - Use plain TF32 (allow_tf32=True / input_precision='tf32') if a max-error probe passes the tolerance (~1e-2); avoid tf32x3 here unless timed.
  - Inside YOUR OWN kernel (or a custom pre-pass), convert input to NHWC and order K as (kh,kw,ic) so a row of ic is contiguous per pixel (this is about the custom gather, not about handing NHWC to cuDNN; see (D)). Dilation only changes the tap offsets (kh*dh, kw*dw).
  - Pass KH/KW (and dilation) as tl.constexpr and loop kh/kw as static outer loops with an inner ic block, so there is no per-element div/mod. Or precompute k->offset tables.
  - Pre-permute and cache the weight to match the K order, e.g. [KH,KW,IC,OC] so the B tile is OC-contiguous (key on weight.data_ptr(), keep the conv kernel launching every forward).
  - When IC is small (e.g., 32), a single BK step per tap gives tiny dots; consider grouping several taps per K step (BK spanning kw*ic) to raise arithmetic intensity.
(b) A register-tiled direct conv: each thread computes several oc x several output pixels, the input tile plus halo is staged in shared memory once per block, and padding is handled at staging time (zero-fill), not in the inner loop.
(c) Cheap cuDNN-side levers: enable TF32 for the conv (torch.backends.cudnn.allow_tf32=True around the conv call, within tolerance), keeping NCHW. channels_last / channels_last_3d is NOT a free lever: for dilated convs it measured 64x slower (failure D), so only try it for non-dilated convs and only with a local timing. On their own these leave own-kernel time at 0, so pair them with real custom work if the task requires your own kernels. If every custom variant stays below cuDNN in local timing, fall back to cuDNN TF32 (NCHW) plus a minimal custom piece rather than submitting a sub-1x kernel.

Set a realistic target. cuDNN is already near-optimal in strict FP32, so beating it needs TF32 tensor cores or a cheaper algorithm. Time locally against the reference module before SUBMIT, with bounded warmup and iteration counts. All measured failures used CHECK-then-SUBMIT flows with no timing: they hid a 64x slowdown, a 6x slowdown, two ~3x slowdowns and parity results while turns remained. Spend at least one turn after the first passing CHECK on a full-size timing vs nn.Conv2d.
[str-00028] helpful=1 harmful=0 :: LayerNorm/norm reductions with FEW, very LONG rows (KernelBench norm family, e.g., L1_040: nn.LayerNorm over (64,256,256) with batch 16, so N=16 rows of D≈4.2M each). Eager nn.LayerNorm parallelizes per row, so only a handful of blocks each reduce millions of elements. It is badly underutilized: about 10ms here.

Fix: split each row across many blocks.
- Use a 2D grid with blockIdx.y=row and blockIdx.x as a grid-stride chunk index, about 512 blocks x 256 threads per row.
- Pass 1 (stats kernel): each block accumulates partial sum and sumsq in double, reduces them in shared memory, then thread 0 atomicAdds (double) into per-row accumulators allocated with torch.zeros.
- Pass 2 (apply kernel, same grid): each thread recomputes mean and rstd from the accumulators. Compute var = E[x^2]-mean^2 in double, rstd = 1/sqrt(var+eps). Then write (x-mean)*rstd*gamma[i]+beta[i], with gamma/beta indexed by the in-row offset i and coalesced grid-stride access.
- Use long sizes/indices (str-00058) and the str-00001 load_inline layout; pass ln.weight/ln.bias flattened (.contiguous().view(-1)) and ln.eps.

Measured, CONFIRMED TWICE on L1_040:
- Run 1: 8.99x (10.17 -> 1.13ms), own kernels 99.6%.
- Run 2: 9.252x (10.17 -> 1.10ms, lo95 9.238), 1.108x vs torch.compile, own_kernel_time_frac 0.994. The recipe was copied nearly verbatim as a turn-0 SUBMIT with no CHECK or timing; it worked only because the recipe was proven (str-00002 still applies to anything novel).
Treat ~9.0-9.25x as the SOLVED BASELINE: resubmitting it cannot raise the score. The pass is near memory-bound (2 reads + 1 write of ~268MB), so expect at most ~10-20% more. Double accumulation keeps the E[x^2]-mean^2 form accurate over millions of elements.

Rule: check the rows x row-length shape first. Few long rows call for a multi-block-per-row reduction (atomics or a two-stage partials buffer). Many short rows (e.g., transformer LN with C=96..768) call for several rows per block/program. Eager LN is already efficient for short rows, so swapping it there is a slowdown; see the large transformer networks bullet.

Next attempt: CHECK on turn 0, then time these headroom levers at full size against the ~1.10ms baseline before SUBMIT, keeping the fastest (all untested):
- float4 vectorized loads/stores for x, out, gamma and beta (D=4194304 is divisible by 4 and rows are 16B-aligned; guard D%4==0 with a scalar fallback otherwise).
- fp32 per-thread accumulation in the stats pass with only the block combine and atomics in double (str-00043 found fp32 FMAs much cheaper than double FMAs); print max error vs reference.
- Compute mean/rstd once per block (thread 0 into shared memory, one __syncthreads) instead of each thread recomputing the double math.
- Remove the 2 small torch.zeros launches: have the apply kernel re-zero the accumulators for the next call, or replace zeros+atomics with a per-block partials buffer plus a finalize step.
[str-00029] helpful=27 harmful=0 :: Tolerance pitfall: several bullets assume KernelBench tolerance is about 1e-2. That is NOT universal. Some matmul-family, network-family and conv-family tasks check at atol=rtol=1e-4, e.g., the L3 MLP with K=N=16384, or a Conv2d → elementwise → MaxPool → Mish task. The CHECK error states the value, e.g. 'output mismatch beyond atol=rtol=0.0001'. Read the actual tolerance from the first CHECK failure, or test locally against the reference at 1e-4, BEFORE choosing a precision path.

At 1e-4:
- TF32 tl.dot or allow_tf32 cuBLAS over long K (e.g., 8192 or 16384) may fail. This was confirmed failing on the MLP.
- Algebraic folds that reorder long summations (W.sum(0), pooled weights) may also fail.
- Mitigate by building folded weights in fp64 and accumulating dot products in fp64 or pairwise.
- Even with a cuDNN conv plus only an elementwise tail, write each activation the way PyTorch does. Examples: hardswish as y*min(max(y+3,0),6)/6, and mish as x*tanh(log1p(exp(x))) with a threshold of 20. Avoid fast intrinsics (__expf, --use_fast_math) and reassociated or algebraically rewritten forms unless you have verified them locally at 1e-4. One such attempt failed its first CHECK at 1e-4 for an unrecorded reason and passed once the formulas matched PyTorch exactly.

For a dense GEMM at 1e-4 (or unknown tolerance) with no fold available, options in order:
(a) tl.dot(..., input_precision='tf32x3'). This 3xTF32 split emulation runs on tensor cores at near-FP32 accuracy. It is now confirmed end to end on a dense K=8192 GEMM: correct and 1.59x over cuBLAS fp32 (see the large dense fp32 GEMM bullet). That task's exact tolerance was not probed, so still verify the max error with CHECK at 1e-4 tasks and time it locally.
(b) Keep cuBLAS in strict FP32 and fuse what you can around it, e.g., torch.addmm with the bias folded in, plus an in-place ReLU. Expect about 1.0x, never below 1.
Never submit strict-FP32 tl.dot(allow_tf32=False); it measured 0.60x.

If a fold or precision change fails CHECK, print the max abs and max rel error and the output magnitude (see the debugging-mismatches bullet) and fix the accumulation. When a fold or a tensor-core path is still possible, do not retreat to the full cuBLAS GEMM, because that layout reliably gives only ~1.0x. However, a ~1.0x cuBLAS fallback beats submitting a sub-1x custom GEMM.
[str-00030] helpful=1 harmful=1 :: Strided (stride>1) and/or grouped and/or DILATED ConvTranspose2d/3d with no epilogue (KernelBench conv_transpose family, e.g., ConvTranspose3d 32->32, k=(3,5,7), stride 2, padding (1,2,3), output_padding 1, groups 4; or ConvTranspose3d 32->64, k=3, stride 2, padding 1, dilation 2). The stride-1 forward-conv rewrite (str-00009) does not apply directly, and wrapping cuDNN with a no-op kernel gives ~0.997x = failure (see failure mode B of the dense conv bullet).

Measured failure of the naive alternative: a gather-style direct kernel scored 0.0435x (199.8 vs 8.68ms, own kernel 100% of runtime), about 23x slower than cuDNN. Its design: one thread per output pixel, a float acc[MAX_OC=128] register array (spills to local memory), per-tap `% stride` branches that skip most taps, weight reads strided by K^3 (uncoalesced, no shared-memory staging), and scalar FP32 FMAs with no tensor cores. Do not write this kernel.

Sub-pixel / phase decomposition: a WIN ONLY in the dilation-shortcut case (empty phases). For non-shortcut shapes (every phase populated) it is a MEASURED LOSS; see below. Per spatial dim, output o = i*s - p + k*d (d = dilation). For each output phase r in [0,s) (s^D phases, 8 for 3D stride 2), out[s*q + r] only receives taps k with (r + p - k*d) % s == 0, reading x[(s*q + r + p - k*d)/s]. So each phase is an independent stride-1 FORWARD conv on the un-dilated input with a sub-sampled, spatially flipped weight slice, channel dims swapped per group: [IC, OC/g, ...] -> [OC, IC/g, ...] with the same groups. This avoids the s^D-fold wasted multiplies of the zero-inserted form.

Dilation shortcut: when d is a multiple of s (e.g., s=2, d=2), k*d % s == 0 for every k, so tap validity depends only on (r+p) % s. A phase then gets either ALL K taps per dim, or NO taps at all. Phases with no taps are just zeros (or bias). Example: with s=2, d=2, p=1, only the r=1 phase in each dim is non-zero, so 1 of 8 phases in 3D carries all the work.
MEASURED WIN (shortcut): ConvTranspose3d 32->64, k=3, s=2, p=1, d=2, input 16x32x16x32x32: 2.17x vs eager (8.61 -> 3.97ms) and 2.21x vs torch.compile, custom flip kernel only 0.1% of runtime. It reproduced on CHECK-then-SUBMIT. Recipe that worked:
- Per forward, run the str-00009 flip/swap kernel to get w_eq [OC,IC,K,K,K].
- For each valid phase r (with (r+p)%s==0, per dim): F.conv3d(x, w_eq, None, stride=1, dilation=d/s, padding=(K-1)*(d/s) - (r+p)/s per dim), crop to len(range(r, Out, s)) per dim, and assign into out[..., r::s, r::s, r::s] of a torch.zeros output; add bias afterwards if present.
- Out = (In-1)*s - 2p + d*(K-1) + 1 per dim.
- Fall back to the reference module when d % s != 0 or any computed padding is negative (or groups/output_padding non-default, unless handled).
The gain comes from replacing cuDNN's backward-data path (mostly wasted work on empty phases) with one dense forward conv over 1/s^D of the output.
Untested headroom for the shortcut: avoid the full-size torch.zeros write plus the strided scatter copy (e.g., one custom kernel that writes the conv result into the strided slice and zero-fills/bias-fills the rest in a single pass, or torch.empty + a fill kernel for empty phases only); cuDNN TF32 within tolerance; cudnn.benchmark. Time these locally.

MEASURED LOSS (non-shortcut, all 8 phases populated): ConvTranspose3d 32->32, k=(3,5,7), s=2, p=(1,2,3), output_padding 1, groups 4, input 8x32x12x24x48, atol=rtol=1e-4: correct but 0.757x (5.77 vs 4.36ms), own_kernel_time_frac 0.0007. Layout: flip/swap kernel + 8 separate F.conv3d calls (one per phase, each with its own F.pad copy) + torch.zeros + 8 strided scatter copies. That is effectively a cuDNN wrapper doing the same total MACs as cuDNN's single backward-data call plus 8x launch/pad/scatter overhead. It went CHECK then immediate SUBMIT with no timing. Do NOT use the per-phase F.conv3d decomposition for non-shortcut shapes; at best it is a hypothesis that must time above 1x at full size. Better candidates for this shape (all untested; time each vs nn.ConvTranspose3d): direct F.conv_transpose3d with cudnn.benchmark; channels_last_3d (non-dilated, str-00051); TF32 only if tolerance allows (at 1e-4 it likely fails); or ONE custom TF32/tf32x3 implicit-GEMM kernel writing all phases in a single launch (per group, M = N*OD*OH*OW per phase, N_gemm = OC/g spanning all OC per tile, K = IC/g*valid_taps).

Correctness details that worked for the general (non-shortcut) phase math, if reused inside a single kernel: in flipped-weight coordinates, tap fk is valid for phase r iff (r - K + 1 + P + fk) % S == 0, at input offset (r - K + 1 + P + fk)//S (Python floor division). Output_padding makes some phases one element longer: derive each phase's length as len(range(r, Out, s)) with Out = (In-1)*s - 2p + d*(K-1) + op + 1, and compute left/right input padding independently per side from the first/last offsets. A 47-vs-48 size mismatch and two 1e-4 mismatches came from exactly these boundary details; validate on small shapes against nn.ConvTranspose3d first.

General implementation notes (non-shortcut case):
- Write each phase into its strided slice out[..., r::s] directly (custom kernel or strided store) rather than concatenating.

Verify against nn.ConvTranspose locally on small shapes first (indexing is error-prone). Then TIME locally against nn.ConvTranspose3d before SUBMIT. The 0.0435x and 0.757x attempts both used CHECKs for correctness only and never timed.
[str-00031] helpful=19 harmful=1 :: no_custom_kernel CHECK pitfall: do NOT assume which mode the evaluator's forward check and timer use. Earlier network tasks were observed to run in eval mode. On L3 EfficientNetMBConv, however, the module was run in train() mode, which is the default for a freshly constructed nn.Module. There, custom kernels placed only in the eval branch (with BN folded) failed with 'none of the candidate's kernels [...] ran on the GPU during forward', and the TIMED path was the training branch. Diagnostic: if an eval-only kernel fails no_custom_kernel and a training-branch kernel then passes, the evaluator runs train(). BN cannot then be folded, and the training branch is what gets scored. Rules:
(1) Put REAL work in BOTH branches. The custom kernels must replace eager passes in each mode, not merely satisfy the check.
(2) Never satisfy the check with a dummy kernel, e.g. x = relu6(x.clone()) * 0.0 + x. It adds clone, mul and add passes over a large tensor and guaranteed a slowdown (0.87x, own_kernel_time_frac 0.029).
(3) A low own_kernel_time_frac combined with a sub-1x score signals that the timed branch is not the one holding your real kernels.
(4) Time locally at full size in the mode that actually runs (check both train() and eval()) before SUBMIT.
Torch-only paths such as F.adaptive_avg_pool2d do not count as a custom kernel. Caching parameter-derived tensors (e.g., a BN fold) inside the eval branch is fine, as long as the custom kernel itself still launches on every forward (confirmed on a DenseNet dense block). The failure arises only when the cached step is the sole custom kernel (see str-00009's caveat).
[str-00032] helpful=4 harmful=1 :: Proven recipe for a large fp32 GEMM + elementwise epilogue (KernelBench matmul family, e.g., Linear → /d → erf-GELU at 1024x8192 @ 8192x8192). This is the measured winner for rule (2b) of the GEMM+epilogue bullet. One fused Triton kernel scored 5.5x (7.98 → 1.45 ms), with its own kernel at 100% of runtime.

Kernel setup (the 5.5x FIXED config):
- tl.dot(a, b, allow_tf32=True); correctness held within tolerance at K=8192 (confirmed on two runs).
- 1D pid with grouped tile ordering (GROUP_M=8) for L2 reuse.
- BLOCK_M=BLOCK_N=128, BLOCK_K=64, num_warps=8, num_stages=3, with an fp32 accumulator.
- Read the nn.Linear weight as W.t() by passing its strides (stride_bk=W.stride(1), stride_bn=W.stride(0)). No transpose copy, and never rescale the weight. With PLAIN TF32 this swapped-stride NT read showed no slowdown; the NT-layout penalty seen elsewhere is tied to tf32x3 (str-00075).
- Epilogue on the accumulator in registers: + bias[offs_n], then * a precomputed scalar 1/d passed as a kernel arg, then the activation, e.g. 0.5*x*(1+tl.erf(x*0.70710678)) (top-level tl.erf compiles, str-00012). Finally one masked store, so the output is written once.

Result: this beat cuBLAS F.linear + standalone epilogue (always ~1.0x) by more than 5x, because the TF32 tensor cores speed up the dominant GEMM.

AUTOTUNE IS NOT AUTOMATICALLY BETTER (measured): on the SAME task, an autotuned version (configs BM/BN in {128,256}, BK {32,64}, num_warps=8, stages 3-4, GROUP_M=8, masked K loop every iteration) was correct but scored only 4.415x (8.01 → 1.81ms) vs 5.5x for the fixed 128x128x64 config: ~20% slower, likely from autotune picking 256-wide tiles and/or the masked K loop. It went CHECK-then-SUBMIT with no local timing, so the shortfall was never seen. Default to the fixed 128x128x64/8 warps/3 stages config; if autotuning, restrict to BM=BN=128-class configs and verify locally it is not slower than the fixed one.

Untried headroom (time each locally against the fixed config and F.linear before adopting):
- An unmasked main K loop when K % BK == 0 (8192 % 64 == 0), with the tail peeled otherwise.
- Small autotune subsets around 128x128 (BK 32/64, stages 3-4).

The ~4.4-5.5x result is a known baseline for this task; re-submitting it unchanged cannot improve the score. Still do a bounded local timing run against the reference before SUBMIT, and compare to the known 1.45ms. Also confirm the task tolerance is ~1e-2, not 1e-4 (see the tolerance pitfall bullet), before relying on TF32.
[str-00033] helpful=0 harmful=0 :: Linear (large GEMM, e.g., 1024x8192 @ 8192x8192 fp32) followed by a per-row normalization and an elementwise tail (KernelBench norm family; e.g., Linear → InstanceNorm2d on x.unsqueeze(1).unsqueeze(1) → +y → *y).

Shape trick first: InstanceNorm2d on a (B,1,1,N) tensor normalizes over all N features of each row. It is a per-row LayerNorm with no affine and biased variance; momentum and running stats are unused. The Linear bias is per-column, so it does NOT cancel under per-row normalization. Keep it (contrast with per-channel batch-norm mean cancellation in the BN-subtract bullet).

Measured low-gain layout: cuBLAS F.linear, then ONE Triton kernel with one program per row (BLOCK=1024, num_warps=8). Loop 1 accumulates sum and sumsq; mean = sum/N and var = E[x^2]-mean^2 in fp32 was accurate enough at N=8192 with a mean near 0. Loop 2 computes (x-mean)*rstd, applies the tail (norm+y)*y, and does a single store. It was correct at atol=rtol=1e-4 but scored 1.02x (8.17 vs 8.00ms), with own_kernel_time_frac 0.0093. The GEMM is about 99%, so about 1.02x is the ceiling for this layout.

Precision caveat: this task checks at 1e-4, so a plain TF32 GEMM rewrite (GEMM+epilogue rule 2, proven-recipe bullet) is likely to FAIL. Dividing by the row std preserves the GEMM's relative error (about 1e-3 for TF32 at K=8192), so normalization does not hide it. One CHECK in this attempt failed at 1e-4, plausibly from such an experiment.

Candidate for a real gain (untested): a Triton GEMM with 3xTF32 split precision, via tl.dot input_precision='tf32x3' or manual hi/lo splits (a_hi*b_hi + a_hi*b_lo + a_lo*b_hi). Its epilogue adds the bias and emits per-row partial sum and sumsq (atomics or a partials buffer per N-tile). Then run one light apply pass for normalize + tail. Verify the max error at 1e-4 locally and time it against the reference before relying on it; otherwise submit the cuBLAS + fused row-kernel fallback.
[str-00034] helpful=3 harmful=2 :: Linear (large GEMM, e.g., 1024x8192 @ 8192x8192 fp32) followed by GroupNorm with SMALL groups (e.g., 512 groups of 16 channels, or 256 groups of 32 channels) and a tail or prologue, KernelBench matmul family, checked at atol=rtol=1e-4. The tail can be elementwise (e.g., LeakyReLU, then x+x) or a reduction (e.g., min over all features with keepdim, then + bias). The chain before GN can also be elementwise (e.g., +bias → hardtanh → Mish).

Measured low-gain layouts. Both cap at about 1.0x, because the GEMM is about 94-95% of runtime.
- Elementwise tail: cuBLAS F.linear, then one Triton kernel with one program per (row, group). Each program loads 16 elements, computes mean and biased var, applies the affine, LeakyReLU and doubling, and stores once. It was correct but scored 1.0116x (8.48 vs 8.39ms), with own_kernel_time_frac 0.06. It also launched about 524k tiny programs; processing several groups or whole rows per program would trim that.
- GN → min over features → +bias tail: cuBLAS F.linear, then one program per row that loops serially over all G=512 groups of 16 elements, keeping a running min. It scored 1.0046x (8.44 vs 8.40ms) and 0.96x vs torch.compile, with own_kernel_time_frac 0.055. The serial group loop is a poor layout. Instead, load a 2D tile [groups, CPG] per program, reduce over axis 1 for per-group mean and var, then take a min over the whole tile. Or process several rows per program.

Measured FAILURE of the tf32x3 GEMM rewrite (Linear → +bias → hardtanh → Mish → GN(256), 1024x8192x8192): an autotuned tf32x3 Triton GEMM with the combined bias, hardtanh and Mish fused into the epilogue, then a SEPARATE GN kernel with one program per (row, group) (262k programs of 32 elements). It was correct but scored 0.85x (9.73 vs 8.29ms), with own kernels at 99.97% of runtime. It was submitted on turn 1 right after CHECK with no timing. The GN pass is only ~1ms, so the custom GEMM itself ran ~8.5ms+, far from the ~5.1ms the same recipe reached on equal-FLOP NN/TN GEMMs. Likely causes: the nn.Linear weight read as W.t() through swapped strides gives a [BK,BN] B tile with K contiguous (the non-native orientation, see str-00004's layout note); BM=256/BN=256 tf32x3 configs causing register pressure; and a masked K loop every iteration. Do NOT assume the tf32x3 recipe wins on the nn.Linear NT layout at M=1024. Levers, timing each (GEMM alone vs F.linear first):
(a) load W tiles as [BN,BK] (K contiguous) and use tl.dot(a, tl.trans(b)) instead of a swapped-stride [BK,BN] tile;
(b) restrict autotune to BM=BN=128 configs and run the K loop unmasked when K % BK == 0;
(c) plain tf32 only if a local max-error check passes 1e-4;
(d) fuse GN into the epilogue (below), removing the 32MB intermediate and the huge GN launch.
If every variant times below 1x, fall back to cuBLAS + one fused act+GN row kernel (~1.0x).

Shape pitfall for the min tail: the min with keepdim gives shape [M,1]. Adding bias of shape (1,N,1,1) broadcasts the result to (1,N,M,1). Reproduce this exactly with mins.view(M,1) + self.bias, which passed CHECK.

Precision: turn-0 attempts in both early tasks failed CHECK at 1e-4. Plain TF32 cannot be assumed to pass over K=8192 here. One attempt retreated without printing the max abs or max rel error; print them before abandoning a precision path. tf32x3 passed CHECK in the hardtanh/Mish task.

Candidate for a real gain (untested): one Triton GEMM using 3xTF32 split precision, via input_precision='tf32x3' or manual splits a=a_hi+a_lo, b=b_hi+b_lo, acc += a_hi·b_hi + a_hi·b_lo + a_lo·b_hi. Fuse GroupNorm entirely into the epilogue by choosing BLOCK_N as a multiple of channels-per-group, so each tile holds whole groups. Reshape the acc tile to [BM, BN/CPG, CPG] (tl.reshape with constexpr-folded shapes, str-00070) and reduce over the last axis for per-group mean and var. No cross-tile reduction is needed for the normalization itself (unlike per-row LayerNorm; see the Linear + per-row normalization bullet).
- Elementwise tail: fold x+x into a 2x scale after the activation, because LeakyReLU commutes with positive scaling, or fold the 2 into gamma and beta.
- Min tail: reduce each tile to a per-row min. Combine the mins across N-tiles with tl.atomic_min into an inf-initialized buffer, or write a small [M, N/BN] partials buffer and run a tiny final min. The 32MB GEMM output is then never written.

Verify the max error at 1e-4 and time it locally before SUBMIT. Otherwise submit the cuBLAS + fused GN kernel fallback.
[str-00035] helpful=2 harmful=0 :: Conv → elementwise activation (e.g., HardSwish) → GroupNorm → global spatial mean (KernelBench conv family, e.g., L2_027: Conv3d → hardswish → GroupNorm(G) → x.mean(dim=[2,3,4])). Never materialize the activation or the normalized tensor.

Algebra: the spatial mean of GN(a) for channel c is gamma_c*rstd_g*(mean_nc(a) - mean_ng) + beta_c. Here mean_ng is the average of the group's channel means (equal S per channel). var_ng = E[a^2] - mean_ng^2 over the group's CPG*S elements (biased). This is the GroupNorm analogue of the BN+global-avgpool bullet.

Kernel layout:
(1) One reduction kernel with one 256-thread block per (n,c). It applies the activation in registers (PyTorch-exact hardswish x*fminf(fmaxf(x+3,0),6)/6) and accumulates sum and sumsq of the ACTIVATED values (fp32 per thread, double shared-memory tree).
(2) A tiny finalize kernel with one thread per (n,c). It sums its group's CPG channel stats in double, computes mean and rstd (rsqrt(var+eps)), and writes gamma*rstd*(csum/S - mean_g) + beta to a [N,C] output.
Use the str-00001 load_inline layout and 64-bit sizes (str-00058).

The tail becomes one read of the conv output and no full-tensor writes. Measured, CONFIRMED TWICE on L2_027: 1.157x and 1.158x vs eager (20.09 → 17.35ms), own kernels only ~3.5-4% of runtime, but only ~0.97x vs torch.compile, which finds the same fusion. ~1.16x is the SOLVED BASELINE for this layout: resubmitting it cannot raise the score.

MISALIGNED-ADDRESS PITFALL (measured crash on L2_027): the conv output slab is S = 13*29*29 = 10933, which is odd, so a float4 reinterpret over per-(n,c) slab bases (base = x + bid*S) crashes with 'CUDA error: misaligned address'. Use scalar loads (the tail is only ~3.5% of runtime, so float4 is not worth the risk), or guard float4 behind S%4==0 plus a 16-byte pointer-alignment check with a scalar fallback (str-00068 risk (1), str-00059 pitfall).

Next attempt on L2_027: paste the scalar-load recipe and CHECK it on turn 0. Because a CUDA error occurred in a prior CHECK of this task, also run the full-size forward loop of str-00082 before SUBMIT. Then spend every remaining turn timing conv-side levers at full size against the reference (the conv is ~96% of runtime) and keep the fastest:
- cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv, within tolerance; print max error first, str-00029).
- torch.backends.cudnn.benchmark=True.
- channels_last_3d (timed per str-00051). If you use it, make the reduction read NDHWC coalesced, with threads spanning C or a 2D tile.
- Fusing the per-(n,c) stats into a conv epilogue so the conv output is never written.
The last run went CHECK, CHECK, SUBMIT with no timing and no conv lever, leaving all of this headroom untried.

The same collapse applies to any per-group or per-channel affine norm followed by global average pooling. Only per-(n,c) sum/sumsq of the pre-norm values are needed. A preceding elementwise op of any kind is folded into the reduction load.
[str-00036] helpful=4 harmful=0 :: Conv2d → min over channels (keepdim), with an elementwise op before or after it (KernelBench conv family). Examples: 16->64 ch, 3x3, N=128, 256px input → min → tanh → tanh (L2_025); or Conv2d 64->128, 3x3, N=64, 256px → *scalar scale → min (L2_032). The channel min already shrinks the data by a factor of OC (64-128x here), so the eager tail after it is cheap.

The obvious layout: cuDNN conv, then ONE fused kernel with one thread per (n,h,w), spatial index fastest, and a coalesced register loop over the strided c. Take the min first, apply the elementwise chain once per pixel, and write [N,1,H,W]. Put the wrapper and stream retrieval in cuda_sources (str-00001).

Measured:
- min → tanh → tanh (L2_025), CONFIRMED TWICE: run 1 1.011x (19.21 vs 18.99ms) and 1.09x vs torch.compile, own_kernel_time_frac 0.085. Run 2 1.0119x (19.26 → 19.03ms, lo95 1.0113), own_kernel_time_frac 0.076: grid-stride min+tanhf+tanhf kernel (256 threads, 4096-block cap, long indices per str-00058, str-00001 layout, #pragma unroll 4 over c). Run 2 was a turn-0 SUBMIT with no CHECK or timing, correct only because the recipe was copied, and it tried no conv lever. ~1.01x is the SOLVED BASELINE for L2_025: the tail it replaces is already tiny and the conv (~92%) writes a ~2GB output that is re-read, so resubmitting this layout or tweaking the tail kernel cannot move the score. Next attempt on L2_025: CHECK on turn 0, then time levers (b) at full size (print TF32 max error first), and aim for (a): with IC=16, k=3 the weight is only 64x144 floats and fits in smem, so a fused conv with a running-min epilogue is feasible.
- scale → min, CONFIRMED TWICE: run 1 1.11x (31.26 → 28.13ms), 1.12x vs torch.compile, own kernel 5.5% (scale applied in the min kernel, no negative-scale guard). Run 2 1.1202x (31.33 → 27.97ms, lo95 1.117), own_kernel_time_frac 0.056: scale folded into conv weight/bias (cached in self.<x> keyed on weight.data_ptr(), str-00014/str-00005), F.conv2d, then a grid-stride channel-min kernel (256 threads, 4096-block cap, long indices per str-00058, str-00001 layout; compiled first try), with an eager fallback for scale < 0. The min kernel still launched every forward, so the custom-kernel check passed. Folding the scale into the weights vs applying it in the min kernel is equivalent in speed: removing the eager scale pass is the only saving either way. The gain is larger than the tanh case because the reference runs an eager scale pass, a full read+write over the ~2GB conv output, BEFORE the min (str-00006 principle: the gain grows with the number of collapsed eager passes before the reduction).

~1.12x is the SOLVED BASELINE for the scale → min task: resubmitting this layout or tweaking the min kernel (~5.6% of runtime) cannot raise the score. Both scale→min runs went CHECK (or nothing) then SUBMIT with no timing. In all cases the conv is about 91-94% of runtime and writes a ~2GB output that the reduction then re-reads, so this layout caps near ~1.0-1.1x. A good reduction layout (str-00021) avoids a slowdown but cannot produce a bigger gain.

Sign and order rules:
- Monotone non-decreasing ops (tanh, multiplying by a POSITIVE scalar) commute with min, so compute the min first and apply the op once per pixel.
- A scalar multiply applied after the min is valid ONLY for scale > 0. For scale < 0, min becomes max: take the channel max and then scale, branching on the sign at runtime, or fold the scale into the conv weight and bias (exact for a linear conv, any sign), or fall back to the eager path for s < 0 (done in run 2).

Decide on turn 0 to attack the conv, and spend the spare turns after the first passing CHECK timing these at full size against the reference before SUBMIT:
(a) Write a fused direct or implicit-GEMM conv (TF32 tl.dot if the tolerance allows). Each tile should span all OC for a block of pixels, with the small weight staged in shared memory. The epilogue keeps a running min over OC per pixel and writes only the 1-channel output after applying the tail, so the big conv output is never written or re-read.
(b) Minimum: try cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv, within tolerance per str-00029), cudnn.benchmark, and/or channels_last (non-dilated, timed per str-00051). With NHWC, the min kernel should read the C values of each pixel contiguously.

Notes: x.contiguous() on the cuDNN output is a no-op. Time locally against the reference before SUBMIT. One attempt used CHECK-then-SUBMIT and found the parity result only after submission. Two others submitted on turn 0 with no CHECK or timing and left levers (a) and (b) untried; another went CHECK then turn-1 SUBMIT and likewise tried no lever.
[str-00037] helpful=3 harmful=1 :: Stride-1 ConvTranspose2d -> training-mode BatchNorm2d -> elementwise activation (e.g., tanh) -> MaxPool2d(2) -> GroupNorm (KernelBench conv_transpose family, L2_011: 512x64x32x32 -> 128 ch, 5x5, p=1).

Measured, CONFIRMED TWICE:
- Run 1: 3.81x (27.9 -> 7.34ms), own kernels about 23% of runtime.
- Run 2: 3.4879x (27.81 -> 7.97ms, lo95 3.487), own_kernel_time_frac 0.22. The recipe was copied nearly verbatim, INCLUDING the slow stats layout this bullet warns against: one block per channel (128 blocks), a serial loop over N, double accumulation. The pool kernel also did a per-element 64-bit div/mod. It compiled first try (str-00001 layout), then went CHECK on turn 0 and SUBMIT on turn 1 with no timing. That slow tail is the likely cause of the drop below 3.81x (evaluator variance is also possible).
~3.5-3.8x is the SOLVED BASELINE: resubmitting it cannot raise the score. Nearly all of the gain is the conv-transpose rewrite (5x5 is a listed win in str-00009).

Recipe:
(1) Rewrite the conv-transpose as a cuDNN forward F.conv2d (padding k-1-p) with the flipped and permuted weight (see str-00009).
(2) Run a per-channel BN stats kernel: sum/sumsq, biased var for normalizing. A finalize step updates running_mean/running_var in-kernel (unbiased var*n/(n-1), momentum) by receiving the buffers as tensors. In eval mode, use the running stats instead (eval branch so a custom kernel runs in both modes).
(3) Run ONE fused kernel with one thread per POOLED output. It normalizes each element of its 2x2 window with the per-channel scale/shift, applies the activation, and then takes the max. Normalize and activate BEFORE the max, because gamma can be negative. The large conv output is read once and only the quarter-size tensor is written.
(4) Leave GroupNorm on the small pooled tensor as eager F.group_norm; it is cheap there.

Training-mode BN is NOT folded into the conv.

WARNING: do NOT copy the stats-kernel layout of the measured runs (one block per channel, serial loop over N, double accumulation per element, per-element div/mod). The conv-transpose rewrite masked how slow it is. Reused on a plain conv -> training BN -> scale task, the same layout scored 0.69x (see str-00043).

Use instead (str-00043 layout):
- one 256-thread block per (n,c) (N*C = 65536 blocks here), contiguous float4 fp32 loads plus scalar tail;
- fp32 per-thread accumulation, shared tree per block;
- fp32 atomicAdd into zeroed per-channel sum/sumsq, then the finalize kernel.

Next attempt plan: CHECK on turn 0, then time each lever at full size against the reference and against the ~7.3-8.0ms baseline before SUBMIT, keeping the fastest:
(a) the str-00043 stats layout above;
(b) a 2D grid in the pool kernel (blockIdx.y = n*C+c) so scale/shift load once per block and there is no 64-bit div/mod per element;
(c) fuse the GroupNorm sum/sumsq into the pool kernel (per-(n,g) atomics) plus a tiny apply pass;
(d) conv side (~78% of runtime): cudnn.benchmark, cuDNN TF32 if a max-error probe passes, channels_last (timed, str-00051).
[str-00038] helpful=3 harmful=2 :: SCOPE GUARD FIRST: this recipe applies ONLY to plain sequential conv->BN->ReLU stacks (MobileNetV1-style) where eager runs separate BN and ReLU passes. It is a MEASURED LOSS on ResNet18 (twice: 0.73x and 0.708x, str-00057) and on MobileNetV2/EfficientNet (0.79-0.80x, str-00026). Do not apply it to residual or inverted-residual networks.

Measured modest WIN for an eval-mode small-batch CNN with plain conv→BN→ReLU stacks (KernelBench network family, MobileNetV1 L3_019 at batch 10, 224px, checked at atol=rtol=1e-4). This contrasts with the MobileNetV2/EfficientNet failures in the small-batch CNN bullet. The reference runs separate eager BN and ReLU passes after each conv. The candidate collapses those two passes into ONE pass rather than adding an extra activation launch on top.

CONFIRMED TWICE:
- Run 1: 1.07x (3.78 -> 3.53ms), custom kernel about 10% of runtime.
- Run 2: 1.0643x (3.75 -> 3.52ms, lo95 1.061), 1.08x vs torch.compile, own_kernel_time_frac 0.106. The recipe was copied nearly verbatim (str-00001 layout, long totals per str-00058, plan precomputed in __init__ per str-00026 rule 3, fold cache keyed on data_ptr with no private attributes per str-00014, 256 threads, 4096-block cap). CHECK on turn 0, SUBMIT on turn 1 with no timing and no lever tried.
~1.06-1.07x is the SOLVED BASELINE for L3_019: resubmitting this layout cannot raise the score.

Recipe:
- Fold BN in fp32, exactly as scale = gamma / sqrt(running_var + eps), w' = w * scale[:,None,None,None], b' = beta - running_mean * scale. A turn-0 fold or numerics slip failed CHECK at 1e-4.
- Build the folded plan once and cache it as self.<x> attributes, keyed on a conv weight's data_ptr() (not _version).
- Per layer, run F.conv2d(x, w', None, ...) and then one in-place bias+ReLU CUDA kernel (c = (idx/HW)%C), launched on the current stream inside cuda_sources.
- Keep a training branch (F.batch_norm with batch stats plus a custom ReLU) so custom kernels run in both modes.

Variant measured as a FAILURE on MobileNetV2 (not retested on V1): pass b' straight to F.conv2d so cuDNN adds it, and keep only an in-place custom ReLU6 after each conv. It scored 0.79x (str-00026 failure 4). Do not treat this variant as a cheaper drop-in; on inverted-residual networks it is a guaranteed slowdown.

Scope limit: this works only when the custom pass replaces eager passes, here BN + ReLU collapsed into one, on a plain sequential conv stack. It did NOT transfer to ResNet18: folded bias + custom residual add+ReLU kernel scored 0.73x, and the full recipe (conv without bias + bias_relu kernel + bias+residual+ReLU kernel) scored 0.708x, because the reference's in-place add_+relu_ already costs one pass and the extra launches dominate (see str-00057). Nor to MobileNetV2/EfficientNet (0.79-0.80x, str-00026).

This layout caps around 1.1x; remaining headroom is launch count (~28 conv + ~27 kernel launches) and cuDNN depthwise convs. Next attempt on L3_019: CHECK on turn 0, then spend the remaining turns timing these launch-cutting levers at full size in eval mode against the reference before SUBMIT, keeping the fastest: (a) CUDA-graph the eval forward with static input/output buffers and the custom kernels inside the graph (str-00026 4a); (b) fuse the last pointwise conv + bias + ReLU + AvgPool(7) into one reduction kernel; (c) a fused depthwise+bias+ReLU kernel only at stage shapes where it beats cuDNN (time per shape, str-00095), ideally fused with the following 1x1; (d) cuDNN TF32 only if a max-error probe passes 1e-4 (likely fails).
[str-00039] helpful=1 harmful=0 :: Conv3d -> x/d -> MaxPool3d(2) -> global avg pool (AdaptiveAvgPool to 1x1x1) -> +bias[C,1,1,1] -> sum over C (KernelBench conv family, L2_008: IC=8, OC=16, k=3, N=128, 16x64x64 input). The tail collapses algebraically, so never materialize the divided or pooled tensors.

Algebra:
- Per (n,c), compute the mean over pooled windows of the window max.
- Apply 1/d AFTER the max. This is valid only for d>0; for d<0 the max becomes a min, so take the min instead or apply the scale before the max.
- out[n] = sum_c(mean_max_nc)/d + sum_c(bias_c). The bias sum is added once, not per window.
- The output shape must match the reference's [N,1,1,1] (view(-1,1,1,1)).

Measured layout, CONFIRMED TWICE: cuDNN conv, then one reduction kernel with one 256-thread block per (n,c). Each thread takes the max of its 2x2x2 window and accumulates the maxes, followed by a shared-memory tree into an [N,C] buffer. A second tiny kernel (one block per n) sums over C, multiplies by 1/d and adds the bias sum.
- Run 1: bias.sum().item() on the host: 1.057x (12.67 -> 11.99ms), own kernels ~3% of runtime.
- Run 2: cleanup (a) applied (bias summed in-kernel, no .item()), long sizes (str-00058), str-00001 layout (compiled first try): 1.0723x (12.67 -> 11.82ms, lo95 1.0717), own_kernel_time_frac 0.024. CHECK on turn 0, a turn lost to no_code (str-00008), SUBMIT on turn 2, no timing, no conv lever.
~1.06-1.07x is the SOLVED BASELINE: resubmitting this layout or further tail cleanups cannot move the score (the tail is ~2.4% of runtime). Caveats in the measured kernel: hardcoded 2x2x2 windows, sm[256] matching the 256-thread launch, -1e30f max seed (use -INFINITY), assumes d>0.

Remaining cheap cleanups (negligible gain): (b) merge the channel sum into the reduction (one block per n looping over c, or atomicAdd into zeroed out[n]); (c) 2D/3D thread mapping with w fastest instead of per-element div/mod.

Real gains must come from the conv (~97% of runtime). torch.backends flags (cudnn.allow_tf32, cudnn.benchmark) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Next attempt: CHECK the baseline on turn 0, then time at full size against the reference (~11.8ms) and keep the fastest:
(1) A fused direct conv whose epilogue does the window max and mean: IC=8, k=3, OC=16, so the weight is only 16x216 floats and fits in smem. Each block computes conv values for all 16 OC over a tile of pooled windows (input tile plus halo staged in smem), takes the 2x2x2 max in registers, accumulates per-(n,c) sums (block reduce + atomicAdd), so the ~1GB conv output is never written or re-read. Verify indexing vs nn.Conv3d on a small shape first.
(2) channels_last_3d via your own transpose kernel (str-00102 pattern, timed per str-00051), with the tail reading NDHWC.
Resubmit the baseline only if no lever beats it.
[str-00040] helpful=7 harmful=0 :: Skinny-M, huge-weight Linear + elementwise epilogue (KernelBench matmul family, e.g., M=128 batch with K=N=32768 or K=N=16384 fp32; the weight is read once per forward). The problem is weight-bandwidth bound and parallelism-limited, not compute bound.

Measured WIN, CONFIRMED TWICE at K=N=32768 (Linear -> swish -> scale):
- Run 1: one fused Triton tl.dot kernel scored 1.33x (16.1 -> 12.1ms), correct, own kernel at 100% of runtime. It used input_precision='tf32x3', grouped pid ordering (GROUP_M=8), W.t() read via swapped strides, and a bias + swish (tl.exp) + scale epilogue with one masked store.
- Run 2: 1.3299x (16.05 -> 12.07ms, lo95 1.329), own kernel 100%. Same tf32x3 swapped-stride read, but a plain 2D grid (no grouped ordering), autotune over BM=128 with BN in {64,128} (256-512 tiles), BK {32,64}, num_warps=8, num_stages 3-4, and a K loop with no K mask (only M/N masks; K % BK == 0). The identical score shows that grouped ordering and the K mask are both NEUTRAL here.
Both went CHECK then SUBMIT with no local timing. ~1.33x is the SOLVED BASELINE for this task: resubmitting it cannot raise the score. It is still far from the limit: about 5ms is ideal for 3xTF32, and about 2-3ms (roughly one pass over the weight) for plain TF32.

NT-layout note: in this weight-bandwidth-bound skinny-M regime, the in-place swapped-stride NT tf32x3 read WINS. The NT-layout tf32x3 losses (str-00075/str-00076) were on compute-bound M>=1024 GEMMs and do not transfer here. Do NOT apply str-00075 lever (4) (cached W.t().contiguous()) at this size: the weight is ~4.3GB, so a cached copy doubles memory and adds OOM risk.

Measured FAILURE of the same recipe at a smaller N: Linear M=128, K=N=16384, with a min(x,c)-c epilogue (fused as min(x-c,0), algebra correct). It was correct but scored 0.80x (5.61 vs 4.47ms, own kernel 100%), about 12 effective TFLOPS against about 23 at N=32768. The likely cause was GRID UNDERFILL. Every autotune config had BM=128 (so M fits one tile row) and BN in {64,128,256}, so the grid was only N/BN = 64-128 programs on an ~132-SM GPU. Each program ran a long 256-512-step K loop, leaving SMs idle or giving a single wave with a long tail. Minor contributors: a %M/%N index wrap, a masked K loop every iteration, and a .item() host sync per forward. It went CHECK then immediate SUBMIT with no timing. The tf32x3 recipe is NOT an automatic win for skinny M.

CONFIRMED FIX on that EXACT task (M=128, K=N=16384, min(x,c)-c): 1.17x (4.58 -> 3.91ms, lo95 1.157), correct, own kernel 100%. Changes vs the 0.80x version: autotune configs all with BLOCK_M=128 (whole batch in one tile, weight streamed once) and BN=64 configs included (256 programs; BN=128 gives 128, BN=256 gives 64), num_warps=8, num_stages 3-4, BK in {32,64}; plain 2D grid (no grouped ordering needed with one M tile); no %M/%N index wrap; the scalar nn.Parameter passed as a pointer and read with tl.load(ptr) (no .item()); epilogue acc + b - c then tl.minimum(acc, 0.0); tf32x3; 1D kmask broadcast per operand (str-00042). It still used a masked K loop and went CHECK then SUBMIT with no timing, and 3.9ms is well above the ~1-2ms weight-read bound, so headroom remains.

RULE before picking configs: compute the tile count cdiv(M,BM)*cdiv(N,BN) and make sure it is at least ~2x the SM count. The 1.33x and 1.17x wins had 256+ tiles; the 0.80x loss had 64-128 tiles.

Levers (time each locally against F.linear on the FULL shape before SUBMIT):
(1) Put the whole batch in one M tile: BLOCK_M >= M (e.g., 128). The weight is then streamed from HBM exactly once. BLOCK_M=64 configs read the weight twice, so drop them for small M. Combine this with (2) so the grid still fills the GPU. (Confirmed in the 1.17x fix.)
(2) Fill the GPU when N/BN tiles are too few: use BN=64 (256 programs at N=16384; confirmed), or split-K (untested; e.g., SPLIT_K 2-4 with BN=128, each split a contiguous K chunk, with tl.atomic_add into a zeroed fp32 buffer followed by an epilogue pass, or a small partials buffer plus a fused reduce+epilogue kernel). A persistent or stream-K schedule is another option. A nonlinear epilogue (min, swish, activation) must run after the split-K reduction, not per split.
(3) If the task tolerance is about 1e-2 (confirm it first; see the tolerance pitfall bullet), time input_precision='tf32' against 'tf32x3' (untested) after printing the max abs/rel error. Plain TF32 cuts tensor-core work so the kernel can approach the weight-read bound.
(4) When K % BK == 0, run the K loop unmasked (measured neutral at N=32768, harmless), and drop the %M/%N index wrap when the dims divide the tiles (done in the fix).
(5) Pick BK and stages so enough loads are in flight to saturate bandwidth (e.g., num_stages 4-5, BK 64/128; untested).
(6) Pass a scalar nn.Parameter (e.g., a constant) as a tensor pointer and tl.load it in-kernel instead of calling .item(), which forces a host sync every forward. (Done in the fix.)

Next attempt at K=N=32768: after CHECK, spend the turns on a bounded full-size timing of, in order, (3) tf32, (2) split-K, (5) stages/BK, keeping the fastest; do not resubmit the 1.33x baseline. If every variant stays below 1x on another shape, fall back to cuBLAS plus a fused epilogue (about 1.0x) rather than submitting a sub-1x kernel.
[str-00041] helpful=1 harmful=0 :: Single-step RNN cell (KernelBench rnn family, e.g., VanillaRNN: combined = cat(x, hidden) -> i2h Linear -> tanh -> new hidden -> h2o Linear, with batch 256, input = hidden = 16384, output 8192, fp32). There is no time loop here. It is NOT a recurrence problem, and the multi-layer cuDNN LSTM/GRU bullet (str-00018) does not apply.

Treat it as two dense skinny-M GEMMs with fused epilogues, using the tf32x3 Triton recipe (str-00004):
- GEMM 1: tl.dot(..., input_precision='tf32x3') with + bias and an exp-based tanh (2/(1+tl.exp(-2x))-1; tl.math.tanh is missing, see str-00012) in the epilogue. Its output is the new hidden state; store it to self.hidden.
- GEMM 2: the same kernel with a bias-only epilogue.
- Read nn.Linear weights via swapped strides (no transpose copy), autotune keyed on M,N,K, grouped pid ordering (GM=8).

Measured BASELINE: 1.25x (20.5 -> 16.3ms) and 1.36x vs torch.compile, with own kernels at 99.5% of runtime. That version used torch.cat, BM/BN 128-256 autotune and grouped ordering. It passed correctness at K=32768 with unknown tolerance, and was submitted right after CHECK with no timing.

Measured REGRESSION from applying 'levers' blindly: 1.0094x (20.44 -> 20.25ms), correct, own kernels 99.9%. Changes from the 1.25x version:
- no torch.cat, with two K loops (x, then hidden with a W column offset);
- autotune dominated by BM=256 configs (whole batch in one M tile);
- W loaded as [BN,BK] tiles with tl.dot(a, tl.trans(w));
- no grouped ordering;
- a masked K loop.
Likely cause: GRID UNDERFILL (str-00040). At BM=256 the grid is only N/BN programs: 128 for i2h and 64 for h2o at BN=128, on ~132 SMs. Each program runs a ~1024-step K loop, and tf32x3 at 256x128 adds register pressure. The NT-layout tf32x3 weakness (str-00075/00076) may also contribute. The rewrite was submitted on turn 2 with no timing. Turn 0 also lost a CHECK to an output mask written as offs_n[:, None] instead of offs_n[None, :] (the str-00042 broadcast error).

RULE: before choosing configs, compute cdiv(M,BM)*cdiv(N,BN) and keep it at least ~2x the SM count. Never submit a rewrite of this task without timing it against the known ~16.3ms (1.25x) baseline.

Headroom levers (time each locally against 16.3ms):
(1) Dropping torch.cat saves only an ~32MB copy against ~1.5GB of weight reads. It is fine to do, but it is NOT the lever. Pass x and hidden as two A pointers and split the K loop at input_size.
(2) Weight-bandwidth bound (see str-00040). Streaming each weight about once (BM >= M, e.g. 256) only helps if the grid still fills the GPU. Use either:
  - BM=128 with BN=64 (512/256 tiles for i2h/h2o), or
  - BM=256 with BN=64 plus split-K (2-4 contiguous K chunks, tl.atomic_add into a zeroed fp32 buffer). The bias+tanh epilogue must then run AFTER the reduction, since tanh is nonlinear.
  Drop BM=64 configs, because they re-read the weight M/64 times. Keep grouped pid ordering when there are several M tiles.
(3) Run the K loop unmasked, since 16384 % BK == 0.
(4) If the tolerance is ~1e-2, time plain tf32 against tf32x3. Given the NT layout, also consider a once-cached contiguous W.t() fed to the NN recipe (str-00075 lever 4).

General rule: check whether an 'rnn'-family task actually loops over time. A single cell step reduces to GEMM+epilogue fusion.
[str-00042] helpful=22 harmful=0 :: Triton broadcast/shape pitfall: a mask built as a 2D tensor, e.g. valid_k = k[None, :] < K with shape [1,BK], must not be re-indexed with [:, None]. That makes it 3D and fails at compile time with 'Cannot make_shape_compatible: incompatible dimensions'. Keep 1D index/mask vectors (e.g., kmask = k < K) and broadcast them separately per operand orientation: kmask[None, :] for an [BM,BK] A tile and kmask[:, None] for a [BK,BN] B tile.
[str-00043] helpful=4 harmful=0 :: Conv → training-mode BatchNorm2d → scalar multiply (KernelBench conv family, e.g., Conv2d 8->64, 3x3, N=128, 128px; conv output about 130MB, about 2M elements per channel). The eager BN is already an efficient cuDNN kernel. Collapsing the tail from 3 passes to 2 helps only if your stats reduction is at least as fast as cuDNN's, so the reduction layout decides the outcome.

Measured WIN: 1.58x vs eager (5.61 → 3.54ms) and 1.40x vs torch.compile, with own kernels at ~42% of runtime. It went CHECK then SUBMIT on the exact task where the layout below had failed. Keep the cuDNN conv and use THREE kernels, all launched from cuda_sources (str-00001), with 64-bit sizes (str-00058):
(1) Stats: one 256-thread block per (n,c), so N*C = 8192 blocks over the contiguous S slab. Use float4 loads plus a scalar tail, fp32 per-thread sum/sumsq, a shared-memory tree, then thread 0 does an fp32 atomicAdd into zeroed per-channel sum/sumsq. fp32 atomics were within tolerance, so double accumulation is not required here.
(2) Finalize: one thread per channel. mean, biased var = E[x^2]-m^2, inv = rsqrtf(var+eps). Fold gamma, beta AND the scalar into scale = inv*gamma*s and shift = (beta - mean*inv*gamma)*s. Update running_mean/running_var in-kernel (unbiased var*n/(n-1), momentum), with the buffers passed as tensors (str-00014).
(3) Apply in place on the conv output, using a 2D grid (blockIdx.y = n*C+c, blockIdx.x capped at 256 blocks, grid-stride over S), so scale/shift load once per block with no per-element div/mod.
Also provide an eval branch whose apply kernel uses the running stats, so a custom kernel runs in both modes (str-00031).

Measured FAILURE of the earlier layout: 0.69x (8.16 vs 5.61ms), with own kernels at 75% of runtime. The algebra was the same, but it lost because:
(1) The stats kernel launched only C=64 blocks x 256 threads to reduce the whole tensor, leaving the GPU badly underutilized.
(2) It used double FMAs per element, which are slow relative to fp32.
(3) It decoded n and spatial index per element with 64-bit idx/S and idx%S.
(4) The apply pass allocated a new output and did a 64-bit div/mod per element.
Switching from 64 blocks with double FMAs to N*C blocks with float4 fp32 reads is what flipped 0.69x into 1.58x.

Untested headroom; time each locally before SUBMIT:
- float4 loads/stores in the apply pass (it is still scalar);
- dropping the torch.zeros launches for the accumulators;
- cuDNN TF32 for the conv, which still dominates, within tolerance (str-00029);
- fusing the stats into a conv epilogue so the conv output is read only once more.
A sub-1x local timing means the reduction layout is wrong.
[str-00044] helpful=1 harmful=0 :: ConvTranspose3d (stride 2) → MaxPool3d(2) → softmax over C → subtract per-channel vector → swish → max over C (KernelBench conv_transpose family, e.g., IC=3, OC=16, k=3, input 128x3x16x32x32, conv output ~2GB at 128x16x32x64x64).

Baseline layout, CONFIRMED TWICE at ~1.06x: cuDNN conv, eager nn.MaxPool3d, then ONE register-resident channel kernel.
- Kernel: one thread per (n,s) with s fastest so the strided-c loads coalesce, and a templated/#define C=16 with a `float vals[C]` array.
- In registers it computes max/__expf/sum softmax, subtracts sub[c], applies swish y/(1+__expf(-y)), and keeps a running max over c.
- Swish must run BEFORE the channel max, because the subtract makes it non-monotone across channels.
- Run 1: 1.063x (19.77 → 18.60ms). Run 2: 1.066x (19.75 → 18.53ms), own_kernel_time_frac 0.005.
Both went CHECK-then-SUBMIT with no timing, and run 2 skipped the levers below even though a turn remained. Do NOT resubmit this layout expecting more. The time sits in the cuDNN conv, which writes the ~2GB output, and in the eager pool that re-reads it.

Next levers, in order of expected payoff. All are untested; time each locally at FULL size before SUBMIT.
(1) A FULLY fused direct transposed-conv kernel, which is plausible because IC and OC are tiny.
- The input is only ~25MB, and the weights are IC*OC*27 = 1296 floats, which fit in constant or shared memory.
- Use one thread per POOLED output. For each of its 2x2x2 conv positions and all OC, accumulate only the stride-2-valid taps: per dim, output o = i*2 - p + k, so at most 2 taps per dim × IC.
- Take the window max per channel in registers, then run softmax, subtract, swish and the channel max, and write only the [N,Do,Ho,Wo] output.
- The 2GB conv tensor is never written, so the task drops from ~2GB+ of traffic to roughly 25MB of reads plus the output write.
- Verify the indexing against nn.ConvTranspose3d on a small shape first.
(2) At minimum, fold the MaxPool3d(2) window into the channel kernel.
- Index the kernel by POOLED position, and for each channel take the max of the 2x2x2 window straight from the conv output, then run the softmax chain over C.
- Max-first is exact here because the reference pools BEFORE the softmax. This contrasts with str-00006/str-00060, where the activation or softmax must precede the pool.
- This removes the pool launch and the pooled-tensor write and re-read.
(3) Conv-side levers:
- cuDNN TF32 within tolerance (__expf already passed tolerance here);
- channels_last_3d, timed per str-00051;
- the stride-2 phase decomposition (str-00030).

When a turn remains after CHECK, spend it timing these levers, not submitting the known baseline.
[str-00045] helpful=9 harmful=0 :: Conv -> monotone non-decreasing elementwise chain -> MaxPool2d(P), KernelBench conv family. Example (L2_082): Conv2d 8->64, k=3 -> tanh -> *scale -> +bias[C,1,1] -> MaxPool2d(4) on a 128x64x254x254 conv output.

Recipe:
- Keep the cuDNN conv.
- Write ONE fused load_inline kernel with one thread per POOLED output. It reads the PxP window of the RAW conv output and writes only the pooled tensor.
- When every step in the chain is monotone non-decreasing, max commutes with the whole chain. Qualifying steps include tanh, sigmoid, ReLU-family ops, multiplying by a POSITIVE scalar, and adding a per-channel constant. So take the window max of the raw values first, then apply the chain once per output, e.g. tanhf(m)*scale + bias[c]. This cuts transcendental work by P^2.
- About 4 eager passes over the large conv output become 1 read plus a 1/P^2-size write.

Measured, CONFIRMED THREE TIMES on L2_082 (~1.93-1.94x band):
- Run 1: 1.94x (19.4 -> 9.99ms) and 2.03x vs torch.compile.
- Run 2: 1.9347x (19.44 -> 10.05ms, lo95 1.933), own_kernel_time_frac 0.173. Grid-stride kernel (256 threads, 4096-block cap), long indices (str-00058), str-00001 layout (compiled first try), eager fallback when scale <= 0. CHECK on turn 0, SUBMIT on turn 1, no timing, no lever.
- Run 3: 1.9268x (19.44 -> 10.09ms, lo95 1.925), own_kernel_time_frac 0.175. Same recipe copied verbatim as a turn-0 SUBMIT with no CHECK or timing; correct only because the recipe was proven. It gained no information.
~1.93x is the SOLVED BASELINE for this task: resubmitting this layout CANNOT raise the score (three runs prove it). The conv is ~82-83% of runtime, so the next gain must come from the conv side.

Correctness guard: max-first holds only when the scale is > 0. If scale < 0, take the window MIN instead, apply the chain per element before the max, or fall back to the eager path (done in runs 2-3). Use the per-element form whenever any step is non-monotone or has a possibly negative per-channel multiplier (as in str-00006, str-00037 and str-00052).

Output shape: floor(H/P) x floor(W/P), which matches MaxPool2d's defaults (stride = kernel, no padding).

Next attempt: CHECK on turn 0, then spend the remaining turns on bounded full-size timings against the reference and the ~10.05ms baseline, keeping the fastest. NOTE: torch.backends flags (cudnn.benchmark, allow_tf32) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Realistic levers (all untested):
(a) channels_last via your own smem-tiled NCHW->NHWC transpose kernel plus a channels_last weight cached on data_ptr (str-00102 non-dilated forward-conv win pattern), with the pool kernel made NHWC-aware (threads span C per pixel window). Time per str-00051.
(b) A fused direct conv with a pool epilogue: with IC=8, k=3 the weight is only 64x72 floats and fits in smem. Each block computes a tile of pooled outputs for all OC (or an OC group), takes the window max of conv values in registers, applies the chain, and writes only the pooled tensor, so the ~2GB 128x64x254x254 conv output is never written or re-read. TF32 tl.dot / implicit GEMM only if a max-error probe passes the tolerance (str-00029). Verify indexing vs nn.Conv2d on a small shape first.
- Minor kernel side (~17% of runtime): a 2D grid with ow fastest instead of per-element 64-bit div/mod; float4 window-row loads for P=4, but guard alignment, because rows with W=254 are not 16B-aligned (str-00059/str-00035 misaligned-address pitfall).
[str-00046] helpful=3 harmful=0 :: Conv2d -> GroupNorm -> per-channel scale -> MaxPool2d(P) -> clamp (KernelBench conv family, L2_085: Conv2d 8->64, 3x3, N=128, 128px, so the conv output is 128x64x126x126; G=16, P=4).

Measured, CONFIRMED TWICE with the recipe below:
- Run 1: 1.54x vs eager (4.74 -> 3.08ms) and 1.78x vs torch.compile, own kernels about 32% of runtime.
- Run 2: 1.5137x (4.73 -> 3.13ms, lo95 1.5128), 1.75x vs compile, own_kernel_time_frac 0.30. The recipe was reproduced nearly verbatim (str-00001 layout, long sizes per str-00058) and compiled first try. It went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no lever tried.
~1.51-1.54x is the SOLVED BASELINE: resubmitting this layout cannot raise the score.

Recipe: keep the cuDNN conv and replace the whole tail with TWO kernels, both launched on the current stream from cuda_sources.
(1) GN stats kernel with one block per (n,g). Each thread accumulates fp32 sum/sumsq over the group's contiguous CPG*S range. Combine them in a double shared-memory tree. Compute biased var = E[x^2]-m^2 and rstd = 1/sqrt(var+eps).
(2) ONE fused kernel with one thread per POOLED output. For each element of its PxP window, compute ((v-m)*rstd*gamma_c + beta_c)*scale_c and take the max. Then clamp once per output.

Ordering rules:
- Normalize BEFORE the max, because gamma_c*scale_c can be negative. Keep this as the general path.
- A max-first fast path on raw values is valid only behind a runtime check that every gamma_c*scale_c > 0.
- The clamp is monotone, so it commutes with max and runs once per output.

Result: about 5 eager passes become 2 reads of the conv output plus a 1/P^2-size write.

The custom kernels are ~30% of runtime, which is much more than in most conv-dominated tasks, so kernel-side levers are worth timing. After the first passing CHECK, time these at full size against the reference, in priority order, and keep the fastest (all untested):
(1) Split each (n,g) reduction across several blocks. There are only N*G = 2048 blocks of ~254k elements each. Use float4 loads, fp32 per-thread sums, and either atomics or a partials buffer plus a finalize step.
(2) Use a 2D grid in the pool kernel with blockIdx.y = n*C+c, and stage the per-channel scale = rstd*gamma*scale and shift = (beta - mean*rstd*gamma)*scale once per block. The measured kernel recomputes g, mean and rstd per output and runs a per-element 64-bit div/mod chain; this change removes both and cuts FLOPs.
(3) Use float4 window-row loads for P=4, guarded for alignment, because W=126 is not a multiple of 4 (str-00059/str-00035 misaligned-address pitfall).
(4) Conv side: cuDNN TF32 within tolerance (str-00029), cudnn.benchmark, and channels_last (non-dilated, timed per str-00051).
[str-00047] helpful=2 harmful=0 :: Conv/ConvTranspose -> Softmax(dim=1) over channels -> elementwise activation (e.g., ConvTranspose3d stride 2 -> Softmax(dim=1) -> Sigmoid, C=64 output channels, KernelBench conv_transpose family, L2_049).

Measured, two kernel versions:
- 3-loop version: 1.18x (12.28 -> 10.43ms), own kernel about 12% of runtime. It did 3 channel loops (max, sum of __expf, then normalize + sigmoid 1/(1+__expf(-sm))) and wrote once.
- Register-resident single-read version (headroom items (a)+(b) applied), CONFIRMED TWICE:
  - Run 1: 1.2264x (11.97 -> 9.76ms, lo95 1.226), own_kernel_time_frac 0.083, but only 1.008x vs torch.compile. It compiled first try, went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no conv lever tried.
  - Run 2: 1.2197x (12.08 -> 9.90ms, lo95 1.2185), own_kernel_time_frac 0.074. Recipe copied verbatim as a turn-0 SUBMIT with no CHECK or timing; correct only because the recipe was proven (str-00002). It gained NO new information; the difference from run 1 is evaluator variance.
  Layout: template<int C> with float vals[C] (C=64 caused no visible spill problem); one global read per value into vals while tracking the max; then __expf(v-m) stored back into vals and summed (each exp computed once); then sigmoid via __expf; one thread per (n,s) with s fastest so the strided-c loads coalesce (str-00021 layout); grid-stride with a 4096-block cap, 256 threads; long indices (str-00058); empty_like output; TORCH_CHECK fallback for unsupported C (plus an eager torch.softmax/sigmoid branch in forward for other C); wrapper and launch in cuda_sources (str-00001).

Treat ~1.22-1.23x as the SOLVED BASELINE for this task (three measurements across both versions). The cuDNN conv is ~92% of runtime, so resubmitting this layout or kernel micro-tweaks cannot meaningfully raise the score. Remaining kernel items touch <8% of runtime: (c) write in place on the conv output instead of empty_like; (d) a 2D grid (blockIdx.y=n) to avoid 64-bit div/mod per thread.

Next attempt: CHECK the baseline on turn 0, then spend the spare turns on a bounded full-size timing of conv-side levers against the reference (~9.8-9.9ms baseline) before SUBMIT, keeping the fastest. NOTE: torch.backends flags (cudnn.allow_tf32, cudnn.benchmark) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Realistic levers (all untested; never submit any of them untimed):
- channels_last_3d via your own transpose kernel, with the softmax kernel then reading each pixel's C values contiguously. RISKY: for transposed convs channels_last measured a loss (0.28x on ConvTranspose2d, str-00051/str-00102), so time it at full size.
- A fused or implicit-GEMM transposed conv (TF32 tl.dot within tolerance; __expf already passed, so tolerance is likely ~1e-2, confirm per str-00029) with the softmax+sigmoid epilogue spanning all 64 OC per tile, so the conv output is never written and re-read.
- In-place output (minor).
Do NOT use the per-phase F.conv3d decomposition here: this shape (k=3, s=2, p=1, dilation 1) is non-shortcut, where it is a measured loss (str-00030). Resubmit the baseline only if no lever beats it.
[str-00048] helpful=1 harmful=0 :: Huge-K, small-MN fp32 GEMM (KernelBench matmul family, e.g., M=N=256, K=524288, checked at atol=rtol=1e-4). The output is only a few tiles, so parallelism must come from split-K. Plain TF32 FAILED CHECK at 1e-4 over this K, as did other reorder/precision experiments (twice more in a later attempt), so tf32x3 or strict FP32 is mandatory here.

Measured, both correct but sub-1x vs cuBLAS fp32 (3.72ms):
- Attempt 1: 0.45x (8.30ms), own kernel 99.85%. It autotuned over BK 64-256 (mostly 128/256), BM/BN 64-128, SPLIT_K 32-128 and num_stages 3-4, wrote partials to a [128,M,N] buffer, then reduced them with torch .sum(0).
- Attempt 2, following levers (1)-(3) below: 0.86x (4.31ms), own kernel 99.96%. Config: BM=BN=64, BK=32, SPLIT_K=128 (grid 16x128 = 2048 programs), each split walking K in BK strides interleaved across splits (k = pid_sk*BK, step BK*SPLIT_K), tl.dot(input_precision='tf32x3'), and tl.atomic_add into a torch.zeros fp32 output. The fp32 atomic ordering passed 1e-4.
Small tiles plus atomics roughly halved the runtime, but tf32x3 (3 MMAs per tile) still loses here: cuBLAS fp32 is near the ~1GB operand-read bound (~0.5-0.7ms of HBM traffic; ideal tf32x3 compute is ~1.3ms). Both attempts submitted after CHECK with no local timing.

The dense-GEMM tf32x3 win (str-00004, square shapes) does NOT transfer to this shape.

Levers; time each locally against torch.matmul:
(1) Stay in the proven tile range: BK 32-64, BM/BN 64-128, num_stages 3-5. BK 128/256 with tf32x3 hurt (done in attempt 2).
(2) Choose SPLIT_K so the grid fills the SMs about 2-4x over (~128-256 with 4-16 output tiles; done).
(3) Reduce partials with tl.atomic_add into a zeroed fp32 output, not a large partials buffer plus torch .sum (done, passes 1e-4).
(4) Untested next variants: give each split a CONTIGUOUS K chunk instead of interleaved BK strides (better DRAM page locality and prefetching); use BM=BN=128 so A/B tiles are re-read fewer times; or a manual hi/lo split (a_hi*b_hi + a_hi*b_lo + a_lo*b_hi) or bf16x3.
(5) MANDATORY: if local timing of every variant stays below 1x, submit strict-FP32 cuBLAS plus minimal custom work (~1.0x) rather than a sub-1x kernel (str-00029). Attempt 2 skipped this and submitted at 0.86x.
[str-00049] helpful=2 harmful=0 :: Memory-bound broadcast or elementwise op disguised as matmul (KernelBench matmul family, e.g., diag(A) @ B written as A.unsqueeze(1) * B, N=M=4096 fp32, ref ~0.148-0.156ms). Despite the 'matmul' label this is NOT a GEMM, so the GEMM bullets (str-00004, tf32/tf32x3) do not apply. It moves 64MB read + 64MB write, and eager PyTorch's vectorized broadcast-mul kernel already runs near HBM bandwidth. The realistic target is parity to ~1.0-1.05x, won by trimming overhead and raising memory-level parallelism, not by extra compute.

Measured FAILURE 1: a naive Triton kernel was correct (own kernel 100% of runtime) but scored 0.76x (0.205 vs 0.156ms). Its design: grid (N, cdiv(M,1024)), BLOCK=1024, default num_warps, so about 16k tiny programs. At a ~150us runtime, Triton's per-call Python launch overhead and small per-program tiles are a large fraction of the total. It went CHECK, CHECK, SUBMIT with no local timing.

Measured FAILURE 2 (the previously recommended CUDA layout): load_inline (str-00001 layout, compiled first try), float4 loads/stores with __restrict__, 2D grid with blockIdx.y = row and A[row] in a register, 256 threads, per-row block cap of 64 (in practice 4 blocks per row, 1 float4 per thread, ~16k blocks), long indices, host contiguous/empty_like/TORCH_CHECK. Correct, own_kernel_time_frac 1.0, but 0.9394x (0.158 vs 0.148ms). It went CHECK then immediate SUBMIT with no timing. Following items (1)-(3) below is NOT sufficient by itself; a single float4 per thread gives too little memory-level parallelism, and 1024-element blocks are too small.

Candidate levers (all untested; time each locally and submit only one above 1x):
(a) Each thread handles 2-4 independent float4s, unrolled, issuing all loads before any stores.
(b) A flat 1D grid over all float4s (row = idx / M4, or rows processed per block) with blocks capped near SMs x 8-16 and grid-stride; compare against 1 thread per float4 uncapped (str-00054 found a fixed 1024 cap ~9% slower, so time both).
(c) Streaming cache hints: __ldcs/__stcs on B/C, __ldg on A.
(d) Fewer, larger blocks (several thousand elements per block) or 512 threads.
(e) Minimize host overhead: skip redundant contiguous()/checks on the fast path.
Baseline items still worth keeping: (1) prefer load_inline CUDA over Triton (lower launch overhead), launched on the current stream from cuda_sources (str-00001); (2) float4 when M % 4 == 0, with __restrict__; (3) row scale held in a register, no per-element div/mod where avoidable; (4) if staying in Triton, BLOCK 4096+ or several rows per program, num_warps=8.

MANDATORY: after CHECK, time 2-3 variants locally against the eager op (bounded warmup/iterations, cuda synchronize) before SUBMIT. Expect parity at best; a sub-1x result must trigger trying the next lever, not a submit. General rule: before applying GEMM recipes to a matmul-family task, check whether the op is actually elementwise or broadcast (diagonal, scalar or row/column scaling).
[str-00050] helpful=0 harmful=0 :: ConvTranspose3d → +scalar → LayerNorm(norm_shape) → AvgPool3d(2) → GELU (KernelBench conv_transpose family). Example: out_channels=64, so norm_shape=(64,), applied to a 32x64x32x64x64 conv output.

Read the shapes literally. nn.LayerNorm(norm_shape) normalizes over the LAST dim(s), which here is the spatial W (=64), not the channels. That gives millions of SHORT rows of 64 elements each.

Algebra: a scalar added before LayerNorm cancels exactly. The mean absorbs it and the variance is unchanged. Drop it, together with any sum_weight.item() call, which forces a host sync every forward.

Measured FAILURE: 0.848x (34.1 vs 28.9ms), with own kernels at 49% of runtime. The tail was slower than eager LN + pool + GELU. It used two kernels:
(1) add_ln: one 256-thread block per 64-element row, about 4.2M blocks with 75% of threads idle. Each row was read 3 times through shared-memory tree reductions with many __syncthreads, and a full-size (~1GB) LN output was written.
(2) pool_gelu: re-read that full-size tensor.
The attempt went CHECK then SUBMIT with no local timing.

Recommended layout: ONE kernel that reads the raw conv output once and writes only the 1/8-size output.
- Use one warp per W-row, with W/32 elements per lane held in registers. Compute mean and var with __shfl_xor reductions; no shared memory or syncthreads are needed for the stats.
- Each block (or warp group) handles the 2x2 rows (d,h pairs) of a pooling window.
- Normalize and apply gamma[w]/beta[w] per element, then average the 2x2x2 window across the 4 rows and pairs of adjacent w. Pairs of w can be exchanged via a shuffle, or with float2 loads per lane.
- Apply erf-GELU and store.
- Pool AFTER normalizing, because LN is per-row nonlinear.

Never materialize the full-size LN output, and never split LN and pool+GELU into separate kernels: that adds a full write and re-read of the huge conv-sized tensor.

Further levers: cuDNN TF32 or channels_last_3d for the conv, within tolerance. Time locally against the reference before SUBMIT.
[str-00051] helpful=2 harmful=0 :: cuDNN memory-format pitfall: several bullets list 'try channels_last / channels_last_3d' as a cheap conv lever. It is shape-, algorithm- and op-dependent, not free, and it goes BOTH ways.

- DILATED fp32 conv: CATASTROPHIC. Feeding cuDNN NHWC input + channels_last weight measured 0.0157x (1285 vs 20ms) on Conv2d 32->64, 5x9, dilation (2,3), 512x512. The custom transpose kernel was only 0.4% of runtime, so cuDNN picked a terribly slow dilated NHWC path (torch.compile matched this slowness too, speedup_vs_compile ~1.0).
- NON-dilated fp32 forward conv: measured WIN. Conv2d 64->128, 5x7, N=8, 512x256, bias=False: a custom NCHW->NHWC transpose kernel + F.conv2d with a channels_last weight scored 1.4454x (31.4 -> 21.7ms), correct, and returning the conv output in channels_last strides passed the correctness check (see str-00102).
- NON-dilated fp32 TRANSPOSED conv: measured LOSS. ConvTranspose2d 32->32, 3x3, padding 0, 8x32x512x1024: naive transpose + F.conv_transpose2d with a channels_last weight scored 0.2823x. cuDNN's NHWC transposed path alone (~4.8ms) was slower than the NCHW reference (3.33ms), even before the ~7ms transpose cost.

Rules: gate any channels_last path on dilation == 1 and keep an NCHW fallback for dilated convs. Do not assume a forward-conv win carries over to transposed convs. Treat any memory-format or algorithm switch (channels_last, conv-transpose rewrite, TF32) as a hypothesis that must be timed locally against the reference module before SUBMIT; a single bounded timing run catches order-of-magnitude regressions.
[str-00052] helpful=1 harmful=0 :: Conv2d → scalar subtract → HardSwish → MaxPool2d(P) → Mish (or any pre-pool chain containing a NON-monotone activation, followed by an activation after the pool), KernelBench conv family (L2_035).

Measured, CONFIRMED TWICE with the recipe below:
- Run 1: 1.255x (17.41 → 13.87ms) and 1.22x vs torch.compile. The own kernel was about 8.5% of runtime.
- Run 2: 1.2205x (17.61 → 14.43ms, lo95 1.215), 1.169x vs compile, own_kernel_time_frac 0.11. It used a grid-stride kernel with a 4096-block cap, 256 threads, long indices (str-00058), the str-00001 layout and PyTorch-exact numerics. It compiled first try, then went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no conv lever tried.
Treat ~1.22-1.26x as the SOLVED BASELINE for this task: resubmitting it cannot raise the score. The cuDNN conv is ~89-92% of runtime, so the next gain must come from the conv.

Recipe:
- Keep the cuDNN conv.
- Write ONE load_inline kernel with one thread per POOLED output; put the wrapper and stream retrieval in cuda_sources (str-00001).
- Each thread applies subtract+HardSwish to EVERY element of its PxP window and takes the max.
- It then applies the post-pool activation (Mish) once per output, with the x>20 softplus guard.

Ordering rules:
- HardSwish is NOT monotone: it decreases on about [-3,-1.5]. So max-first (str-00045) is invalid, and the activation must run per element before the max.
- Any op that the reference applies AFTER the pool runs only once per pooled output, so expensive transcendentals there (Mish) cost 1/P^2 as much.

Numerics at atol=rtol=1e-4:
- The first CHECK of run 1 failed at 1e-4 and the next passed. The trajectory does not record the cause.
- The passing versions wrote the formulas exactly like PyTorch: y*fminf(fmaxf(y+3,0),6)/6 for HardSwish, and log1pf(expf(x)) with a threshold of 20, then tanhf, for Mish.
- They used no __expf or other fast intrinsics, and no reassociated forms.

After the first passing CHECK, spend the spare turns on a bounded full-size timing of conv-side levers against the reference before SUBMIT, keeping the fastest (all untested):
- torch.backends.cudnn.benchmark=True;
- channels_last (non-dilated, timed per str-00051), with the pool kernel made NHWC-aware if adopted;
- cuDNN TF32 only if the 1e-4 tolerance allows (likely not; print max error first);
- fusing the pool into a conv epilogue.
Minor kernel headroom (kernel ~11% of runtime): a 2D grid with ow fastest instead of per-element 64-bit div/mod, float2 loads for P=2, and a grid cap near SMs x 8-16 or uncapped instead of 4096 (str-00054).
[str-00053] helpful=0 harmful=0 :: Large 3D conv (e.g., Conv3d 32->64, k=3, N=16, 32x64x64 input, about 25ms, 64-channel fp32 output) followed by a SHORT chain of built-in activations (e.g., F.mish → torch.tanh; KernelBench L2_047), conv family. Note that F.mish (like F.hardswish) is already ONE native kernel, so the eager chain is only 2 passes. Collapsing it into one fused pass saves just one pass, unlike a hand-written multi-op Mish (str-00007).

Measured, CONFIRMED TWICE with the same layout: cuDNN nn.Conv3d plus one in-place grid-stride kernel computing tanhf(v*tanhf(log1pf(expf(v)))) with a v>20 guard (PyTorch-exact numerics, no fast intrinsics; 256 threads, 4096-block cap, long n per str-00058, str-00001 layout, compiled first try).
- Run 1: 1.022x (24.96 vs 24.42ms), own_kernel_time_frac 0.029. It went CHECK then SUBMIT on turn 1 with no timing.
- Run 2: 1.0218x (24.92 -> 24.38ms, lo95 1.0205), 0.978x vs torch.compile, own_kernel_time_frac 0.032. Turn-0 SUBMIT with no CHECK, no timing and no conv lever; correct only because the recipe was copied.
~1.02x is the SOLVED BASELINE and is below torch.compile: resubmitting this layout cannot raise the score. The conv is about 97% of runtime.

Next attempt: CHECK on turn 0, then spend the remaining turns timing conv-side levers at full size against the reference (bounded warmup/iterations, cuda synchronize), and pair the fastest with the fused activation kernel. Levers, in order:
(1) cuDNN TF32 (torch.backends.cudnn.allow_tf32=True around the conv), only if a max-error probe fits the tolerance (str-00029).
(2) channels_last_3d. This is plausible for non-dilated convs but must be timed (str-00051). If you adopt it, keep the elementwise kernel layout-agnostic (a flat pass over a contiguous buffer) or make it read NDHWC.
(3) torch.backends.cudnn.benchmark=True for algorithm selection.

Kernel cleanups for small extra gains (kernel only ~3% of runtime): cap the grid-stride grid at about SMs x 4-8 blocks rather than a fixed 4096 or 65535, and use float4 loads and stores when numel % 4 == 0 (scalar tail).
[str-00054] helpful=1 harmful=0 :: Standalone multi-op eager elementwise chain with no conv or GEMM around it. Example: KernelBench L1 MinGPT tanh-GELU, 0.5*x*(1+tanh(sqrt(2/pi)*(x+0.044715*pow(x,3)))), on 8192x8192 fp32. This applies even when the task's op-family label says 'attention' or 'matmul'. It is a pure elementwise task, so the GEMM and attention bullets do not apply. Eager runs each op (pow, mul, add, tanh, mul, add, mul) as a separate full read+write pass, about 7+ passes over a 256MB tensor. Collapsing all of them into ONE pass is the whole win. This contrasts with a single eager elementwise or broadcast op, which is already near bandwidth and caps at about parity (str-00049).

Measured WIN, CONFIRMED TWICE:
- Run 1, 1 thread per float4 (no grid cap): 8.09x vs eager and 1.28x vs torch.compile (3.58 to 0.44ms, own kernel 100% of runtime). That is near the HBM minimum for 512MB of traffic.
- Run 2, grid-stride loop with blocks capped at a fixed 1024 (256 threads), long indices (str-00058), separate scalar tail kernel: 7.45x (3.58 to 0.480ms, lo95 7.43), 1.18x vs torch.compile, own kernel 100%. The 1024-block cap was ~9% SLOWER than 1 thread per float4. 'Fewer blocks' is NOT a lever on this bandwidth-bound kernel; it already answered the untested grid-stride idea.

Recipe:
- Use a load_inline CUDA kernel. Put the wrapper, stream retrieval and launch in cuda_sources; cpp_sources holds only a declaration (str-00001).
- Use float4 vectorized loads and stores with __restrict__ pointers.
- Handle numel%4 with a scalar tail (same kernel or a tiny second kernel).
- Launch on at::cuda::getCurrentCUDAStream(), with 256 threads and 1 thread per float4 (preferred). If using grid-stride, cap near SMs*8-16 blocks with ~4 float4s per thread, never as low as 1024, and time it against the uncapped version.
- Keep n and indices as long (str-00058).
- Write x*x*x instead of pow, and precompute sqrt(2/pi) as a constant.
- Use accurate tanhf, not __expf or fast-math forms, to stay safe on tolerance.
- Call x.contiguous() before the kernel.

~8x (0.44ms) is the solved baseline; remaining headroom is small because the kernel is bandwidth-bound. Both runs were turn-0 SUBMITs with no CHECK, which succeeded only because the recipe was copied verbatim. Still CHECK and briefly time before SUBMIT (str-00002); timing would have caught the grid-cap regression.
[str-00055] helpful=1 harmful=0 :: ConvTranspose3d (stride 2) -> MaxPool3d(2) -> MaxPool3d(3) -> sum over C (KernelBench conv_transpose family, L2_078, e.g., C=64, conv output about 1GB, pooled output only N*Do*Ho*Wo ~ 16k positions). The pool algebra is sound but gives little here. Fusing the tail is a measured loss in TWO layouts.

Algebra:
- The two pools compose into one 6x6x6, stride-6 max when the dims divide evenly.
- Otherwise, match the floor sizes: D1 = D/2, then Do = D1/3, and likewise for H and W.

Measured FAILURE 1: 0.967x (75.2 vs 72.7ms), own_kernel_time_frac about 0.05. The cuDNN conv was about 95% of runtime. The fused tail kernel used one thread per output (n,od,oh,ow), and each thread serially looped over all 64 channels x 216 window elements. Problems:
- Only about 16k threads for a ~1GB read.
- Poor memory-level parallelism: long serial load chains, and adjacent lanes sit 6 floats apart.
- The result (~3.7ms) was slower than the eager tail it replaced.
- The eager double-pool is already cheap, because pool1 is a well-parallelized coalesced pass that shrinks the data 8x. Collapsing the pools therefore saves little.

Measured FAILURE 2, WORSE: 0.9163x (79.3 vs 72.6ms), own_kernel_time_frac ~0.10, so the tail took ~7.9ms, about 2x the eager tail. Layout: one block per output (n,od,oh,ow) (~16k blocks of 64 threads), threads = channels, each thread serially scanning its own 6x6x6 window, then a shared-memory tree sum over channels. Adjacent threads read addresses D*H*W floats apart (channel stride), so EVERY warp load hits 32 different channel planes (fully uncoalesced), and each thread still runs a 216-long dependent load chain. Mapping threads to the channel axis is the WRONG axis for coalescing. Both attempts went CHECK then immediate SUBMIT with no timing, which hid the slowdown.

Do NOT fuse the composed 6x6x6 window per channel per thread (with threads over outputs OR over channels). If fusing at all, map consecutive threads to consecutive W addresses:
- e.g., one block per (n,c,od,oh) slab, with threads spanning the 6 rows x W (w fastest), a warp-shuffle max per output window, then atomicAdd of the channel max into a zeroed out[n,0,od,oh,ow]; or
- keep eager pool1 (or a coalesced custom pool1 kernel) and fuse only pool2 + channel sum on the 8x smaller tensor.
Time any such variant against the eager tail before adopting it.

The real lever is the conv (~90-95% of runtime):
- cuDNN TF32 within tolerance (str-00029);
- torch.backends.cudnn.benchmark=True;
- channels_last_3d (time it, per str-00051);
- the stride-2 phase decomposition (str-00030; non-shortcut shapes measured as a loss there, so time it).

General rule: when a pool or reduction tail shrinks the data by a large factor (here 216 x C), check BOTH that the fused kernel launches enough threads to saturate bandwidth AND that consecutive threads read consecutive (innermost-dim) addresses, before assuming fusion wins. Time locally before SUBMIT (str-00002).
[str-00056] helpful=1 harmful=0 :: NetVLAD / NetVLAD-with-ghost-clusters (KernelBench L3, labeled matmul family). This covers both variants: K=32 clusters + G=16 ghosts (L3_046), and the no-ghost variant (L3_047) with G=0. Example shapes: B=2048, N=100, D=512, so x is BN x D, about 420MB fp32. The reference pipeline is: x@clusters, then training-mode BatchNorm1d, softmax, drop the ghost columns, a_sum, A^T·x per batch element, subtract a_sum*clusters2, intra-cluster L2 normalize, flatten, and a global L2 normalize.

Runtime is dominated by two skinny, memory-bound passes over the large x (GEMM1 with N_out=K+G, and the batched A^T·x), plus about 6-8 eager passes in the tail.

SOLVED LOSS, measured THREE times: every attempt replaced only BN+softmax with a Triton row kernel (one program per row) and left both cuBLAS matmuls and the eager subtract/normalize tail untouched.
- Ghost variant: batch stats and running-stat updates were done in torch. The custom kernel was about 5% of runtime (own_kernel_time_frac 0.0495). It touched only the ~39MB logits tensor, so ~1.0x was the ceiling. It also failed the official run with 'candidate produced no output' after a correctness-only CHECK, and was submitted with no local full-size run. Separately, its turn-0 CHECK failed no_custom_kernel because bn_softmax_kernel never ran during forward (str-00031).
- No-ghost variant (G=0), run 1: 0.973x (3.83 vs 3.73ms), own_kernel_time_frac 0.043.
- No-ghost variant (G=0), run 2: 0.9537x (3.83 vs 3.66ms, lo95 0.951), own_kernel_time_frac 0.043. An identical copy of run 1's layout, scoring even lower.
In both no-ghost runs the kernel was correct and ran in both modes, yet lost. The per-forward BN fold in eager torch (mean, var, rsqrt, scale, shift and running-stat updates, about 6-8 tiny launches) plus an extra Triton launch replaced only one eager BN and one softmax over the 8MB logits. Both went CHECK then immediate SUBMIT on turn 1 with no timing, while this bullet was shown. This mirrors the 'cuBLAS + standalone epilogue' pattern (str-00011/str-00016). NEVER submit or resubmit this layout in either variant; it caps at ~1.0x or below, and a correct BN+softmax kernel is NOT progress.

GHOST VARIANT 'NO OUTPUT' FAILURE, NOW TWICE (second time with a partial minimum-fallback layout): load_inline eval-only BN-fold+softmax kernel (runtime-computed power-of-2 block size, extern + static shared arrays, BN=204800 blocks), per-forward torch fold ops still present, both cuBLAS matmuls kept, and a one-block-per-batch tail kernel fusing subtract + both normalizes. Turn-0 CHECK failed 'CUDA error: invalid argument' (launch-config class error); it was patched without a CUDA_LAUNCH_BLOCKING diagnosis, turn 1 CHECK passed, turn 2 SUBMIT with no full-size run, no timing and no train/eval sweep. Official: 'candidate produced no output', own_kernel_time_frac 0.1986 (kernels did run, so the failure is at full scale / under repeated timing, not compile). Tail-kernel flaws to avoid: idx -> (d = idx/K, k = idx%K) makes consecutive threads read bmm_out D floats apart (uncoalesced); a 128MB B x D*K vtmp written then re-read twice; shared-memory atomicAdd per element. Index with d as the fastest thread index (read bmm_out [b,k,d] coalesced, write out[b, d*K+k] via smem transpose), or keep each batch's K x D tile in smem/registers (K*D*4 = 64KB; needs care or split per k) so no temp tensor exists. Remove launch fragility: fixed 64- or 128-thread blocks for KG=48 (no runtime-computed sizes), static shared sizes, TORCH_CHECKs on shapes/contiguity/dtype, and a cudaGetLastError/C10_CUDA_KERNEL_LAUNCH_CHECK after each launch.

Decide on turn 0 to restructure the data flow, so x is read exactly twice with no intermediate full-size tensors:
(1) GEMM1: a Triton tf32x3 tl.dot kernel (str-00004). clusters is stored [D,K], so this is the native NN layout (not the nn.Linear NT trap of str-00075). Pad N_out up to a power of 2: 48 becomes 64, while 32 needs no padding when G=0. Its epilogue writes the small BN x (K+G) logits plus per-column sum/sumsq partials, which training-mode BN needs.
(2) A tiny finalize kernel folds gamma/beta/batch stats into per-column scale/shift. It updates running_mean/running_var in-kernel through the buffer tensors, using unbiased var*n/(n-1) and momentum (str-00014). In eval mode it uses the running stats instead. Do not do this fold with eager torch ops per forward.
(3) ONE kernel per batch element does the rest:
- apply scale/shift and softmax on the N x (K+G) logits, keeping only the first K columns (with G=0 the ghost slice is a no-op);
- compute a_sum;
- compute vlad = A^T·x_b (K x D) with tl.dot while the x_b tile streams once;
- subtract a_sum*clusters2;
- per-cluster L2 norm over D, then the global L2 norm over D*K. Both are cheap once the per-batch K x D vlad tile is in registers or SMEM; never materialize the full B x K x D vlad (~128MB) in global memory;
- one store of the D*K output in the reference's B x (D*K) layout, i.e. a D-major then K order after the transpose(1,2).
Match F.normalize's eps (1e-12 clamp on the norm).
Minimum fallback if (3) is infeasible: keep the bmm but fuse the a_sum*clusters2 subtract and BOTH normalizes into one COALESCED kernel over the B x K x D bmm output (one program per batch element, no full-size temp), plus the in-kernel BN finalize, so the eager tail passes and torch fold launches disappear.

Before SUBMIT (mandatory here, ignored twice): CHECK, then run the forward at the full get_inputs shape under CUDA_LAUNCH_BLOCKING=1, many iterations, in BOTH model.train() and model.eval(), printing peak memory, time and max error vs the reference (str-00002, str-00082). A sub-1x local timing or any CUDA error must block SUBMIT.
[str-00057] helpful=0 harmful=0 :: ResNet-style residual CNNs in eval mode at small batch (KernelBench network family, e.g., ResNet18 L3_009 at batch 2, 224px, ref ~3.2ms). These are launch- and dispatch-bound (see str-00026). The reference's residual tail `out += identity; relu(out)` is already in place and cheap (one pass), so a custom add+ReLU kernel saves no memory pass.

SOLVED LOSS, measured TWICE at ~0.71-0.73x. Do NOT resubmit either layout.

Measured FAILURE 1: 0.73x (4.34 vs 3.18ms), 0.69x vs torch.compile, own_kernel_time_frac 0.017. The candidate did three things:
- folded BN once into conv weights, lazily cached in a dict plan;
- replaced module calls with per-block F.conv2d calls carrying the folded bias;
- added one standalone in-place add+ReLU CUDA kernel per block.

Measured FAILURE 2: 0.708x (4.56 vs 3.23ms, lo95 0.707), own_kernel_time_frac 0.05. It followed the str-00038 MobileNetV1 recipe despite that bullet's scope warning:
- BN folded into conv weights, cached in self.<x> attributes keyed on data_ptr;
- F.conv2d WITHOUT bias, then a separate in-place bias+ReLU kernel after conv1 (and the stem);
- a bias+residual+ReLU kernel replacing the in-place add_+relu_;
- the downsample conv given its folded bias via cuDNN.
The training branch also allocated torch.zeros(C) for every kernel call (extra launches). It went CHECK then immediate SUBMIT with no timing.

In both runs the custom kernels were a small share of runtime; the slowdown came from extra launches and Python dispatch, not from the kernels. This is the same failure as str-00026 failure 2 (EfficientNetB0, 0.80x). The MobileNetV1 ~1.07x recipe (str-00038) does NOT transfer here. It works only where the reference runs separate eager BN and ReLU passes that can be collapsed; in ResNet, conv(no bias) + separate bias/ReLU kernel has the same or MORE launches than eager conv+BN+relu_.

What can actually beat eager here is cutting launches, keeping the launch count at or below eager's:
(a) CUDA-graph the whole eval forward with static input and output buffers. The custom kernel must still run inside the graphed path. (Check first that the static checker allows it; str-00014 (e) suggests graphs may be rejected.)
(b) Fuse conv + bias + residual + ReLU in one conv epilogue (e.g., a TF32 implicit-GEMM conv within tolerance), so no separate elementwise launch exists.
(c) Fuse the tail: avgpool + fc in one kernel.

Also avoid these:
- A training-path `out = out + identity` followed by a kernel call with torch.zeros_like / torch.zeros(C) per call. That is an extra allocation and an extra launch/pass.
- Delegating turns to a weaker model. One such turn dropped the _make_layer method and wasted a CHECK.

Route the custom kernel through the eval branch (str-00031). MANDATORY: time the evaluated mode at full size against the reference before SUBMIT (str-00002); both failures were CHECK-then-SUBMIT flows that hid a ~30% slowdown.
[str-00058] helpful=50 harmful=0 :: Index-width pitfall in load_inline kernels: pass element counts (numel/total) and loop bounds as int64 (long / int64_t), not int. A kernel that casts total to int while using a long loop index works at ~126M elements but silently overflows (wrong or missing output) once numel exceeds 2^31, which large conv outputs (>8GB-scale or batched 3D) can hit. Keep the kernel signature, the host-side computation and the cast all 64-bit; only per-channel indices like c may be int.
[str-00059] helpful=3 harmful=1 :: Depthwise conv (groups == in_channels; KernelBench L1 conv_depthwise family) is purely memory-bound, and cuDNN's depthwise path is weak. This holds both for FEW channels (vertical (K,1) kernel, N=64, C=8, 512x512 fp32) and for MANY channels (C=128, asymmetric 3x7 kernel, N=32, 128x256 fp32). Unlike dense convs (str-00027), a hand-written load_inline CUDA kernel reliably beats it.

Measured:
- WIN (vectorized recipe below), C=8, (K,1) kernel: 3.44x vs eager (2.93 -> 0.85ms) and 15.3x vs torch.compile, own kernel 100% of runtime. It went CHECK, CHECK, SUBMIT.
- CONFIRMED on the same C=8 (K,1) task (L1_083): 3.068x (2.93 -> 0.954ms, lo95 3.058), 13.7x vs torch.compile, own kernel 100%. Same float4 recipe (one thread per float4 of 4 adjacent outputs, register weights templated on KH=3/5/7, float4 input-row loads and stores, scalar fallback), BUT with a FLAT 1D grid-stride loop capped at 4096 blocks and per-element 64-bit div/mod decode of (ow4, oh, c, n) instead of the 3D grid. The fast path was gated on stride==1 && OW==W && W%4==0 (this also routes int padding>0, which widens OW, to the scalar path, avoiding the latent bug below). It went CHECK on turn 0 then SUBMIT on turn 1 with no timing. The ~11% shortfall vs 3.44x is likely the grid cap (str-00054 measured a cap ~9% slower) plus the flat div/mod decode, or evaluator variance. Treat ~3.1-3.4x as the SOLVED BASELINE for this shape: resubmitting cannot raise the score. Next attempt: use the 3D grid (x=OW/4, y=oh, z=n*C+c) with 1 thread per float4 and NO grid cap, try several output rows per thread reusing the KH input rows from registers, and time each against the reference after the first passing CHECK instead of submitting immediately.
- Modest WIN (plain SCALAR kernel), C=128, 3x7 kernel: 1.21x (3.74 -> 3.09ms), own kernel 100% of runtime. It used one thread per output pixel, the same 3D grid, template<KH,KW> unrolled taps, per-tap bounds checks, weights re-read from global memory every tap, and a generic fallback. It went CHECK then SUBMIT with no timing. Do not settle for the scalar kernel.
- Better WIN (4 outputs per thread, SCALAR loads/stores), same C=128, 3x7, N=32, 128x256 task: 1.347x (3.73 -> 2.77ms, own kernel 100%). Same 3D grid with x = OW/4 chunks, template<KH,KW>, weights in a register array, 4 accumulators per thread, per-output bounds checks on each store. It went CHECK then SUBMIT with no timing, so levers (1)-(3) below remain untried. Treat ~1.35x as the baseline for this shape.

MISALIGNED-ADDRESS PITFALL (measured crash): a float4 reinterpret on output rows or input taps crashed with 'CUDA error: misaligned address' at W=256, KW=7 (OW=250). The row base offset oh*OW is not 16-byte aligned when OW%4 != 0, and the input tap offsets ow0+kw (kw != 0) are never 16-byte aligned. Use real float4 only behind an explicit alignment guard: OW%4==0 for output stores, W%4==0 plus 16B-aligned base pointers for input row reads, and only at aligned offsets. Read the shifted taps as scalars, or as one aligned float4 load plus register shuffling. Otherwise keep the 4-outputs-per-thread structure with scalar loads and stores, which still beats one output per thread.

Recipe:
- One thread per 4 adjacent outputs (float4 store only when aligned per the pitfall above; else scalar stores with per-output bounds checks).
- 3D grid: x = ceil(OW/4) chunks (128 threads per block), y = output row oh, z = n*C + c. This needs no per-element div/mod; c = blockIdx.z % C once per thread. Prefer this over a flat grid-stride loop with a block cap (measured ~3.07x vs 3.44x on the C=8 task).
- Template on the kernel size (template<int K> dispatching K=3/5/7, or template<int KH,int KW> for asymmetric kernels), with the per-channel weights loaded into a register array once per thread (or into shared memory per block), and the taps unrolled.
- Initialize the accumulator with the bias.
- Read the input rows directly with a bounds check on ih = oh*stride - pad + k*dil. Neighbouring output rows share these rows through L1/L2, so no shared-memory staging is needed.
- Keep a scalar generic kernel as the fallback for stride != 1 or other kernel sizes.
- Use 64-bit base offsets (str-00058). Put the wrapper and stream in cuda_sources (str-00001). Read the weight and bias from the nn.Conv2d submodule, and pass an empty tensor when bias=None.

Untested levers for wide (KW>1) kernels; time each locally:
(1) For the 4 outputs of a thread, load the KW+3 input values of each row once into registers and slide the window over them, so input loads are shared across outputs.
(2) A padding-0 fast path with no bounds checks in the inner loop.
(3) Each block computing several output rows, so the KH input rows are reused from registers or shared memory.

CORRECTNESS CAVEAT (latent bug in the 3.44x kernel): nn.Conv2d with an int padding/stride applies it to BOTH dims, even for a (K,1) kernel. W_out = (W + 2p - dil*(KW-1) - 1)/stride + 1, not W. A vec4 path that writes width W is valid only for stride 1 and no width shrink. Always compute OW generally, and route stride != 1 or unaligned cases to scalar code with proper OW indexing. A guard of stride==1 && OW==W && W%4==0 handles this correctly (padding>0 makes OW != W).

The same layout (register weights, a thread per output vector, 2D/3D grid) should extend to square KxK depthwise kernels. In that case, stage an input tile plus halo in shared memory if L1 reuse proves insufficient (untested).
[str-00060] helpful=4 harmful=0 :: Conv3d → Softmax(dim=1) over a small C (e.g., 16) → MaxPool3d(2) → MaxPool3d(2), KernelBench conv family (L2_006: IC=3, OC=16, k=3, input 128x3x16x32x32).

Measured WIN, CONFIRMED THREE TIMES with this exact recipe:
- Run 1: 1.32x vs eager (1.89 → 1.43ms) and 1.17x vs torch.compile. Went CHECK then SUBMIT with no timing.
- Run 2: 1.2934x (1.80 → 1.39ms, lo95 1.293), own_kernel_time_frac 0.077. Turn-0 SUBMIT with no CHECK or timing; it compiled and was correct only because the recipe was copied verbatim (str-00001 layout, long indices per str-00058).
- Run 3: 1.2888x (1.81 → 1.40ms, lo95 1.287), own_kernel_time_frac 0.072. Recipe copied verbatim (template<OC=16>, 4096-block grid-stride cap, 256 threads, __expf); CHECK on turn 0, SUBMIT on turn 1, no timing, no conv lever tried. It gained NO new information.
~1.29-1.32x is the SOLVED BASELINE for this layout: resubmitting it cannot raise the score. The own kernel is only ~7-8% of runtime, so the cuDNN conv (~92-93%) sets the ceiling and kernel tweaks cannot move the score. Do NOT go CHECK -> immediate SUBMIT again.

Recipe:
- Keep the cuDNN conv.
- Write ONE load_inline kernel that does the softmax and both pools; put the wrapper, stream retrieval and launch in cuda_sources with a declaration-only cpp_sources (str-00001).
- Compose the two 2x2x2 pools into a single 4x4x4, stride-4 window.
- Size the output with floor composition: Do = (D/2)/2, likewise for H and W. Bounds-check window elements with `continue`.
- Use one thread per FINAL pooled output (n,od,oh,ow), with ow fastest; keep total and base offsets as long (str-00058).
- For each window element, load the C channel values into a `float vals[OC]` register array (#define OC). Compute max, exp (__expf passed tolerance), sum and normalize in registers. Then update a per-channel running max `runmax[OC]`.
- Write the OC outputs at the end.

Ordering: softmax must run per element BEFORE the max. Softmax is not monotone per channel, so max-first on raw values is invalid. Contrast str-00044, where the pool precedes the softmax.

Why fusion paid off here but not in str-00055: the eager tail includes a full softmax read+write pass over the conv output, and the per-thread work is small (64 elements × 16 channels). About 25k threads were enough. With larger windows or more channels, re-check the thread count, per str-00055.

Plan from turn 0: paste this tail kernel, CHECK it, then spend EVERY remaining turn timing conv-side levers at full size against the reference (~1.40ms baseline; bounded warmup/iterations, cuda synchronize) and keep the fastest. NOTE: torch.backends flags (cudnn.benchmark, cudnn.allow_tf32) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Realistic levers (all untested):
(a) A fused direct conv with the softmax+4x4x4-max epilogue: IC=3, k=3, so the weight is only 16x81 floats and fits in smem/constant memory. Each thread (or block tile) computes the conv values for all 16 OC at its window positions, applies softmax per position in registers and keeps the running max, so the ~0.7GB 128x16x14x30x30 conv output is never written or re-read. Stage the input tile plus halo in smem; verify indexing against nn.Conv3d on a small shape first.
(b) channels_last_3d via your own transpose kernel (str-00102 pattern, timed per str-00051); if adopted, make the tail kernel read each pixel's C values contiguously.
- Minor: float4 loads or a 2D grid in the tail kernel (small, since the tail is ~7-8% of runtime).
Resubmit the baseline only if no lever beats ~1.40ms.
[str-00061] helpful=3 harmful=0 :: Conv3d → MaxPool3d(2) → logsumexp over C (dim=1, keepdim) → ReLU (KernelBench conv family, e.g., Conv3d 32->64, k=3, N=4, 32x128x128 input, conv ~97-98% of runtime).

The tail-only layout is CONFIRMED TWICE at ~1.00-1.016x. It keeps the cuDNN conv and adds ONE load_inline kernel: one thread per pooled output (ow fastest); for each channel, the max of the raw 2x2x2 window; then an online LSE over C in registers; then ReLU; writing the [N,1,Do,Ho,Wo] output. This replaces the eager pool, the multi-pass logsumexp and ReLU with one read of the conv output plus a 1/8-size write.
- Run 1: 1.0046x vs eager (27.54 → 27.41ms), 0.9996x vs torch.compile, own_kernel_time_frac 0.032. It went CHECK then immediate SUBMIT with no timing.
- Run 2: 1.0157x (27.57 → 27.15ms, lo95 1.014), 1.009x vs compile, own_kernel_time_frac 0.018. It was a turn-0 SUBMIT with no CHECK and no timing. It was correct only by luck, since a compile_error would have scored 0 (str-00002). It tried no conv lever.
Do NOT resubmit this layout on turn 0 expecting more. It is the baseline and the custom piece to keep. The gain has to come from the conv.

Correctness notes, all validated:
- Max-first on raw values is exact here, because the reference pools BEFORE the channel reduction (same as str-00044, contrast str-00006/str-00060).
- Do not keep a float cmax[C] array (C=64 or 128) per thread, since it risks register spills. Use an online LSE in one channel loop instead: if mx > m, then s = s*exp(m-mx)+1 and m = mx; else s += exp(mx-m). Finally lse = m + log(s).
- A templated-C dispatch (if C==64 / 32 / ...) must have a generic fallback, either a runtime-C kernel or TORCH_CHECK. Otherwise an unsupported C silently returns the uninitialized torch::empty output.
- Keep total and base offsets as long (str-00058). Use the str-00001 layout. x.contiguous() on the cuDNN output is a no-op.

Plan from turn 0:
- Write the tail kernel.
- CHECK it.
- Spend the remaining turns timing conv-side levers locally against the reference module, with bounded warmup and iterations and cuda synchronize, and keep whichever is fastest.
The levers:
(1) cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv), only after confirming the tolerance (str-00029).
(2) channels_last_3d, timed per str-00051. If you adopt it, make the tail kernel read the C values of each pixel contiguously.
(3) torch.backends.cudnn.benchmark=True.
(4) A conv whose epilogue does the pool and LSE, so the large conv output is never written.
[str-00062] helpful=2 harmful=0 :: Training- or eval-mode BatchNorm followed by AvgPool (one or a chain of AvgPool3d/2d), KernelBench conv_transpose/conv family. Example (L2_072): ConvTranspose3d(stride 2) → BatchNorm3d → AvgPool3d(2) → AvgPool3d(2).

Algebra: once the batch stats are known, BN is a per-channel affine map (scale_c*x + shift_c). AvgPool is linear, so the two commute. Pool the RAW conv output, then apply scale_c*pooled + shift_c only on the small pooled tensor. The normalized full-size tensor is never written. This generalizes the BN + global-avgpool collapse (str-00010) to non-global average pools.
- Two AvgPool(2) layers compose into one 4x4x4, stride-4 window.
- floor(floor(D/2)/2) == floor(D/4), so Do = D/4 with no bounds checks.
- This does NOT apply to MaxPool when gamma can be negative (see str-00037/str-00046).

Measured WIN, CONFIRMED THREE TIMES with the layout below:
- Run 1: 1.65x vs eager (36.3 → 22.0ms), only 1.005x vs torch.compile (which finds the same reordering).
- Run 2: 1.6491x (36.40 → 22.07ms, lo95 1.648), own_kernel_time_frac 0.079. Compiled first try (str-00001 layout, long sizes per str-00058), CHECK on turn 0 then SUBMIT on turn 1 with no timing and no lever tried.
- Run 3: 1.6438x (36.45 → 22.17ms, lo95 1.6387), own_kernel_time_frac 0.079. Recipe copied verbatim (grid-stride pool kernel, 4096-block cap, double-math finalize, separate eval finalize); CHECK on turn 0, SUBMIT on turn 1, no timing, no lever. It gained NO new information.
The cuDNN conv is ~92% of runtime, so ~1.64-1.65x is the SOLVED BASELINE for this task: resubmitting this layout cannot raise the score. Do NOT go CHECK -> immediate SUBMIT; spend the spare turns after the first passing CHECK timing the levers below at full size against the ~22.1ms baseline, keeping the fastest, and resubmit the baseline only if none beats it.

Layout used (3 kernels, launched from cuda_sources per str-00001):
(1) Stats kernel: one block per (n,c), fp32 accumulation, shared-memory tree reduction, atomicAdd into zeroed per-channel sum/sumsq. (If shared arrays are hardcoded, e.g. [256], they must match the launch block size; size them from a compile-time constant shared with the launch.)
(2) Finalize kernel: biased var for scale/shift; unbiased var*n/(n-1) with momentum for running_var. Update running_mean/running_var in-kernel, with the buffers passed as tensors (str-00014). Use the running stats in eval mode (separate eval finalize kernel so a custom kernel runs in both modes).
(3) Fused pool + affine kernel: one thread per pooled output, summing the 64-element window × 1/64, then *scale[c] + shift[c].

Untried headroom; time each. NOTE: torch.backends flags (cudnn.benchmark, cudnn.allow_tf32) are REJECTED by the static checker (str-00014 (e)), so they are NOT available. Realistic levers:
(a) Fold sum/sumsq accumulation into the pool kernel, so the conv output is read once instead of twice (removes ~1/3 of custom-kernel traffic). With non-divisible sizes (e.g., 63 → 15*4 = 60 covered), the pool windows miss the edge slabs, and those elements must still be added to the stats.
(b) Conv side (the main payoff, needed to beat compile): a fused direct transposed conv with a pool(+stats) epilogue, plausible when IC/OC are tiny (L2_072: IC=3, OC=16, weight 1296 floats in smem/constant; per dim only stride-2-valid taps, see str-00044 lever 1), so the full conv output is never written; channels_last_3d via your own transpose kernel (str-00051, transposed convs measured a loss there, so time it); the stride-2 phase decomposition (str-00030) is non-shortcut for k=3,s=2,p=1, likely a loss.
(c) Use a 2D grid or float4 loads in the pool kernel instead of per-element 64-bit div/mod.
(d) Drop the torch.zeros launches for the accumulators.
Kernel levers (a), (c), (d) touch only ~8% of runtime.

General rule: any per-channel affine (eval BN, or training BN after its stats) followed by a LINEAR reduction (avg pool, mean, sum) can be applied after the reduction.
[str-00063] helpful=1 harmful=0 :: Time-looped single-layer RNN with a SMALL hidden state (KernelBench rnn family, e.g., L3 VanillaRNNHidden: per step cat(x[t], h) -> i2h Linear -> tanh -> h2o Linear, with H=256, B=8, T=256, in=1024; ref ~28-29ms of per-step eager launches). This is a real recurrence, unlike the single-step cell in str-00041. It is also not the large multi-layer or gated case of str-00018.

Measured WIN, CONFIRMED TWICE with this recipe, both via CHECK then SUBMIT and own kernel at ~95% of runtime:
- 9.02x (27.97 -> 3.10ms).
- 11.03x (29.13 -> 2.64ms, lo95 10.73).
The 'excessive_speedup>10x: audit' flag is expected for this legitimate restructuring.

Recipe:
(1) Split the i2h weight into Wx = W[:, :in] and Wh = W[:, in:], so torch.cat is never needed.
(2) Precompute the input projection for ALL timesteps as one cuBLAS GEMM, with the bias folded in: F.linear(x.reshape(T*B, in), Wx, b).
(3) Run the whole recurrence in ONE persistent single-block load_inline kernel (wrapper and stream in cuda_sources, declaration-only cpp_sources, per str-00001; compiled first try), e.g. template<int H,int B> launched <<<1, H>>>:
- Use H threads, with thread j holding W_hh row j in a float wrow[H] register array for all T steps.
- Keep h for all B batches in shared memory (__shared__ float hs[B*H]), so the h[k] reads are broadcasts.
- Each step: acc[b] = xproj[t,b,j] + sum_k w[k]*h[b,k] for all B, then __syncthreads, then tanhf, then write h to both shared memory and hout[t], then __syncthreads.
(4) Run h2o as one batched GEMM (F.linear) over all T*B hidden states afterwards.

HARD SCOPE GUARD: a single block suffices ONLY when the recurrent weights are held on-chip (registers or smem) for all T steps. That requires gates*H*H*4 bytes to fit within one SM: ~256KB for vanilla H=256, which is borderline. Do NOT reuse this recipe for GRU/LSTM or multi-layer stacks at H=256 (768KB-1MB per layer). Never let threads stream their W_hh rows from global memory inside a single block.

Measured: a <<<1, H>>> GRU (L=6, B=10, T=512) streaming W_hh from global scored 0.0071x (4480 vs 31.7ms). Use the multi-block hidden-unit-partitioned grid-sync design of str-00018/str-00074 for larger H, multiple layers or gated cells.

Remaining time is inside the serial recurrence kernel. Untested headroom:
- A float wrow[256] per thread risks register spills. Split k across warps, or use 2+ threads per row with a shuffle reduction.
- Template or dispatch over H and B, with a generic fallback, instead of a hardcoded TORCH_CHECK (fragile, though sufficient for the fixed test shape).
- Cache the contiguous Wx/Wh slices, keyed on weight.data_ptr(), instead of re-slicing every forward.
- Write the h2o output directly instead of materializing hout, only if that is cheaper.

Time locally before SUBMIT if turns allow (str-00002). With the exact proven shape, CHECK-then-SUBMIT has reproduced the result. On any other shape or cell type, a full-size timing run is mandatory.
[str-00064] helpful=5 harmful=0 :: Matmul tasks whose operands carry a structural label that implies NO zero work, e.g., KernelBench L1 'Matmul_for_symmetric_matrices' computing C = A @ B with A and B both symmetric (N=4096 fp32). Symmetry does NOT reduce FLOPs for a general product: C = A@B is not symmetric, so no output tiles can be skipped. Since B = B^T, B could be read in K-contiguous orientation, but the NN layout is already native for tl.dot, so this gains nothing. Contrast this with triangular, banded or block-sparse operands, where skipping structural zeros wins big (str-00019). Before writing the kernel, decide whether the label implies zeros (exploit them) or only redundancy that the product does not inherit (treat it as a plain dense GEMM).

Treat the symmetric case as plain dense A@B and apply the str-00004 tf32x3 recipe unchanged:
- tl.dot(..., input_precision='tf32x3');
- autotune over BM/BN in {128,256}, BK in {32,64}, num_warps=8, num_stages 3-4;
- grouped pid ordering (GM=8);
- 1D kmask broadcast per operand (str-00042).

Measured, CONFIRMED TWICE with this exact recipe (masked K loop, %M/%N index wrap), own kernel 100% of runtime, correct:
- Run 1: 1.43x (7.31 -> 5.11ms).
- Run 2: 1.4261x (7.30 -> 5.12ms, lo95 1.425).
Both were turn-0 SUBMITs with no CHECK or timing; they worked only because the recipe was copied verbatim. ~1.43x is the SOLVED BASELINE for this layout: resubmitting it cannot raise the score. For any novel kernel, still CHECK first (str-00002).

Next attempt on this task: CHECK, then time these headroom levers locally against torch.matmul (bounded warmup/iterations, cuda synchronize) before SUBMIT, keeping the fastest:
(1) plain tf32 (input_precision='tf32') after printing the max abs/rel error against the reference;
(2) an unmasked K loop and no %M/%N index wrap, since 4096 divides every tile size;
(3) a FIXED BM=BN=128, BK=64, 8-warp, 3-stage config instead of the autotune set: on the str-00032 task, autotune including 256-wide tiles measured ~20% slower than this fixed config;
(4) more stages (4) or larger-BK variants around the 128x128 config.
[str-00065] helpful=0 harmful=0 :: UNet-style networks with repeated conv→BatchNorm2d→Softmax(dim=-1) pairs (KernelBench network family, e.g., UNetSoftmax: about 18 BN→softmax sites, tensors up to 8x64x64x512 fp32, checked at atol=rtol=1e-4).

Measured FAILURE: 0.9966x (23.43 vs 23.35ms), own_kernel_time_frac 0.0068. The candidate kept the whole eager reference and replaced only the tiny final 1x1 conv with a naive pointwise kernel. This is the network-family 'eager/cuDNN + trivial kernel' layout (see str-00026 failure 3 and str-00027 failure B): it removes no memory passes and lands at parity. It was submitted right after its first passing CHECK with no timing. The three earlier CHECKs had failed at 1e-4 and were abandoned without printing the max abs/rel error (str-00024).

The real lever is the BN→softmax chain. Each site is 2 eager full passes over a large tensor; collapse each one into ONE row kernel.
- Eval mode (verify the mode first; str-00003/str-00031): softmax over the last dim W is per (n,c,h) row, so BN's per-channel shift (beta_c - mean_c*scale_c) is constant along the row and cancels exactly. Only scale_c = gamma_c*rsqrt(running_var_c+eps) matters, and softmax(scale_c*x). Fold scale once and cache it as self.<x> attributes; never recompute it per forward (str-00015).
- Training mode: the shift still cancels, but scale_c needs batch variance, so a per-channel stats reduction is required first (use the many-block layout of str-00043). Update running stats in-kernel if needed.
- Kernel layout: one warp (W<=1024) or one block per W-length row, with the row held in registers (e.g., 512/32 = 16 floats per lane). Do a shuffle max and sum, then a single write: 1 read + 1 write instead of the eager BN read/write + softmax read/write.
- Further pass removal: fuse each MaxPool2d(2) into the softmax kernel that feeds it. Have each block handle the 2 rows of a pooling window, writing both the full-res output (needed for the skip connection) and the pooled output. Also preallocate each decoder's concat buffer and write the upconv output and the encoder skip output directly into their channel slices, so no torch.cat copies are made.
- Numerics at 1e-4: use expf, not __expf or fast-math, and keep the max-subtracted form. If CHECK fails, print the max abs/rel error and compare a single site against F.softmax(bn(x), -1) before retreating.

Keep the cuDNN 3x3 convs and transposed convs. Time locally against the reference before SUBMIT (str-00002).
[str-00066] helpful=1 harmful=0 :: Small-feature-map CNNs at LARGE batch (KernelBench network family, e.g., L3_004 LeNet5 with batch 4096, 1x32x32 input, conv->ReLU->MaxPool2d(2) twice, then tiny FCs 400->120->84->classes; ref ~2.3ms).

Measured, CONFIRMED TWICE with the recipe below:
- Run 1: 1.0855x vs eager (2.30 -> 2.12ms) and 1.064x vs torch.compile, own_kernel_time_frac 0.052. It also added standalone in-place ReLU kernels after each F.linear.
- Run 2: 1.0885x (2.30 -> 2.11ms, lo95 1.087), own_kernel_time_frac 0.045. It dropped the FC ReLU kernels and kept eager F.relu.
Both went CHECK on turn 0 then SUBMIT on turn 1 with no local timing and no lever tried. ~1.09x is the SOLVED BASELINE for this task: resubmitting it cannot raise the score.

Recipe:
- Keep the cuDNN convs and cuBLAS linears.
- Collapse each conv->ReLU->MaxPool2d(2) into ONE load_inline kernel with one thread per pooled output: grid-stride, 256 threads, 4096-block cap, long totals and indices (str-00058). Take the max of the RAW 2x2 window, then apply ReLU once, since ReLU is monotone and commutes with max (str-00045). This removes the full-size ReLU write and re-read: 1 read plus a quarter-size write.
- Put the wrappers and stream retrieval in cuda_sources (str-00001). It compiled first try both times.
- The kernel assumes even H and W (28 and 10 here), which matches MaxPool2d's floor only for even dims. Use OH = H/2, OW = W/2 and index generally if reusing on odd shapes.

What added nothing (now measured): standalone in-place ReLU kernels after F.linear. Run 2 removed them and scored the same, so they are confirmed NEUTRAL. They replace eager relu one-for-one with no pass reduction (str-00003/str-00026).

The own kernel is only ~4.5-5% of runtime; the rest is cuDNN/cuBLAS plus launch overhead, so further gains need fewer launches or a changed conv. Next attempt: CHECK on turn 0, then time each of these at full size against the reference (bounded warmup/iterations, cuda synchronize) and keep the fastest, instead of submitting the baseline:
(a) Fuse the whole FC stack (fc1->relu->fc2->relu->fc3) into one kernel. These widths are tiny, so the weights (~58K floats) fit in shared memory and each block can process a batch of rows.
(b) Fold bias+ReLU into the GEMM: torch.addmm with the bias, then an in-place relu_, or a cuBLASLt RELU_BIAS epilogue.
(c) CUDA-graph the forward with static input/output buffers, keeping the custom kernels inside the graph (str-00026 4a).
(d) cuDNN TF32 (within tolerance; print the max error first, str-00029) and torch.backends.cudnn.benchmark=True.
(e) A fused direct conv+ReLU+pool kernel, so the conv output is never written. This is plausible because the channel counts are small (1->6, 6->16) and the weights are tiny.

Time locally before SUBMIT (str-00002) to see whether the convs or the FC/launch overhead dominate.
[str-00067] helpful=0 harmful=1 :: Linear -> NONLINEAR op -> reduction over output features (KernelBench matmul family, skinny-M huge weight: M=128, K=N=32768 fp32, ~4.3GB weight, checked at atol=rtol=1e-4). Examples: Linear -> sigmoid -> torch.sum(dim=1, keepdim=True) (L2_056), and Linear -> MaxPool1d(k) over features -> torch.sum(dim=1) -> *scale (L2_055).

The linearity fold of str-00013 does NOT apply here. sum_j sigmoid((xW^T+b)_j) is not linear in W, and neither is MAX pooling (only AVG pooling folds), so W cannot be pre-summed or pre-pooled. The fold is valid only when every op between the GEMM and the reduction is linear (scale, add-constant, avg-pool).

Measured parity, twice. Both used the 'cuBLAS + standalone epilogue' ~1.0x layout of str-00011:
- sigmoid+sum: cuBLAS nn.Linear plus a Triton per-row sigmoid+sum kernel. It was correct but scored 1.0006x (16.04 vs 16.05ms), own_kernel_time_frac 0.0011. The turn-0 CHECK had failed at 1e-4 for an unrecorded reason. The attempt then retreated without printing the max abs/rel error (str-00024) and submitted with no timing.
- maxpool(2)+sum+scale: cuBLAS nn.Linear plus a Triton per-row even/odd max+sum kernel. It was correct but scored 1.0007x, own_kernel_time_frac 0.0012. Earlier fused-GEMM attempts behaved as follows:
  - One failed to compile on a non-constexpr tl.arange bound (see the Triton constexpr pitfall bullet).
  - Three failed CHECK at 1e-4 with no error printout, so it is unknown whether precision or a pooling/indexing bug was the cause.
  - The attempt then retreated and submitted with no timing.

Measured CATASTROPHIC FAILURE of the fused path (sigmoid+sum): a fused tf32x3 Triton GEMM with bias+sigmoid+row-sum in the epilogue passed CHECK but scored 0.034x (475.9 vs 16.1ms, own_kernel_time_frac 1.0), about 30x slower than cuBLAS. Its config: fixed BM=BN=128, BK=64, num_warps=8, num_stages=3, no autotune. W.t() was read through swapped strides as a [BK,BN] tile with K contiguous, and the K loop was masked every iteration. Partials went to an [M, N/BN] buffer, then torch .sum. The grid was 1x256 programs, so underfill does not explain the loss. Likely causes are register spill and shared-memory pressure from tf32x3 at 128x128x64 with 3 stages. It went CHECK then immediate SUBMIT with no timing. Turn 0 had failed CHECK at 1e-4 with no error printed.

Measured SUB-1x of a fused path (sigmoid+sum, second try): fused tf32x3 GEMM with the CORRECT swapped-stride [BK,BN] W read (W.stride(1), W.stride(0)), FIXED BM=128, BN=64, BK=64, num_warps=8, num_stages=3, no autotune, K loop with only M/N masks, 1x512 tiles (no underfill), epilogue bias -> sigmoid -> tl.where mask -> tl.sum(axis=1) -> tl.atomic_add into a torch.zeros [M] buffer. Correct at 1e-4 (tf32x3 at K=32768 and fp32 atomics both pass), but 0.618x (25.90 vs 16.01ms, own kernel 100%). Process: turn 0 failed CHECK at 1e-4 (error not printed), turn 2 lost to no_code (str-00008), final SUBMIT with no timing.

Measured SUB-1x of a fused path (maxpool(2)+sum+scale, L2_055): fused tf32x3 GEMM, FIXED BM=128, BN=64, BK=64, num_warps=8, num_stages=4, no autotune, unmasked K loop, W loaded as [BN,BK] tiles with tl.dot(x, tl.trans(w)), epilogue bias -> tl.max(tl.reshape(acc,(BM,BN//2,2)),2) -> tl.sum(axis=1) -> *scale -> tl.atomic_add into a torch.zeros [M] buffer. Correct at 1e-4 (the reshape-based pool compiles because its shape is constexpr-folded, str-00070; fp32 atomics passed), but 0.9159x (17.52 vs 16.05ms, own kernel 100%). Grid was 512 tiles, so not underfill. Process: turn 0 failed CHECK at 1e-4 (cause not printed), turn 1 lost to no_code (str-00008), turn 2 passed, turn 3 SUBMIT with no timing; one timing run would have caught it while a turn remained.

KEY RULE for this M/K/N (revised): three fused GEMM+reduction-epilogue kernels have all lost (0.034x, 0.618x, 0.916x), while the ONLY >1x kernel at this shape (str-00040, 1.33x) used autotune AND a light elementwise epilogue with a plain store. It is UNPROVEN whether the GEMM config or the reduction epilogue (tl.sum + atomics) is to blame; do not assume the epilogue is free. Every loss also used a FIXED config instead of str-00040's autotune set. Staged plan:
(1) Copy the str-00040 kernel VERBATIM (tf32x3, swapped-stride [BK,BN] W read, BM=128, autotune BN {64,128}, BK {32,64}, num_warps=8, num_stages 3-4, unmasked K loop, plain 2D grid) with a LIGHT epilogue (bias + sigmoid, or bias only for maxpool) and a plain store into an [M,N] buffer (16MB here; no reduction, no atomics).
(2) Follow with a tiny Triton per-row reduction kernel over that buffer (sum, or pair-max+sum*scale).
(3) Time the GEMM alone against F.linear, then the full pipeline against the ~16ms reference, before SUBMIT.
(4) Only then try moving the row-sum into the epilogue, and keep it only if re-timing shows no regression.
(5) If every variant stays below 1x, submit cuBLAS + the row kernel (~1.0x).
Do NOT carry the [BN,BK]+tl.trans lever from the compute-bound NT bullets (str-00034/str-00075) into this weight-bandwidth-bound regime: it measured 0.916x here.

Rules:
(1) After the first passing CHECK of any fused variant, ALWAYS do a bounded full-size timing against the reference (nn.Linear + epilogue) before SUBMIT. The cuBLAS + Triton row-kernel layout reliably gives ~1.0x and is the fallback. Never submit a sub-1x fused kernel when that fallback exists. Use the turn after a passing CHECK for timing, and include the full ```python block in every reply (str-00008), since three attempts on these tasks lost a turn or submitted untimed.
(2) If the fused path is slow, do not tune blindly. First time the GEMM alone against F.linear, comparing:
  - The str-00040 autotune set (BN=64/128, BK=32/64, stages 3-4), with BLOCK_M >= M so the weight streams once, and a tile count of at least ~2x the SM count.
  - Fewer stages (2-3) to cut shared-memory and register pressure under tf32x3.
  - An unmasked K loop when K % BK == 0 (32768 is divisible by any BK).
  - Split-K or plain tf32 only per str-00040's lever list (plain tf32 likely fails 1e-4).

Fused epilogue details:
- Use tl.dot(input_precision='tf32x3') for the 1e-4 tolerance, never plain tf32.
- Epilogue for sigmoid: add bias, apply sigmoid (1/(1+tl.exp(-v))), and zero the invalid columns with tl.where.
- Epilogue for MaxPool1d(k):
  - Add bias, then pool groups of k adjacent columns in registers, e.g. tl.max(tl.reshape(acc, [BM, BN//k, k]), axis=2) (validated correct at 1e-4).
  - Make BN a multiple of k so windows never straddle tiles. MaxPool1d floors, so drop any tail columns beyond (N//k)*k.
  - Mask invalid windows to 0 before the sum, then multiply by the scale.
- Reduce each tile over axis 1 with tl.sum. Then either tl.atomic_add into a zeroed [M] fp32 output (validated at 1e-4), or write a small [M, cdiv(N,BN)] partials buffer and run a tiny final sum (deterministic).

Debug before retreating. Validate the fused epilogue on a small shape against the reference first. The summed outputs are large (~N/2 ~ 16k for sigmoid), so rtol=1e-4 is loose in absolute terms. When CHECK fails, print the max abs/rel error and the output magnitude: a large or structured error means an indexing or pool bug, while a small one points to precision.
[str-00068] helpful=1 harmful=0 :: Conv -> single built-in activation (e.g., erf-GELU) -> global average pool (adaptive_avg_pool2d(x,1) + squeeze), KernelBench conv family (e.g., L2_067). A standalone single-activation kernel gives nothing (str-00003). Fusing the activation INTO the pooling reduction does win. Eager does an activation write, a re-read, then a pool pass; the fused kernel does 1 read and writes only [N,C].

Measured, CONFIRMED TWICE on L2_067 (128x64x254x254 conv output):
- Run 1: 1.31x vs eager (12.86 -> 9.84ms) and 1.79x vs torch.compile, own kernel ~15% of runtime.
- Run 2: 1.3073x (12.84 -> 9.82ms, lo95 1.3067), own_kernel_time_frac 0.154. The recipe was copied almost verbatim, with the shared array templated on BS and S/NC kept long. It compiled first try with the str-00001 layout, then went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no conv lever tried.
~1.31x is the SOLVED BASELINE for this task: resubmitting it cannot raise the score.

Recipe:
- Keep the cuDNN conv.
- Write ONE load_inline reduction kernel with one 256-thread block per (n,c) over the contiguous H*W slab.
- Use float4 loads, plus a scalar tail loop for S%4.
- Apply the activation per element in registers with accurate erff (0.5*v*(1+erff(v*0.70710678))), accumulating in fp32.
- Do a shared-memory tree reduction and write sum/S to out[n*C+c].
- Put the wrapper and stream in cuda_sources (str-00001).

Risks to handle when reusing:
(1) The float4 reinterpret needs every slab base to be 16-byte aligned, i.e. S%4==0. It is safe on L2_067 only because S=254*254=64516 is divisible by 4. On other shapes, add an S%4==0 plus 16-byte pointer-alignment guard with a scalar fallback. An odd S crashed with 'misaligned address' in str-00035.
(2) Size the shared array with a template block size (done in run 2); do not hardcode sm[256].
(3) Keep NC and S as 64-bit (str-00058; done in run 2).

Next attempt: after the first passing CHECK, spend the spare turns timing conv-side levers at full size against the reference before SUBMIT, and keep the fastest. The conv is ~85% of runtime. Levers:
- cuDNN TF32 within tolerance (print the max error first, str-00029);
- cudnn.benchmark;
- channels_last (non-dilated only, str-00051), with the reduction then reading NHWC with threads spanning C;
- an epilogue-fused conv that accumulates per-(n,c) activated sums, so the conv output is never written.
[str-00069] helpful=1 harmful=0 :: Eval-mode CNN stages built from conv3x3->BN->ReLU->conv3x3->BN->ReLU->MaxPool2d(2) at moderate batch and full resolution (KernelBench network family, e.g., L3 RegNet L3_027 with batch 8, 224px, widths 64/128/256; ref ~8.35-8.40ms). The cuDNN 3x3 convs are about 95-99% of runtime, so the str-00038 recipe gives only parity here.

The layout, measured THREE times:
- fold BN once (str-00038 formula), cached and keyed on data_ptr, guarded by `not self.training`;
- pass the folded bias to F.conv2d, then run an in-place custom ReLU after conv1;
- after conv2, run one fused relu+2x2 maxpool kernel (max first, then ReLU, since ReLU is monotone; str-00045).
Results:
- 0.9975x (8.37 vs 8.35ms), own_kernel_time_frac 0.0006.
- 1.029x (8.35 -> 8.12ms, lo95 1.028), own_kernel_time_frac 0.048, 1.56x vs torch.compile.
- 1.0416x (8.40 -> 8.06ms, lo95 1.040), own_kernel_time_frac 0.04, 1.57x vs compile (str-00001 layout, long sizes, kernels in both branches, grid-stride 4096-block cap).
All went CHECK then SUBMIT with no local timing and tried no conv lever. This layout is a SOLVED ~1.00-1.04x baseline: resubmitting it cannot raise the score. The standalone ReLU saves no memory pass (str-00003); the relu+pool fusion removes one pass, negligible next to the convs.

no_custom_kernel pitfall seen here: the turn-0 CHECK of the 1.029x attempt failed because neither kernel ran during the checked forward. The passing version launches both kernels in BOTH the training branch (eager conv+BN, then the custom relu / relu_maxpool) and the eval branch (str-00031).

FLAG LEVERS ARE BLOCKED: the 1.0416x run lost turn 0 because the static checker rejects any 'torch.backends' reference (cudnn.benchmark / allow_tf32), see str-00014 (e). The same checker message lists CUDA graphs as forbidden, so do not count on graph capture either without a CHECK.

The gain must come from the convs. Decide on turn 0, and spend the turns after the first passing CHECK timing each locally against the reference instead of submitting:
(1) channels_last (these convs are non-dilated): your own NCHW->NHWC transpose kernel (ideally the smem-tiled one) plus channels_last weights cached on data_ptr, following the str-00102 1.45x pattern; make the relu_maxpool kernel NHWC-aware so the whole stack stays NHWC (transpose only once at the input). Time per str-00051.
(2) A conv whose epilogue does bias+ReLU(+pool), e.g. a TF32 Triton implicit GEMM within tolerance (str-00027 option a; print max error first, str-00029).
(3) CUDA-graph the eval forward only if a CHECK confirms the checker allows it.
Keep the fused relu_maxpool kernel as the custom piece that runs in both branches. A local timing at or near 1x should trigger these levers, not a SUBMIT.
[str-00070] helpful=3 harmful=0 :: Triton constexpr pitfall: values derived inside a @triton.jit kernel from constexpr args are NOT constexpr. For example, half = BLOCK_N // 2 followed by tl.arange(0, half) fails at JIT compile time with 'arange's arguments must be of type tl.constexpr', which scores runtime_error. The same applies to shapes passed to tl.reshape, tl.zeros and similar calls.

Fixes:
(1) Pass the derived size as its own tl.constexpr kernel argument, e.g. HALF: tl.constexpr, with HALF=BLOCK_N//2 computed on the host. If you use autotune, derive it per config.
(2) Avoid the derived range entirely by reshaping a full-size tile, e.g. tl.reshape(acc, [BM, BN//2, 2]), whose shape expressions are folded from the constexpr args.

Always launch the kernel locally (CHECK) before SUBMIT, because compilation happens only on the first launch.
[str-00071] helpful=0 harmful=0 :: Batched-LHS matmul with a SHARED 2D RHS (KernelBench matmul family, e.g., L1 3D_tensor_matrix_multiplication: A [N,M,K] @ B [K,L], with N=16, M=1024, K=2048, L=768 fp32). Collapse the batch into the row dimension. A.reshape(N*M, K) is a free view when A is contiguous, so no copy is made. Allocate C as [N*M, L], run ONE plain NN dense GEMM, then return C.reshape(N, M, L). Do NOT write a batched/3D-grid kernel or loop over N. This trick applies only when B is shared across the batch. A per-batch B ([N,K,L], i.e. true bmm) cannot be collapsed and needs a batch program axis.

Apply the str-00004 tf32x3 recipe unchanged:
- tl.dot(..., input_precision='tf32x3');
- autotune keyed on M,N,K over BM/BN in {128,256} and BK in {32,64}, num_warps=8, num_stages 3-4;
- grouped pid ordering (GM=8);
- 1D kmask broadcast per operand (str-00042).

Measured: 1.335x (2.90 -> 2.17ms, own kernel 100% of runtime, correct) at an effective 16384x2048x768 GEMM. This is lower than the 1.43-1.59x seen on 4096-class square or TN shapes. The likely cause is the narrow output width L=768, which gives fewer N-tiles and a relatively larger epilogue cost. The tile count is still ~768 at 128x128, so the grid is not underfilled (contrast str-00040).

The attempt went CHECK then SUBMIT with no timing, so the following headroom is untested. Time each lever locally against torch.matmul:
- an unmasked K loop (K % BK == 0);
- dropping the %M/%N index wrap when the dims divide the tiles;
- plain tf32 if the tolerance allows;
- BN choices that divide L exactly (768 = 3x256 = 6x128), comparing BN=256-heavy configs against BN=128.
[str-00072] helpful=1 harmful=0 :: Conv -> scalar subtract -> NONLINEAR activation (e.g., tanh) -> scalar subtract -> AvgPool2d(P), KernelBench conv family (e.g., 128x128x126x126 conv output, P=2).

Measured, CONFIRMED TWICE with this exact recipe:
- Run 1: 1.37x vs eager (19.21 -> 14.02ms) and 1.20x vs torch.compile, own_kernel_time_frac ~0.10.
- Run 2: 1.379x (19.32 -> 14.01ms, lo95 1.377), own_kernel_time_frac 0.086.
Both were turn-0 SUBMITs with no CHECK or timing, and both succeeded only because a proven recipe was copied. Treat ~1.38x as the BASELINE (a ceiling for this layout), not the target. Do NOT resubmit it on turn 0 expecting more. The custom kernel is under ~10% of runtime, so the cuDNN conv dominates and the next gain must come from the conv.

There are two exact folds:
(1) Fold the pre-activation scalar into the conv bias: F.conv2d(x, w, b - s1). The per-forward bias subtract is a tiny tensor and costs nothing measurable.
(2) AvgPool is LINEAR, so a post-activation scalar (or per-channel) add/subtract commutes with it. Apply it once per pooled output: mean(act(v)) - s2.

Kernel: ONE load_inline kernel with one thread per POOLED output, grid-stride with a 4096-block cap, long totals and indices (str-00058). Put the wrapper and stream retrieval in cuda_sources with a declaration-only cpp_sources (str-00001). Each thread reads its PxP window of the RAW conv output, applies tanhf to EVERY element, averages with a precomputed 1/P^2, subtracts s2, and writes. About 4 eager passes become 1 read plus a 1/P^2-size write.

Ordering rule: do NOT average first. The activation is nonlinear, so avg does not commute with it. Contrast max-pool, which commutes with monotone ops (str-00045). Only linear ops after the activation (scalar or per-channel add, positive or negative scale) move past the avg pool. The same holds for BN-affine before an AvgPool (str-00062).

Plan from turn 0: write this kernel, CHECK it, then spend the remaining turns on a bounded full-size local timing of conv-side levers against the reference, keeping the fastest (all untested):
- cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv) within tolerance (str-00029);
- torch.backends.cudnn.benchmark=True;
- channels_last (non-dilated only, timed per str-00051), with the pool kernel then made NHWC-aware;
- an epilogue-fused conv+tanh+pool, so the conv output is never written.
Minor kernel headroom (small, since the kernel is <10% of runtime): a 2D grid with ow fastest to remove per-element 64-bit div/mod, and float2 loads for P=2.
[str-00073] helpful=1 harmful=0 :: Conv2d → GroupNorm → elementwise chain (e.g., tanh → HardSwish) → residual add of the RAW conv output → logsumexp over C (dim=1, keepdim), KernelBench conv family (e.g., 128x64x126x126 conv output, G=16).

Measured, CONFIRMED TWICE with the recipe below:
- Run 1: 2.34x vs eager (8.50 → 3.63ms), own kernels ~42% of runtime.
- Run 2: 2.3863x (8.51 → 3.57ms, lo95 2.384), own_kernel_time_frac 0.41. It compiled first try (str-00001 layout, long sizes per str-00058), then went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no lever tried.
Treat ~2.35-2.39x as the SOLVED BASELINE: resubmitting this layout cannot meaningfully raise the score. The win is large because the reference materializes x_norm, x_tanh, x_hs and x_res and then runs a multi-pass logsumexp, about 7 full passes over a ~512MB tensor. The 2D conv is a smaller share of runtime than in the 3D cases.

Recipe: keep the cuDNN conv and use exactly TWO kernels, both launched from cuda_sources (str-00001).
(1) GN stats kernel (str-00046 layout): one block per (n,g), fp32 sum/sumsq per thread, a float shared-memory tree, then a double finalize for mean and rstd.
(2) Fused kernel: one thread per (n,s), with s fastest so the strided-c loads coalesce (str-00021). In registers, each thread does the following for every channel: normalize with gamma/beta, apply the chain (HardSwish as t*clamp(t+3,0,6)/6), add the raw conv value v, and update an online logsumexp (str-00061: if res>m, s = s*exp(m-res)+1 and m = res; else s += exp(res-m)). It writes only the [N,1,H,W] output.
Result: the conv output is read twice, and the full-size output is never written. __expf passed tolerance here, so check the task tolerance first (str-00029).

The custom kernels are ~41% of runtime, so kernel-side levers matter MORE here than in conv-dominated tasks. After the first passing CHECK, time these at full size against the baseline before SUBMIT (all untested):
(a) Split each (n,g) stats reduction across several blocks with float4 loads. There are only N*G≈2048 blocks of ~254k elements each, plus a serial shared-memory tree. Use double atomics or a partials buffer plus a finalize step.
(b) A 2D grid (blockIdx.y=n) with gamma/beta and the per-(n,g) mean/rstd staged in shared memory. The measured kernel still recomputes g = c/cpg, does per-element 64-bit div/mod, and reloads mean/rstd/gamma/beta from global memory per channel. Precomputing per-channel scale/shift (rstd*gamma, beta - mean*rstd*gamma) also removes FLOPs.
(c) Fuse the GN sum/sumsq into a pass that is already reading the conv output.
(d) cuDNN TF32 or channels_last for the conv, within tolerance.

General rule: a norm → elementwise → residual → channel reduction tail collapses to stats + one register-resident channel loop, whatever the elementwise chain is.
[str-00074] helpful=2 harmful=0 :: Persistent hidden-unit-partitioned GRU/LSTM recurrence kernel (str-00018 option a). Being correct is not enough: the thread mapping inside each block decides speed. Measured TWICE, and lost to cuDNN both times.

Measurement 1: bidirectional GRU returning h_n (B=10, T=512, H=256, L=6). The design beat the cuDNN-wrapper layout (~0.77x) but still LOST to cuDNN: correct, 0.879x (69.7 vs 61.2ms), own_kernel_time_frac 0.955. Nearly all of the time was inside the recurrence kernel, so the inner loop is the lever.
Measured design (do NOT copy as is):
- 96 blocks x 256 threads (48 per direction); each block owns ~6 hidden units and their 3 W_hh rows (18KB) in shared memory.
- Per step, one thread per (j,b) ran a serial 256-long dot product, reading h[k] directly from global/L2. That is only 6x10=60 work items for 256 threads, so ~76% of threads sat idle.
- Two atomic-counter grid barriers per step (write out, then copy out into hstate), across 3072 steps.
- Layers ran sequentially, and torch.stack made extra copies of xproj and the weights.

Measurement 2 (WORSE): bidirectional LSTM -> fc(out[:, -1, :]) (B=10, T=512, H=256, L=6): correct, 0.3675x (215 vs 79ms), own_kernel_time_frac 0.98. Design: min(H,132) blocks x 128 threads, one block per hidden unit u, one thread per batch element b (only 10 of 128 threads working, ~92% idle); each active thread streamed its 4 W_hh rows (4x256 floats) from global/L2 every step inside a serial 256-long k loop (never smem-resident; lanes read different rows at the same k, uncoalesced); h_prev was staged in smem and h double-buffered with ONE atomic-counter barrier per step (those parts were fine); but forward and backward directions were SEPARATE sequential launches (2*L*T = 6144 serial barrier steps), and the last-layer backward pruning of str-00018 was not used. Turn 2 was lost to no_code (str-00008), and turn 3 submitted with no timing.
Lesson: smem-staged h plus a single barrier is NOT sufficient. Without resident weights, a warp k-split and concurrent directions, expect well below 1x.

Required design (apply ALL together, timing against nn.LSTM/nn.GRU):
(1) Load each block's slice of W_hh rows (several units x all gates, e.g. 16 rows = 16KB for LSTM) into shared memory ONCE before the T loop. Never stream W_hh from global inside the step loop.
(2) Map warps to (row or unit, b) work items, split k across the 32 lanes, and reduce with __shfl_xor, so every thread is busy. Do not use one thread per (j,b) or per b with a serial k loop. Size the units per block so work items >= threads.
(3) At the start of each step, stage h_prev (B*H floats, ~10KB) into shared memory with coalesced loads. Never read h from global memory inside the dot product.
(4) Double-buffer h in global memory: read buf[t%2] and write buf[(t+1)%2]. This needs ONE grid barrier per step and removes the copy pass.
(5) Run BOTH directions concurrently in ONE launch on disjoint block sets, with grid <= SMs x occupancy so all blocks are co-resident.
(6) When only out[:, -1] is read, apply the str-00018 dead-work pruning: the last layer's backward direction needs only one step.
(7) Use a layer wavefront to overlap the L sequential layer launches. Write xproj directly in the layout the kernel reads (e.g., one GEMM producing [T,B,4H] per direction, or both directions at once), instead of torch.stack or permute().contiguous() copies.

Access the weights via self.lstm/self.gru.all_weights[2*l+dir] (static-check safe). Both attempts went CHECK then SUBMIT with no timing (str-00002); a single local timing vs the reference module would have exposed each slowdown while turns remained. A sub-1x local timing should trigger the fixes above before SUBMIT.
[str-00075] helpful=11 harmful=2 :: nn.Linear (NT layout: weight stored [N,K]) + fused elementwise epilogue or BN/softmax tail at large size (e.g., Linear 2048x8192x8192 -> *scale -> hardtanh -> erf-GELU, Linear 1024x8192x8192 -> training BN -> scale -> softmax, Linear 1024x8192x8192 -> *mult -> LeakyReLU, or Linear 1024x8192x8192 -> +bias -> ReLU; KernelBench matmul family).

Do NOT CHECK-then-SUBMIT a tf32x3 Triton kernel that reads W in place here. It has failed FIVE times on NT/TT layouts, always correct but slower than cuBLAS:
- 0.90x on A.T@B.T (str-00004).
- 0.85x at M=1024 NT with a separate GN pass (str-00034).
- 0.944x at M=2048 NT (16.84 vs 15.90ms, own kernel 100%). It used autotune over BM/BN {128,256}, BK {32,64}, GM=8, a masked K loop, W.t() read via swapped strides as a [BK,BN] tile with K contiguous, and the bias/scale/clamp/tl.erf epilogue. It went CHECK then immediate SUBMIT on turn 1, with turns left and no timing.
- 0.9228x at M=1024 NT with a lean BN+softmax tail (8.77 vs 8.09ms, str-00016). The tail was only ~0.5-1ms, so the GEMM alone ran ~7.5-8ms. It went CHECK, CHECK, SUBMIT with no timing.
- 0.9529x at M=1024 NT, Linear -> *mult -> LeakyReLU at atol=rtol=1e-4 (8.42 vs 8.02ms, own kernel 100%). This one ALREADY applied lever (2): W loaded as [BN,BK] tiles with tl.dot(a, tl.trans(w), input_precision='tf32x3'), an unmasked K loop with %M/%N wrap, and a fused bias+mult+LeakyReLU epilogue, but autotune still included 256-wide configs. A plain-TF32 turn-0 CHECK failed at 1e-4, so TF32 was not an option. It submitted on turn 3 with no timing. Lever (2) narrows the NT gap (0.85-0.94x -> ~0.95x) but does NOT close it.

The same op family has a proven 5.5x win with PLAIN TF32 (str-00032: allow_tf32=True, BM=BN=128, BK=64, 8 warps, 3 stages, GROUP_M=8, Linear -> /d -> erf-GELU), which passed that task's tolerance at K=8192.

For Linear+epilogue tasks, choose the precision path this way:
(1) Start from the str-00032 plain-TF32 recipe, and probe the tolerance by printing the max abs/rel error against the reference. At 1e-4 it fails (confirmed again on Linear+bias+ReLU 1024x8192x8192); go to (4).
(2) In-place NT tf32x3: load the W tile as [BN,BK] (K contiguous) and use tl.dot(a, tl.trans(b)) instead of a swapped-stride [BK,BN] tile. Measured: helps, but stays sub-1x. Superseded by (4).
(3) Restrict to BM=BN=128 configs and run the K loop unmasked when K % BK == 0.
(4) MEASURED WIN, the 1e-4 default: convert to the NN layout. Linear 1024x8192x8192 -> +bias -> ReLU at atol=rtol=1e-4 scored 1.2403x (7.95 -> 6.41ms, lo95 1.2395), own kernel 100%. Recipe:
  - Build Wt = W.t().contiguous() ONCE, cached in a self.<x> attribute keyed on weight.data_ptr() and filled lazily on the first forward. Never rebuild it per forward, and never key on ._version (str-00014). See the str-00081 caveat on stale data_ptr keys.
  - Run a Triton NN GEMM with tl.dot(a, w, acc, input_precision='tf32x3') and a FIXED config: BM=BN=128, BK=64, GROUP_M=8, num_warps=8, num_stages=3.
  - Use an unmasked K loop (K % BK == 0).
  - Put the bias/scale/activation in the epilogue, followed by one masked store.
  - Keep the GEMM kernel launching every forward, so the cached transpose is never the only custom work (str-00009 caveat).
  Treat ~1.24x as the BASELINE, not the target: equal-FLOP NN GEMMs reach 1.43-1.59x with str-00004. That run submitted with no timing. Spend the spare turns after the first passing CHECK timing num_stages 4, BK 32, and a small BM=BN=128 autotune subset.

Always use the spare turns after the first passing CHECK for a bounded full-size timing against the reference, and time the GEMM alone against F.linear as well. A sub-1x timing triggers the next variant, not a SUBMIT. If every variant stays below 1x, submit cuBLAS F.linear + one in-place epilogue kernel (~1.01x, str-00011), never a sub-1x custom GEMM.
[str-00076] helpful=11 harmful=0 :: Scope guard for the tf32x3 dense-GEMM default (str-00004): when the task's GEMM is an nn.Linear (weight stored [N,K], the NT layout) at compute-bound sizes (M >= ~1024), do NOT use the str-00004 tf32x3 recipe reading W in place as the default. It has lost to cuBLAS on every such NT/TT attempt so far: 0.85x, 0.90x, 0.92x, 0.944x, and 0.953x even with [BN,BK] W tiles + tl.trans + an unmasked K loop. This held even with lean surrounding passes.

For these nn.Linear shapes, follow str-00075 instead:
- Use plain TF32 (str-00032) when a max-error probe shows a ~1e-2 tolerance.
- At 1e-4, go straight to a once-cached contiguous W.t() (keyed on data_ptr, filled lazily on the first forward) fed to the NN tf32x3 kernel with a fused epilogue. This is MEASURED at 1.24x on Linear+bias+ReLU 1024x8192x8192 with fixed 128x128x64, 8 warps, 3 stages.
- Time it against F.linear, and keep cuBLAS + an in-place epilogue as the fallback.

EXCEPTION (skinny-M, weight-bandwidth bound, e.g., M=128, K=N=32768): the in-place swapped-stride NT tf32x3 read WINS (1.33x, confirmed twice; see str-00040), so this guard does not apply there. Also skip the cached W.t() copy for multi-GB weights (memory doubling / OOM risk).

The str-00004 tf32x3 wins (1.43-1.59x) apply directly only to NN/TN layouts such as A @ B and A.T @ B, to NT converted to NN via the cached transpose, or to the skinny-M exception above.
[str-00077] helpful=0 harmful=0 :: Transient toolchain/harness failure pitfall: a CHECK or official run can fail for environmental reasons that are not bugs in your kernel. Three forms have been seen:
(a) CalledProcessError from /usr/bin/gcc compiling Triton's own cuda_utils.c (paths like /tmp/.../cuda_utils.c, -lcuda, triton/backends/nvidia/include). This is a hiccup in Triton's driver-stub build.
(b) FileNotFoundError on an evaluator temp log, e.g. '/tmp/sandbox/work/tmp/tmpXXXX.log'. This is a harness hiccup.
(c) FileNotFoundError: '/tmp/sandbox/work/triton' on the OFFICIAL run, after CHECK had passed (L2_097, Triton GEMM + load_inline tail). This is most likely a missing Triton cache directory. It scored 0 (runtime_error), and the official run allows no retry.
In cases (a)/(b) during CHECK, an unchanged retry passed on the next turn. Retry the same code once before rewriting or abandoning the kernel. Only debug the kernel if the error recurs or names your own source. Budget one turn for this when planning CHECK and timing turns, and still keep a turn for local timing after the passing CHECK.
Mitigation for (c), untested: any submission that JIT-compiles Triton carries this risk. At module import, before any Triton launch, create the cache dir: `import os; d = os.environ.get('TRITON_CACHE_DIR'); os.makedirs(d, exist_ok=True) if d else None` (optionally also os.makedirs('/tmp/sandbox/work/triton', exist_ok=True) inside a try/except). Where a load_inline/cuBLAS path performs equally well, it avoids Triton JIT entirely.
[str-00078] helpful=1 harmful=0 :: Pure dense Conv2d/Conv3d with no epilogue and bias=False (KernelBench L1 conv family). The zero-bias wrapper is a SOLVED LOSS, measured TWICE. This is a repeat of str-00027 failure mode B in its worst form: cuDNN F.conv2d/F.conv3d followed by a custom add_bias kernel that, because bias=False, adds a zero vector. That is an extra full read+write pass over the output (plus an allocation or a .to(dtype) on the zero buffer every forward), so it is strictly slower than parity. Never submit a no-op, zero-add or identity-copy wrapper around cuDNN for a pure conv.

Measured:
- Dilated Conv2d 32->64, 5x9, dilation (2,3), padding (2,4), N=8, 512x512 (~0.5GB output): a torch.zeros(OC) bias allocated every forward, written into empty_like. CHECK passed, but the official run scored 0 as 'incorrect / candidate produced no output' (own_kernel_time_frac 0.046), after an immediate turn-1 SUBMIT with no full-size run (str-00002).
- Non-dilated Conv3d 3->64, k=(3,5,7), input 8x3x16x128x128 (~0.5GB output, L1_066): an in-place grid-stride add_bias over a registered zero buffer. It was correct but scored 0.9069x (8.57 vs 7.77ms), own_kernel_time_frac 0.077, so the extra pass alone cost ~7.7% of runtime. It went CHECK then turn-1 SUBMIT with no timing.

On turn 0, commit to changing the conv itself. Routes:
(a) A TF32 implicit-GEMM kernel with constexpr KH/KW(/KD)/dilation, K ordered (kh,kw,ic) over an NHWC copy made inside your OWN kernel, a tile spanning all OC, and the weight pre-permuted and cached.
(b) For TINY IC (e.g., 3), a register-tiled direct conv is plausible: the whole weight (e.g., 64*3*105 floats) fits in shared/constant memory. Each thread computes several W-adjacent outputs for all OC, the input tile plus halo is staged in smem once per block, padding is handled at staging time, and the kernel sizes are constexpr. Verify it on a small shape against nn.Conv3d first.
(c) cuDNN levers, timed: allow_tf32 (print the max error against the tolerance first, str-00029) and cudnn.benchmark. For dilated convs keep NCHW, since dilated channels_last is ~64x slower (str-00051 / failure D). channels_last / channels_last_3d is only an option for NON-dilated convs, and must be timed. Pair any cuDNN lever with real custom work.
After CHECK, run the forward at the FULL get_inputs size with time and peak memory printed, and compare against nn.Conv2d/nn.Conv3d. A sub-1x result must block SUBMIT.
[str-00079] helpful=0 harmful=0 :: Dense conv failure mode (E), a companion to str-00027: a MATERIALIZED im2col loses badly for large-M convs. Measured on AlexNet conv1 (Conv2d 3->96, 11x11, stride 4, pad 2, N=256, 224px; cuDNN ref ~5.35ms). The layout was a custom im2col kernel + torch.matmul(col, W.t()) + an eager bias add + reshape(N,OH,OW,OC).permute(0,3,1,2).contiguous(). It was correct but scored 0.363x (14.72ms). own_kernel_time_frac was 0.69, so the im2col kernel alone took ~10ms, about 2x the whole cuDNN conv.

Causes:
(1) The col matrix is [N*OH*OW=774400, K=363] fp32, ~1.1GB, written once and then re-read by the GEMM.
(2) One thread per output pixel writing its own contiguous 363-float row means lanes in a warp store 1452 bytes apart. The writes are fully uncoalesced.
(3) The eager bias add and the NHWC->NCHW permute+contiguous are two more full passes over the ~0.2GB output.
This attempt went CHECK then turn-1 SUBMIT with no timing (str-00002). That happened while str-00027 option (a) explicitly said to gather on the fly.

Rules:
- Do not materialize im2col when N*OH*OW*K is in the GB range.
- If you must use an im2col-style kernel, write col TRANSPOSED as [K, M] (per image: [K, OH*OW]), so consecutive threads (consecutive output pixels) write consecutive addresses.
- Then use a GEMM that writes NCHW directly with the bias fused, e.g. per image out[n] = W[OC,K] @ col_n[K, OH*OW] via torch.baddbmm(bias.view(1,OC,1), W.expand(...), col). This needs no permute copy and no separate bias pass. Untested; time it.
- Better: a TF32 implicit GEMM that gathers patches on the fly, with the tile spanning all OC and constexpr KH/KW (str-00027 option a).
- Or use cuDNN TF32 (NCHW) if the tolerance allows.
- Always time at full size against nn.Conv2d before SUBMIT.
[str-00080] helpful=0 harmful=0 :: Conv3d -> x*scale[c] -> tanh -> *bias[c] -> sigmoid (per-channel elementwise chain after a cuDNN 3D conv, ~1GB 128x16x14x62x62 output). This is the per-channel variant in str-00006, now CONFIRMED TWICE as a SOLVED baseline.

Measured:
- Run 1: 1.287x vs eager (7.83 -> 6.09ms), 0.88x vs torch.compile.
- Run 2: 1.2345x (7.81 -> 6.32ms, lo95 1.234), 0.8466x vs compile, own_kernel_time_frac 0.19. It used one fused kernel with a 2D grid (blockIdx.y = n*C + c, per-channel params loaded once per block, blockIdx.x capped at 256, grid-stride over S), an empty_like output, accurate tanhf/expf, long sizes (str-00058) and the str-00001 layout. It compiled first try, then went CHECK on turn 0 and SUBMIT on turn 1 with no timing and no lever tried.

Resubmitting this plain fused kernel cannot raise the score. torch.compile finds the same fusion, so ~0.85-0.88x vs compile is the ceiling of this layout.

After the first passing CHECK, spend the remaining turns on a bounded full-size timing against the reference before SUBMIT, and keep the fastest variant.
(1) Conv-side levers (str-00053):
- cuDNN TF32 (torch.backends.cudnn.allow_tf32 around the conv), only within tolerance (str-00029);
- torch.backends.cudnn.benchmark=True;
- channels_last_3d, timed per str-00051. The flat per-channel kernel must then index c from NDHWC or stay layout-aware.
(2) Cheap kernel tweaks. These matter a bit more here, because the kernel is ~19% of runtime:
- write in place on the conv output instead of empty_like;
- use float4 loads/stores when S % 4 == 0, with alignment checked (S = 14*62*62 = 53816 is divisible by 4);
- use a grid cap near SMs x 8 instead of 256 blocks per (n,c);
- use __expf or exp-based tanh only if the tolerance allows.
[str-00081] helpful=2 harmful=0 :: SwinV2 cosine window attention (KernelBench L3_030 SwinTransformerV2, attention family, batch 10, 224px, ref ~18.9-19.7ms). Removing parameter-only work is a measured WIN, CONFIRMED TWICE:
- Run 1: 1.2027x (18.93 -> 15.74ms, lo95 1.195).
- Run 2: 1.2175x (19.66 -> 16.15ms, lo95 1.2127), own_kernel_time_frac 0.011, but only 0.56x vs torch.compile.
Almost all of the gain comes from CACHING, not from the custom kernel.

What worked: per WindowAttention module, lazily cache two tensors in self.<x> attributes, each keyed on a parameter's data_ptr() (str-00014):
- the relative-position bias: 16*sigmoid(cpb_mlp(relative_coords_table)), gathered by relative_position_index, permuted and reshaped to a contiguous [nH, N*N] (or [1,nH,N,N]);
- clamp(logit_scale, max=log(100)).exp().
This removes ~10 tiny launches per block across all blocks.

The custom kernel in both runs was a trivial grid-stride bias add (4096-block cap, long indices, str-00001 layout; compiled first try), 1.1-1.7% of runtime. It adds NOTHING and likely costs a pass: in run 2 the proj bias was moved out of F.linear into this kernel, so cuBLAS's fused addmm became F.linear(no bias) + .contiguous() + a separate add pass. Keep proj (and qkv) WITH their biases in cuBLAS; a custom kernel must remove passes (fused attention chain), not re-implement a bias add.

Treat ~1.2x as the SOLVED BASELINE: resubmitting this layout cannot improve the score, and both runs went CHECK then SUBMIT with no timing, learning nothing new. torch.compile at ~2x the speed shows large headroom in the attention chain. Next attempt: CHECK on turn 0, then spend the remaining turns timing real fusions at full size against the reference before SUBMIT (str-00017):
(a) Precombine the cached bias and shift mask into ONE cached [nW, nH, N, N] additive tensor per block.
(b) One fused per-(window, head) Triton kernel doing normalize(q), normalize(k), QK^T, *scale, +bias/mask, softmax and @v, with N=49 padded to 64 and masked (-inf on padded columns), reading q/k/v from the qkv output via strides.
(c) Fold roll/window_partition/window_reverse into index math.
(d) Fuse the residual add with the post-norm LayerNorm (x = shortcut + LN(y)) in one kernel.

Cache caveat (applies to any data_ptr-keyed parameter cache): the key does NOT change on in-place weight updates (param.copy_, load_state_dict into existing storage), so a cache filled BEFORE the weights are loaded goes stale silently. A cache filled lazily on the first forward after the weights are loaded is safe, and that is what passed here (twice). Keep a real custom kernel launching every forward, so the cache is never the only custom work (str-00009/str-00031).
[str-00082] helpful=0 harmful=0 :: CUDA-error-then-pass pitfall: if ANY CHECK raised a CUDA runtime error (e.g., 'invalid argument', 'illegal memory access', 'misaligned address'), a later passing CHECK does NOT prove the launch path is safe; the official run can still end with 'candidate produced no output' (measured on NetVLAD-ghost, str-00056, and the dilated-conv wrapper, str-00078). Before SUBMIT: (1) reproduce/diagnose with CUDA_LAUNCH_BLOCKING=1 to find the failing launch instead of blind patching; (2) run the full get_inputs-size forward in a loop (e.g., 20-50 iterations, like the timing harness) in both train() and eval() modes, with torch.cuda.synchronize(), peak memory (torch.cuda.max_memory_allocated) and time printed, and compare against the reference output each mode; (3) harden launches: fixed compile-time block sizes rather than runtime-computed ones, static shared-memory sizes with blockDim asserted <= the static array size, grid dims within limits, TORCH_CHECKs on shape/contiguity/dtype/device, and C10_CUDA_KERNEL_LAUNCH_CHECK() (or cudaGetLastError) after every launch so errors surface at the right call; (4) avoid large per-forward temporaries (hundreds of MB) that add OOM risk under repeated timing.
[str-00083] helpful=0 harmful=0 :: Fit-check before reusing a single-block persistent recurrence kernel (str-00063) on GRU/LSTM. Compute the recurrent weight bytes per layer first: gates*H*H*4. A vanilla RNN at H=256 is 256KB, which just fits in one SM's registers. A GRU at H=256 is 3x = 768KB and an LSTM is 4x = 1MB, and neither fits in registers plus smem (~256KB regs + ~228KB smem per SM).

Measured CATASTROPHIC FAILURE: multi-layer GRU (B=10, T=512, H=256, L=6) with one persistent kernel per layer launched <<<1, H>>>. Thread j computed hidden unit j for all 3 gates. Each step it looped over b and ran a serial 256-long dot product, reading its 3 W_hh rows straight from global/L2. It was correct but scored 0.0071x (4480 vs 31.7ms), own_kernel_time_frac 0.9998. That is 10x worse than the 10-block batch-parallel variant (0.07x, str-00018).

Why it failed:
- Only 1 SM was busy.
- 768KB of weights were re-read per batch element per step.
- Lanes read different rows at the same k, so the loads were uncoalesced.

It went CHECK then turn-1 SUBMIT with no timing. A single full-size timing run would have shown the 140x slowdown.

When the weights do not fit on one SM, use the hidden-unit-partitioned multi-block design (str-00018 option a, with the inner-loop fixes of str-00074):
- Split the W_hh rows across many blocks (each block owns a slice of units j with all its gate rows), and keep each slice resident in smem or registers for all T steps.
- Process all B batch elements per weight load.
- Split k across warp lanes with a shuffle reduction.
- Stage h_prev into smem each step.
- Double-buffer h in global memory, with ONE grid barrier per step and all blocks co-resident.

Keep the parts that worked: the input-projection GEMM per layer with b_ih (+ b_hr, b_hz) folded in, and b_hn kept separate.
[str-00084] helpful=1 harmful=1 :: Linear (1024x8192x8192, nn.Linear NT weight) -> +bias -> hardtanh -> Mish -> GroupNorm(256 groups, CPG=32), checked at atol=rtol=1e-4. This is a companion to str-00034.

Measured SOLVED fallback baseline: 1.0112x (8.37 -> 8.28ms, lo95 1.010), 0.98x vs torch.compile, own_kernel_time_frac 0.033. Layout:
- Fold self.bias into the Linear bias (exact; str-00011 rule 1), then cuBLAS F.linear.
- A stats kernel with grid (G, M) and 32 threads per block. It uses PyTorch-exact hardtanh+Mish (log1pf(expf), tanhf, >20 guard; str-00029/str-00052), writes a full act tensor, and writes per-(row, group) mean/rstd.
- A separate grid-stride apply kernel doing a 64-bit div/mod per element.
One CHECK in this run failed at 1e-4 for an unrecorded cause; the PyTorch-exact version passed. The run then submitted with no timing.

The tail is only ~3% of runtime, so tail cleanups cannot move the score. Resubmitting this layout cannot improve the result. Optional cleanup if the tail is kept: ONE kernel per (row, several groups) that holds the 32 activated values in registers. This removes the act-tensor write and re-read and the second launch.

The only lever with real upside is the GEMM. Do NOT use the in-place swapped-stride tf32x3 GEMM: it measured 0.85x on this exact task (str-00034).

WARNING: fusing GroupNorm into the epilogue of the fixed 128x128x64 tf32x3 NN GEMM (reshape acc to [BM, BN/CPG, CPG] and reduce) is a MEASURED CATASTROPHE. On the sibling task L2_062 (GN with CPG=16) it scored 0.0371x (str-00090 signature).

Staged candidate instead (untested here):
(1) Use the str-00075 lever-(4) NN conversion: a once-cached contiguous W.t(), keyed on data_ptr and filled lazily on the first forward.
(2) Run a fixed BM=BN=128, BK=64, 8-warp, 3-stage tf32x3 GEMM with an unmasked K loop and a LIGHT epilogue only: bias+hardtanh+Mish, with no reductions. This config gave 1.24x on Linear+bias+ReLU.
(3) Then run one GN kernel per row or per several groups on a 2D [G,CPG] tile.
Time the GEMM alone against F.linear first, then the full pipeline against ~8.28ms. Only after that may you try GN-in-epilogue, and only if re-timing shows no 10x jump. Submit only a variant that beats 8.28ms; otherwise submit the cuBLAS fallback.
[str-00085] helpful=0 harmful=0 :: Linear (1024x8192x8192, nn.Linear NT weight) -> +add_value[out] -> swish -> tanh -> erf-GELU -> hardtanh(-1,1) (KernelBench L2 matmul family).

Measured SAFE BASELINE: 1.1851x (8.23 -> 6.94ms, lo95 1.1845), correct, own kernel ~100%. It used the str-00075 lever-(4) NN conversion, which confirms that lever (4) transfers to other epilogues at ~1.2x (1.24x on bias+ReLU, 1.185x here).

Setup:
- Fold bias + add_value into one vector per forward (exact and cheap).
- Lazily cache W.t().contiguous(), keyed on data_ptr. It is filled on the first forward, and the GEMM still launches every forward.
- Use a tf32x3 NN Triton GEMM with a FIXED config: 128x128x64, 8 warps, 3 stages, GROUP_M=8, and an unmasked K loop.
- Fuse the whole epilogue: sigmoid via tl.exp, overflow-safe exp-based tanh 2/(1+exp(-2x))-1, top-level tl.erf GELU (compiles inside the tf32x3 epilogue, str-00012), then clamp. Finish with one masked store.

This run went CHECK on turn 0 and SUBMIT on turn 1, with no tolerance probe and no timing. That left the main lever untested. The tanh/hardtanh tail squashes outputs into [-1,1], which shrinks absolute error, and str-00011 rule (2b) lists this exact chain for plain TF32 (the str-00032 5.5x recipe).

Next attempt on this task, after the first passing CHECK:
(1) Run the same kernel with allow_tf32=True (input_precision='tf32'), either in place via swapped strides per str-00032 or on the cached NN Wt. Print the max abs/rel error against the reference.
(2) If it fits the tolerance (likely ~1e-2), time it against the 6.94ms tf32x3 version and submit the faster one.
(3) Only if the error exceeds the tolerance, keep tf32x3 and time num_stages 4 / BK 32.
Resubmitting the tf32x3 baseline cannot raise the score.
[str-00086] helpful=0 harmful=0 :: Scope guard for the str-00004 tf32x3 dense-GEMM recipe: IRREGULAR / ODD-STRIDE shapes (KernelBench L1 'Matmul_with_irregular_shapes', A[8205,2949] @ B[2949,5921] fp32, NN layout). Before applying any GEMM recipe, check M, N and K for divisibility by 16 and by the tile sizes.

Measured FAILURE: the unchanged str-00004 recipe was correct but scored 0.5989x (25.92 vs 15.52ms, own kernel 100%). That is about 11 effective TFLOPS, against ~27 on a 4096-square NN GEMM. The recipe used tf32x3, autotune over BM/BN {128,256} and BK {32,64}, GM=8, a masked K loop and a %M/%N index wrap. Strict FP32 had scored 0.56x on the same task, so precision is NOT the bottleneck. The run went CHECK then immediate SUBMIT with no timing, even though str-00004 already listed this shape as a failure.

Likely causes:
- K=2949 and N=5921 are odd, so the row strides of A and B are not 16-byte multiples. Triton cannot prove alignment and emits scalar loads instead of 16B vectorized/cp.async loads.
- K % BK != 0 forces a masked load on every K iteration.
- The %M/%N index wrap.
- Possible register pressure from 256-wide tf32x3 configs.

Untested fixes; time each at full size against torch.matmul before SUBMIT:
(1) Copy A and B into zero-padded buffers, with K padded to a multiple of BK (e.g., 64) and N padded to a multiple of 16 (or of BN). The copy is ~330MB, ~0.2ms, negligible next to ~15ms. Zero K-padding is exact. Then run an unmasked-K NN kernel on the aligned strides, writing C through a masked epilogue store (or into a padded C that is then sliced).
(2) Use a FIXED BM=BN=128, BK=32/64, 8-warp, 3-stage config instead of the 256-wide autotune set (compare str-00032's finding that autotune was ~20% slower).
(3) Use plain tf32 if a max-error probe passes the tolerance.
If every variant stays below 1x, submit cuBLAS (torch.matmul) plus minimal custom work (~1.0x) rather than a sub-1x kernel (str-00029).
[str-00087] helpful=0 harmful=0 :: Linear -> torch.sum(dim=1, keepdim) over all output features (often followed by max/mean/logsumexp over the now size-1 dim, which is an exact identity; see str-00013 Case 2), checked at atol=rtol=1e-4.

The NO-FOLD layout is a SOLVED ~1.01x dead end. It has been measured TWICE on the same 1024x8192x8192 sum task:
- 1.0099x with full cuBLAS plus a Triton row-sum kernel.
Both runs had two earlier CHECKs fail at 1e-4 (probably fold attempts). Both abandoned the fold without printing any error, then retreated.
The fold (cached fp64 w_sum = W.double().sum(0) plus the bias-sum constant, then a per-row GEMV) is the only route to a large win. The mean variant of the same fold reached ~60x.

When a fold CHECK fails at 1e-4, debug it; do not retreat. Steps:
(1) Print max abs/rel error and the output magnitude range (str-00024). A row sum over 8192 mixed-sign features can be small relative to the element sizes, so atol may be the binding limit.
(2) Build both w_sum and the bias sum in fp64.
(3) Do the per-row dot in fp64: torch.mv(x.double(), w_sum64) as a probe, or a Triton GEMV loading x as fp32, casting to fp64 and accumulating in fp64. Cast to fp32 only at the final store.
(4) If it still mismatches, compare BOTH the candidate and the fp32 reference against a float64 reference, i.e. (x.double() @ W.double().T + b.double()).sum(1). This tells you whether the reference's own fp32 rounding (cuBLAS GEMM then torch.sum) dominates the gap.
(5) Only if the fold cannot match, fall back to a tf32x3 GEMM with a per-tile row-sum epilogue: NN layout via a cached W.t(), str-00075 lever 4, with partials or atomics. Do not fall back to cuBLAS plus a copy.

Caching rules:
- Cache w_sum lazily on the first forward, keyed on weight.data_ptr(). Never cache it before the weights are loaded (str-00081).
- Keep the GEMV kernel launching every forward (str-00009/str-00031).

Time the result against the reference before SUBMIT.
[str-00088] helpful=0 harmful=0 :: A.T @ B with A stored [K,M] (KernelBench L1 Matmul_with_transposed_A, M=2048, K=8192, N=4096 fp32, TN layout). The str-00004 tf32x3 recipe is now CONFIRMED TWICE, but scores vary from run to run:
- Run 1: 1.59x (8.08 -> 5.09ms).
- Run 2: 1.3216x (8.49 -> 6.43ms, lo95 1.319), own kernel 100%, correct.
Run 2 used essentially the same recipe: tf32x3; autotune over BM/BN {128,256}, BK {32,64}, 8 warps, stages 3-4; GM=8; A read through swapped strides as a [BM,BK] tile (M contiguous); %M/%N index wrap; unmasked K loop (K % BK == 0). Both runs went CHECK on turn 0 and SUBMIT on turn 1 with no local timing.
The ~20% spread with an identical recipe suggests the cause is autotune picking 256-wide tf32x3 tiles (str-00032 measured these ~20% slower than a fixed 128x128x64 config), or evaluator variance. Treat ~1.3-1.6x as the SOLVED baseline: resubmitting it unchanged cannot reliably raise the score.
Next attempt: after the first passing CHECK, time these at full size against torch.matmul (bounded warmup/iterations, cuda synchronize) and keep the fastest:
(1) A FIXED BM=BN=128, BK=64, 8-warp, 3-stage config instead of autotune.
(2) Plain tf32 (input_precision='tf32'), after printing the max abs/rel error against the reference.
(3) Loading A as a [BK,BM] tile (its native M-contiguous storage order) and using tl.dot(tl.trans(a_tile), b).
(4) Dropping the %M/%N wrap, since 2048 and 4096 divide the tiles.
(5) num_stages 4 around the 128x128 config.
[str-00089] helpful=0 harmful=0 :: Per-step ceiling of a LAYER-SEQUENTIAL persistent LSTM/GRU recurrence (companion to str-00018/str-00074). A design that applies ALL of the str-00074 per-block fixes still only reaches near-parity with cuDNN when layers run one at a time. Treat ~0.89x as that layout's ceiling, not a near-win.

Measured on unidirectional LSTM returning h_n (B=10, T=512, H=256, L=6): correct but 0.894x (36.0 vs 32.2ms), own_kernel_time_frac 0.965. Design (it compiled first try with the str-00001 layout and all_weights access per str-00014):
- one cooperative launch per layer, 64 blocks x 256 threads, 4 hidden units per block;
- that block's W_hh slice (16KB) resident in smem, h_prev (10KB) staged in smem each step;
- warp per (unit,b) with a k-split over lanes and __shfl_down reduction;
- double-buffered h and ONE grid.sync per step;
- input projection via F.linear with b_ih+b_hh folded in, then transposed to [T,B,4H].

Why it caps: L*T = 3072 serial grid-synced steps at ~11.7us each, versus cuDNN's ~10.5us per step. Only 64 of ~132 SMs were busy. Per-step tuning cannot close a gap set by the serial barrier count.

The dominant lever is structural: a LAYER WAVEFRONT in ONE cooperative launch. Layer l at step t runs concurrently with layer l+1 at step t-1, on disjoint block sets (L x 64 blocks must fit co-resident; use fewer units per block or more units per block to fit). This cuts serial barriers from L*T to T+L-1 (~517) and keeps ~6x more SMs busy. It requires computing layer l+1's input projection per step INSIDE the kernel from layer l's fresh h, i.e. W_ih slices resident alongside W_hh, instead of a precomputed per-layer GEMM. Layer 0 can still use the precomputed x-projection GEMM.

Cheaper secondary levers (time each):
- skip seq_out writes for the last layer when only h_n/c_n is returned;
- drop the gate.transpose(0,1).contiguous() and seq_out.transpose().contiguous() copies by indexing [B,T,*] directly in the kernel and GEMM;
- upb=2 (128 blocks) to fill more SMs, provided all blocks stay co-resident.

That attempt went CHECK then immediate SUBMIT with no timing. A local timing against nn.LSTM after the first passing CHECK would have exposed the sub-1x result while a turn remained. A result below 1x here should trigger the wavefront rewrite, not a SUBMIT.
[str-00090] helpful=2 harmful=0 :: Scope guard and CATASTROPHE SIGNATURE for the fixed-config tf32x3 NN GEMM (str-00075 lever 4, str-00016 item 4). The combination is a fixed tf32x3 Triton GEMM at BM=BN=128, BK=64, 8 warps, 3 stages, plus a HEAVY epilogue: reshape+reductions, row-sum, or atomic stats. Treat it as a ~30x slowdown signature (likely register spill or a forced layout conversion), not a tuning issue. Measured THREE times, each correct with own kernels at 100% of runtime:
(a) L2_039 (Linear 16384x4096x4096 -> *scale -> training BatchNorm1d): 0.0336x (893 vs 30ms, ~0.6 TFLOPS). Epilogue was bias+scale, a store, then per-column tl.sum of sum/sumsq with tl.atomic_add, followed by a finalize CUDA kernel and an apply kernel. The BN tail costs only ~1-2ms, so nearly all the time was in the GEMM.
(b) str-00067 sigmoid+row-sum: 0.034x.
(c) L2_062 (Linear 1024x8192x8192 -> GroupNorm(512, CPG=16) -> LeakyReLU -> x+x): 0.0371x (232 vs 8.6ms). It used a cached W.t(), NN tf32x3, fixed 128x128x64 with num_warps=8/num_stages=3 correctly passed at launch, M/N masks on loads every K iteration, and an epilogue doing tl.reshape(acc,(BM,8,16)) with mean/var reductions, affine, LeakyReLU and x2.
The only measured success of this config (1.24x, str-00075) had a trivial bias/activation epilogue and unmasked loads. Runs (a) and (c) differ from it only by per-iteration load masks and the heavy epilogue. Treat the config as UNPROVEN with masked loads or any reduction epilogue, at any M.

Mandatory staged process (all three runs went CHECK then SUBMIT with no timing):
(1) First build a PLAIN-epilogue NN tf32x3 GEMM (cached W.t(), unmasked loads when the dims divide the tiles, bias-only epilogue). Time it at full size; it should be at or below F.linear.
(2) Then add the heavy epilogue and re-time. A ~10x jump means the epilogue broke codegen.
(3) Escape routes, in order:
  - Keep the GEMM epilogue light and do the norm/stats in a separate lightweight pass. For GN, use one program per row or per several groups on a 2D [G,CPG] tile. For BN stats, use the str-00016 2D-grid stats kernel.
  - Smaller tiles (BN=64), num_stages=2, or BK=32.
  - Plain tf32 only if a max-error probe passes. BN and GN do not hide relative error.
  - Fall back to cuBLAS F.linear plus a fused tail kernel (~1.01x, str-00016/str-00034).
Never submit a fused GEMM+reduction epilogue without a full-size timing.

Parts that worked and are worth reusing: the str-00001 load_inline layout; the BN finalize kernel (double mean/var, running stats updated in-kernel with unbiased var*M/(M-1) and momentum, gamma/beta/scale folded into sc/sh, eval branch using running stats); the data_ptr-keyed W.t() cache; and derived tile shapes (e.g. NG=BN//CPG) passed as tl.constexpr args from the host (str-00070).
[str-00091] helpful=6 harmful=0 :: Triton launch meta-parameter pitfall: num_warps and num_stages are NOT kernel signature args. With a FIXED config and no @triton.autotune, they must be passed explicitly at launch, e.g. kernel[grid](..., BM, BN, BK, GM, num_warps=8, num_stages=3). If they are omitted, Triton silently uses its defaults (4 warps). That is too few for a 128x128x64 tf32x3 tile, and the kernel still compiles and passes CHECK.

Measured on Linear 1024x8192x8192 -> sigmoid -> *s -> +residual (the GEMM output). The str-00075 lever-(4) NN conversion was used: a cached W.t().contiguous() keyed on data_ptr, tf32x3, BM=BN=128, BK=64, GM=8, an unmasked K loop, and a fused bias+sigmoid*s+residual epilogue. With the launch kwargs omitted, it was correct but scored 0.9838x (8.25 vs 8.11ms, own kernel 100%). The same recipe with 8 warps and 3 stages reached ~1.19-1.24x on other epilogues. The run went CHECK then SUBMIT with no timing; one timing run would have exposed it.

Rules:
(1) When copying a fixed-config recipe, check the launch line for num_warps/num_stages before CHECK.
(2) An epilogue of the form GEMM value + bounded term (sigmoid*s + residual) preserves the GEMM's relative error. Print the max abs/rel error. If the tolerance is ~1e-2, time plain tf32 (str-00032) against tf32x3.
(3) After the first passing CHECK, time the GEMM against F.linear before SUBMIT (str-00002/str-00075).
[str-00092] helpful=0 harmful=0 :: Single MBConv / inverted-residual block at large size (KernelBench L3 EfficientNetMBConv: N=10, 112->672 hidden, 224px, 5x5 depthwise stride 2, ref ~15.7ms), evaluated in train() mode (see str-00031).

Measured FAILURE: 0.873x. The real kernel (fused dw-conv+folded BN+ReLU6) sat only in the eval path. The training branch ran eager modules plus a dummy relu6 on a clone, which added ~2.3ms of extra passes over the ~0.34GB hidden tensor. It went CHECK then SUBMIT with no timing.

Training-mode plan (untested; time at full size in train() against the reference before SUBMIT):
(1) Replace each eager BN+ReLU6 with the str-00043 layout:
- many-block float4 per-channel sum/sumsq stats kernel;
- finalize kernel computing scale/shift and updating running_mean/running_var in-kernel (buffers passed as tensors, str-00014);
- ONE in-place apply+ReLU6 pass.
(2) Better: a custom depthwise conv kernel (str-00059 layout) that accumulates the per-channel BN sum/sumsq in its epilogue, so the dw output is not re-read for stats.
(3) Keep the cuDNN 1x1 convs. The project BN (no activation) can use the same stats/finalize/apply kernels.
(4) Provide an eval branch (folded BN plus the fused dw+BN+ReLU6 kernel) so custom kernels run in both modes.
[str-00093] helpful=0 harmful=0 :: ShuffleNet (KernelBench L3 network family, e.g., L3_026: batch 10, 224px, groups=3, stages [3,7,3], ref ~20.2ms; mode not yet probed). Measured FAILURE: the candidate kept the entire eager network and replaced only ChannelShuffle (view -> transpose -> contiguous) with a custom 3D-grid copy kernel (grid = (HW chunks, C, N); in_c = (oc % g)*cpg + oc/g; str-00001 layout, compiled first try). It was correct but scored 0.9729x (20.76 vs 20.20ms, lo95 0.970), with own_kernel_time_frac 0.028. The custom copy swaps one memory pass for another, and it adds module and launch overhead on a launch-bound network. This is the 'eager + one trivial kernel' dead end of str-00026. Do NOT submit it. It also wasted turns: turn 0 defined a class named Model, which the static check forbids; turn 2 was a SUBMIT with no python block (str-00008); and the final SUBMIT had no local timing.

To win, the custom kernels must REMOVE passes or launches. Untested levers, in order; time each at full size against the reference after the first passing CHECK:
(1) A custom depthwise 3x3 kernel (str-00059 layout; cuDNN's depthwise path is weak) that writes each output channel directly to its shuffled channel index. This fuses conv2 with the shuffle and deletes the separate shuffle pass. BN2 is per-channel, so its gamma/beta/stats simply permute with the channels. In eval mode, fold BN2 into the dw weights and bias. In train mode, accumulate the per-channel sum/sumsq in the dw epilogue (str-00092 item 2), then use the str-00043 finalize/apply.
(2) Fuse BN3 + ReLU + the residual add into ONE pass, replacing eager BN, relu and the out-of-place add.
(3) Fuse the tail: conv5's BN5 + ReLU + global average pool in one reduction kernel (str-00068 layout), so the 1024xHxW activation is never re-read.
(4) CUDA-graph the eval forward with the custom kernels inside the graph (str-00026 4a).
Probe which mode the evaluator runs (str-00031) before folding any BN, and put real kernels in both branches.
[str-00094] helpful=0 harmful=0 :: Static-check pitfall, reference class names: the checker rejects any class named Model ('do not define a class named Model (the reference); define ModelNew only'). When starting from a copied reference file, rename or remove the reference Model class before the first CHECK. Helper classes (e.g., ShuffleNetUnit, ChannelShuffle) may keep their names, but giving them distinct names (e.g., ...New) avoids confusion. Only 'Model' itself is forbidden. One ShuffleNet attempt lost its turn-0 CHECK this way.
[str-00095] helpful=0 harmful=0 :: Scope guard for the str-00059 depthwise layout and str-00026 lever (4b), for small feature maps in networks such as EfficientNet/MobileNet later stages (7x7-56x56 maps with hundreds of channels, batch ~10). str-00059 was proven only on LARGE maps (128x256, 512x512). On small maps it is a measured LOSS.

Measured FAILURE, EfficientNetB0 at batch 10 (L3_022): 0.7161x (8.05 vs 5.76ms, lo95 0.715), own_kernel_time_frac 0.66, so the custom kernel itself was slow.
- Layout: eager cuDNN 1x1 expand/project convs kept as they were.
- Only the depthwise conv+BN+ReLU6 was replaced with the str-00059 kernel: grid (ceil(OW/4)/128, OH, N*C), 128 threads per block, 4 outputs per thread, scalar loads, per-tap bounds checks, template<KH,KW>.
- In eval mode BN was folded and cached keyed on data_ptr; in train mode the kernel ran with a zero bias followed by eager BN.
- Root cause: at OW=7/14/28 only 2-7 of 128 threads per block are active, so the launch consists of N*C*OH near-empty blocks. Strided scalar reads and per-tap branches add to the cost. cuDNN's depthwise path wins at these shapes.
- The run went CHECK then SUBMIT with no timing.

Rules:
(1) A 1:1 depthwise swap is not a network-level win on launch-bound small-batch CNNs, even with BN+ReLU6 folded in. A win needs fusion (expand->dw->project, or dw+project, so hidden tensors are not round-tripped) or a CUDA graph of the whole forward (str-00026 4a).
(2) When the spatial width is small, map threads over channels and rows rather than OW/4 chunks so warps stay full. Examples: one block per (n, channel-group) covering the whole small plane, or several rows/channels per block, with the input plane plus halo staged in shared memory.
(3) Before building a network around a custom dw kernel, time it per stage shape against F.conv2d(groups=C) for every (C, H, W, k, stride) in the network. Use it only where it wins.
(4) Then time the full network in the evaluated mode before SUBMIT (str-00002, str-00031).
[str-00096] helpful=0 harmful=0 :: Conv2d -> ReLU -> hand-written HardSwish (x * clamp((x+3)/6, 0, 1)), KernelBench conv family (L2_057, Conv2d 8->64, 3x3, N=128, 128px; output 128x64x126x126). This is the exact chain behind str-00006's ~2.17x figure, and it is now CONFIRMED TWICE.

Measured: 2.1677x (6.31 -> 2.91ms, lo95 2.166), 1.55x vs torch.compile, own_kernel_time_frac 0.24.

Layout:
- cuDNN nn.Conv2d.
- ONE fused grid-stride kernel computing r = max(v,0); out = r*min(r+3,6)*(1/6). The lower clamp at 0 is unnecessary because r >= 0, and 1/6 is a multiply, not a divide.
- 256 threads, 4096-block cap, long n and indices (str-00058), empty_like output, str-00001 layout. It compiled first try.

This replaces about 5 eager passes with 1. ~2.17x is the SOLVED BASELINE: resubmitting it cannot raise the score. The confirming run was a turn-0 SUBMIT with no CHECK or timing; it worked only because the recipe was copied (str-00002).

Next attempt: CHECK on turn 0, then time these levers at full size against the reference and keep the fastest.
- Kernel side (~24% of runtime): write in place on the conv output instead of empty_like; use float4 loads/stores with a scalar tail; set the grid cap near SMs x 8, or launch 1 thread per float4 (str-00054 found a fixed 1024 cap ~9% slower).
- Conv side (~76% of runtime): cuDNN TF32 within tolerance (str-00029), cudnn.benchmark, and channels_last (non-dilated, timed per str-00051).
[str-00097] helpful=0 harmful=0 :: Persistent GRU/LSTM recurrence, measurement 3 (companion to str-00074/str-00089). Applying only PART of the str-00074 checklist is a measured loss. Case: bidirectional GRU returning `output` (B=10, T=512, H=256, L=6). Result: correct but 0.7011x (85.5 vs 59.9ms), own_kernel_time_frac 0.965.

What that design already had: W_hh slice resident in smem, h_prev staged in smem, h double-buffered, ONE grid.sync per step.

What it lacked:
(a) One cooperative launch per (layer, direction), run sequentially. That is 2*L*T = 6144 serial grid-synced steps at ~13.9us each.
(b) Thread mapping of one thread per (j,b) with a serial 256-long k loop. With 4 units x B=10 that is 40 work items for 256 threads, so 84% of threads sat idle.
(c) The hp = hprev + b*H access pattern. Lanes with different b read the same k at a 256-float stride, which hits the same smem bank (~10-way conflicts).
(d) Only 64 of ~132 SMs busy.
(e) Separate per-direction addmm, plus an add_ of b_hr/b_hz on a narrowed slice.
(f) An extra hout -> lout narrow copy per direction.

The fixes are structural and must be applied TOGETHER:
(1) Run both directions concurrently in ONE cooperative launch on disjoint block sets (2x blocks, all co-resident; check cudaOccupancyMaxActiveBlocksPerMultiprocessor). This halves the serial barriers to L*T.
(2) One warp per (unit, b) work item, with k split across lanes and a __shfl_xor reduction. Size the units per block so warps x work items fill every thread.
(3) Lay out the smem h so lanes read consecutive addresses. Lanes should vary k, not b. Otherwise pad the row stride to H+1 to break bank conflicts.
(4) One GEMM for both directions' input projections: W_ih stacked [6H, in], with b_ih and b_hr/b_hz folded into one bias vector (b_hn kept separate). Have the kernel write straight into its half of the [T,B,2H] layer output, which is the next layer's input, so there is no copy.
(5) Check the cudaLaunchCooperativeKernel return value.
(6) Longer term, use the layer wavefront (str-00089).

Process: this run went CHECK on turn 0 and SUBMIT on turn 1 with no timing, while turns remained. On rnn tasks, a local timing against nn.GRU/nn.LSTM after the first passing CHECK is mandatory. A sub-1x timing must trigger these fixes, not a SUBMIT.
[str-00098] helpful=0 harmful=0 :: A.T @ B.T with A stored [K,M], B stored [N,K] (KernelBench L1_018 Matmul_with_transposed_both, M=2048, K=8192, N=4096 fp32, TT layout). MEASURED: the operand-swap kernel (C^T = B_stored @ A_stored: B tile [BN,BK] K-contiguous, A tile [BK,BM] M-contiguous, acc [BN,BM] = tl.dot(b, a, acc, input_precision='tf32x3'), stored through swapped C strides; autotune BM/BN 128-256, BK 32/64, 8 warps, stages 3-4, GM=8 grouped ordering over the N-tiles; masked K loop with 1D kmask broadcast per operand per str-00042; %M/%N wrap) scored 1.4638x (7.53 -> 5.15ms, lo95 1.461), correct, own kernel 100%, vs 0.90x for the in-place swapped-stride version. Treat ~1.46x as the SOLVED BASELINE: resubmitting it cannot raise the score. It went CHECK then SUBMIT with no timing. Next attempt: after the first passing CHECK, time at full size against torch.matmul and keep the fastest: (1) fixed BM=BN=128, BK=64, 8 warps, 3 stages instead of autotune; (2) unmasked K loop (K % BK == 0) and no %M/%N wrap; (3) plain tf32 after printing max abs/rel error; (4) num_stages 4 around 128x128.
[str-00099] helpful=1 harmful=0 :: Linear 1024x8192x8192 -> GroupNorm(512 groups, CPG=16) -> LeakyReLU -> x+x (KernelBench L2_062, matmul family, atol=rtol=1e-4).

Measured, two layouts:
- Fallback: cuBLAS F.linear plus one fused GN+LeakyReLU+2x kernel gives ~1.01x. That kernel processed one (row, group) per program; use whole rows or several groups per program with a 2D [G,CPG] tile instead.
- GN-epilogue fusion: CATASTROPHIC at 0.0371x (232 vs 8.6ms, str-00090). Its kernel was a fixed 128x128x64 tf32x3 NN GEMM on a cached W.t(), with masked loads every K iteration and the tl.reshape(acc,(BM,NG,CPG)) GN reduction in the epilogue.

Turn 0 of that run also failed to compile. NG = BN//CPG was computed inside the kernel and passed to tl.reshape; pass NG as a tl.constexpr from the host instead (str-00070).

Plan for the next attempt:
(1) Write the NN tf32x3 GEMM with an unmasked K loop (1024 and 8192 divide 128) and a bias-only epilogue.
(2) CHECK it, then time it alone against F.linear. It should be about 6.4ms or better (str-00075).
(3) Add a separate lightweight GN+LeakyReLU+2x pass and time the full pipeline against the reference. Fold the 2x into the activation; LeakyReLU commutes with positive scaling.
(4) Only then try GN in the epilogue, with BN=64 or num_stages=2, and re-time it.
Submit only a variant that times above about 1.01x. Otherwise submit the cuBLAS fallback.
[str-00100] helpful=0 harmful=0 :: DenseNet121DenseBlock single block (KernelBench L3_014: batch 10, 224px, 6 layers, C_in 32, growth 32, ref ~15.5ms). Assume the evaluator times train() mode, since that is the nn.Module default (str-00031). The custom TRAINING-BN path then decides the score. The str-00015 eval-mode 1.098x recipe does NOT carry over unless the training branch is also lean.

Measured FAILURE: 0.7669x (20.22 vs 15.51ms, lo95 0.766), own_kernel_time_frac 0.53, correct. The layout was a preallocated block buffer with an eval fold cached on data_ptr, which matches str-00015 (i)-(iv). Every training layer then did the following:
(1) buf.narrow(1,0,cur).contiguous(), a full copy of the growing prefix.
(2) A stats kernel with only C (32-192) blocks x 256 threads, double FMAs and per-element 64-bit div/mod. This is the underfilled layout that measured 0.69x in str-00043, and it re-reduced the WHOLE prefix every layer (O(L^2)).
(3) About 10 eager torch ops for mean/var/fold/running-stat updates.
(4) bn_relu into a fresh tensor with per-element div/mod, then conv to a temp plus copy_.
The run went CHECK then SUBMIT with no timing.

Key fact: buffer channels never change once written, so each channel's BN sum/sumsq is the same for every later layer that reads it. Recipe (untested):
(a) Compute the stats ONCE per channel slice, at the moment it is written. Reduce the input slice up front, then each conv-output slice right after its conv. Cache them in a [C_total] sum/sumsq buffer zeroed once per forward.
(b) Use the str-00043 reduction layout: N*C_slice blocks, float4 fp32 loads, a shared tree, and fp32 atomicAdd. Do not use one block per channel with double FMAs.
(c) Per layer, run only a tiny finalize kernel over the first cur channels. It computes mean and biased var, folds gamma/beta into scale/shift, and updates that BN's running_mean/running_var in-kernel (unbiased var*M/(M-1), momentum). Pass the buffers as tensors (str-00014). No eager torch ops.
(d) bn_relu reads the strided buffer view directly (batch stride C_total*HW), with no .contiguous(). Use a 2D grid (blockIdx.y = n*cur+c) so there is no per-element 64-bit div/mod, and write into a REUSED preallocated pre-activation buffer.
(e) Keep the eval branch: cached fold plus the same bn_relu kernel.
Before SUBMIT, time BOTH train() and eval() at full size against the reference. A recipe copied from an eval-mode 'measured win' is not evidence when the training branch differs.
[str-00101] helpful=1 harmful=0 :: Skinny-M multi-layer Linear+ReLU MLP stacks with huge weights (KernelBench network family, e.g., L3_002 ShallowWideMLP: M=128, 16384->32768->32768->16384 fp32, ~8.6GB of weights). The str-00040 skinny-M kernel transfers VERBATIM when run once per layer. Fuse ReLU into the epilogue of every layer except the last, using an APPLY_RELU constexpr flag. Pass the bias by pointer.

Recipe settings:
- tf32x3.
- W read in place via swapped strides as a [BK,BN] tile (pass w.stride(1), w.stride(0)).
- BM=128, so the whole batch fits in one M tile.
- Autotune BN {64,128} x BK {32,64}, num_warps=8, stages 3-4.
- Plain 2D grid.
- Unmasked K loop with only M/N masks.
- One kernel launch per layer, no W.t() cache.

MEASURED, CONFIRMED TWICE, correct, own kernel 100%:
- Run 1: 1.3175x (33.6 -> 25.5ms, lo95 1.317), about 21.6 effective fp32 TFLOPS, matching the ~1.33x of str-00040 at K=N=32768. It went CHECK then SUBMIT with no timing.
- Run 2: 1.353x (33.75 -> 24.94ms, lo95 1.352), 1.458x vs torch.compile. Same recipe copied as a turn-0 SUBMIT with no CHECK or timing; it worked only because the recipe was proven, and it produced no new information. The small gain over run 1 is likely run variance.
Treat ~1.32-1.35x (~25ms) as the SOLVED BASELINE: resubmitting it cannot raise the score.

Checks before choosing configs:
- Compute the tile count per layer. BN=64 gives 512 tiles at N=32768 and 256 at N=16384, both >= 2x the SM count.
- Do NOT cache W.t() copies; that would double ~8.6GB of weights (OOM risk).
- Use tf32x3, not plain TF32. The related MLP task checked at 1e-4, and TF32 failed there (str-00029).

Next attempt: CHECK on turn 0, then time these untested levers at full size against the reference and keep a variant only if it beats ~25ms (otherwise submit the baseline):
(1) num_stages 4-5 with BK 64/128;
(2) BN=64-only configs;
(3) split-K on the N=16384 last layer (256 tiles at BN=64), with the bias epilogue applied after the reduction.
[str-00102] helpful=0 harmful=1 :: Pure NON-dilated dense fp32 FORWARD Conv2d with no epilogue (KernelBench L1 conv family, e.g. L1_056: Conv2d 64->128, asymmetric 5x7, N=8, 512x256, bias=False). MEASURED WIN: 1.4454x (31.40 -> 21.73ms, lo95 1.4437), correct, own_kernel_time_frac 0.248.

Layout:
(1) Your own load_inline NCHW->NHWC transpose kernel (str-00001 layout, int64 sizes per str-00058) returning out.as_strided([N,C,H,W], [C*HW, 1, W*C, C]).
(2) w = conv.weight.contiguous(memory_format=torch.channels_last).
(3) F.conv2d(x_nhwc, w, bias, stride, padding, dilation, groups), returning the channels_last-strided output as is (passed correctness).
(4) self.use_cl = dilation == 1, else plain nn.Conv2d (dilated NHWC is ~64x slower, str-00051 / str-00027 failure D).
The transpose is real custom work, unlike the zero-bias wrapper (str-00078). Treat ~1.45x as the BASELINE, not the target.

SCOPE LIMIT (measured loss): this recipe is proven ONLY for a forward Conv2d. It does NOT transfer to ConvTranspose2d. On ConvTranspose2d 32->32, 3x3, padding 0, input 8x32x512x1024 (L1_071), the same naive transpose + F.conv_transpose2d with a channels_last weight scored 0.2823x (3.33 -> 11.81ms). The transpose took ~59% of runtime (~7ms for ~537MB, writes scattered at stride C=32), and cuDNN's NHWC transposed path (~4.8ms) was itself slower than the NCHW backward-data path (3.33ms). It was submitted with no timing. Small C (e.g., 32) makes the naive transpose's scattered writes especially costly. Never port it to another conv type or shape without a full-size timing against the reference module.

The naive transpose (one thread per input element, coalesced reads but writes scattered at stride C) takes ~5ms for a ~270MB tensor, several times the bandwidth bound.

Untested headroom; time each at full size after the first passing CHECK instead of submitting immediately:
(a) A shared-memory tiled transpose: per n, 32x32 tiles of the [C, HW] matrix in a [32][33] smem array, so both reads and writes are coalesced.
(b) Cache the channels_last weight in self.<x> keyed on weight.data_ptr() (str-00014/str-00081) instead of converting every forward, with the transpose kernel still launching every forward.
(c) cuDNN TF32 within tolerance (print max error first, str-00029) and torch.backends.cudnn.benchmark=True.
Together these may push past ~1.7x. The recipe is a candidate for other non-dilated dense forward convs, but time it per shape (str-00051).
[str-00103] helpful=0 harmful=0 :: Linear 2048x8192x8192 (nn.Linear NT weight) -> *scale -> hardtanh(-2,2) -> erf-GELU (KernelBench L2_053, matmul family). This task has now been measured with three layouts:
- cuBLAS F.linear + one in-place epilogue kernel: 1.014x (str-00011 attempt 5).
- In-place swapped-stride tf32x3 Triton GEMM with an autotune set that included 256-wide tiles: 0.944x (str-00011 attempt 6, str-00075).
- str-00075 lever-(4) NN conversion: 1.2162x (15.90 -> 13.08ms, lo95 1.2157), 1.14x vs torch.compile, correct, own kernel 100%.

The winning lever-(4) recipe:
- Wt = W.t().contiguous(), cached lazily in self.<x> and keyed on weight.data_ptr() (str-00014).
- An NN tf32x3 Triton GEMM with a FIXED config: BM=BN=128, BK=64, GM=8. Pass num_warps=8 and num_stages=3 explicitly at launch (str-00091).
- An unmasked K loop.
- Epilogue: +bias, *scale, tl.minimum/tl.maximum clamp, then top-level tl.erf GELU (compiles, str-00012), followed by one masked store.

This confirms that lever (4) gives ~1.2x on NT Linear + light-epilogue tasks at M=2048 as well as at M=1024 (1.24x bias+ReLU, 1.185x for the str-00085 chain). Treat ~1.22x as the SOLVED BASELINE for L2_053: resubmitting it cannot raise the score.

The run went CHECK on turn 0 and SUBMIT on turn 1, with no tolerance probe and no timing. It also hardcoded the clamp bounds (-2.0, 2.0). Pass hardtanh_min/max as kernel args instead, so the code is safe to reuse.

Next attempt, after the first passing CHECK:
(1) Run the same kernel with input_precision='tf32' (plain TF32; str-00032 reached 5.5x on a sibling Linear+GELU task), either in place via swapped strides or on the cached Wt. Print the max abs/rel error against the reference. The clamp bounds outputs to GELU([-2,2]), which shrinks absolute error.
(2) If the error fits the tolerance, time it against 13.08ms and submit the faster version.
(3) Otherwise keep tf32x3 and time num_stages 4 and BK 32 around the 128x128 config.
[str-00104] helpful=0 harmful=0 :: ConvTranspose2d 32->32, square 3x3, padding 0, stride 1, input 8x32x512x1024 (KernelBench L1_071; companion to str-00009). This shape now has THREE measured losses, all submitted without local timing:
- The flip+F.conv2d rewrite: 0.75x and 0.7377x.
- channels_last (naive NCHW->NHWC kernel + F.conv_transpose2d with a channels_last weight): 0.2823x (3.33 -> 11.81ms).
The NCHW cuDNN backward-data path (~3.33ms) beats both cuDNN's forward algorithm and its NHWC transposed path here.

On this shape, do NOT submit any memory-format or algorithm switch untimed. After the correctness CHECK, time these remaining untested routes at full size against nn.ConvTranspose2d (bounded warmup/iterations, cuda synchronize), and submit only a variant above 1x:
(1) Direct NCHW F.conv_transpose2d with torch.backends.cudnn.allow_tf32 (if a max-error probe passes the tolerance) plus cudnn.benchmark, paired with a custom kernel doing real work.
(2) A TF32 implicit-GEMM or register-tiled direct kernel spanning all 32 OC per tile (K = 32*9 = 288), reading NCHW directly with coalesced W-contiguous loads.
If every variant times below 1x, do not expect a trivial wrapper to rescue it: that lands at parity, which also fails (str-00027 failure B).
[str-00105] helpful=0 harmful=0 :: Linear 1024x8192x8192 -> AvgPool1d(16) over features -> erf-GELU -> *2.0 -> row max (KernelBench L2_098, matmul family). This is a companion to str-00013 Case 1.

The pooling fold is now CONFIRMED TWICE:
- Run 1: 9.37x (8.00 -> 0.853ms).
- Run 2: 9.5966x (7.95 -> 0.828ms, lo95 9.587), correct, own_kernel_time_frac 0.007.

Recipe, copied from str-00013:
- Build W' and b' in fp64 via reshape(M/k, k, K).mean(1), cast to fp32, and cache them keyed on weight.data_ptr() with no private attributes (str-00014).
- Run F.linear down to 512 outputs.
- Run one Triton program per row: BLOCK_N = next_pow2(N), -inf masking, top-level tl.erf GELU (str-00012), *scale, tl.max, num_warps=8.

Treat ~9.4-9.6x as the SOLVED BASELINE: resubmitting it cannot raise the score. Run 2 went CHECK then SUBMIT with no timing.

The remaining ~99% of runtime is the reduced cuBLAS GEMM (1024x8192x512).

Next attempt: CHECK the fold on turn 0. Then spend the spare turns timing GEMM-side levers at full size against ~0.83ms, and submit only a faster variant.
(1) A Triton GEMM on a cached contiguous W'.t() (NN layout, str-00075 lever 4).
  - Precision: tf32x3 by default. Use plain tf32 only after printing the max abs/rel error against the reference.
  - Use a LIGHT epilogue: bias + GELU + scale, a per-tile row max, then combine the row maxes across N-tiles with tl.atomic_max into a -inf buffer or a small [M, N/BN] partials buffer.
  - Keep it light to avoid the str-00090 heavy-epilogue catastrophe. Time the GEMM alone first, then add the row-max epilogue and re-time.
(2) Mind grid fill. N=512 gives only 4 N-tiles at BN=128 (32 tiles total at BM=128), so try BN=64, smaller BM, or split-K over K=8192. Split-K needs partials plus a reduce, with GELU and max applied after the reduction because GELU is nonlinear.
(3) Pass num_warps and num_stages explicitly at launch (str-00091).
[str-00106] helpful=0 harmful=0 :: Plain A @ B NN layout (KernelBench L1_002 Standard_matrix_multiplication, M=2048, K=8192, N=4096 fp32; same family as the square/symmetric N=4096 cases in str-00004/str-00064). The str-00004 tf32x3 recipe copied verbatim scored 1.4288x (7.36 -> 5.15ms, lo95 1.426), correct, own kernel 100%. Recipe details: tl.dot(input_precision='tf32x3'); a 5-config autotune including 128x256 and 256x128 tiles; BK 32/64, 8 warps, stages 3-4; GM=8 grouped pid; masked K loop with a 1D kmask broadcast per operand (str-00042); %M/%N wrap. This matches the 1.43x figures in str-00004 and both str-00064 runs, so NN dense GEMM at this scale is a SOLVED ~1.43x baseline (now seen three times). Resubmitting it cannot raise the score, and that run learned nothing because it went CHECK on turn 0 and SUBMIT on turn 1 with no timing. Next attempt: after the first passing CHECK, time these at full size against torch.matmul (bounded warmup/iterations, cuda synchronize) and keep the fastest: (1) a fixed BM=BN=128, BK=64, num_warps=8, num_stages=3 config passed explicitly at launch (str-00091), since autotune with 256-wide tiles measured ~20% slower in str-00032; (2) an unmasked K loop with no %M/%N wrap, since 2048/4096/8192 divide every tile; (3) plain tf32 (input_precision='tf32') after printing the max abs/rel error against the reference; (4) num_stages 4 or BK 32 around 128x128. Submit the baseline only if no variant beats ~5.15ms.
[str-00107] helpful=0 harmful=0 :: Linear 1024x8192x8192 (nn.Linear NT weight) -> GroupNorm(512 groups, CPG=16) -> min over features (keepdim) -> +bias(1,N,1,1) (KernelBench L2_075, matmul family, atol=rtol=1e-4). This is a companion to str-00034, str-00090 and str-00099.

MEASURED WIN: 1.1115x (8.47 -> 7.62ms, lo95 1.1109), 1.06x vs torch.compile, correct, own_kernel_time_frac 0.996. It beats the ~1.00x cuBLAS + serial-group-loop fallback (1.0046x). It compiled first try and went CHECK, CHECK, SUBMIT with no timing.

Layout, three parts:
(1) Wt = W.t().contiguous(), cached lazily in self.<x> and keyed on weight.data_ptr() (str-00014/str-00081).
(2) A FIXED tf32x3 NN Triton GEMM: BM=BN=128, BK=64, GM=8, with num_warps=8 and num_stages=3 passed explicitly at launch (str-00091). Use an unmasked K loop and a bias-only epilogue (str-00075 lever 4). Keep GN OUT of the epilogue (str-00090 catastrophe).
(3) One Triton program per row (grid=M, num_warps=8). It loads the whole 8192-wide row (BLOCK_N=next_pow2(N)) plus gamma/beta, then runs tl.reshape(x,(NG,CPG)) with NG and CPG passed as tl.constexpr from the host (str-00070). It reduces axis 1 for the per-group mean and biased var, applies rstd and the affine, takes tl.min with +inf masking, and stores one value per row. Return out.view(M,1) + self.bias to reproduce the (1,N,M,1) broadcast (str-00034 shape pitfall).

Treat ~1.11x as the BASELINE: resubmitting it cannot raise the score. It sits below the ~1.24x the same GEMM reached with a pure bias+ReLU epilogue (~6.4ms), so the GN/min pass plus the 32MB intermediate likely costs about 1.2ms, far more than a 32MB read should take.

Next attempt, after the first passing CHECK, time each lever in isolation at full size:
(a) Time the GEMM alone and the row kernel alone to locate the gap.
(b) Row kernel: process several rows per program, or try num_warps 4/16, so the 8192-element register tile does not spill.
(c) GEMM: drop the %M/%N index wrap, since the dims divide the tiles; try num_stages 4.
(d) Only after (a)-(c), try a per-tile GN + row-min epilogue with tl.atomic_min into an inf-initialized buffer. Re-time it, and revert on any ~10x jump (str-00090).
Submit only a variant that beats ~7.62ms.
[str-00108] helpful=0 harmful=0 :: EfficientNetB0 (KernelBench L3_022, batch 10, 224px, ref ~5.7ms): keeping the WHOLE eager network and fusing only the tail is a measured LOSS. This is the fifth ~0.8x result in the str-00026 family. The tail fusion was conv2 output -> bn2 + ReLU + global avgpool in ONE custom per-(n,c) reduction kernel, applying scale/shift in registers.

Result: correct but 0.7974x (7.18 vs 5.72ms, lo95 0.7965), 0.74x vs torch.compile, own_kernel_time_frac 0.011.

Why it lost:
(1) In train mode, the BN stats were computed with eager conv2_out.mean and .var. That is 2 extra full reductions over the 1280xHxW tensor, plus several running-stat mul_/add_ launches and scale/shift torch ops on every forward. These cost more than the one saved pass.
(2) In eval mode, the fold (sqrt/div/mul/sub) was recomputed per forward, about 4 tiny launches, instead of being cached once.
(3) Every MBConv was re-wrapped as a new module, and the residual was changed to the non-inplace `x = x + identity`. That violates str-00026 rule (2).
(4) The process was CHECK on turn 0, then an untimed SUBMIT on turn 1 (str-00002, str-00026 rule 5).

The str-00010 BN+avgpool algebra is correct, but on a launch-bound network it pays only if the training-mode stats are computed IN-KERNEL. Use per-(n,c) sum/sumsq in the same reduction, a tiny finalize kernel that updates running stats via buffer tensors (str-00014), and an eval fold cached once. Even done perfectly, a tail kernel touches about 1% of runtime and cannot beat the reference alone.

A win must cut launches across the blocks:
- fused expand->dw->project per MBConv (respect the small-map guard in str-00095);
- fold work cached, never recomputed per forward;
- stats/finalize kernels instead of eager torch reductions in train mode;
- the residual kept in place (add_).

Time the evaluated mode (str-00031) at full size against the reference before SUBMIT. Treat any layout that leaves all MBConvs eager as a guaranteed ~0.8x.
[str-00109] helpful=0 harmful=0 :: Linear 16384x4096x4096 (nn.Linear NT weight) -> training BatchNorm1d -> erf-GELU -> ReLU (KernelBench L2_041, matmul family, atol=rtol=1e-4). This is a companion to str-00016 and str-00090.

Three tails have been measured on the same 'strict-FP32 cuBLAS F.linear + custom tail' layout:
- Slow tail (one thread per column, double FMAs): 0.9517x.
- First tail: 1.014x.
- LEAN tail: 1.0224x (30.55 -> 29.88ms, lo95 1.0214), 1.002x vs torch.compile, own_kernel_time_frac 0.020. It compiled and passed CHECK on turn 0 and was submitted on turn 1 with no timing.

The lean tail used the str-00001 layout, long sizes (str-00058), and three kernels:
(1) A 2D-grid stats kernel: 256-thread column tiles x 128 row chunks, fp32 per-thread sum/sumsq, and atomicAdd into torch.zeros buffers.
(2) A finalize kernel in double math. It computes biased var for scale/shift and updates running_mean/running_var in-kernel (unbiased var*M/(M-1), momentum), with the buffers passed as tensors (str-00014).
(3) An in-place 2D-grid apply (256 row blocks) with sc/sh in registers, no per-element idx%N, and accurate erff GELU followed by ReLU.
It also had an eval branch that computes sc/sh once per thread from running stats (str-00031).

The custom kernels are now only ~2% of runtime. ~1.02x is therefore the SOLVED CEILING for 'cuBLAS + any tail' on this task. Resubmitting it, or tweaking the tail further (float4, dropping the torch.zeros launches), can move the score by at most ~1%.

The only lever with real upside is the GEMM itself. Next attempt:
- CHECK the lean-tail baseline on turn 0.
- Build a cached-W.t() NN tf32x3 GEMM (fixed 128x128x64, 8 warps / 3 stages passed explicitly at launch per str-00091, unmasked K loop, bias-only light epilogue). Keep the BN stats in the separate 2D-grid kernel, never in the epilogue (str-00090 0.0336x signature at this exact M=16384).
- Time that GEMM ALONE against F.linear at full size.
- Adopt it with the lean tail only if the full pipeline beats ~29.9ms. Otherwise resubmit the lean-tail baseline.
- Plain TF32 likely fails at 1e-4, because BN followed by GELU/ReLU does not hide relative error. Print the max error before trying it.

## verified_technique_library
[kf-E-H1] helpful=0 harmful=0 :: ### Verified technique (E-H1, admitted by the KFORGE engine): Dense Conv2d/Conv3d as a Triton implicit GEMM with an in-register bf16 3-product split and a fused epilogue
For dense convolutions with reduction length K = Ci*kh*kw (*kd) >= ~256 and Co >= 64, a Triton implicit-GEMM kernel should beat strict-fp32 cuDNN by 1.5-3x on the conv itself, and more once bias/activation/pool epilogues are fused. The kernel splits each loaded fp32 operand into bf16 hi+lo in registers and issues three bf16 tl.dot calls into one fp32 accumulator.
Physics: cuDNN under strict fp32 runs on the 19.5 TFLOP/s SIMT path. A100 bf16 mma peaks at 312 TFLOP/s, so even at 40% efficiency the three products deliver about 40 TFLOP/s of fp32-equivalent work. P1 already shows the split stays within about 0.34 of the tolerance for K <= 8192, and conv K is almost always below that. This is not the tf32x3 path that measured only 1.1-1.2x: bf16 mma has 2x the TF32 rate, and the hi/lo split is exact in fp32 adds. Implicit GEMM never materializes im2col, so it avoids the bandwidth loss of the AlexNet failure.
How: Map GEMM M = N*OH*OW (output pixels), GEMM N = Co, and K = (r, s, ci), with ci innermost. Use one cheap Triton prep kernel to convert x to channels_last and emit xh = bf16(x) and xl = bf16(x - float(xh)); these two bf16 tensors take the same bytes as fp32. Cache the weight reorganised as [kh*kw*Ci, Co] in hi/lo bf16 form in self.<x>, keyed on weight.data_ptr(). In the kernel, loop over (r, s), compute shifted input coordinates with padding masks, and load BK = 32 channel slabs of xh/xl and wh/wl. Accumulate acc += dot(xl, wh) + dot(xh, wl) first, then acc += dot(xh, wh), with fp32 accumulation. Use BM = 128, BN = 64 or 128, num_stages = 3-4, and a small autotune list. In the epilogue, add bias, apply the following elementwise ops (and pool if the window fits the tile), and store straight to NCHW with int64 offsets. Consecutive M indices are consecutive ow, so the stores coalesce. Validate against F.conv2d at atol = rtol = 1e-4 before timing, and fall back to cuDNN when K < 128 or Co < 32.
- Verified on training problems (protected evaluator): L1_080 conv_standard_2D_square_input_asymmetric_kernel___dilated____padded__: 2.29x vs 0.44x with the previous library; L2_032 Conv2d_Scaling_Min: 2.50x vs 1.11x with the previous library.
- Verified example implementing it (L2_032, 2.50x; adapt what applies):
```python
import torch
import torch.nn as nn
import triton
import triton.language as tl


@triton.autotune(
    configs=[
        triton.Config({"BM": 128, "BK": 32}, num_stages=3, num_warps=8),
        triton.Config({"BM": 64, "BK": 32}, num_stages=4, num_warps=4),
        triton.Config({"BM": 128, "BK": 64}, num_stages=2, num_warps=8),
    ],
    key=["M", "Ci", "Co"],
)
@triton.jit
def conv_min_kernel(x, wh, wl, b, out, M, Ci, H, W, OH, OW, Co, scale,
                    KH: tl.constexpr, KW: tl.constexpr, BN: tl.constexpr,
                    BM: tl.constexpr, BK: tl.constexpr):
    pid = tl.program_id(0)
    rm = pid * BM + tl.arange(0, BM)
    mm = rm < M
    rmc = tl.where(mm, rm, 0)
    ow = rmc % OW
    t = rmc // OW
    oh = t % OH
    n = t // OH
    base = n * (Ci * H * W) + oh * W + ow
    rk = tl.arange(0, BK)
    rn = tl.arange(0, BN)
    nm = rn < Co
    HW = H * W
    nci = tl.cdiv(Ci, BK)
    total = KH * KW * nci
    acc = tl.zeros((BM, BN), dtype=tl.float32)
    for kk in range(0, total):
        tap = kk // nci
        c0 = (kk - tap * nci) * BK
        r = tap // KW
        s = tap - r * KW
        ci = c0 + rk
        cm = ci < Ci
        a = tl.load(x + base[:, None] + (r * W + s) + ci[None, :] * HW,
                    mask=cm[None, :], other=0.0)
        krow = tap * Ci + ci
        boff = krow[:, None] * Co + rn[None, :]
        bmask = cm[:, None] & nm[None, :]
        bh = tl.load(wh + boff, mask=bmask, other=0.0)
        bl = tl.load(wl + boff, mask=bmask, other=0.0)
        ah = a.to(tl.bfloat16)
        al = (a - ah.to(tl.float32)).to(tl.bfloat16)
        acc = tl.dot(al, bh, acc)
        acc = tl.dot(ah, bl, acc)
        acc = tl.dot(ah, bh, acc)
    bias = tl.load(b + rn, mask=nm, other=0.0)
    v = (acc + bias[None, :]) * scale
    v = tl.where(nm[None, :], v, float("inf"))
    mn = tl.min(v, axis=1)
    tl.store(out + rm, mn, mask=mm)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, scale_factor):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.scale_factor = scale_factor

    def forward(self, x):
        x = x.contiguous()
        w = self.conv.weight
        Co, Ci, KH, KW = w.shape
        N, _, H, W = x.shape
        OH = H - KH + 1
        OW = W - KW + 1
        wr = w.permute(2, 3, 1, 0).reshape(KH * KW * Ci, Co).contiguous()
        whi = wr.to(torch.bfloat16)
        wlo = (wr - whi.float()).to(torch.bfloat16)
        out = torch.empty((N, 1, OH, OW), device=x.device, dtype=torch.float32)
        M = N * OH * OW
        BN = max(16, triton.next_power_of_2(Co))
        grid = lambda META: (triton.cdiv(M, META["BM"]),)
        conv_min_kernel[grid](x, whi, wlo, self.conv.bias, out, M, Ci, H, W, OH, OW, Co,
                              float(self.scale_factor), KH=KH, KW=KW, BN=BN)
        return out
```
