```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 nchw_to_channels_last_cuda(torch::Tensor input);
"""

_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 nchw_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 tile_index = blockIdx.x;
    const int64_t spatial_tile = tile_index % spatial_tiles;
    const int64_t channel_batch = tile_index / 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;

    // Coalesced reads from the original NCHW storage.
    #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();

    // Coalesced writes into NHWC storage, preserving logical NCHW shape.
    #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];
        }
    }
}

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

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

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

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

    const int64_t channels = input.size(1);
    const int64_t spatial = input.size(2) * input.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 =
        input.size(0) * spatial_tiles * channel_tiles;

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

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

    return output;
}
"""

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


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

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


def _channels_last(x):
    if x.is_contiguous(memory_format=torch.channels_last):
        return x
    if x.is_cuda and x.dtype == torch.float32 and x.is_contiguous():
        return _ChannelsLast.apply(x)
    return x.contiguous(memory_format=torch.channels_last)


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 or x.ndim != 4:
            return conv(x)

        # Keep the tensor-core convolution and its output in channels-last
        # format, eliminating the large output conversion back to NCHW.
        return F.conv_transpose2d(
            _channels_last(x),
            _channels_last(conv.weight),
            conv.bias,
            stride=conv.stride,
            padding=conv.padding,
            output_padding=conv.output_padding,
            groups=conv.groups,
            dilation=conv.dilation,
        )
```