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


_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__ input,
    scalar_t* __restrict__ output,
    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 spatial_base = int64_t(blockIdx.x) * 32;
    const int64_t channel_base = int64_t(blockIdx.y) * 32;
    const int64_t batch_base = int64_t(blockIdx.z) * channels * spatial;

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

    __syncthreads();

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

template <typename scalar_t>
__global__ void transform_weights(
    const scalar_t* __restrict__ weight,
    scalar_t* __restrict__ output,
    int64_t count,
    int64_t inputs_per_group,
    int64_t outputs_per_group,
    int64_t kernel_h,
    int64_t kernel_w)
{
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        // Output physical layout: [out_channels, kernel_h, kernel_w,
        //                          inputs_per_group].
        int64_t remaining = index;
        const int64_t input_channel = remaining % inputs_per_group;
        remaining /= inputs_per_group;
        const int64_t kw = remaining % kernel_w;
        remaining /= kernel_w;
        const int64_t kh = remaining % kernel_h;
        const int64_t output_channel = remaining / kernel_h;

        const int64_t group = output_channel / outputs_per_group;
        const int64_t group_output = output_channel % outputs_per_group;
        const int64_t original_input =
            group * inputs_per_group + input_channel;

        const int64_t source =
            ((original_input * outputs_per_group + group_output) *
             kernel_h + (kernel_h - 1 - kh)) *
             kernel_w + (kernel_w - 1 - kw);

        output[index] = weight[source];
    }
}

std::vector<torch::Tensor> prepare_transposed_convolution(
    torch::Tensor input,
    torch::Tensor weight,
    int64_t groups)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must share a device");
    TORCH_CHECK(input.scalar_type() == weight.scalar_type(),
                "Input and weight must share a dtype");
    TORCH_CHECK(input.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid groups");
    TORCH_CHECK(input.size(1) == weight.size(0),
                "Input channel mismatch");

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

    torch::Tensor packed_input;
    if (input.is_contiguous(at::MemoryFormat::ChannelsLast)) {
        packed_input = input;
    } else {
        auto contiguous_input = input.contiguous();
        packed_input = torch::empty(
            input.sizes(),
            input.options().memory_format(at::MemoryFormat::ChannelsLast));

        const int64_t channels = input.size(1);
        const int64_t spatial = input.size(2) * input.size(3);

        if (input.numel() > 0) {
            const dim3 threads(32, 8);
            const dim3 blocks(
                (spatial + 31) / 32,
                (channels + 31) / 32,
                input.size(0));

            AT_DISPATCH_FLOATING_TYPES_AND2(
                at::ScalarType::Half,
                at::ScalarType::BFloat16,
                input.scalar_type(),
                "pack_transposed_conv_input",
                [&] {
                    nchw_to_channels_last<scalar_t>
                        <<<blocks, threads, 0, stream>>>(
                            contiguous_input.data_ptr<scalar_t>(),
                            packed_input.data_ptr<scalar_t>(),
                            channels,
                            spatial);
                });
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        }
    }

    auto contiguous_weight = weight.contiguous();
    const int64_t inputs_per_group = weight.size(0) / groups;
    const int64_t outputs_per_group = weight.size(1);
    const int64_t kernel_h = weight.size(2);
    const int64_t kernel_w = weight.size(3);

    auto transformed_weight = torch::empty(
        {outputs_per_group * groups, inputs_per_group, kernel_h, kernel_w},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

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

        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half,
            at::ScalarType::BFloat16,
            weight.scalar_type(),
            "transform_transposed_conv_weight",
            [&] {
                transform_weights<scalar_t>
                    <<<blocks, 256, 0, stream>>>(
                        contiguous_weight.data_ptr<scalar_t>(),
                        transformed_weight.data_ptr<scalar_t>(),
                        count,
                        inputs_per_group,
                        outputs_per_group,
                        kernel_h,
                        kernel_w);
            });
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return {packed_input, transformed_weight};
}
"""

_extension = load_inline(
    name="transposed_conv_channels_last_cuda_v1",
    cpp_sources=r"""
#include <torch/extension.h>
#include <vector>
std::vector<torch::Tensor> prepare_transposed_convolution(
    torch::Tensor input, torch::Tensor weight, int64_t groups);
""",
    cuda_sources=_cuda_source,
    functions=["prepare_transposed_convolution"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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

    @staticmethod
    def backward(ctx, grad_x, grad_transformed_weight):
        grad_weight = None
        if grad_transformed_weight is not None and ctx.needs_input_grad[1]:
            in_channels, outputs_per_group, kh, kw = ctx.weight_shape
            inputs_per_group = in_channels // ctx.groups
            grad_weight = (
                grad_transformed_weight.reshape(
                    ctx.groups,
                    outputs_per_group,
                    inputs_per_group,
                    kh,
                    kw,
                )
                .permute(0, 2, 1, 3, 4)
                .reshape(ctx.weight_shape)
                .flip((-2, -1))
            )

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

        # At unit stride, transpose convolution is an ordinary convolution
        # with exchanged channel axes and spatially reversed weights.
        if (
            x.is_cuda
            and x.ndim == 4
            and x.numel() > 0
            and x.size(0) <= 65535
            and x.size(1) <= 65535 * 32
            and x.dtype == conv.weight.dtype
            and x.device == conv.weight.device
            and x.dtype
            in (torch.float16, torch.bfloat16, torch.float32, torch.float64)
            and conv.stride == (1, 1)
            and conv.dilation == (1, 1)
            and conv.output_padding == (0, 0)
            and ph >= 0
            and pw >= 0
        ):
            packed_x, transformed_weight = _PrepareConvolution.apply(
                x, conv.weight, conv.groups
            )
            return F.conv2d(
                packed_x,
                transformed_weight,
                conv.bias,
                stride=1,
                padding=(ph, pw),
                groups=conv.groups,
            )

        return conv(x)
```