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


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

std::vector<torch::Tensor> transpose_conv_prepare(
    torch::Tensor x, torch::Tensor weight, 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 <vector>

template <typename scalar_t>
__global__ void nchw_to_nhwc(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int channels,
    int spatial)
{
    __shared__ scalar_t tile[32][33];

    const int spatial_base = blockIdx.x * 32;
    const int channel_base = blockIdx.y * 32;
    const int batch = blockIdx.z;
    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int64_t batch_offset =
        static_cast<int64_t>(batch) * channels * spatial;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int c = channel_base + ty + j;
        const int s = spatial_base + tx;
        if (c < channels && s < spatial) {
            tile[ty + j][tx] =
                input[batch_offset + static_cast<int64_t>(c) * spatial + s];
        }
    }

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int c = channel_base + tx;
        const int s = spatial_base + ty + j;
        if (c < channels && s < spatial) {
            output[batch_offset + static_cast<int64_t>(s) * channels + c] =
                tile[tx][ty + j];
        }
    }
}

template <typename scalar_t>
__global__ void transpose_and_reverse_weight(
    const scalar_t* __restrict__ source,
    scalar_t* __restrict__ destination,
    int input_per_group,
    int output_per_group,
    int kernel_h,
    int kernel_w,
    int64_t count)
{
    for (int64_t index =
             static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        int64_t remaining = index;
        const int input_local = remaining % input_per_group;
        remaining /= input_per_group;
        const int kw = remaining % kernel_w;
        remaining /= kernel_w;
        const int kh = remaining % kernel_h;
        const int output_channel = remaining / kernel_h;

        const int group = output_channel / output_per_group;
        const int output_local = output_channel % output_per_group;
        const int input_channel = group * input_per_group + input_local;

        const int64_t source_index =
            ((static_cast<int64_t>(input_channel) * output_per_group
              + output_local) * kernel_h + (kernel_h - 1 - kh))
              * kernel_w + (kernel_w - 1 - kw);

        destination[index] = source[source_index];
    }
}

std::vector<torch::Tensor> transpose_conv_prepare(
    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(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == weight.scalar_type(), "Dtype mismatch");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4, "Expected 4D 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");

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

    const int n = x.size(0);
    const int ci = x.size(1);
    const int h = x.size(2);
    const int w = x.size(3);
    const int ig = ci / groups;
    const int og = weight.size(1);
    const int co = og * groups;
    const int kh = weight.size(2);
    const int kw = weight.size(3);

    torch::Tensor converted_x;
    const bool convert_input =
        !x.is_contiguous(at::MemoryFormat::ChannelsLast);

    if (convert_input) {
        TORCH_CHECK(x.is_contiguous(), "Expected contiguous input");
        converted_x = torch::empty({n, h, w, ci}, x.options())
                          .permute({0, 3, 1, 2});
    } else {
        converted_x = x;
    }

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

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        x.scalar_type(),
        "transpose_conv_prepare",
        [&] {
            if (convert_input && x.numel() != 0) {
                const dim3 block(32, 8);
                const dim3 grid((h * w + 31) / 32, (ci + 31) / 32, n);
                nchw_to_nhwc<scalar_t><<<grid, block, 0, stream>>>(
                    x.data_ptr<scalar_t>(),
                    converted_x.data_ptr<scalar_t>(),
                    ci, h * w);
            }

            const int64_t count = weight.numel();
            const int blocks = static_cast<int>(
                std::min<int64_t>((count + 255) / 256, 4096));

            transpose_and_reverse_weight<scalar_t>
                <<<blocks, 256, 0, stream>>>(
                    weight.data_ptr<scalar_t>(),
                    converted_weight.data_ptr<scalar_t>(),
                    ig, og, kh, kw, count);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return {converted_x, converted_weight};
}
"""

_extension = None


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


class _PrepareConvolution(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, groups):
        ctx.groups = groups
        ctx.weight_shape = tuple(weight.shape)
        converted_x, converted_weight = (
            _get_extension().transpose_conv_prepare(
                x, weight.contiguous(), groups
            )
        )
        return converted_x, converted_weight

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        restored_weight = None
        if grad_weight is not None and ctx.needs_input_grad[1]:
            ci, og, kh, kw = ctx.weight_shape
            groups = ctx.groups
            restored_weight = (
                grad_weight.reshape(groups, og, ci // groups, kh, kw)
                .transpose(1, 2)
                .reshape(ci, og, kh, kw)
                .flip(-2, -1)
                .contiguous()
            )

        return (
            grad_x if ctx.needs_input_grad[0] else None,
            restored_weight,
            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:
        layer = self.conv_transpose2d
        effective_padding = tuple(
            d * (k - 1) - p
            for d, k, p in zip(
                layer.dilation, layer.kernel_size, layer.padding
            )
        )

        use_custom = (
            x.is_cuda
            and x.ndim == 4
            and x.numel() > 0
            and x.shape[0] <= 65535
            and x.dtype == layer.weight.dtype
            and x.dtype in (torch.float32, torch.float16, torch.bfloat16)
            and layer.stride == (1, 1)
            and layer.output_padding == (0, 0)
            and min(effective_padding) >= 0
            and (
                x.is_contiguous()
                or x.is_contiguous(memory_format=torch.channels_last)
            )
        )

        if not use_custom:
            return layer(x)

        # A stride-one transposed convolution is a regular convolution
        # with reversed spatial weights and exchanged channel axes.
        prepared_x, prepared_weight = _PrepareConvolution.apply(
            x, layer.weight, layer.groups
        )
        return F.conv2d(
            prepared_x,
            prepared_weight,
            layer.bias,
            stride=1,
            padding=effective_padding,
            dilation=layer.dilation,
            groups=layer.groups,
        )
```