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

std::vector<torch::Tensor> prepare_conv3d_cuda(
    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>

template <typename scalar_t>
__global__ void channels_last_kernel(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int64_t spatial,
    int channels)
{
    __shared__ scalar_t tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int64_t spatial_start = int64_t(blockIdx.x) * 32;
    const int channel_start = blockIdx.y * 32;
    const int64_t batch_offset =
        int64_t(blockIdx.z) * channels * spatial;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int c = channel_start + ty + j;
        const int64_t s = spatial_start + tx;
        if (c < channels && s < spatial) {
            tile[ty + j][tx] =
                input[batch_offset + int64_t(c) * spatial + s];
        }
    }

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int c = channel_start + tx;
        const int64_t s = spatial_start + ty + j;
        if (c < channels && s < spatial) {
            output[batch_offset + s * channels + c] =
                tile[tx][ty + j];
        }
    }
}

template <typename scalar_t>
__global__ void reverse_weight_kernel(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int64_t count,
    int input_channels_per_group,
    int output_channels_per_group,
    int64_t kernel_volume)
{
    for (int64_t index =
             int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x)
    {
        int64_t remaining = index;
        const int ci = remaining % input_channels_per_group;
        remaining /= input_channels_per_group;

        const int64_t kernel_index = remaining % kernel_volume;
        const int co = remaining / kernel_volume;
        const int group = co / output_channels_per_group;
        const int local_co = co % output_channels_per_group;
        const int global_ci = group * input_channels_per_group + ci;

        const int64_t source =
            (int64_t(global_ci) * output_channels_per_group + local_co)
                * kernel_volume
            + kernel_volume - 1 - kernel_index;

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

std::vector<torch::Tensor> prepare_conv3d_cuda(
    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(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == weight.scalar_type(), "Dtype mismatch");
    TORCH_CHECK(x.dim() == 5 && weight.dim() == 5, "Expected 5D tensors");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid groups");
    TORCH_CHECK(x.size(1) == weight.size(0), "Channel mismatch");

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

    const int64_t n = x.size(0);
    const int64_t ci = x.size(1);
    const int64_t d = x.size(2);
    const int64_t h = x.size(3);
    const int64_t w = x.size(4);
    const int64_t spatial = d * h * w;

    torch::Tensor channels_last_x;
    if (x.is_contiguous(at::MemoryFormat::ChannelsLast3d)) {
        channels_last_x = x;
    } else {
        x = x.contiguous();
        channels_last_x = torch::empty({n, d, h, w, ci}, x.options())
                              .permute({0, 4, 1, 2, 3});

        if (x.numel() > 0) {
            const dim3 block(32, 8);
            const dim3 grid((spatial + 31) / 32, (ci + 31) / 32, n);

            AT_DISPATCH_FLOATING_TYPES_AND2(
                at::ScalarType::Half, at::ScalarType::BFloat16,
                x.scalar_type(), "channels_last_3d", [&] {
                    channels_last_kernel<scalar_t>
                        <<<grid, block, 0, stream>>>(
                            x.data_ptr<scalar_t>(),
                            channels_last_x.data_ptr<scalar_t>(),
                            spatial, static_cast<int>(ci));
                });
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        }
    }

    weight = weight.contiguous();
    const int64_t co_per_group = weight.size(1);
    const int64_t ci_per_group = ci / groups;
    const int64_t co = co_per_group * groups;
    const int64_t kd = weight.size(2);
    const int64_t kh = weight.size(3);
    const int64_t kw = weight.size(4);

    auto forward_weight =
        torch::empty({co, kd, kh, kw, ci_per_group}, weight.options())
            .permute({0, 4, 1, 2, 3});

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

    if (count > 0) {
        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half, at::ScalarType::BFloat16,
            weight.scalar_type(), "reverse_conv3d_weight", [&] {
                reverse_weight_kernel<scalar_t>
                    <<<blocks, 256, 0, stream>>>(
                        weight.data_ptr<scalar_t>(),
                        forward_weight.data_ptr<scalar_t>(),
                        count,
                        static_cast<int>(ci_per_group),
                        static_cast<int>(co_per_group),
                        kd * kh * kw);
            });
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return {channels_last_x, forward_weight};
}
"""

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


class _PrepareConv3d(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, groups):
        ctx.groups = groups
        ctx.weight_shape = tuple(weight.shape)
        prepared_x, prepared_weight = _extension.prepare_conv3d_cuda(
            x, weight, groups
        )
        return prepared_x, prepared_weight

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        original_grad_weight = None
        if grad_weight is not None and ctx.needs_input_grad[1]:
            ci, co_per_group, kd, kh, kw = ctx.weight_shape
            original_grad_weight = (
                grad_weight.reshape(
                    ctx.groups, co_per_group, ci // ctx.groups, kd, kh, kw
                )
                .transpose(1, 2)
                .flip((-3, -2, -1))
                .reshape(ctx.weight_shape)
            )
        return (
            grad_x if ctx.needs_input_grad[0] else None,
            original_grad_weight,
            None,
        )


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: tuple,
        stride: tuple = (1, 1, 1),
        padding: tuple = (0, 0, 0),
        output_padding: tuple = (0, 0, 0),
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv_transpose3d = nn.ConvTranspose3d(
            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_transpose3d
        forward_padding = tuple(
            dilation * (kernel - 1) - padding
            for kernel, dilation, padding in zip(
                conv.kernel_size, conv.dilation, conv.padding
            )
        )

        if (
            not x.is_cuda
            or x.ndim != 5
            or x.shape[0] == 0
            or x.dtype != conv.weight.dtype
            or x.dtype not in (
                torch.float16, torch.bfloat16, torch.float32, torch.float64
            )
            or conv.stride != (1, 1, 1)
            or conv.output_padding != (0, 0, 0)
            or min(forward_padding) < 0
        ):
            return conv(x)

        prepared_x, prepared_weight = _PrepareConv3d.apply(
            x, conv.weight, conv.groups
        )
        return F.conv3d(
            prepared_x,
            prepared_weight,
            conv.bias,
            stride=1,
            padding=forward_padding,
            dilation=conv.dilation,
            groups=conv.groups,
        )
```