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

void pack_input_cuda(torch::Tensor x, torch::Tensor out);
void pack_weight_cuda(torch::Tensor w, torch::Tensor out, 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.h>
#include <cuda_runtime.h>

template <typename scalar_t>
__global__ void pack_input_kernel(
    const scalar_t* __restrict__ x,
    scalar_t* __restrict__ out,
    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 channel_tiles = (channels + 31) / 32;

    const int64_t batch = blockIdx.y / channel_tiles;
    const int64_t channel_base = (blockIdx.y % channel_tiles) * 32;
    const int64_t spatial_base = int64_t(blockIdx.x) * 32;

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int64_t c = channel_base + ty + j;
        const int64_t s = spatial_base + tx;
        if (c < channels && s < spatial) {
            tile[ty + j][tx] =
                x[(batch * channels + c) * spatial + s];
        }
    }

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 32; j += 8) {
        const int64_t c = channel_base + tx;
        const int64_t s = spatial_base + ty + j;
        if (c < channels && s < spatial) {
            out[(batch * spatial + s) * channels + c] =
                tile[tx][ty + j];
        }
    }
}

template <typename scalar_t>
__global__ void pack_weight_kernel(
    const scalar_t* __restrict__ w,
    scalar_t* __restrict__ out,
    int64_t input_per_group,
    int64_t output_per_group,
    int64_t kernel_volume,
    int64_t count)
{
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        // Destination physical layout: [output, depth, height, width, input].
        const int64_t input_local = index % input_per_group;
        const int64_t k = (index / input_per_group) % kernel_volume;
        const int64_t output_channel =
            index / (input_per_group * kernel_volume);
        const int64_t group = output_channel / output_per_group;
        const int64_t output_local = output_channel % output_per_group;
        const int64_t input_channel =
            group * input_per_group + input_local;

        const int64_t source =
            (input_channel * output_per_group + output_local) * kernel_volume
            + kernel_volume - 1 - k;

        out[index] = w[source];
    }
}

void pack_input_cuda(torch::Tensor x, torch::Tensor out)
{
    const c10::cuda::CUDAGuard guard(x.device());
    const int64_t batch = x.size(0);
    const int64_t channels = x.size(1);
    const int64_t spatial = x.size(2) * x.size(3) * x.size(4);

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

    const dim3 block(32, 8);
    const dim3 grid(
        static_cast<unsigned int>((spatial + 31) / 32),
        static_cast<unsigned int>(batch * ((channels + 31) / 32)));

    const auto stream = at::cuda::getCurrentCUDAStream();

    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>(),
                out.data_ptr<scalar_t>(),
                channels,
                spatial);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
}

void pack_weight_cuda(
    torch::Tensor w, torch::Tensor out, int64_t groups)
{
    const c10::cuda::CUDAGuard guard(w.device());
    const int64_t input_per_group = w.size(0) / groups;
    const int64_t output_per_group = w.size(1);
    const int64_t kernel_volume = w.size(2) * w.size(3) * w.size(4);
    const int64_t count = w.numel();

    if (count == 0) return;

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

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

    C10_CUDA_KERNEL_LAUNCH_CHECK();
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="transpose3d_forward_conv_pack_v1",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["pack_input_cuda", "pack_weight_cuda"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _extension


class _PackInput(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x):
        out = torch.empty(
            x.shape,
            dtype=x.dtype,
            device=x.device,
            memory_format=torch.channels_last_3d,
        )
        _get_extension().pack_input_cuda(x, out)
        return out

    @staticmethod
    def backward(ctx, grad_out):
        return grad_out.contiguous()


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

        input_channels, output_per_group, kd, kh, kw = weight.shape
        out = torch.empty(
            (output_per_group * groups, input_channels // groups, kd, kh, kw),
            dtype=weight.dtype,
            device=weight.device,
            memory_format=torch.channels_last_3d,
        )
        _get_extension().pack_weight_cuda(weight, out, groups)
        return out

    @staticmethod
    def backward(ctx, grad_out):
        input_channels, output_per_group, kd, kh, kw = ctx.weight_shape
        groups = ctx.groups
        grad_weight = (
            grad_out.reshape(
                groups, output_per_group, input_channels // groups, kd, kh, kw
            )
            .permute(0, 2, 1, 3, 4, 5)
            .flip((3, 4, 5))
            .reshape(ctx.weight_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
        forward_padding = tuple(
            k - 1 - p for k, p in zip(conv.kernel_size, conv.padding)
        )

        supported = (
            x.is_cuda
            and x.ndim == 5
            and x.numel() > 0
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
            and conv.stride == (1, 1, 1)
            and conv.dilation == (1, 1, 1)
            and conv.output_padding == (0, 0, 0)
            and min(forward_padding) >= 0
            and x.shape[0] * ((x.shape[1] + 31) // 32) <= 65535
        )
        if not supported:
            return conv(x)

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

        weight = _PackWeight.apply(conv.weight.contiguous(), conv.groups)

        # A unit-stride transposed convolution equals a forward convolution
        # with spatially reversed, channel-transposed weights.
        return F.conv3d(
            x,
            weight,
            conv.bias,
            stride=1,
            padding=forward_padding,
            dilation=1,
            groups=conv.groups,
        )
```