```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 pack_input_cuda(torch::Tensor x);
torch::Tensor pack_weight_cuda(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 <algorithm>

__global__ void pack_input_float4_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t channels,
    int64_t spatial)
{
    __shared__ float tile[32][33];

    const int t = threadIdx.x;
    const int row = t >> 3;
    const int col = (t & 7) << 2;

    const int64_t s0 = int64_t(blockIdx.x) * 32;
    const int64_t c0 = int64_t(blockIdx.y) * 32;
    const int64_t base = int64_t(blockIdx.z) * channels * spatial;

    const float4 v = *reinterpret_cast<const float4*>(
        input + base + (c0 + row) * spatial + s0 + col);

    tile[row][col]     = v.x;
    tile[row][col + 1] = v.y;
    tile[row][col + 2] = v.z;
    tile[row][col + 3] = v.w;

    __syncthreads();

    const float4 result = make_float4(
        tile[col][row],
        tile[col + 1][row],
        tile[col + 2][row],
        tile[col + 3][row]);

    *reinterpret_cast<float4*>(
        output + base + (s0 + row) * channels + c0 + col) = result;
}

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

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int64_t s0 = int64_t(blockIdx.x) * 32;
    const int64_t c0 = int64_t(blockIdx.y) * 32;
    const int64_t base = int64_t(blockIdx.z) * channels * spatial;

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

    __syncthreads();

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

template <typename scalar_t>
__global__ void pack_weight_kernel(
    const scalar_t* __restrict__ weight,
    scalar_t* __restrict__ packed,
    int64_t count,
    int64_t input_per_group,
    int64_t output_per_group,
    int64_t kernel_volume)
{
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        const int64_t ic = i % input_per_group;
        const int64_t q = i / input_per_group;
        const int64_t k = q % kernel_volume;
        const int64_t oc = q / kernel_volume;
        const int64_t group = oc / output_per_group;
        const int64_t oc_local = oc % output_per_group;

        const int64_t source =
            ((group * input_per_group + ic) * output_per_group + oc_local)
                * kernel_volume
            + kernel_volume - 1 - k;

        packed[i] = weight[source];
    }
}

torch::Tensor pack_input_cuda(torch::Tensor x)
{
    TORCH_CHECK(x.is_cuda() && x.dim() == 5,
                "Expected a five-dimensional CUDA tensor");
    const c10::cuda::CUDAGuard guard(x.device());

    if (x.is_contiguous(at::MemoryFormat::ChannelsLast3d))
        return x;

    x = x.contiguous();
    auto output = at::empty(
        x.sizes(), x.options(), at::MemoryFormat::ChannelsLast3d);

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

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

    TORCH_CHECK(x.size(0) <= 65535 && (channels + 31) / 32 <= 65535,
                "Input exceeds CUDA grid limits");

    const dim3 grid(
        (spatial + 31) / 32,
        (channels + 31) / 32,
        x.size(0));
    const auto stream = at::cuda::getCurrentCUDAStream();

    if (x.scalar_type() == at::ScalarType::Float &&
        channels % 32 == 0 && spatial % 32 == 0) {
        pack_input_float4_kernel<<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), output.data_ptr<float>(),
            channels, spatial);
    } else {
        const dim3 block(32, 8);
        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half, at::ScalarType::BFloat16,
            x.scalar_type(), "pack_transpose3d_input", [&] {
                pack_input_kernel<scalar_t><<<grid, block, 0, stream>>>(
                    x.data_ptr<scalar_t>(), output.data_ptr<scalar_t>(),
                    channels, spatial);
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

torch::Tensor pack_weight_cuda(torch::Tensor weight, int64_t groups)
{
    TORCH_CHECK(weight.is_cuda() && weight.dim() == 5,
                "Expected five-dimensional CUDA weights");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "Invalid convolution groups");
    const c10::cuda::CUDAGuard guard(weight.device());

    weight = weight.contiguous();

    const int64_t input_per_group = weight.size(0) / groups;
    const int64_t output_per_group = weight.size(1);
    const int64_t output_channels = output_per_group * groups;
    const int64_t kernel_volume =
        weight.size(2) * weight.size(3) * weight.size(4);

    auto packed = at::empty(
        {output_channels, input_per_group,
         weight.size(2), weight.size(3), weight.size(4)},
        weight.options(), at::MemoryFormat::ChannelsLast3d);

    const int64_t count = weight.numel();
    if (count == 0)
        return packed;

    const int blocks = static_cast<int>(
        std::min<int64_t>((count + 255) / 256, 65535));
    const auto stream = at::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        weight.scalar_type(), "pack_transpose3d_weight", [&] {
            pack_weight_kernel<scalar_t><<<blocks, 256, 0, stream>>>(
                weight.data_ptr<scalar_t>(), packed.data_ptr<scalar_t>(),
                count, input_per_group, output_per_group, kernel_volume);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return packed;
}
"""

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


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

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


class _PackWeight(torch.autograd.Function):
    @staticmethod
    def forward(ctx, weight, groups):
        ctx.groups = groups
        ctx.original_shape = tuple(weight.shape)
        return _extension.pack_weight_cuda(weight, groups)

    @staticmethod
    def backward(ctx, grad_output):
        cin, cout_per_group, kd, kh, kw = ctx.original_shape
        grad_weight = (
            grad_output.reshape(
                ctx.groups, cout_per_group, cin // ctx.groups, kd, kh, kw
            )
            .transpose(1, 2)
            .flip((3, 4, 5))
            .reshape(ctx.original_shape)
            .contiguous()
        )
        return 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
        padding = tuple(
            dilation * (kernel - 1) - p
            for kernel, dilation, p in zip(
                conv.kernel_size, conv.dilation, conv.padding
            )
        )

        if (
            not x.is_cuda
            or x.ndim != 5
            or x.numel() == 0
            or conv.stride != (1, 1, 1)
            or conv.output_padding != (0, 0, 0)
            or any(p < 0 for p in padding)
            or x.shape[1] != conv.in_channels
            or x.device != conv.weight.device
            or x.dtype != conv.weight.dtype
            or x.dtype not in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
            or x.shape[0] > 65535
            or (x.shape[1] + 31) // 32 > 65535
        ):
            return conv(x)

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

        with torch.backends.cudnn.flags(
            enabled=torch.backends.cudnn.enabled,
            benchmark=True,
            deterministic=torch.backends.cudnn.deterministic,
            allow_tf32=torch.backends.cudnn.allow_tf32,
        ):
            return F.conv3d(
                packed_x,
                packed_weight,
                conv.bias,
                stride=1,
                padding=padding,
                dilation=conv.dilation,
                groups=conv.groups,
            )
```