```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> prepare_transpose_conv_cuda(
    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 prepare_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ weight,
    scalar_t* __restrict__ x_channels_last,
    scalar_t* __restrict__ rotated_weight,
    int64_t channels,
    int64_t spatial,
    int64_t spatial_tiles,
    int64_t channel_tiles,
    int64_t input_blocks,
    int64_t weight_elements,
    int64_t input_per_group,
    int64_t output_per_group,
    int64_t kernel_h,
    int64_t kernel_w)
{
    __shared__ scalar_t tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int64_t block = blockIdx.x;

    if (block < input_blocks) {
        const int64_t spatial_tile = block % spatial_tiles;
        const int64_t channel_tile =
            (block / spatial_tiles) % channel_tiles;
        const int64_t batch =
            block / (spatial_tiles * channel_tiles);

        const int64_t s0 = spatial_tile * 32;
        const int64_t c0 = channel_tile * 32;

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

        __syncthreads();

        #pragma unroll
        for (int j = 0; j < 32; j += 8) {
            const int64_t s = s0 + ty + j;
            const int64_t c = c0 + tx;
            if (s < spatial && c < channels) {
                x_channels_last[(batch * spatial + s) * channels + c] =
                    tile[tx][ty + j];
            }
        }
        return;
    }

    const int64_t index =
        (block - input_blocks) * 256 + ty * 32 + tx;
    if (index >= weight_elements) {
        return;
    }

    // Destination physical layout: [Cout, Kh, Kw, Cin/groups].
    int64_t q = index;
    const int64_t ci = q % input_per_group;
    q /= input_per_group;
    const int64_t kw = q % kernel_w;
    q /= kernel_w;
    const int64_t kh = q % kernel_h;
    const int64_t co = q / kernel_h;

    const int64_t group = co / output_per_group;
    const int64_t co_local = co % output_per_group;
    const int64_t ci_global = group * input_per_group + ci;

    // Original transposed-convolution layout:
    // [Cin, Cout/groups, Kh, Kw].
    const int64_t source =
        ((ci_global * output_per_group + co_local) * kernel_h
         + kernel_h - 1 - kh) * kernel_w
         + kernel_w - 1 - kw;

    rotated_weight[index] = weight[source];
}

std::vector<torch::Tensor> prepare_transpose_conv_cuda(
    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 be on the same device");
    TORCH_CHECK(x.scalar_type() == weight.scalar_type(),
                "Input and weight must have the same dtype");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(groups > 0 && x.size(1) % groups == 0,
                "Invalid groups");
    TORCH_CHECK(x.size(1) == weight.size(0),
                "Input channel mismatch");

    const c10::cuda::CUDAGuard device_guard(x.device());

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

    if (!already_channels_last) {
        x = x.contiguous();
    }
    weight = weight.contiguous();

    const int64_t batch = x.size(0);
    const int64_t channels = x.size(1);
    const int64_t spatial = x.size(2) * x.size(3);
    const int64_t input_per_group = channels / groups;
    const int64_t output_per_group = weight.size(1);
    const int64_t output_channels = output_per_group * groups;
    const int64_t kernel_h = weight.size(2);
    const int64_t kernel_w = weight.size(3);

    auto prepared_x = already_channels_last
        ? x
        : at::empty(
            x.sizes(), x.options(), at::MemoryFormat::ChannelsLast);

    auto prepared_weight = at::empty(
        {output_channels, input_per_group, kernel_h, kernel_w},
        weight.options(), at::MemoryFormat::ChannelsLast);

    const int64_t spatial_tiles = (spatial + 31) / 32;
    const int64_t channel_tiles = (channels + 31) / 32;
    const int64_t input_blocks = already_channels_last
        ? 0
        : batch * spatial_tiles * channel_tiles;
    const int64_t weight_elements = weight.numel();
    const int64_t blocks =
        input_blocks + (weight_elements + 255) / 256;

    const auto stream =
        at::cuda::getCurrentCUDAStream(x.get_device());

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        x.scalar_type(),
        "prepare_transpose_conv_cuda",
        [&] {
            prepare_kernel<scalar_t>
                <<<static_cast<unsigned int>(blocks),
                   dim3(32, 8), 0, stream>>>(
                    x.data_ptr<scalar_t>(),
                    weight.data_ptr<scalar_t>(),
                    prepared_x.data_ptr<scalar_t>(),
                    prepared_weight.data_ptr<scalar_t>(),
                    channels,
                    spatial,
                    spatial_tiles,
                    channel_tiles,
                    input_blocks,
                    weight_elements,
                    input_per_group,
                    output_per_group,
                    kernel_h,
                    kernel_w);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return {prepared_x, prepared_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=["prepare_transpose_conv_cuda"],
            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)
        prepared_x, prepared_weight = (
            _get_extension().prepare_transpose_conv_cuda(
                x, weight, groups
            )
        )
        return prepared_x, prepared_weight

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        original_weight_grad = None

        if grad_weight is not None and ctx.needs_input_grad[1]:
            cin, cout_per_group, kh, kw = ctx.weight_shape
            groups = ctx.groups
            original_weight_grad = (
                grad_weight.reshape(
                    groups, cout_per_group, cin // groups, kh, kw
                )
                .permute(0, 2, 1, 3, 4)
                .flip(-2, -1)
                .reshape(ctx.weight_shape)
                .contiguous()
            )

        return (
            grad_x if ctx.needs_input_grad[0] else None,
            original_weight_grad,
            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

        forward_padding = tuple(
            d * (k - 1) - p
            for d, k, p in zip(
                conv.dilation, conv.kernel_size, conv.padding
            )
        )

        supported = (
            x.is_cuda
            and x.ndim == 4
            and x.numel() > 0
            and conv.weight.device == x.device
            and conv.weight.dtype == x.dtype
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
            and conv.stride == (1, 1)
            and conv.output_padding == (0, 0)
            and min(forward_padding) >= 0
        )

        if not supported:
            return conv(x)

        # Unit-stride transposed convolution equals a forward convolution
        # with spatially reversed, channel-transposed weights.
        prepared_x, prepared_weight = _PrepareConvolution.apply(
            x, conv.weight, conv.groups
        )

        return F.conv2d(
            prepared_x,
            prepared_weight,
            conv.bias,
            stride=1,
            padding=forward_padding,
            dilation=conv.dilation,
            groups=conv.groups,
        )
```