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

torch::Tensor to_channels_last_cuda(torch::Tensor x);
torch::Tensor prepare_conv_weight_cuda(torch::Tensor weight);
"""

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

constexpr int TILE_C = 64;
constexpr int TILE_S = 128;
constexpr int THREADS = 256;

__global__ void to_channels_last_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t channels,
    int64_t spatial,
    int64_t spatial_tiles,
    int64_t channel_tiles)
{
    __shared__ float tile[TILE_C][TILE_S + 1];

    const int64_t spatial_tile = blockIdx.x % spatial_tiles;
    const int64_t channel_batch = blockIdx.x / spatial_tiles;
    const int64_t channel_tile = channel_batch % channel_tiles;
    const int64_t batch = channel_batch / channel_tiles;

    const int64_t channel_start = channel_tile * TILE_C;
    const int64_t spatial_start = spatial_tile * TILE_S;
    const int64_t batch_offset = batch * channels * spatial;

    #pragma unroll
    for (int index = threadIdx.x;
         index < TILE_C * TILE_S;
         index += THREADS) {
        const int c = index / TILE_S;
        const int s = index % TILE_S;
        const int64_t channel = channel_start + c;
        const int64_t position = spatial_start + s;

        float value = 0.0f;
        if (channel < channels && position < spatial) {
            value = input[batch_offset + channel * spatial + position];
        }
        tile[c][s] = value;
    }

    __syncthreads();

    #pragma unroll
    for (int index = threadIdx.x;
         index < TILE_C * TILE_S;
         index += THREADS) {
        const int s = index / TILE_C;
        const int c = index % TILE_C;
        const int64_t channel = channel_start + c;
        const int64_t position = spatial_start + s;

        if (channel < channels && position < spatial) {
            output[batch_offset + position * channels + channel] =
                tile[c][s];
        }
    }
}

// Convert [Cin, Cout, 3, 3] transposed-convolution weights into
// spatially reversed [Cout, Cin, 3, 3] channels-last convolution weights.
__global__ void prepare_conv_weight_kernel(
    const float* __restrict__ weight,
    float* __restrict__ output,
    int64_t in_channels,
    int64_t out_channels,
    int64_t count)
{
    const int64_t index =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;

    if (index < count) {
        const int64_t ci = index % in_channels;
        const int64_t tap = (index / in_channels) % 9;
        const int64_t co = index / (in_channels * 9);

        output[index] =
            weight[(ci * out_channels + co) * 9 + (8 - tap)];
    }
}

torch::Tensor to_channels_last_cuda(torch::Tensor x)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 tensor");
    TORCH_CHECK(x.dim() == 4 && x.is_contiguous(),
                "Expected contiguous NCHW input");

    c10::cuda::CUDAGuard guard(x.device());

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

    if (x.numel() == 0) {
        return output;
    }

    const int64_t channels = x.size(1);
    const int64_t spatial = x.size(2) * x.size(3);
    const int64_t spatial_tiles = (spatial + TILE_S - 1) / TILE_S;
    const int64_t channel_tiles = (channels + TILE_C - 1) / TILE_C;
    const int64_t blocks = x.size(0) * spatial_tiles * channel_tiles;

    TORCH_CHECK(blocks <= 2147483647LL, "Input is too large");

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    to_channels_last_kernel<<<
        static_cast<unsigned int>(blocks), THREADS, 0, stream>>>(
            x.data_ptr<float>(),
            output.data_ptr<float>(),
            channels,
            spatial,
            spatial_tiles,
            channel_tiles);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}

torch::Tensor prepare_conv_weight_cuda(torch::Tensor weight)
{
    TORCH_CHECK(weight.is_cuda() && weight.scalar_type() == at::kFloat,
                "Expected CUDA float32 weights");
    TORCH_CHECK(
        weight.dim() == 4 && weight.is_contiguous() &&
        weight.size(2) == 3 && weight.size(3) == 3,
        "Expected contiguous [Cin, Cout, 3, 3] weights");

    c10::cuda::CUDAGuard guard(weight.device());

    const int64_t in_channels = weight.size(0);
    const int64_t out_channels = weight.size(1);

    auto output = torch::empty(
        {out_channels, in_channels, 3, 3},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

    const int64_t count = weight.numel();
    if (count == 0) {
        return output;
    }

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    prepare_conv_weight_kernel<<<
        static_cast<unsigned int>((count + 255) / 256),
        256, 0, stream>>>(
            weight.data_ptr<float>(),
            output.data_ptr<float>(),
            in_channels,
            out_channels,
            count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

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


class _ChannelsLast(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        return _extension.to_channels_last_cuda(x)

    @staticmethod
    def backward(ctx, grad_output):
        return grad_output


class _PrepareConvWeight(torch.autograd.Function):
    @staticmethod
    def forward(ctx, weight):
        return _extension.prepare_conv_weight_cuda(weight)

    @staticmethod
    def backward(ctx, grad_output):
        return grad_output.flip((2, 3)).transpose(0, 1).contiguous()


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

        if not (
            x.is_cuda
            and x.ndim == 4
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.device == conv.weight.device
            and x.shape[1] == conv.in_channels
            and x.shape[2] > 0
            and x.shape[3] > 0
            and conv.weight.is_contiguous()
            and conv.kernel_size == (3, 3)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.output_padding == (0, 0)
            and conv.dilation == (1, 1)
            and conv.groups == 1
            and not torch.is_autocast_enabled()
        ):
            return conv(x)

        if not x.is_contiguous(memory_format=torch.channels_last):
            if x.is_contiguous():
                x = _ChannelsLast.apply(x)
            else:
                x = x.contiguous(memory_format=torch.channels_last)

        weight = _PrepareConvWeight.apply(conv.weight)

        # A stride-one transposed convolution is a forward convolution
        # with exchanged channel axes, reversed spatial weights, and
        # padding equal to kernel_size - 1.
        return F.conv2d(
            x,
            weight,
            conv.bias,
            stride=1,
            padding=2,
            dilation=1,
            groups=1,
        )
```