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


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

std::vector<torch::Tensor> pack_transposed_conv(
    torch::Tensor x, torch::Tensor weight, int64_t groups);

torch::Tensor run_convolution(
    torch::Tensor x,
    torch::Tensor weight,
    std::vector<int64_t> padding,
    std::vector<int64_t> dilation,
    int64_t groups,
    bool deterministic,
    bool allow_tf32);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/ops/cudnn_convolution.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <vector>

// A larger tile reduces block scheduling overhead on large feature maps.
__global__ void pack_input_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int64_t spatial)
{
    __shared__ float tile[64][65];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int c0 = int(blockIdx.y) * 64;
    const int64_t s0 = int64_t(blockIdx.x) * 64;
    const int64_t base = int64_t(blockIdx.z) * channels * spatial;

    #pragma unroll
    for (int j = 0; j < 64; j += 8) {
        const int c = c0 + ty + j;
        if (c < channels) {
            if (s0 + tx < spatial)
                tile[ty + j][tx] =
                    input[base + int64_t(c) * spatial + s0 + tx];
            if (s0 + tx + 32 < spatial)
                tile[ty + j][tx + 32] =
                    input[base + int64_t(c) * spatial + s0 + tx + 32];
        }
    }
    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 64; j += 8) {
        const int64_t s = s0 + ty + j;
        if (s < spatial) {
            if (c0 + tx < channels)
                output[base + s * channels + c0 + tx] =
                    tile[tx][ty + j];
            if (c0 + tx + 32 < channels)
                output[base + s * channels + c0 + tx + 32] =
                    tile[tx + 32][ty + j];
        }
    }
}

__global__ void pack_weight_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int ci_per_group,
    int co_per_group,
    int kh,
    int kw,
    int64_t count)
{
    const int64_t index =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
    if (index >= count)
        return;

    int64_t q = index;
    const int ic = q % ci_per_group;
    q /= ci_per_group;
    const int s = q % kw;
    q /= kw;
    const int r = q % kh;
    const int oc = q / kh;

    const int group = oc / co_per_group;
    const int local_oc = oc % co_per_group;
    const int global_ic = group * ci_per_group + ic;

    const int64_t source =
        ((int64_t(global_ic) * co_per_group + local_oc) * kh
         + kh - 1 - r) * kw + kw - 1 - s;

    output[index] = input[source];
}

std::vector<torch::Tensor> pack_transposed_conv(
    torch::Tensor x, torch::Tensor weight, int64_t groups)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.device() == weight.device(),
                "Input and weight must share a device");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(weight.is_contiguous(),
                "Weight must be contiguous");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid groups");
    TORCH_CHECK(x.size(1) == weight.size(0),
                "Input channel count does not match weight");

    const c10::cuda::CUDAGuard guard(x.device());
    const auto stream = at::cuda::getCurrentCUDAStream();

    const int64_t n = x.size(0);
    const int64_t ci = x.size(1);
    const int64_t h = x.size(2);
    const int64_t w = x.size(3);
    const int64_t kh = weight.size(2);
    const int64_t kw = weight.size(3);
    const int64_t cig = ci / groups;
    const int64_t cog = weight.size(1);
    const int64_t co = cog * groups;

    torch::Tensor packed_x;
    if (x.is_contiguous(at::MemoryFormat::ChannelsLast)) {
        packed_x = x;
    } else {
        TORCH_CHECK(x.is_contiguous(), "Input must be contiguous");
        packed_x = torch::empty({n, h, w, ci}, x.options())
                       .permute({0, 3, 1, 2});

        if (x.numel() != 0) {
            const dim3 block(32, 8);
            const dim3 grid((h * w + 63) / 64, (ci + 63) / 64, n);
            pack_input_kernel<<<grid, block, 0, stream>>>(
                x.data_ptr<float>(),
                packed_x.data_ptr<float>(),
                static_cast<int>(ci),
                h * w);
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        }
    }

    auto packed_weight =
        torch::empty({co, kh, kw, cig}, weight.options())
            .permute({0, 3, 1, 2});

    const int64_t count = weight.numel();
    if (count != 0) {
        pack_weight_kernel<<<(count + 255) / 256, 256, 0, stream>>>(
            weight.data_ptr<float>(),
            packed_weight.data_ptr<float>(),
            static_cast<int>(cig),
            static_cast<int>(cog),
            static_cast<int>(kh),
            static_cast<int>(kw),
            count);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return {packed_x, packed_weight};
}

torch::Tensor run_convolution(
    torch::Tensor x,
    torch::Tensor weight,
    std::vector<int64_t> padding,
    std::vector<int64_t> dilation,
    int64_t groups,
    bool deterministic,
    bool allow_tf32)
{
    const c10::cuda::CUDAGuard guard(x.device());
    const std::vector<int64_t> stride{1, 1};

    // cuDNN caches the selected algorithm for subsequent calls.
    return at::cudnn_convolution(
        x, weight, padding, stride, dilation, groups,
        true, deterministic, allow_tf32);
}
"""

_extension = load_inline(
    name="asymmetric_transpose_conv_tuned_v2",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pack_transposed_conv", "run_convolution"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _TransposedConvFunction(torch.autograd.Function):
    @staticmethod
    def forward(
        ctx, x, weight, bias, padding, dilation, groups,
        deterministic, allow_tf32
    ):
        packed_x, packed_weight = _extension.pack_transposed_conv(
            x, weight.contiguous(), groups
        )
        output = _extension.run_convolution(
            packed_x, packed_weight, list(padding), list(dilation),
            groups, deterministic, allow_tf32
        )
        if bias is not None:
            output.add_(bias.reshape(1, -1, 1, 1))

        ctx.padding = padding
        ctx.dilation = dilation
        ctx.groups = groups
        ctx.input_shape = tuple(x.shape)
        ctx.weight_shape = tuple(packed_weight.shape)
        ctx.original_weight_shape = tuple(weight.shape)
        ctx.has_bias = bias is not None

        saved_x = packed_x if ctx.needs_input_grad[1] else None
        saved_weight = packed_weight if ctx.needs_input_grad[0] else None
        ctx.save_for_backward(saved_x, saved_weight)
        return output

    @staticmethod
    def backward(ctx, grad_output):
        packed_x, packed_weight = ctx.saved_tensors
        grad_x = grad_weight = grad_bias = None

        if ctx.needs_input_grad[0]:
            grad_x = torch.nn.grad.conv2d_input(
                ctx.input_shape,
                packed_weight,
                grad_output,
                stride=1,
                padding=ctx.padding,
                dilation=ctx.dilation,
                groups=ctx.groups,
            )

        if ctx.needs_input_grad[1]:
            packed_grad = torch.nn.grad.conv2d_weight(
                packed_x,
                ctx.weight_shape,
                grad_output,
                stride=1,
                padding=ctx.padding,
                dilation=ctx.dilation,
                groups=ctx.groups,
            )
            ci, cog, kh, kw = ctx.original_weight_shape
            grad_weight = (
                packed_grad.reshape(
                    ctx.groups, cog, ci // ctx.groups, kh, kw
                )
                .transpose(1, 2)
                .flip((-2, -1))
                .reshape(ctx.original_weight_shape)
            )

        if ctx.has_bias and ctx.needs_input_grad[2]:
            grad_bias = grad_output.sum(dim=(0, 2, 3))

        return (
            grad_x, grad_weight, grad_bias,
            None, None, None, None, None
        )


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

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose2d
        padding = tuple(
            d * (k - 1) - p
            for d, k, p in zip(
                conv.dilation, conv.kernel_size, conv.padding
            )
        )

        eligible = (
            x.is_cuda
            and x.ndim == 4
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.device == conv.weight.device
            and conv.stride == (1, 1)
            and conv.output_padding == (0, 0)
            and min(padding) >= 0
            and 0 < x.shape[0] <= 65535
            and x.shape[2] > 0
            and x.shape[3] > 0
            and (x.shape[1] + 63) // 64 <= 65535
            and torch.backends.cudnn.is_available()
            and torch.backends.cudnn.enabled
            and not torch.is_autocast_enabled()
        )
        if not eligible:
            return conv(x)

        if not (
            x.is_contiguous()
            or x.is_contiguous(memory_format=torch.channels_last)
        ):
            x = x.contiguous()

        return _TransposedConvFunction.apply(
            x,
            conv.weight,
            conv.bias,
            padding,
            conv.dilation,
            conv.groups,
            torch.backends.cudnn.deterministic
            or torch.are_deterministic_algorithms_enabled(),
            torch.backends.cudnn.allow_tf32,
        )
```