```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> transpose_as_forward_cuda(
    torch::Tensor input,
    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>

__global__ void pack_channels_last(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int64_t spatial,
    int64_t spatial_tiles,
    int channel_tiles,
    int64_t total_tiles)
{
    __shared__ float tile[64][65];
    const int tx = threadIdx.x;
    const int ty = threadIdx.y;

    for (int64_t tile_id = blockIdx.x;
         tile_id < total_tiles;
         tile_id += gridDim.x) {
        int64_t q = tile_id;
        const int64_t s0 = (q % spatial_tiles) * 64;
        q /= spatial_tiles;
        const int c0 = (q % channel_tiles) * 64;
        const int64_t batch = q / channel_tiles;
        const int64_t base = batch * channels * spatial;

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

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

// Convert [Cin, Cout/groups, Kh, Kw] into channels-last
// [Cout, Cin/groups, Kh, Kw], reversing both spatial axes.
__global__ void reverse_transpose_filter(
    const float* __restrict__ input,
    float* __restrict__ output,
    int input_per_group,
    int output_per_group,
    int kernel_spatial,
    int64_t count,
    int64_t stride0,
    int64_t stride1,
    int64_t stride2,
    int64_t stride3,
    int kernel_width)
{
    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 ci = q % input_per_group;
        q /= input_per_group;
        const int k = q % kernel_spatial;
        const int co = q / kernel_spatial;
        const int group = co / output_per_group;
        const int local_co = co % output_per_group;
        const int source_ci = group * input_per_group + ci;
        const int reversed_k = kernel_spatial - 1 - k;
        const int r = reversed_k / kernel_width;
        const int s = reversed_k % kernel_width;

        output[index] = input[
            int64_t(source_ci) * stride0 +
            int64_t(local_co) * stride1 +
            int64_t(r) * stride2 +
            int64_t(s) * stride3];
    }
}

static torch::Tensor channels_last_input(
    const torch::Tensor& input,
    cudaStream_t stream)
{
    if (input.is_contiguous(at::MemoryFormat::ChannelsLast))
        return input;

    TORCH_CHECK(input.is_contiguous(), "Expected contiguous input");

    const int64_t n = input.size(0);
    const int64_t c = input.size(1);
    const int64_t h = input.size(2);
    const int64_t w = input.size(3);
    const int64_t spatial = h * w;

    auto output = torch::empty({n, h, w, c}, input.options())
                      .permute({0, 3, 1, 2});

    const int64_t spatial_tiles = (spatial + 63) / 64;
    const int channel_tiles = (c + 63) / 64;
    const int64_t tiles = n * channel_tiles * spatial_tiles;
    const unsigned blocks =
        static_cast<unsigned>(tiles < 65535 ? tiles : 65535);

    pack_channels_last<<<blocks, dim3(32, 8), 0, stream>>>(
        input.data_ptr<float>(),
        output.data_ptr<float>(),
        static_cast<int>(c),
        spatial,
        spatial_tiles,
        channel_tiles,
        tiles);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

std::vector<torch::Tensor> transpose_as_forward_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    std::vector<int64_t> padding,
    std::vector<int64_t> dilation,
    int64_t groups,
    bool deterministic,
    bool allow_tf32)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must have the same device");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(input.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid groups");

    const c10::cuda::CUDAGuard guard(input.device());
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    auto packed_input = channels_last_input(input, stream);

    const int64_t ci = weight.size(0) / groups;
    const int64_t co = weight.size(1) * groups;
    const int64_t kh = weight.size(2);
    const int64_t kw = weight.size(3);

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

    const int64_t count = weight.numel();
    const int64_t requested_blocks = (count + 255) / 256;
    const int blocks = static_cast<int>(
        requested_blocks < 65535 ? requested_blocks : 65535);

    reverse_transpose_filter<<<blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        forward_weight.data_ptr<float>(),
        static_cast<int>(ci),
        static_cast<int>(weight.size(1)),
        static_cast<int>(kh * kw),
        count,
        weight.stride(0),
        weight.stride(1),
        weight.stride(2),
        weight.stride(3),
        static_cast<int>(kw));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    std::vector<int64_t> forward_padding = {
        dilation[0] * (kh - 1) - padding[0],
        dilation[1] * (kw - 1) - padding[1]
    };
    TORCH_CHECK(forward_padding[0] >= 0 && forward_padding[1] >= 0,
                "Unsupported padding");

    const std::vector<int64_t> stride = {1, 1};

    auto output = at::cudnn_convolution(
        packed_input,
        forward_weight,
        forward_padding,
        stride,
        dilation,
        groups,
        true,
        deterministic,
        allow_tf32);

    return {output, packed_input};
}
"""

_extension = load_inline(
    name="asymmetric_transpose_conv_forward_nhwc_v5",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["transpose_as_forward_cuda"],
    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
    ):
        output, packed_x = _extension.transpose_as_forward_cuda(
            x,
            weight,
            list(padding),
            list(dilation),
            groups,
            deterministic,
            allow_tf32,
        )

        if bias is not None:
            output.add_(bias.reshape(1, -1, 1, 1))

        ctx.save_for_backward(packed_x, weight)
        ctx.padding = padding
        ctx.dilation = dilation
        ctx.groups = groups
        ctx.bias_sizes = list(bias.shape) if bias is not None else None
        return output

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        mask = (
            ctx.needs_input_grad[0],
            ctx.needs_input_grad[1],
            ctx.bias_sizes is not None and ctx.needs_input_grad[2],
        )

        grad_x, grad_weight, grad_bias = (
            torch.ops.aten.convolution_backward.default(
                grad_output,
                x,
                weight,
                ctx.bias_sizes,
                [1, 1],
                list(ctx.padding),
                list(ctx.dilation),
                True,
                [0, 0],
                ctx.groups,
                mask,
            )
        )
        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
        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 x.numel() > 0
            and x.shape[1] == conv.in_channels
            and x.shape[1] < 2**31 - 64
            and conv.out_channels < 2**31
            and conv.kernel_size[0] * conv.kernel_size[1] < 2**31
            and conv.stride == (1, 1)
            and conv.output_padding == (0, 0)
            and conv.padding[0]
            <= conv.dilation[0] * (conv.kernel_size[0] - 1)
            and conv.padding[1]
            <= conv.dilation[1] * (conv.kernel_size[1] - 1)
            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,
            conv.padding,
            conv.dilation,
            conv.groups,
            torch.backends.cudnn.deterministic
            or torch.are_deterministic_algorithms_enabled(),
            torch.backends.cudnn.allow_tf32,
        )
```