```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(
    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_channels_last(
    const scalar_t* __restrict__ src,
    scalar_t* __restrict__ dst,
    int64_t channels,
    int64_t spatial)
{
    __shared__ scalar_t tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int64_t batch = blockIdx.z;
    const int64_t spatial_base = int64_t(blockIdx.x) * 32;
    const int64_t channel_base = int64_t(blockIdx.y) * 32;

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

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

template <typename scalar_t>
__global__ void transform_weight(
    const scalar_t* __restrict__ src,
    scalar_t* __restrict__ dst,
    int64_t count,
    int64_t input_per_group,
    int64_t output_per_group,
    int64_t kernel_h,
    int64_t kernel_w)
{
    for (int64_t idx = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         idx < count;
         idx += int64_t(blockDim.x) * gridDim.x)
    {
        // Destination physical layout: [output, kh, kw, input/group].
        int64_t rem = idx;
        const int64_t ic = rem % input_per_group;
        rem /= input_per_group;
        const int64_t kw = rem % kernel_w;
        rem /= kernel_w;
        const int64_t kh = rem % kernel_h;
        const int64_t oc = rem / kernel_h;

        const int64_t group = oc / output_per_group;
        const int64_t source_ic = group * input_per_group + ic;
        const int64_t source_oc = oc % output_per_group;

        const int64_t source_idx =
            ((source_ic * output_per_group + source_oc) * kernel_h
             + kernel_h - 1 - kh) * kernel_w
             + kernel_w - 1 - kw;

        dst[idx] = src[source_idx];
    }
}

std::vector<torch::Tensor> prepare_transpose_conv(
    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(), "Devices must match");
    TORCH_CHECK(x.scalar_type() == weight.scalar_type(), "Dtypes must match");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4, "Expected 4D tensors");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid groups");
    TORCH_CHECK(x.size(1) == weight.size(0), "Input channels mismatch");
    TORCH_CHECK(weight.is_contiguous(), "Weight must be contiguous");

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

    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 kernel_h = weight.size(2);
    const int64_t kernel_w = weight.size(3);

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

    TORCH_CHECK(already_channels_last || x.is_contiguous(),
                "Input must be contiguous or channels-last");

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

    auto packed_weight = torch::empty(
        {output_per_group * groups, input_per_group, kernel_h, kernel_w},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        x.scalar_type(), "prepare_transpose_conv", [&] {
            if (!already_channels_last && x.numel() != 0) {
                const dim3 block(32, 8);
                const dim3 grid(
                    (spatial + 31) / 32,
                    (channels + 31) / 32,
                    batch);

                nchw_to_channels_last<scalar_t><<<grid, block, 0, stream>>>(
                    x.data_ptr<scalar_t>(),
                    packed_x.data_ptr<scalar_t>(),
                    channels, spatial);
                C10_CUDA_KERNEL_LAUNCH_CHECK();
            }

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

                transform_weight<scalar_t><<<blocks, 256, 0, stream>>>(
                    weight.data_ptr<scalar_t>(),
                    packed_weight.data_ptr<scalar_t>(),
                    count, input_per_group, output_per_group,
                    kernel_h, kernel_w);
                C10_CUDA_KERNEL_LAUNCH_CHECK();
            }
        });

    return {packed_x, packed_weight};
}
"""

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


class _PrepareTransposeConv(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, groups):
        ctx.groups = groups
        ctx.weight_shape = tuple(weight.shape)
        packed_x, packed_weight = _extension.prepare_transpose_conv(
            x, weight, groups
        )
        return packed_x, packed_weight

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        if grad_weight is not None:
            cin, cout_per_group, kh, kw = ctx.weight_shape
            grad_weight = (
                grad_weight.reshape(
                    ctx.groups, cout_per_group, cin // ctx.groups, kh, kw
                )
                .transpose(1, 2)
                .flip(-2, -1)
                .reshape(ctx.weight_shape)
                .contiguous()
            )
        return grad_x, grad_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:
        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.dtype == conv.weight.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
            and 0 < x.shape[0] <= 65535
            and 0 < x.shape[1] <= 65535 * 32
            and x.shape[2] > 0
            and x.shape[3] > 0
        )

        if not supported:
            return conv(x)

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

        packed_x, packed_weight = _PrepareTransposeConv.apply(
            x, conv.weight.contiguous(), conv.groups
        )

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