```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 pad_h, int64_t pad_w, int64_t groups);

torch::Tensor run_convolution(
    torch::Tensor x, torch::Tensor weight,
    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>

// Fuse spatial padding and NCHW -> NHWC conversion. Explicit padding
// allows the subsequent convolution to operate entirely on valid windows.
__global__ void pad_pack_nchw(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels, int height, int width,
    int padded_height, int padded_width,
    int pad_h, int pad_w)
{
    __shared__ float tile[64][65];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int c0 = blockIdx.y * 64;
    const int spatial = padded_height * padded_width;
    const int s0 = blockIdx.x * 64;
    const int64_t input_base =
        int64_t(blockIdx.z) * channels * height * width;
    const int64_t output_base =
        int64_t(blockIdx.z) * channels * spatial;

    const int s1 = s0 + tx;
    const int s2 = s1 + 32;
    const int y1 = s1 / padded_width - pad_h;
    const int x1 = s1 % padded_width - pad_w;
    const int y2 = s2 / padded_width - pad_h;
    const int x2 = s2 % padded_width - pad_w;

    const bool valid1 =
        s1 < spatial &&
        static_cast<unsigned>(y1) < static_cast<unsigned>(height) &&
        static_cast<unsigned>(x1) < static_cast<unsigned>(width);
    const bool valid2 =
        s2 < spatial &&
        static_cast<unsigned>(y2) < static_cast<unsigned>(height) &&
        static_cast<unsigned>(x2) < static_cast<unsigned>(width);

    #pragma unroll
    for (int j = 0; j < 64; j += 8) {
        const int c = c0 + ty + j;
        float a = 0.0f;
        float b = 0.0f;
        if (c < channels) {
            const int64_t base =
                input_base + int64_t(c) * height * width;
            if (valid1)
                a = input[base + int64_t(y1) * width + x1];
            if (valid2)
                b = input[base + int64_t(y2) * width + x2];
        }
        tile[ty + j][tx] = a;
        tile[ty + j][tx + 32] = b;
    }
    __syncthreads();

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

__global__ void pad_pack_nhwc(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels, int height, int width,
    int padded_height, int padded_width,
    int pad_h, int pad_w, int64_t count)
{
    for (int64_t index =
             int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        int64_t q = index;
        const int c = q % channels;
        q /= channels;
        const int x = q % padded_width - pad_w;
        q /= padded_width;
        const int y = q % padded_height - pad_h;
        const int64_t n = q / padded_height;

        float value = 0.0f;
        if (static_cast<unsigned>(y) < static_cast<unsigned>(height) &&
            static_cast<unsigned>(x) < static_cast<unsigned>(width)) {
            value = input[
                ((n * height + y) * width + x) * channels + c];
        }
        output[index] = value;
    }
}

__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 pad_h, int64_t pad_w, int64_t groups)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    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), "Channel mismatch");
    TORCH_CHECK(pad_h >= 0 && pad_w >= 0, "Invalid padding");

    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 hp = h + 2 * pad_h;
    const int64_t wp = w + 2 * pad_w;
    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;

    const bool channels_last =
        x.is_contiguous(at::MemoryFormat::ChannelsLast);

    torch::Tensor packed_x;
    if (channels_last && pad_h == 0 && pad_w == 0) {
        packed_x = x;
    } else {
        packed_x = torch::empty({n, hp, wp, ci}, x.options())
                       .permute({0, 3, 1, 2});

        if (channels_last) {
            const int64_t count = packed_x.numel();
            const int blocks = static_cast<int>(
                std::min<int64_t>((count + 255) / 256, 65535));
            pad_pack_nhwc<<<blocks, 256, 0, stream>>>(
                x.data_ptr<float>(), packed_x.data_ptr<float>(),
                ci, h, w, hp, wp, pad_h, pad_w, count);
        } else {
            TORCH_CHECK(x.is_contiguous(), "Input must be contiguous");
            const dim3 block(32, 8);
            const dim3 grid((hp * wp + 63) / 64, (ci + 63) / 64, n);
            pad_pack_nchw<<<grid, block, 0, stream>>>(
                x.data_ptr<float>(), packed_x.data_ptr<float>(),
                ci, h, w, hp, wp, pad_h, pad_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();
    pack_weight_kernel<<<(count + 255) / 256, 256, 0, stream>>>(
        weight.data_ptr<float>(), packed_weight.data_ptr<float>(),
        cig, cog, kh, 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> dilation, int64_t groups,
    bool deterministic, bool allow_tf32)
{
    const c10::cuda::CUDAGuard guard(x.device());
    const std::vector<int64_t> padding{0, 0};
    const std::vector<int64_t> stride{1, 1};
    return at::cudnn_convolution(
        x, weight, padding, stride, dilation, groups,
        true, deterministic, allow_tf32);
}
"""

_extension = load_inline(
    name="asymmetric_transpose_conv_fused_padding_v3",
    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(), padding[0], padding[1], groups
        )
        output = _extension.run_convolution(
            packed_x, packed_weight, 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.padded_shape = tuple(packed_x.shape)
        ctx.weight_shape = tuple(packed_weight.shape)
        ctx.original_weight_shape = tuple(weight.shape)
        ctx.has_bias = bias is not None
        ctx.save_for_backward(
            packed_x if ctx.needs_input_grad[1] else None,
            packed_weight if ctx.needs_input_grad[0] else None,
        )
        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]:
            padded_grad = torch.nn.grad.conv2d_input(
                ctx.padded_shape,
                packed_weight,
                grad_output,
                stride=1,
                padding=0,
                dilation=ctx.dilation,
                groups=ctx.groups,
            )
            ph, pw = ctx.padding
            h, w = ctx.input_shape[-2:]
            grad_x = padded_grad[:, :, ph:ph + h, pw:pw + w]

        if ctx.needs_input_grad[1]:
            packed_grad = torch.nn.grad.conv2d_weight(
                packed_x,
                ctx.weight_shape,
                grad_output,
                stride=1,
                padding=0,
                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 0 < x.shape[1] <= 65535 * 64
            and x.shape[2] > 0
            and x.shape[3] > 0
            and (
                (x.shape[2] + 2 * padding[0])
                * (x.shape[3] + 2 * padding[1])
                < 2**31 - 64
            )
            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,
        )
```