```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, op=0)
// using TF32 tensor cores (mma.m16n8k8) with a 3-stage cp.async pipeline.
// out[co, y, x] = sum_{ky,kx,ci} w[ci, co, ky, kx] * in[ci, y-ky, x-kx]
// GEMM view: C[m=co, n=(b,y,x)] = sum_k A[m,k] * B[k,n], k = (ky,kx,ci)
#define BK 32

__device__ __forceinline__ void cpa16(void* dst, const void* src) {
    unsigned d = (unsigned)__cvta_generic_to_shared(dst);
    asm volatile("cp.async.cg.shared.global [%0], [%1], 16;\n" :: "r"(d), "l"(src));
}
__device__ __forceinline__ void cpa4(void* dst, const void* src, int srcsize) {
    unsigned d = (unsigned)__cvta_generic_to_shared(dst);
    asm volatile("cp.async.ca.shared.global [%0], [%1], 4, %2;\n" :: "r"(d), "l"(src), "r"(srcsize));
}
__device__ __forceinline__ void cp_commit() { asm volatile("cp.async.commit_group;\n"); }
__device__ __forceinline__ void cp_wait1() { asm volatile("cp.async.wait_group 1;\n"); }
__device__ __forceinline__ void cp_wait0() { asm volatile("cp.async.wait_group 0;\n"); }
__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] permuted weights (k-pair order baked)
    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;                  // 3 stages x 64 x 32
    float* Bs = smem + 3 * 64 * BK;    // 3 stages x 128 x 32

    const int b = blockIdx.z / mtiles;
    const int mt = blockIdx.z % mtiles;
    const int m0 = mt * 64;
    const int y0 = blockIdx.y * 8;
    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;
    const int t = lane & 3;
    const int warp_m = warp & 1;
    const int warp_n = warp >> 1;      // 0..3
    const int m_warp = warp_m * 32;
    const int n_warp = warp_n * 32;

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

    // ---- staging setup (loop-invariant parts) ----
    // A: thread handles 2 16B chunks
    const int a_id0 = tid, a_id1 = tid + 256;
    const int a_m0 = a_id0 >> 3, a_kc0 = a_id0 & 7;
    const int a_m1 = a_id1 >> 3, a_kc1 = a_id1 & 7;
    const float* a_src0 = wp + (long long)(m0 + a_m0) * K + a_kc0 * 4;
    const float* a_src1 = wp + (long long)(m0 + a_m1) * K + a_kc1 * 4;
    const int a_dst0 = (a_m0 * BK + ((a_kc0 ^ ((a_m0 & 3) << 1)) << 2));
    const int a_dst1 = (a_m1 * BK + ((a_kc1 ^ ((a_m1 & 3) << 1)) << 2));

    // B: thread handles fixed n, 16 k values
    const int n_idx = tid & 127;
    const int k_tid = tid >> 7;               // 0..1
    const int ty = n_idx >> 4, tx = n_idx & 15;
    const int b_row_off = n_idx * BK;
    const int b_swz = (n_idx & 3) << 3;

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

    auto load_stage = [&](int s) {
        const int buf = s % 3;
        const int k0 = s << 5;
        // A chunks
        cpa16(As + buf * (64 * BK) + a_dst0, a_src0 + k0);
        cpa16(As + buf * (64 * BK) + a_dst1, a_src1 + k0);
        // B gather
        const int tap = k0 / Cin;             // uniform within stage (Cin % 32 == 0)
        const int cibase = k0 - tap * Cin;
        const int ky = tap / 3;
        const int kx = tap - ky * 3;
        const int gy = y0 + ty - ky;
        const int gx = x0 + tx - kx;
        const bool valid = (gy >= 0) & (gy < H) & (gx >= 0) & (gx < W);
        const float* src = inb + (long long)cibase * inHW + (long long)gy * W + gx;
        const int ssz = valid ? 4 : 0;
        float* bd = Bs + buf * (128 * BK) + b_row_off;
        #pragma unroll
        for (int j = 0; j < 16; ++j) {
            const int k = k_tid + 2 * j;      // 0..31
            const int col = ((k & 24) | ((k & 3) << 1) | ((k >> 2) & 1)) ^ b_swz;
            cpa4(bd + col, src + (long long)k * inHW, ssz);
        }
    };

    // prologue: stages 0 and 1 in flight
    load_stage(0); cp_commit();
    load_stage(1); cp_commit();

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

    // precompute lds.64 column offsets per k8 substep
    int acol[4], bcol[4];
    #pragma unroll
    for (int kk = 0; kk < 4; ++kk) {
        acol[kk] = (kk * 8 + 2 * t) ^ ((g & 3) << 3);
        bcol[kk] = acol[kk];
    }
    const int a_row = (m_warp + g) * BK;
    const int a_row8 = (m_warp + g + 8) * BK;
    const int b_row = (n_warp + g) * BK;

    for (int s = 0; s < total_stages; ++s) {
        cp_wait1();
        __syncthreads();
        const int buf = s % 3;
        const float* Asb = As + buf * (64 * BK);
        const float* Bsb = Bs + buf * (128 * BK);

        #pragma unroll
        for (int kk = 0; kk < 4; ++kk) {
            // load A frags (2 x m16k8)
            unsigned af[2][4];
            #pragma unroll
            for (int mi = 0; mi < 2; ++mi) {
                float2 v0 = lds64(Asb + a_row + mi * (16 * BK) + acol[kk]);
                float2 v1 = lds64(Asb + a_row8 + mi * (16 * BK) + acol[kk]);
                af[mi][0] = __float_as_uint(v0.x);
                af[mi][2] = __float_as_uint(v0.y);
                af[mi][1] = __float_as_uint(v1.x);
                af[mi][3] = __float_as_uint(v1.y);
            }
            // load B frags (4 x n8k8) and mma
            #pragma unroll
            for (int ni = 0; ni < 4; ++ni) {
                float2 bv = lds64(Bsb + b_row + ni * (8 * BK) + bcol[kk]);
                unsigned b0 = __float_as_uint(bv.x);
                unsigned b1 = __float_as_uint(bv.y);
                #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);
                }
            }
        }

        if (s + 2 < total_stages) load_stage(s + 2);
        cp_commit();
    }
    cp_wait0();
    __syncthreads();

    // ---- epilogue ----
    float bias_g[2], bias_g8[2];
    if (bias != nullptr) {
        #pragma unroll
        for (int mi = 0; mi < 2; ++mi) {
            bias_g[mi] = bias[m0 + m_warp + mi * 16 + g];
            bias_g8[mi] = bias[m0 + m_warp + mi * 16 + g + 8];
        }
    }
    const int px_y = y0;  // base
    #pragma unroll
    for (int mi = 0; mi < 2; ++mi) {
        const int cm0 = m0 + m_warp + mi * 16 + g;
        const int cm1 = cm0 + 8;
        #pragma unroll
        for (int ni = 0; ni < 4; ++ni) {
            const int nbase = n_warp + ni * 8 + t * 2;
            #pragma unroll
            for (int half = 0; half < 2; ++half) {   // r = 2*half .. 2*half+1 pair rows g / g+8
                const int cm = half ? cm1 : cm0;
                const float bv = (bias != nullptr) ? (half ? bias_g8[mi] : bias_g[mi]) : 0.f;
                const int n0p = nbase;
                const int yy0 = px_y + (n0p >> 4);
                const int xx0 = x0 + (n0p & 15);
                const float v0 = acc[mi][ni][half * 2 + 0] + bv;
                const float v1 = acc[mi][ni][half * 2 + 1] + bv;
                float* op = out + (((long long)(b * Cout + cm) * Ho + yy0) * Wo + xx0);
                if (yy0 < Ho) {
                    if (xx0 + 1 < Wo) {
                        *reinterpret_cast<float2*>(op) = make_float2(v0, v1);
                    } else if (xx0 < Wo) {
                        op[0] = v0;
                    }
                }
            }
        }
    }
}

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 + 7) / 8),
              (unsigned)(Bn * mtiles));
    const int smem_bytes = (3 * 64 * BK + 3 * 128 * BK) * 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_v3",
        cpp_sources=cpp_src,
        cuda_sources=cuda_src,
        functions=["tconv2d_forward"],
        verbose=False,
        extra_cuda_cflags=["-O3", "--use_fast_math"],
    )
except Exception:
    _mod = None

_KPAIR_PERM = [0, 4, 1, 5, 2, 6, 3, 7]


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)
            wt = wre.t().contiguous().view(co, (k * k * ci) // 8, 8)
            perm = torch.tensor(_KPAIR_PERM, device=w.device)
            self._wprep = wt[:, :, perm].reshape(co, k * k * ci).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)
```