```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>

#define TH 8
#define TW 8
#define TILE 64
#define KSTEP 16
#define LDS 20

// Implicit-GEMM transposed 3x3 conv (stride=1, pad=0) using TF32 tensor cores
// via mma.sync.m16n8k8. C[cout, pixel] = sum_k W[k, cout] * B[k, pixel],
// where k enumerates (ky, kx, cin) and B is the im2col gather of the input.
__global__ void __launch_bounds__(256) tconv2d_mma_tf32_kernel(
    const float* __restrict__ in,
    const float* __restrict__ w,   // [K][Cout], k = (ky*3+kx)*Cin + cin
    float* __restrict__ out,
    int Bn, int Cin, int Cout, int H, int W, int Ho, int Wo, int K)
{
    __shared__ float Ws[TILE][LDS];  // [m][kk]
    __shared__ float Bs[TILE][LDS];  // [n][kk]

    const int mtiles = Cout / TILE;
    const int b = blockIdx.z / mtiles;
    const int m0 = (blockIdx.z % mtiles) * TILE;
    const int y0 = blockIdx.y * TH;
    const int x0 = blockIdx.x * TW;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int lane = tid & 31;
    const int warp_m = warp >> 2;   // 0..1
    const int warp_n = warp & 3;    // 0..3
    const int gid = lane >> 2;
    const int tig = lane & 3;

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

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

    for (int k0 = 0; k0 < K; k0 += KSTEP) {
        // stage weights, coalesced global reads (consecutive tid -> consecutive m)
        #pragma unroll
        for (int i = 0; i < 4; ++i) {
            int idx = tid + i * 256;
            int kk = idx >> 6;
            int m = idx & 63;
            int k = k0 + kk;
            float v = 0.f;
            if (k < K) v = w[(long long)k * Cout + m0 + m];
            Ws[m][kk] = v;
        }
        // stage im2col input, coalesced global reads (consecutive tid -> consecutive pixel/x)
        #pragma unroll
        for (int i = 0; i < 4; ++i) {
            int idx = tid + i * 256;
            int kk = idx >> 6;
            int n = idx & 63;
            int k = k0 + kk;
            float v = 0.f;
            if (k < K) {
                int tap = k / Cin;
                int cin = k - tap * Cin;
                int ky = tap / 3;
                int kx = tap - ky * 3;
                int ty = n >> 3;
                int tx = n & 7;
                int gy = y0 + ty - ky;   // out(y,x) = sum w[ky,kx]*in(y-ky, x-kx)
                int gx = x0 + tx - kx;
                if (gy >= 0 && gy < H && gx >= 0 && gx < W)
                    v = inb[(long long)cin * inHW + (long long)gy * W + gx];
            }
            Bs[n][kk] = v;
        }
        __syncthreads();

        #pragma unroll
        for (int ks = 0; ks < KSTEP; ks += 8) {
            unsigned bf[2][2];
            #pragma unroll
            for (int ni = 0; ni < 2; ++ni) {
                int nb = warp_n * 16 + ni * 8;
                bf[ni][0] = __float_as_uint(Bs[nb + gid][ks + tig]);
                bf[ni][1] = __float_as_uint(Bs[nb + gid][ks + tig + 4]);
            }
            #pragma unroll
            for (int mi = 0; mi < 2; ++mi) {
                int mb = warp_m * 32 + mi * 16;
                unsigned a0 = __float_as_uint(Ws[mb + gid][ks + tig]);
                unsigned a1 = __float_as_uint(Ws[mb + gid + 8][ks + tig]);
                unsigned a2 = __float_as_uint(Ws[mb + gid][ks + tig + 4]);
                unsigned a3 = __float_as_uint(Ws[mb + gid + 8][ks + tig + 4]);
                #pragma unroll
                for (int ni = 0; ni < 2; ++ni) {
                    float* c = acc[mi][ni];
                    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"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
                        : "r"(a0), "r"(a1), "r"(a2), "r"(a3),
                          "r"(bf[ni][0]), "r"(bf[ni][1]));
                }
            }
        }
        __syncthreads();
    }

    #pragma unroll
    for (int mi = 0; mi < 2; ++mi) {
        #pragma unroll
        for (int ni = 0; ni < 2; ++ni) {
            int cb = m0 + warp_m * 32 + mi * 16;
            int pb = warp_n * 16 + ni * 8;
            #pragma unroll
            for (int r = 0; r < 4; ++r) {
                int cm = cb + gid + ((r >> 1) << 3);
                int pn = pb + tig * 2 + (r & 1);
                int y = y0 + (pn >> 3);
                int x = x0 + (pn & 7);
                if (y < Ho && x < Wo)
                    out[((long long)(b * Cout + cm) * Ho + y) * Wo + x] = acc[mi][ni][r];
            }
        }
    }
}

void tconv2d_forward(torch::Tensor in, torch::Tensor w, torch::Tensor out,
                     int64_t Bn, int64_t Cin, int64_t Cout, int64_t H, int64_t W,
                     int64_t Ho, int64_t Wo, int64_t K)
{
    dim3 grid((unsigned)((Wo + TW - 1) / TW), (unsigned)((Ho + TH - 1) / TH),
              (unsigned)(Bn * (Cout / TILE)));
    tconv2d_mma_tf32_kernel<<<grid, 256>>>(
        in.data_ptr<float>(), w.data_ptr<float>(), out.data_ptr<float>(),
        (int)Bn, (int)Cin, (int)Cout, (int)H, (int)W, (int)Ho, (int)Wo, (int)K);
}
"""

cpp_src = ("void tconv2d_forward(torch::Tensor in, torch::Tensor w, torch::Tensor out, "
           "int64_t Bn, int64_t Cin, int64_t Cout, int64_t H, int64_t W, "
           "int64_t Ho, int64_t Wo, int64_t K);")

_mod = None
try:
    _mod = load_inline(
        name="tconv2d_mma_tf32_v2",
        cpp_sources=cpp_src,
        cuda_sources=cuda_src,
        functions=["tconv2d_forward"],
        verbose=False,
        extra_cuda_cflags=["-O3"],
    )
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 not bias
        )
        self._wre = 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._wre is None or self._wre.device != x.device:
            k = self.kernel_size
            ci = w.shape[0]
            co = w.shape[1]
            self._wre = w.permute(2, 3, 0, 1).reshape(k * k * ci, co).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
        out = torch.empty((Bn, Cout, Ho, Wo), device=x.device, dtype=x.dtype)
        _mod.tconv2d_forward(x, self._wre, out, Bn, Cin, Cout, H, W, Ho, Wo, K)
        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)
```