```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_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    std::vector<int64_t> padding,
    std::vector<int64_t> output_padding,
    std::vector<int64_t> stride,
    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_transpose.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <vector>

__device__ __forceinline__ void store_shared4(
    float* destination, const float4 value)
{
    destination[0] = value.x;
    destination[1] = value.y;
    destination[2] = value.z;
    destination[3] = value.w;
}

// Each warp covers four rows and eight float4 columns.
// The padded shared-memory pitch avoids bank conflicts during
// both phases of the transpose.
template <bool SingleChannelTile>
__global__ __launch_bounds__(256) void pack_channels_last_vectorized(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t channels,
    int64_t spatial,
    int64_t spatial_tiles,
    int64_t channel_tiles,
    int64_t total_tiles)
{
    __shared__ float tile[64][65];

    const int x = threadIdx.x * 4;
    const int y = threadIdx.y;

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

        #pragma unroll
        for (int row = 0; row < 64; row += 32) {
            #pragma unroll
            for (int col = 0; col < 64; col += 32) {
                const float4 value = *reinterpret_cast<const float4*>(
                    input + base + (c0 + y + row) * spatial
                    + s0 + x + col);
                store_shared4(&tile[y + row][x + col], value);
            }
        }
        __syncthreads();

        #pragma unroll
        for (int row = 0; row < 64; row += 32) {
            #pragma unroll
            for (int col = 0; col < 64; col += 32) {
                const float4 value = make_float4(
                    tile[x + col    ][y + row],
                    tile[x + col + 1][y + row],
                    tile[x + col + 2][y + row],
                    tile[x + col + 3][y + row]);
                *reinterpret_cast<float4*>(
                    output + base + (s0 + y + row) * channels
                    + c0 + x + col) = value;
            }
        }
        __syncthreads();
    }
}

__global__ void pack_channels_last_generic(
    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();
    }
}

__global__ void pack_filter(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int spatial,
    int64_t count)
{
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        const int c = index % channels;
        const int64_t q = index / channels;
        const int s = q % spatial;
        const int64_t n = q / spatial;
        output[index] = input[(n * channels + c) * spatial + s];
    }
}

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 int64_t channel_tiles = (c + 63) / 64;
    const int64_t tiles = n * channel_tiles * spatial_tiles;
    const unsigned blocks =
        static_cast<unsigned>(tiles < 65535 ? tiles : 65535);

    const bool vectorizable =
        c % 64 == 0 &&
        spatial % 64 == 0 &&
        reinterpret_cast<uintptr_t>(input.data_ptr<float>()) % 16 == 0;

    if (vectorizable) {
        if (c == 64) {
            pack_channels_last_vectorized<true>
                <<<blocks, dim3(8, 32), 0, stream>>>(
                    input.data_ptr<float>(), output.data_ptr<float>(),
                    c, spatial, spatial_tiles, channel_tiles, tiles);
        } else {
            pack_channels_last_vectorized<false>
                <<<blocks, dim3(8, 32), 0, stream>>>(
                    input.data_ptr<float>(), output.data_ptr<float>(),
                    c, spatial, spatial_tiles, channel_tiles, tiles);
        }
    } else {
        pack_channels_last_generic<<<blocks, dim3(32, 8), 0, stream>>>(
            input.data_ptr<float>(), output.data_ptr<float>(),
            static_cast<int>(c), spatial, spatial_tiles,
            static_cast<int>(channel_tiles), tiles);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

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

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

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

    auto output = torch::empty({ci, kh, kw, co}, 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);

    pack_filter<<<blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(), output.data_ptr<float>(),
        static_cast<int>(co), static_cast<int>(kh * kw), count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

std::vector<torch::Tensor> transpose_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    std::vector<int64_t> padding,
    std::vector<int64_t> output_padding,
    std::vector<int64_t> stride,
    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(), "Device mismatch");
    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");

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

    auto packed_input = channels_last_input(input, stream);
    auto packed_weight = channels_last_weight(weight, stream);

    auto output = at::cudnn_convolution_transpose(
        packed_input, packed_weight,
        padding, output_padding, stride, dilation, groups,
        true, deterministic, allow_tf32);

    return {output, packed_input, packed_weight};
}
"""

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


class _TransposedConvFunction(torch.autograd.Function):
    @staticmethod
    def forward(
        ctx, x, weight, bias, stride, padding, output_padding,
        dilation, groups, deterministic, allow_tf32
    ):
        output, packed_x, packed_weight = _extension.transpose_conv_cuda(
            x,
            weight,
            list(padding),
            list(output_padding),
            list(stride),
            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, packed_weight)
        ctx.stride = stride
        ctx.padding = padding
        ctx.output_padding = output_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,
                list(ctx.stride),
                list(ctx.padding),
                list(ctx.dilation),
                True,
                list(ctx.output_padding),
                ctx.groups,
                mask,
            )
        )
        return (
            grad_x, grad_weight, grad_bias,
            None, None, 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.weight.shape[1] < 2**31
            and conv.kernel_size[0] * conv.kernel_size[1] < 2**31
            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()

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

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