```python
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

torch.backends.cudnn.benchmark = True
torch.backends.cudnn.allow_tf32 = True

cuda_src = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

// Direct NCHW implicit-GEMM transposed 3x3 conv (stride=1, pad=0, opad=0).
//   out[co, y, x] = sum_{ky,kx,ci} w[ci, co, ky, kx] * in[ci, y-ky, x-kx]
// GEMM: C[m=co, n=(y,x)] += A[m=co][k] * B[k=(ky,kx,ci)][n], TF32 tensor cores.
// 256 threads (8 warps, 2x4), tile M=64 x N=256 (16x16 spatial), BK=32,
// 4-stage cp.async pipeline. A loaded as 16B chunks, B as 16B x-contiguous
// chunks into a k-major padded smem layout (bank-conflict-free frag loads).
// Natural k order everywhere (consistent A/B fragment placement makes the
// mma result independent of the within-group k permutation).

#define BK 32
#define S_B 264
#define STAGES 4
#define A_STAGE_FLOATS (64 * BK)
#define B_STAGE_FLOATS (BK * S_B)
#define SMEM_FLOATS (STAGES * (A_STAGE_FLOATS + B_STAGE_FLOATS))

__device__ __forceinline__ void cpA16(void* dst, const void* src) {
    unsigned d = (unsigned)__cvta_generic_to_shared(dst);
    asm volatile("cp.async.ca.shared.global [%0], [%1], 16;\n" :: "r"(d), "l"(src));
}
__device__ __forceinline__ void cpB16(void* dst, const void* src, int srcsize) {
    unsigned d = (unsigned)__cvta_generic_to_shared(dst);
    asm volatile("cp.async.cg.shared.global [%0], [%1], 16, %2;\n" :: "r"(d), "l"(src), "r"(srcsize));
}
__device__ __forceinline__ void cp_commit() { asm volatile("cp.async.commit_group;\n"); }
__device__ __forceinline__ void cp_wait2() { asm volatile("cp.async.wait_group 2;\n"); }
__device__ __forceinline__ float lds32(const void* p) {
    float v;
    unsigned a = (unsigned)__cvta_generic_to_shared(p);
    asm volatile("ld.shared.f32 %0, [%1];\n" : "=f"(v) : "r"(a));
    return v;
}
__device__ __forceinline__ float2 lds64(const void* p) {
    float2 v;
    unsigned a = (unsigned)__cvta_generic_to_shared(p);
    asm volatile("ld.shared.v2.f32 {%0,%1}, [%2];\n" : "=f"(v.x), "=f"(v.y) : "r"(a));
    return v;
}
__device__ __forceinline__ void mma_tf32(float& c0, float& c1, float& c2, float& c3,
                                         unsigned a0, unsigned a1, unsigned a2, unsigned a3,
                                         unsigned b0, unsigned b1) {
    asm volatile(
        "mma.sync.aligned.m16n8k8.row.col.f32.tf32.tf32.f32 "
        "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};\n"
        : "+f"(c0), "+f"(c1), "+f"(c2), "+f"(c3)
        : "r"(a0), "r"(a1), "r"(a2), "r"(a3), "r"(b0), "r"(b1));
}

__global__ void __launch_bounds__(256) tconv2d_pipe_kernel(
    const float* __restrict__ in,    // [B, Cin, H, W]
    const float* __restrict__ wp,    // [Cout, K] k=(ky,kx,ci), ci fastest
    const float* __restrict__ bias,  // [Cout] or nullptr
    float* __restrict__ out,         // [B, Cout, Ho, Wo]
    int Bn, int Cin, int Cout, int H, int W, int Ho, int Wo, int K, int mtiles)
{
    extern __shared__ float smem[];
    float* As = smem;
    float* Bs = smem + STAGES * A_STAGE_FLOATS;

    const int b = blockIdx.z / mtiles;
    const int mt = blockIdx.z % mtiles;
    const int m0 = mt * 64;
    const int y0 = blockIdx.y * 16;
    const int x0 = blockIdx.x * 16;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int lane = tid & 31;
    const int g = lane >> 2;      // 0..7
    const int t = lane & 3;       // 0..3
    const int m_warp = (warp & 1) * 32;
    const int n_warp = (warp >> 1) * 64;

    const long long inHW = (long long)H * W;
    const float* __restrict__ inb = in + (long long)b * Cin * inHW;

    // A staging: thread handles chunks tid and tid+256
    const int a_m0 = tid >> 3, a_c0 = tid & 7;
    const int a_m1 = (tid + 256) >> 3, a_c1 = (tid + 256) & 7;
    const float* a_src0 = wp + (long long)(m0 + a_m0) * K + a_c0 * 4;
    const float* a_src1 = wp + (long long)(m0 + a_m1) * K + a_c1 * 4;
    const int a_dst0 = a_m0 * BK + a_c0 * 4;
    const int a_dst1 = a_m1 * BK + a_c1 * 4;

    // B staging: thread handles 8 chunks: k = (tid>>2) + 64j, c = tid&3
    const int b_k = tid >> 2;         // 0..63
    const int b_c = tid & 3;          // 0..3
    const int b_tx = b_c * 4;         // x offset 0,4,8,12

    const int NS = K >> 5;            // K / 32

    auto load_stage = [&](int s) {
        const int buf = s % STAGES;
        const int k0 = s << 5;
        cpA16(As + buf * A_STAGE_FLOATS + a_dst0, a_src0 + k0);
        cpA16(As + buf * A_STAGE_FLOATS + a_dst1, a_src1 + k0);

        const int tap = k0 / Cin;                 // uniform (Cin % 32 == 0)
        const int cibase = k0 - tap * Cin;
        const int ky = tap / 3;
        const int kx = tap - 3 * ky;
        float* Bbuf = Bs + buf * B_STAGE_FLOATS;
        #pragma unroll
        for (int j = 0; j < 8; ++j) {
            const int k = b_k + (j << 6);         // 0..511
            const int ty = k >> 5;                // 0..15
            const int kk = k & 31;
            const int gy = y0 + ty - ky;
            const int gx = x0 + b_tx - kx;
            const bool yv = (gy >= 0) & (gy < H);
            const bool xv_full = yv & (gx >= 0) & (gx + 3 < W);
            const bool xv_part = yv & (gx < W) & (gx + 3 >= W) & (gx >= 0);
            const float* src = inb + (long long)(cibase + kk) * inHW
                             + (long long)gy * W + gx;
            int ssz = 0;
            if (xv_full) ssz = 16;
            else if (xv_part) ssz = (W - gx) * 4;
            else src = inb;
            cpB16(Bbuf + kk * S_B + ty * 16 + b_tx, src, ssz);
        }
    };

    load_stage(0); cp_commit();
    load_stage(1); cp_commit();
    load_stage(2); cp_commit();

    float acc[2][8][4];
    #pragma unroll
    for (int mi = 0; mi < 2; ++mi)
        #pragma unroll
        for (int ni = 0; ni < 8; ++ni)
            #pragma unroll
            for (int r = 0; r < 4; ++r) acc[mi][ni][r] = 0.f;

    const int a_row = (m_warp + g) * BK + 2 * t;
    const int a_row8 = (m_warp + g + 8) * BK + 2 * t;
    const int b_col = n_warp + g;

    for (int s = 0; s < NS; ++s) {
        cp_wait2();
        __syncthreads();
        const int buf = s % STAGES;
        const float* Abuf = As + buf * A_STAGE_FLOATS;
        const float* Bbuf = Bs + buf * B_STAGE_FLOATS;

        #pragma unroll
        for (int kk = 0; kk < 4; ++kk) {
            unsigned af[2][4];
            #pragma unroll
            for (int mi = 0; mi < 2; ++mi) {
                float2 v0 = lds64(Abuf + a_row + mi * (16 * BK) + kk * 8);
                float2 v1 = lds64(Abuf + a_row8 + mi * (16 * BK) + kk * 8);
                af[mi][0] = __float_as_uint(v0.x);
                af[mi][1] = __float_as_uint(v1.x);
                af[mi][2] = __float_as_uint(v0.y);
                af[mi][3] = __float_as_uint(v1.y);
            }
            const int krow = 8 * kk + 2 * t;
            #pragma unroll
            for (int ni = 0; ni < 8; ++ni) {
                const float* bp = Bbuf + krow * S_B + b_col + ni * 8;
                unsigned b0 = __float_as_uint(lds32(bp));
                unsigned b1 = __float_as_uint(lds32(bp + S_B));
                #pragma unroll
                for (int mi = 0; mi < 2; ++mi) {
                    mma_tf32(acc[mi][ni][0], acc[mi][ni][1], acc[mi][ni][2], acc[mi][ni][3],
                             af[mi][0], af[mi][1], af[mi][2], af[mi][3], b0, b1);
                }
            }
        }
        __syncthreads();
        if (s + 3 < NS) load_stage(s + 3);
        cp_commit();
    }

    // epilogue: bias + float2 stores along x
    #pragma unroll
    for (int mi = 0; mi < 2; ++mi) {
        const int cm0 = m0 + m_warp + mi * 16 + g;
        const int cm1 = cm0 + 8;
        const float bv0 = (bias != nullptr) ? bias[cm0] : 0.f;
        const float bv1 = (bias != nullptr) ? bias[cm1] : 0.f;
        #pragma unroll
        for (int ni = 0; ni < 8; ++ni) {
            const int nbase = n_warp + ni * 8 + 2 * t;
            const int yy = y0 + (nbase >> 4);
            const int xx = x0 + (nbase & 15);
            if (yy < Ho) {
                float* op0 = out + (((long long)(b * Cout + cm0) * Ho + yy) * Wo + xx);
                float* op1 = out + (((long long)(b * Cout + cm1) * Ho + yy) * Wo + xx);
                const float v00 = acc[mi][ni][0] + bv0;
                const float v01 = acc[mi][ni][1] + bv0;
                const float v10 = acc[mi][ni][2] + bv1;
                const float v11 = acc[mi][ni][3] + bv1;
                if (xx + 1 < Wo) {
                    *reinterpret_cast<float2*>(op0) = make_float2(v00, v01);
                    *reinterpret_cast<float2*>(op1) = make_float2(v10, v11);
                } else if (xx < Wo) {
                    op0[0] = v00;
                    op1[0] = v10;
                }
            }
        }
    }
}

void tconv2d_forward(torch::Tensor in, torch::Tensor wp, torch::Tensor out,
                     c10::optional<torch::Tensor> bias,
                     int64_t Bn, int64_t Cin, int64_t Cout, int64_t H, int64_t W,
                     int64_t Ho, int64_t Wo, int64_t K, int64_t mtiles)
{
    dim3 grid((unsigned)((Wo + 15) / 16), (unsigned)((Ho + 15) / 16),
              (unsigned)(Bn * mtiles));
    const int smem_bytes = SMEM_FLOATS * sizeof(float);
    static bool attr_set = false;
    if (!attr_set) {
        cudaFuncSetAttribute(tconv2d_pipe_kernel,
                             cudaFuncAttributeMaxDynamicSharedMemorySize, smem_bytes);
        attr_set = true;
    }
    const float* bias_ptr = nullptr;
    if (bias.has_value() && bias->defined()) bias_ptr = bias->data_ptr<float>();
    tconv2d_pipe_kernel<<<grid, 256, smem_bytes>>>(
        in.data_ptr<float>(), wp.data_ptr<float>(), bias_ptr, out.data_ptr<float>(),
        (int)Bn, (int)Cin, (int)Cout, (int)H, (int)W, (int)Ho, (int)Wo,
        (int)K, (int)mtiles);
}
"""

cpp_src = ("void tconv2d_forward(torch::Tensor in, torch::Tensor wp, torch::Tensor out, "
           "c10::optional<torch::Tensor> bias, "
           "int64_t Bn, int64_t Cin, int64_t Cout, int64_t H, int64_t W, "
           "int64_t Ho, int64_t Wo, int64_t K, int64_t mtiles);")

_mod = None
try:
    _mod = load_inline(
        name="tconv2d_pipe_tf32_v5",
        cpp_sources=cpp_src,
        cuda_sources=cuda_src,
        functions=["tconv2d_forward"],
        verbose=False,
        extra_cuda_cflags=["-O3", "--use_fast_math"],
    )
except Exception:
    _mod = None


class ModelNew(nn.Module):
    def __init__(self, in_channels: int, out_channels: int, kernel_size: int, 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.kernel_size = kernel_size
        self.stride = stride
        self.padding = padding
        self.output_padding = output_padding
        self.groups = groups

        # stride-1, pad-0 transposed conv == regular conv with 180-degree
        # flipped kernel, channel-transposed weights and padding k-1.
        self._can_conv2d = (stride == 1 and padding == 0 and output_padding == 0 and groups == 1)
        self._use_custom = (
            _mod is not None and self._can_conv2d and kernel_size == 3
            and (out_channels % 64 == 0) and (in_channels % 32 == 0)
            and ((9 * in_channels) % 32 == 0)
        )
        self._wprep = None
        self._wflip = None
        self._wflip_cl = None
        self._path = None  # 0: custom mma kernel, 1: conv2d NCHW, 2: native, 3: conv2d channels-last

    def _custom_run(self, x):
        w = self.conv_transpose2d.weight
        if self._wprep is None or self._wprep.device != x.device:
            ci, co = w.shape[0], w.shape[1]
            k = self.kernel_size
            wre = w.permute(2, 3, 0, 1).reshape(k * k * ci, co)
            self._wprep = wre.t().contiguous()
        x = x.contiguous()
        Bn, Cin, H, W = x.shape
        Cout = w.shape[1]
        Ho = (H - 1) * self.stride - 2 * self.padding + self.kernel_size + self.output_padding
        Wo = (W - 1) * self.stride - 2 * self.padding + self.kernel_size + self.output_padding
        K = self.kernel_size * self.kernel_size * Cin
        mtiles = Cout // 64
        out = torch.empty((Bn, Cout, Ho, Wo), device=x.device, dtype=x.dtype)
        _mod.tconv2d_forward(x, self._wprep, out, self.conv_transpose2d.bias,
                             Bn, Cin, Cout, H, W, Ho, Wo, K, mtiles)
        return out

    def _conv2d_run(self, x):
        w = self.conv_transpose2d.weight
        if self._wflip is None or self._wflip.device != x.device:
            self._wflip = w.flip([2, 3]).permute(1, 0, 2, 3).contiguous()
        return F.conv2d(x, self._wflip, self.conv_transpose2d.bias, 1, self.kernel_size - 1)

    def _conv2d_cl_run(self, x):
        w = self.conv_transpose2d.weight
        if self._wflip_cl is None or self._wflip_cl.device != x.device:
            self._wflip_cl = w.flip([2, 3]).permute(1, 0, 2, 3).contiguous(
                memory_format=torch.channels_last)
        y = F.conv2d(x.contiguous(memory_format=torch.channels_last), self._wflip_cl,
                     self.conv_transpose2d.bias, 1, self.kernel_size - 1)
        return y.contiguous()

    def _autotune(self, x):
        if not x.is_cuda:
            self._path = 2
            return
        ref = self.conv_transpose2d(x)
        candidates = [(2, lambda t: self.conv_transpose2d(t))]
        if self._can_conv2d:
            try:
                o = self._conv2d_run(x)
                if o.shape == ref.shape and torch.allclose(o, ref, atol=1e-2, rtol=1e-2):
                    candidates.append((1, lambda t: self._conv2d_run(t)))
            except Exception:
                pass
            try:
                o = self._conv2d_cl_run(x)
                if o.shape == ref.shape and torch.allclose(o, ref, atol=1e-2, rtol=1e-2):
                    candidates.append((3, lambda t: self._conv2d_cl_run(t)))
            except Exception:
                pass
        if self._use_custom and x.dtype == torch.float32:
            try:
                o = self._custom_run(x)
                if o.shape == ref.shape and torch.allclose(o, ref, atol=1e-2, rtol=1e-2):
                    candidates.append((0, lambda t: self._custom_run(t)))
            except Exception:
                pass
        best_path, best_t = 2, float("inf")
        for path, fn in candidates:
            try:
                fn(x)
                torch.cuda.synchronize()
                start = torch.cuda.Event(enable_timing=True)
                end = torch.cuda.Event(enable_timing=True)
                start.record()
                for _ in range(3):
                    fn(x)
                end.record()
                torch.cuda.synchronize()
                t = start.elapsed_time(end)
                if t < best_t:
                    best_t, best_path = t, path
            except Exception:
                continue
        self._path = best_path

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if self._path is None:
            self._autotune(x)
        if self._path == 0 and x.is_cuda and x.dtype == torch.float32:
            return self._custom_run(x)
        elif self._path == 1:
            return self._conv2d_run(x)
        elif self._path == 3:
            return self._conv2d_cl_run(x)
        return self.conv_transpose2d(x)
```