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


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

torch::Tensor transpose_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride,
    int64_t padding,
    int64_t output_padding,
    int64_t groups);
"""

_CUDA_SOURCE = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace {

constexpr int BM = 32;
constexpr int BN = 64;
constexpr int BK = 32;

union SharedStorage {
    struct {
        float ah[BM * BK];
        float al[BM * BK];
        float bh[BK * BN];
        float bl[BK * BN];
    } tiles;
    float result[BM * BN];
};

template<int KH_STATIC, int KW_STATIC, int STRIDE_STATIC>
__global__ void transpose_conv_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int cin,
    int cout,
    int height,
    int width,
    int hout,
    int wout,
    int kh_runtime,
    int kw_runtime,
    int stride_runtime,
    int padding,
    int groups)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    const int kh = KH_STATIC ? KH_STATIC : kh_runtime;
    const int kw = KW_STATIC ? KW_STATIC : kw_runtime;
    const int stride = STRIDE_STATIC ? STRIDE_STATIC : stride_runtime;
    const int ci = cin / groups;
    const int co = cout / groups;
    const int kernel_area = kh * kw;
    const int reduction_size = ci * kernel_area;
    const int spatial_size = hout * wout;

    const int batch = blockIdx.z / groups;
    const int group = blockIdx.z % groups;
    const int channel_base = blockIdx.y * BM;
    const int spatial_base = blockIdx.x * BN;
    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int warp_row = (warp / 2) * 16;
    const int warp_col = (warp % 2) * 32;

    __shared__ __align__(32) SharedStorage shared;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc[2];
    wmma::fill_fragment(acc[0], 0.0f);
    wmma::fill_fragment(acc[1], 0.0f);

    for (int reduction_base = 0;
         reduction_base < reduction_size;
         reduction_base += BK)
    {
        for (int index = tid; index < BM * BK; index += 128) {
            const int oc = channel_base + index / BK;
            const int k = reduction_base + index % BK;
            float value = 0.0f;

            if (oc < co && k < reduction_size) {
                const int ic = k / kernel_area;
                const int tap = k % kernel_area;
                const int64_t offset =
                    (int64_t(group * ci + ic) * co + oc) *
                    kernel_area + tap;
                value = weight[offset];
            }

            const float hi = __float_to_tf32(value);
            shared.tiles.ah[index] = hi;
            shared.tiles.al[index] = __float_to_tf32(value - hi);
        }

        for (int index = tid; index < BK * BN; index += 128) {
            const int k = reduction_base + index / BN;
            const int position = spatial_base + index % BN;
            float value = 0.0f;

            if (k < reduction_size && position < spatial_size) {
                const int ic = k / kernel_area;
                const int tap = k % kernel_area;
                const int ky = tap / kw;
                const int kx = tap % kw;
                const int oy = position / wout;
                const int ox = position % wout;
                const int iy_scaled = oy + padding - ky;
                const int ix_scaled = ox + padding - kx;

                if (iy_scaled >= 0 && ix_scaled >= 0 &&
                    iy_scaled % stride == 0 &&
                    ix_scaled % stride == 0)
                {
                    const int iy = iy_scaled / stride;
                    const int ix = ix_scaled / stride;
                    if (iy < height && ix < width) {
                        const int64_t offset =
                            ((int64_t(batch) * cin + group * ci + ic) *
                             height + iy) * width + ix;
                        value = x[offset];
                    }
                }
            }

            const float hi = __float_to_tf32(value);
            shared.tiles.bh[index] = hi;
            shared.tiles.bl[index] = __float_to_tf32(value - hi);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < BK; kk += 8) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> ah, al;

            wmma::load_matrix_sync(
                ah, shared.tiles.ah + warp_row * BK + kk, BK);
            wmma::load_matrix_sync(
                al, shared.tiles.al + warp_row * BK + kk, BK);

            #pragma unroll
            for (int j = 0; j < 2; ++j) {
                wmma::fragment<
                    wmma::matrix_b, 16, 16, 8,
                    wmma::precision::tf32, wmma::row_major> bh, bl;

                const int offset = kk * BN + warp_col + j * 16;
                wmma::load_matrix_sync(bh, shared.tiles.bh + offset, BN);
                wmma::load_matrix_sync(bl, shared.tiles.bl + offset, BN);

                // Residual products recover substantially more precision
                // than a single TF32 multiplication.
                wmma::mma_sync(acc[j], al, bh, acc[j]);
                wmma::mma_sync(acc[j], ah, bl, acc[j]);
                wmma::mma_sync(acc[j], ah, bh, acc[j]);
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int j = 0; j < 2; ++j) {
        wmma::store_matrix_sync(
            shared.result + warp_row * BN + warp_col + j * 16,
            acc[j], BN, wmma::mem_row_major);
    }
    __syncthreads();

    for (int index = tid; index < BM * BN; index += 128) {
        const int oc = channel_base + index / BN;
        const int position = spatial_base + index % BN;
        if (oc < co && position < spatial_size) {
            const int global_oc = group * co + oc;
            float value = shared.result[index];
            if (bias != nullptr) {
                value += bias[global_oc];
            }
            output[(int64_t(batch) * cout + global_oc) *
                   spatial_size + position] = value;
        }
    }
#endif
}

} // namespace

torch::Tensor transpose_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride,
    int64_t padding,
    int64_t output_padding,
    int64_t groups)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "CUDA tensors required");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat, "float32 required");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Contiguous tensors required");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4, "Expected 4D tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(groups > 0 && stride > 0, "Invalid convolution parameters");
    TORCH_CHECK(x.size(1) == weight.size(0) &&
                x.size(1) % groups == 0, "Invalid channel dimensions");

    c10::cuda::CUDAGuard guard(x.device());

    const int batch = x.size(0);
    const int cin = x.size(1);
    const int height = x.size(2);
    const int width = x.size(3);
    const int co = weight.size(1);
    const int cout = co * groups;
    const int kh = weight.size(2);
    const int kw = weight.size(3);
    const int64_t hout64 =
        (height - 1) * stride - 2 * padding + kh + output_padding;
    const int64_t wout64 =
        (width - 1) * stride - 2 * padding + kw + output_padding;

    TORCH_CHECK(hout64 > 0 && wout64 > 0, "Invalid output dimensions");
    TORCH_CHECK(hout64 * wout64 <= INT_MAX, "Output spatial size too large");
    TORCH_CHECK(int64_t(batch) * groups <= 65535, "Batch/groups too large");

    const int hout = hout64;
    const int wout = wout64;
    auto output = torch::empty({batch, cout, hout, wout}, x.options());
    if (batch == 0) {
        return output;
    }

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == at::kFloat &&
                    bias->is_contiguous() &&
                    bias->numel() == cout, "Invalid bias");
        bias_ptr = bias->data_ptr<float>();
    }

    const dim3 grid(
        (hout * wout + BN - 1) / BN,
        (co + BM - 1) / BM,
        batch * groups);
    const auto stream = at::cuda::getCurrentCUDAStream();

    if (kh == 3 && kw == 3 && stride == 1) {
        transpose_conv_kernel<3, 3, 1><<<grid, 128, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(), cin, cout, height, width,
            hout, wout, kh, kw, stride, padding, groups);
    } else {
        transpose_conv_kernel<0, 0, 0><<<grid, 128, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(), cin, cout, height, width,
            hout, wout, kh, kw, stride, padding, groups);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="transpose_conv_tf32_residual_v1",
            cpp_sources=_CPP_SOURCE,
            cuda_sources=_CUDA_SOURCE,
            functions=["transpose_conv_cuda"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _extension


class _TransposeConvFunction(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias, stride, padding, output_padding, groups):
        ctx.save_for_backward(x, weight)
        ctx.has_bias = bias is not None
        ctx.params = (stride, padding, output_padding, groups)
        return _get_extension().transpose_conv_cuda(
            x, weight, bias, stride, padding, output_padding, groups
        )

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        stride, padding, output_padding, groups = ctx.params
        mask = (
            ctx.needs_input_grad[0],
            ctx.needs_input_grad[1],
            ctx.has_bias and ctx.needs_input_grad[2],
        )
        grad_x, grad_weight, grad_bias = (
            torch.ops.aten.convolution_backward.default(
                grad_output.contiguous(),
                x,
                weight,
                [weight.shape[1] * groups] if ctx.has_bias else None,
                [stride, stride],
                [padding, padding],
                [1, 1],
                True,
                [output_padding, output_padding],
                groups,
                mask,
            )
        )
        return grad_x, grad_weight, grad_bias, None, None, None, 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().__init__()
        self.conv_transpose2d = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
            groups=groups,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose2d
        supported = (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.ndim == 4
            and x.shape[0] * conv.groups <= 65535
            and conv.stride[0] == conv.stride[1]
            and conv.padding[0] == conv.padding[1]
            and conv.output_padding[0] == conv.output_padding[1]
            and conv.dilation == (1, 1)
            and not torch.is_autocast_enabled()
        )
        if not supported or torch.cuda.get_device_capability(x.device)[0] < 8:
            return conv(x)

        return _TransposeConvFunction.apply(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias,
            conv.stride[0],
            conv.padding[0],
            conv.output_padding[0],
            conv.groups,
        )
```