## Technique library (KFORGE: techniques verified by the protected evaluator on training problems; use what fits this problem, ignore the rest)

### Verified technique (E-H3obs, observed in a verified solution; admitted by the KFORGE engine): Split-bf16 implicit-GEMM transposed conv via padded channels-last taps
Mechanism: A stride-1 ConvTranspose3d is rewritten as an ordinary convolution, using a flipped and transposed weight and an input zero-padded by K−1−p. A prep kernel writes the input once into a channels-last buffer that already contains the spatial halo, with channels padded to a power of two (at least 16). The main kernel is then an implicit GEMM: each program owns BM output voxels, loops over the KD·KH·KW taps, and for each tap does an unmasked contiguous [BM×CIP] activation load against a [CIP×COP] weight slice. It accumulates with tensor-core tl.dot and writes straight to NCDHW with the bias fused.

Numerics: Every fp32 operand is split into a bf16 high part and a bf16 residual (lo = x − bf16(x)). Three bf16 MMAs are computed (hi·hi, hi·lo, lo·hi), dropping lo·lo, all accumulated in fp32. The small terms go into a separate accumulator that is added at the end, which recovers roughly 16 mantissa bits and passes atol/rtol 1e-4 while running on bf16 tensor cores. This avoids TF32, which would fail the tolerance, and avoids slow FP32 FMA.

Where it applies: The pattern suits any conv or transposed conv with modest channel counts (Cin and Cout small enough to fit one tile) and many spatial positions. It also fits any fp32 GEMM-like op that needs near-fp32 accuracy but wants tensor-core throughput.

What to watch:
- The pre-padded buffers cost extra memory and a separate prep pass.
- Channel padding wastes MMA work when C is far from a power of two.
- Large C or Co would need a K-loop and N-tiling rather than whole-channel tiles.
- Stride > 1, groups, or output_padding fall back to PyTorch.
- Error grows with the reduction length (taps × channels), so recheck the tolerance for larger kernels or inputs with a large dynamic range.
- Verified on a training problem (protected evaluator): L1_058 conv_transposed_3D__asymmetric_input__asymmetric_kernel: 4.93x vs 2.17x with the previous library.
- Verified example (L1_058, 4.93x; adapt what applies):
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl


@triton.jit
def prep_x_kernel(x_ptr, xh_ptr, xl_ptr, P, C, Di, Hi, Wi, Dp, Hp, Wp, pd, ph, pw,
                  CIP: tl.constexpr, BP: tl.constexpr):
    p = tl.program_id(0) * BP + tl.arange(0, BP)
    pm = p < P
    pc = tl.where(pm, p, 0)
    wq = pc % Wp
    t = pc // Wp
    hq = t % Hp
    t = t // Hp
    dq = t % Dp
    n = t // Dp
    idd = dq - pd
    ih = hq - ph
    iw = wq - pw
    valid = pm & (idd >= 0) & (idd < Di) & (ih >= 0) & (ih < Hi) & (iw >= 0) & (iw < Wi)
    ci = tl.arange(0, CIP)
    Sin = Di * Hi * Wi
    sp = (idd * Hi + ih) * Wi + iw
    addr = n[:, None] * (C * Sin) + ci[None, :] * Sin + sp[:, None]
    m = valid[:, None] & (ci[None, :] < C)
    v = tl.load(x_ptr + addr, mask=m, other=0.0)
    h = v.to(tl.bfloat16)
    lo = (v - h.to(tl.float32)).to(tl.bfloat16)
    o = pc[:, None] * CIP + ci[None, :]
    sm = pm[:, None] & (ci[None, :] < CIP)
    tl.store(xh_ptr + o, h, mask=sm)
    tl.store(xl_ptr + o, lo, mask=sm)


@triton.autotune(
    configs=[
        triton.Config({"BM": 128}, num_stages=3, num_warps=4),
        triton.Config({"BM": 128}, num_stages=4, num_warps=4),
        triton.Config({"BM": 256}, num_stages=3, num_warps=8),
        triton.Config({"BM": 64}, num_stages=4, num_warps=4),
    ],
    key=["M", "NT", "CO"],
)
@triton.jit
def convt_kernel(xh_ptr, xl_ptr, wh_ptr, wl_ptr, b_ptr, y_ptr, M, S, CO, NT,
                 Do, Ho, Wo, Dp, Hp, Wp, KH, KW,
                 CIP: tl.constexpr, COP: tl.constexpr, HAS_BIAS: tl.constexpr, BM: tl.constexpr):
    pid = tl.program_id(0)
    m = pid * BM + tl.arange(0, BM)
    mm = m < M
    mc = tl.where(mm, m, 0)
    ow = mc % Wo
    t = mc // Wo
    oh = t % Ho
    t = t // Ho
    od = t % Do
    n = t // Do
    base = ((n * Dp + od) * Hp + oh) * Wp + ow
    ci = tl.arange(0, CIP)
    co = tl.arange(0, COP)
    acc = tl.zeros((BM, COP), dtype=tl.float32)
    acc2 = tl.zeros((BM, COP), dtype=tl.float32)
    for tap in range(NT):
        kw = tap % KW
        tt = tap // KW
        kh = tt % KH
        kd = tt // KH
        off = (kd * Hp + kh) * Wp + kw
        pa = (base + off)[:, None] * CIP + ci[None, :]
        ah = tl.load(xh_ptr + pa)
        al = tl.load(xl_ptr + pa)
        pw_ = tap * (CIP * COP) + ci[:, None] * COP + co[None, :]
        bh = tl.load(wh_ptr + pw_)
        bl = tl.load(wl_ptr + pw_)
        acc = tl.dot(ah, bh, acc)
        acc2 = tl.dot(ah, bl, acc2)
        acc2 = tl.dot(al, bh, acc2)
    v = acc + acc2
    if HAS_BIAS:
        v = v + tl.load(b_ptr + co, mask=co < CO, other=0.0)[None, :]
    s = mc - n * S
    addr = n[:, None] * (CO * S) + co[None, :] * S + s[:, None]
    tl.store(y_ptr + addr, v, mask=mm[:, None] & (co[None, :] < CO))


def _p2(v):
    r = 16
    while r < v:
        r *= 2
    return r


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

    def forward(self, x):
        KD, KH, KW = self.ks
        pd0, ph0, pw0 = self.pad
        ok = (self.st == (1, 1, 1) and self.groups == 1 and self.opad == (0, 0, 0)
              and pd0 <= KD - 1 and ph0 <= KH - 1 and pw0 <= KW - 1)
        if not ok:
            return self.conv_transpose3d(x)
        x = x.contiguous()
        N, C, Di, Hi, Wi = x.shape
        w = self.conv_transpose3d.weight
        Co = w.shape[1]
        CIP = _p2(C)
        COP = _p2(Co)
        pd, ph, pw = KD - 1 - pd0, KH - 1 - ph0, KW - 1 - pw0
        Dp, Hp, Wp = Di + 2 * pd, Hi + 2 * ph, Wi + 2 * pw
        Do, Ho, Wo = Dp - KD + 1, Hp - KH + 1, Wp - KW + 1
        wf = w.flip(2, 3, 4).permute(2, 3, 4, 0, 1)
        wpad = torch.zeros((KD, KH, KW, CIP, COP), device=x.device, dtype=torch.float32)
        wpad[:, :, :, :C, :Co] = wf
        wh = wpad.to(torch.bfloat16)
        wl = (wpad - wh.float()).to(torch.bfloat16)
        P = N * Dp * Hp * Wp
        xh = torch.empty((P, CIP), device=x.device, dtype=torch.bfloat16)
        xl = torch.empty((P, CIP), device=x.device, dtype=torch.bfloat16)
        BP = 64
        prep_x_kernel[(triton.cdiv(P, BP),)](x, xh, xl, P, C, Di, Hi, Wi, Dp, Hp, Wp, pd, ph, pw,
                                             CIP=CIP, BP=BP)
        S = Do * Ho * Wo
        M = N * S
        y = torch.empty((N, Co, Do, Ho, Wo), device=x.device, dtype=torch.float32)
        b = self.conv_transpose3d.bias
        NT = KD * KH * KW
        grid = lambda META: (triton.cdiv(M, META["BM"]),)
        convt_kernel[grid](xh, xl, wh, wl, b if b is not None else y, y, M, S, Co, NT,
                           Do, Ho, Wo, Dp, Hp, Wp, KH, KW,
                           CIP=CIP, COP=COP, HAS_BIAS=b is not None)
        return y
```


## Kernel memory (items verified by the protected evaluator on training problems; may not all apply)

### Lessons about this environment (from official checker verdicts and verified outcomes)
- Pitfall: A reply without a ```python fenced code block was rejected by the checker. Always wrap the full ModelNew source in a ```python block.
- Checker rule: getattr needs a literal, non-private attribute name
- Pitfall: Output mismatch beyond 1e-4 occurred repeatedly. Fix by forcing exact fp32 tl.dot(allow_tf32=False) and exact masking: K bounds at block granularity, tril mask, and zero-initialized output for skipped tiles.
- Checker rule: forward modified its input tensors in place
- Checker rule: do not define a class named Model (the reference); define ModelNew only
- Process: never SUBMIT code that has not passed ACTION: CHECK in this episode. Untested submissions that fail to compile or mismatch the reference score 0; a CHECK costs one reply and shows the compiler output.
- Pitfall: cpp_sources is compiled by plain c++ without CUDA headers: never use cudaStream_t or at::cuda::getCurrentCUDAStream there. Put wrappers needing streams in cuda_sources, keeping only plain-type prototypes in cpp_sources.
- Pitfall: If cpp_sources must reference CUDA types or the current stream, explicitly #include <cuda_runtime.h> and <ATen/cuda/CUDAContext.h>; torch/extension.h does not provide them.
- Checker rule: assignment to attribute 'm.conv1' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.bn1' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.conv2' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.bn2' is not allowed (only self.<x>)
- Checker rule: assignment to attribute 'm.downsample' is not allowed (only self.<x>)
- Checker rule: At least one of your own kernels must actually launch during forward; a pure torch path (even with a compiled but unused kernel) is rejected as no_custom_kernel.

### Pitfalls seen on earlier conv_transpose problems (wrong results or slower than eager)
- Pitfall: Don't rewrite a transposed conv as flipped-weight F.conv3d plus a weight-flip kernel. It still calls cuDNN, adds a launch and allocation per forward, and ran slower than eager ConvTranspose.
- Pitfall: Avoid launching a per-forward weight-transform kernel (flip/transpose) before a library call. Cache the transformed weight once, or fuse the transform, since extra launches and memory traffic make it slower than eager.
- Pitfall: Don'...'s native transposed conv is usually as fast or faster, so this measured slower.
- Pitfall: In Triton kernels over large tensors (e.g. 3D conv outputs), compute offsets in int64 (cast program_id/indices with .to(tl.int64)). int32 products like n*C*S can overflow and crash, leaving no output during timing.
- Pitfall: Never launch a pure copy/identity kernel into a fresh buffer just to have a custom kernel; that adds a full memory round-trip and guarantees slowdown versus eager.
- Pitfall: Don't rewrite conv_transpose as flipped-weight F.conv2d plus a separate elementwise kernel; the per-call weight flip and extra full-output pass make it slower. Fuse bias into a real kernel, or skip pointless copies.

### A verified kernel for a related conv_transpose problem (L1_065 conv_transposed_2D__square_input__asymmetric_kernel, 2.21x): adapt what applies
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl


@triton.jit
def flip_kernel(w_ptr, o_ptr, IC, OC, KH, KW, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = IC * OC * KH * KW
    mask = offs < total
    o = offs
    kw = o % KW
    t = o // KW
    kh = t % KH
    t = t // KH
    oc = t % OC
    ic = t // OC
    src = ((ic * OC + oc) * KH + kh) * KW + kw
    fkh = KH - 1 - kh
    fkw = KW - 1 - kw
    dst = ((oc * IC + ic) * KH + fkh) * KW + fkw
    v = tl.load(w_ptr + src, mask=mask, other=0.0)
    tl.store(o_ptr + dst, v, mask=mask)


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

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

### A verified kernel for a related conv_transpose problem (L2_011 ConvTranspose2d_BatchNorm_Tanh_MaxPool_GroupNorm, 4.06x): adapt what applies
```python
import torch
import torch.nn as nn
import torch.nn.functional as F
import triton
import triton.language as tl


@triton.jit
def flip_weight(w_ptr, o_ptr, IC, OC, KH, KW, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = IC * OC * KH * KW
    m = offs < total
    kw = offs % KW
    t = offs // KW
    kh = t % KH
    t = t // KH
    oc = t % OC
    ic = t // OC
    nkh = KH - 1 - kh
    nkw = KW - 1 - kw
    dst = ((oc * IC + ic) * KH + nkh) * KW + nkw
    v = tl.load(w_ptr + offs, mask=m, other=0.0)
    tl.store(o_ptr + dst, v, mask=m)


@triton.jit
def reduce_nc(x_ptr, sum_ptr, sq_ptr, S, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = tl.arange(0, BLOCK)
    base = pid * S
    acc = tl.zeros((), dtype=tl.float32)
    accsq = tl.zeros((), dtype=tl.float32)
    for i in range(0, S, BLOCK):
        idx = i + offs
        m = idx < S
        v = tl.load(x_ptr + base + idx, mask=m, other=0.0)
        acc += tl.sum(v)
        accsq += tl.sum(v * v)
    tl.store(sum_ptr + pid, acc)
    tl.store(sq_ptr + pid, accsq)


@triton.jit
def bn_tanh_pool(x_ptr, out_ptr, scale_ptr, shift_ptr,
                 N, C, H, W, Ho, Wo, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = N * C * Ho * Wo
    m = offs < total
    wo = offs % Wo
    t = offs // Wo
    ho = t % Ho
    t = t // Ho
    c = t % C
    n = t // C

    sc = tl.load(scale_ptr + c, mask=m, other=0.0)
    sh = tl.load(shift_ptr + c, mask=m, other=0.0)

    hbase = ho * 2
    wbase = wo * 2
    inbase = ((n * C + c) * H + hbase) * W + wbase

    v00 = tl.load(x_ptr + inbase, mask=m, other=-1e30)
    v01 = tl.load(x_ptr + inbase + 1, mask=m, other=-1e30)
    v10 = tl.load(x_ptr + inbase + W, mask=m, other=-1e30)
    v11 = tl.load(x_ptr + inbase + W + 1, mask=m, other=-1e30)

    t00 = tl.extra.cuda.libdevice.tanh(sc * v00 + sh)
    t01 = tl.extra.cuda.libdevice.tanh(sc * v01 + sh)
    t10 = tl.extra.cuda.libdevice.tanh(sc * v10 + sh)
    t11 = tl.extra.cuda.libdevice.tanh(sc * v11 + sh)

    mx = tl.maximum(tl.maximum(t00, t01), tl.maximum(t10, t11))
    tl.store(out_ptr + offs, mx, mask=m)


@triton.jit
def gn_apply(y_ptr, out_ptr, mean_ptr, inv_ptr, w_ptr, b_ptr,
             N, C, So, cpg, GPC, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    total = N * C * So
    m = offs < total
    t = offs // So
    c = t % C
    n = t // C
    g = c // cpg
    ng = n * GPC + g

    mean = tl.load(mean_ptr + ng, mask=m, other=0.0)
    inv = tl.load(inv_ptr + ng, mask=m, other=0.0)
    w = tl.load(w_ptr + c, mask=m, other=0.0)
    b = tl.load(b_ptr + c, mask=m, other=0.0)
    v = tl.load(y_ptr + offs, mask=m, other=0.0)
    r = (v - mean) * inv * w + b
    tl.store(out_ptr + offs, r, mask=m)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, groups, num_groups):
        super(ModelNew, self).__init__()
        self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)
        self.batch_norm = nn.BatchNorm2d(out_channels)
        self.max_pool = nn.MaxPool2d(kernel_size=2, stride=2)
        self.group_norm = nn.GroupNorm(num_groups=num_groups, num_channels=out_channels)
        self.num_groups = num_groups
        self.stride = stride
        self.padding = padding
        self.kernel_size = kernel_size

    def forward(self, x):
        W_t = self.conv_transpose.weight  # (IC, OC, KH, KW)
        IC, OC, KH, KW = W_t.shape
        if self.stride == 1:
            wf = torch.empty((OC, IC, KH, KW), device=W_t.device, dtype=W_t.dtype)
            total = IC * OC * KH * KW
            flip_weight[(triton.cdiv(total, 256),)](W_t, wf, IC, OC, KH, KW, BLOCK=256)
            pad = KH - 1 - self.padding
            c = F.conv2d(x, wf, self.conv_transpose.bias, stride=1, padding=pad)
        else:
            c = self.conv_transpose(x)

        N, C, H, W = c.shape
        c = c.contiguous()
        S = H * W
        c2 = c.view(N * C, S)

        ssum = torch.empty(N * C, device=c.device, dtype=torch.float32)
        ssq = torch.empty(N * C, device=c.device, dtype=torch.float32)
        reduce_nc[(N * C,)](c2, ssum, ssq, S, BLOCK=1024)

        ssum = ssum.view(N, C)
        ssq = ssq.view(N, C)
        M = N * S
        mean_c = ssum.sum(0) / M
        var_c = ssq.sum(0) / M - mean_c * mean_c
        scale_c = (self.batch_norm.weight / torch.sqrt(var_c + self.batch_norm.eps)).contiguous()
        shift_c = (self.batch_norm.bias - mean_c * scale_c).contiguous()

        Ho = H // 2
        Wo = W // 2
        y = torch.empty(N, C, Ho, Wo, device=c.device, dtype=torch.float32)
        total = N * C * Ho * Wo
        grid = ((total + 255) // 256,)
        bn_tanh_pool[grid](c, y, scale_c, shift_c, N, C, H, W, Ho, Wo, BLOCK=256)

        G = self.num_groups
        cpg = C // G
        So = Ho * Wo
        seg = cpg * So
        yv = y.view(N * G, seg)
        gsum = torch.empty(N * G, device=c.device, dtype=torch.float32)
        gsq = torch.empty(N * G, device=c.device, dtype=torch.float32)
        reduce_nc[(N * G,)](yv, gsum, gsq, seg, BLOCK=1024)
        gmean = gsum / seg
        gvar = gsq / seg - gmean * gmean
        ginv = 1.0 / torch.sqrt(gvar + self.group_norm.eps)

        out = torch.empty_like(y)
        totalg = N * C * So
        gridg = ((totalg + 255) // 256,)
        gn_apply[gridg](y, out, gmean.contiguous(), ginv.contiguous(),
                        self.group_norm.weight.contiguous(), self.group_norm.bias.contiguous(),
                        N, C, So, cpg, G, BLOCK=256)
        return out
```