```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_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>
#include <cuda_runtime.h>
#include <vector>

// Convert NCDHW to channels-last storage using a tiled transpose.
__global__ void input_channels_last_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int64_t spatial)
{
    __shared__ float tile[32][33];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int channel_base = blockIdx.y * 32;
    const int64_t spatial_base = int64_t(blockIdx.x) * 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] =
                input[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) {
            output[batch_base + s * channels + c] = tile[tx][ty + j];
        }
    }
}

// Convert transposed-convolution weights to spatially reversed
// forward-convolution weights, directly in channels-last storage.
__global__ void prepare_weight_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int input_channels_per_group,
    int output_channels_per_group,
    int kernel_volume,
    int64_t count)
{
    const int64_t index =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;

    if (index >= count) {
        return;
    }

    int64_t remaining = index;
    const int local_input_channel =
        remaining % input_channels_per_group;
    remaining /= input_channels_per_group;
    const int kernel_index = remaining % kernel_volume;
    const int output_channel = remaining / kernel_volume;

    const int group = output_channel / output_channels_per_group;
    const int local_output_channel =
        output_channel % output_channels_per_group;
    const int input_channel =
        group * input_channels_per_group + local_input_channel;

    const int64_t source =
        (int64_t(input_channel) * output_channels_per_group
         + local_output_channel) * 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(),
                "Input and weight must use the same device");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 5 && weight.dim() == 5,
                "Expected five-dimensional tensors");
    TORCH_CHECK(weight.is_contiguous(),
                "Weight must be contiguous");

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

    torch::Tensor packed_input;

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

        const int channels = x.size(1);
        const int64_t spatial = x.size(2) * x.size(3) * x.size(4);

        const dim3 threads(32, 8);
        const dim3 blocks(
            (spatial + 31) / 32,
            (channels + 31) / 32,
            x.size(0));

        input_channels_last_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(),
            packed_input.data_ptr<float>(),
            channels,
            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 kernel_volume =
        weight.size(2) * weight.size(3) * weight.size(4);

    auto packed_weight = torch::empty(
        {output_channels, input_channels_per_group,
         weight.size(2), weight.size(3), weight.size(4)},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast3d));

    const int64_t count = packed_weight.numel();
    prepare_weight_kernel<<<(count + 255) / 256, 256, 0, stream>>>(
        weight.data_ptr<float>(),
        packed_weight.data_ptr<float>(),
        input_channels_per_group,
        output_channels_per_group,
        kernel_volume,
        count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return {packed_input, packed_weight};
}
"""

_extension = load_inline(
    name="conv_transpose3d_forward_packing_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)
        ctx.set_materialize_grads(False)

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

        packed_x, packed_weight = _extension.prepare_conv3d_cuda(
            x, weight.contiguous(), groups
        )
        return packed_x, packed_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]:
            input_channels, output_channels_per_group, kd, kh, kw = (
                ctx.weight_shape
            )
            input_channels_per_group = input_channels // ctx.groups

            original_grad_weight = (
                grad_weight.flip((2, 3, 4))
                .reshape(
                    ctx.groups,
                    output_channels_per_group,
                    input_channels_per_group,
                    kd,
                    kh,
                    kw,
                )
                .transpose(1, 2)
                .reshape(ctx.weight_shape)
                .contiguous()
            )

        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

        # A unit-stride transposed convolution is a forward convolution
        # with exchanged channel axes and spatially reversed weights.
        use_forward_convolution = (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.ndim == 5
            and x.numel() > 0
            and x.size(1) == conv.in_channels
            and x.size(0) <= 65535
            and conv.stride == (1, 1, 1)
            and conv.output_padding == (0, 0, 0)
            and all(
                0 <= p < k
                for p, k in zip(conv.padding, conv.kernel_size)
            )
        )

        if not use_forward_convolution:
            return conv(x)

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

        forward_padding = tuple(
            k - 1 - p
            for k, p in zip(conv.kernel_size, conv.padding)
        )

        return F.conv3d(
            packed_x,
            packed_weight,
            conv.bias,
            stride=1,
            padding=forward_padding,
            dilation=1,
            groups=conv.groups,
        )
```