```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>

__global__ void nchw_to_channels_last(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int channels,
    int64_t spatial)
{
    __shared__ float tile[32][33];

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

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

    __syncthreads();

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

__global__ void transform_weights(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int64_t count,
    int input_channels_per_group,
    int output_channels_per_group,
    int kernel_h,
    int kernel_w)
{
    const int64_t idx =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
    if (idx >= count) {
        return;
    }

    // Destination physical layout: [Cout, Kh, Kw, Cin/groups].
    int64_t rem = idx;
    const int ic = rem % input_channels_per_group;
    rem /= input_channels_per_group;
    const int kw = rem % kernel_w;
    rem /= kernel_w;
    const int kh = rem % kernel_h;
    const int oc = rem / kernel_h;

    const int group = oc / output_channels_per_group;
    const int source_ic = group * input_channels_per_group + ic;
    const int source_oc = oc % output_channels_per_group;

    const int64_t source_idx =
        ((int64_t(source_ic) * output_channels_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(),
                "Input and weight must share a device");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional 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),
                "Input channel count does not match weight");

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

    torch::Tensor packed_x;
    if (x.is_contiguous(at::MemoryFormat::ChannelsLast)) {
        packed_x = x;
    } else {
        TORCH_CHECK(x.is_contiguous(), "Input must be contiguous");
        packed_x = torch::empty(
            x.sizes(),
            x.options().memory_format(at::MemoryFormat::ChannelsLast));

        const int64_t spatial = x.size(2) * x.size(3);
        if (x.numel() != 0) {
            const dim3 block(32, 8);
            const dim3 grid(
                (spatial + 31) / 32,
                (x.size(1) + 31) / 32,
                x.size(0));

            nchw_to_channels_last<<<grid, block, 0, stream>>>(
                x.data_ptr<float>(),
                packed_x.data_ptr<float>(),
                static_cast<int>(x.size(1)),
                spatial);
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        }
    }

    const int input_channels_per_group = weight.size(0) / groups;
    const int output_channels_per_group = weight.size(1);
    const int output_channels = output_channels_per_group * groups;
    const int kh = weight.size(2);
    const int kw = weight.size(3);

    auto packed_weight = torch::empty(
        {output_channels, input_channels_per_group, kh, kw},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

    const int64_t count = packed_weight.numel();
    if (count != 0) {
        transform_weights<<<(count + 255) / 256, 256, 0, stream>>>(
            weight.data_ptr<float>(),
            packed_weight.data_ptr<float>(),
            count,
            input_channels_per_group,
            output_channels_per_group,
            kh,
            kw);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return {packed_x, packed_weight};
}
"""

_extension = load_inline(
    name="transpose_conv_channels_last_cuda_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

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

        return tuple(
            _extension.prepare_transpose_conv(
                x, weight.contiguous(), groups
            )
        )

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        if grad_weight is not None:
            cout, cin_per_group, kh, kw = grad_weight.shape
            groups = ctx.groups
            cout_per_group = cout // groups

            grad_weight = (
                grad_weight.reshape(
                    groups, cout_per_group, cin_per_group, kh, kw
                )
                .permute(0, 2, 1, 3, 4)
                .flip(-2, -1)
                .reshape(
                    groups * cin_per_group, cout_per_group, kh, kw
                )
                .contiguous()
            )

        return grad_x, grad_weight, None


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: int,
        stride: int = 1,
        padding: int = 0,
        output_padding: int = 0,
        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,
            groups=groups,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose2d
        kh, kw = conv.kernel_size
        ph = kh - 1 - conv.padding[0]
        pw = kw - 1 - conv.padding[1]

        eligible = (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.ndim == 4
            and 0 < x.shape[0] <= 65535
            and x.shape[2] > 0
            and x.shape[3] > 0
            and (x.shape[1] + 31) // 32 <= 65535
            and conv.stride == (1, 1)
            and conv.output_padding == (0, 0)
            and conv.dilation == (1, 1)
            and ph >= 0
            and pw >= 0
        )

        if not eligible:
            return conv(x)

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

        # A stride-one transposed convolution is a regular convolution
        # with exchanged channel axes and spatially reversed weights.
        return F.conv2d(
            packed_x,
            packed_weight,
            conv.bias,
            stride=1,
            padding=(ph, pw),
            dilation=1,
            groups=conv.groups,
        )
```