```python
import functools
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(
    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>

__global__ void pack_input(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int64_t spatial,
    int channels)
{
    __shared__ float tile[32][33];

    const int lane = threadIdx.x;
    const int row = 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 offset = 0; offset < 32; offset += 8) {
        const int channel = channel_base + row + offset;
        const int64_t position = spatial_base + lane;
        if (channel < channels && position < spatial) {
            tile[row + offset][lane] =
                src[batch_base + int64_t(channel) * spatial + position];
        }
    }

    __syncthreads();

    #pragma unroll
    for (int offset = 0; offset < 32; offset += 8) {
        const int channel = channel_base + lane;
        const int64_t position = spatial_base + row + offset;
        if (channel < channels && position < spatial) {
            dst[batch_base + position * channels + channel] =
                tile[lane][row + offset];
        }
    }
}

__global__ void transform_weight(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int64_t count,
    int input_per_group,
    int output_per_group,
    int kernel_volume)
{
    const int64_t index =
        int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
    if (index >= count) return;

    // Destination physical order: output, depth, height, width, input.
    const int input_local = index % input_per_group;
    const int64_t remainder = index / input_per_group;
    const int kernel_index = remainder % kernel_volume;
    const int output_channel = remainder / kernel_volume;

    const int group = output_channel / output_per_group;
    const int output_local = output_channel % output_per_group;
    const int input_channel = group * input_per_group + input_local;

    const int64_t source_index =
        (int64_t(input_channel) * output_per_group + output_local)
            * kernel_volume
        + kernel_volume - 1 - kernel_index;

    dst[index] = src[source_index];
}

std::vector<torch::Tensor> prepare_conv3d(
    torch::Tensor x, torch::Tensor weight, int64_t groups)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.dim() == 5 && weight.dim() == 5, "Expected 5D tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    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 = at::cuda::getCurrentCUDAStream();

    const int64_t n = x.size(0);
    const int64_t c = 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 packed_x;
    if (x.is_contiguous(at::MemoryFormat::ChannelsLast3d)) {
        packed_x = x;
    } else {
        x = x.contiguous();
        packed_x = at::empty_strided(
            {n, c, d, h, w},
            {c * spatial, 1, h * w * c, w * c, c},
            x.options());

        if (x.numel() > 0) {
            const dim3 block(32, 8);
            const dim3 grid(
                (spatial + 31) / 32, (c + 31) / 32, n);
            pack_input<<<grid, block, 0, stream>>>(
                x.data_ptr<float>(), packed_x.data_ptr<float>(),
                spatial, static_cast<int>(c));
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        }
    }

    weight = weight.contiguous();

    const int64_t ci = c / groups;
    const int64_t co_local = weight.size(1);
    const int64_t co = co_local * groups;
    const int64_t kd = weight.size(2);
    const int64_t kh = weight.size(3);
    const int64_t kw = weight.size(4);
    const int64_t kernel_volume = kd * kh * kw;

    auto packed_weight = at::empty_strided(
        {co, ci, kd, kh, kw},
        {kernel_volume * ci, 1, kh * kw * ci, kw * ci, ci},
        weight.options());

    const int64_t count = packed_weight.numel();
    if (count > 0) {
        transform_weight<<<(count + 255) / 256, 256, 0, stream>>>(
            weight.data_ptr<float>(), packed_weight.data_ptr<float>(),
            count, static_cast<int>(ci), static_cast<int>(co_local),
            static_cast<int>(kernel_volume));
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return {packed_x, packed_weight};
}
"""


@functools.lru_cache(maxsize=1)
def _extension():
    return load_inline(
        name="transpose3d_forward_layout_cuda_v1",
        cpp_sources=_CPP_SOURCE,
        cuda_sources=_CUDA_SOURCE,
        functions=["prepare_conv3d"],
        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)
        packed_x, packed_weight = _extension().prepare_conv3d(
            x, weight, groups
        )
        return packed_x, packed_weight

    @staticmethod
    def backward(ctx, grad_x, grad_weight):
        original_grad_weight = None
        if grad_weight is not None:
            ci, co_local, kd, kh, kw = ctx.weight_shape
            groups = ctx.groups
            original_grad_weight = (
                grad_weight.reshape(
                    groups, co_local, ci // groups, kd, kh, kw
                )
                .transpose(1, 2)
                .flip((3, 4, 5))
                .reshape(ctx.weight_shape)
            )
        return grad_x, 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(
            k - 1 - p for k, p in zip(conv.kernel_size, conv.padding)
        )

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and conv.weight.dtype == torch.float32
            and x.ndim == 5
            and 0 < x.shape[0] <= 65535
            and conv.stride == (1, 1, 1)
            and conv.dilation == (1, 1, 1)
            and conv.output_padding == (0, 0, 0)
            and min(forward_padding) >= 0
        ):
            packed_x, packed_weight = _PrepareConv3d.apply(
                x, conv.weight, conv.groups
            )
            return F.conv3d(
                packed_x,
                packed_weight,
                conv.bias,
                stride=1,
                padding=forward_padding,
                dilation=1,
                groups=conv.groups,
            )

        return conv(x)
```