```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 reverse_weights_cuda(torch::Tensor weight);
torch::Tensor expand_output_cuda(
    torch::Tensor compact,
    torch::Tensor bias);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/Dispatch.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

template <typename scalar_t>
__global__ void reverse_weights_kernel(
    const scalar_t* __restrict__ source,
    scalar_t* __restrict__ destination,
    int64_t count,
    int input_channels,
    int output_channels)
{
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        const int kernel_index = index % 27;
        const int input_channel = (index / 27) % input_channels;
        const int output_channel = index / (27 * input_channels);

        const int64_t source_index =
            (int64_t(input_channel) * output_channels + output_channel) * 27
            + 26 - kernel_index;

        destination[index] = source[source_index];
    }
}

template <typename scalar_t, bool HasBias>
__global__ void expand_output_kernel(
    const scalar_t* __restrict__ compact,
    const scalar_t* __restrict__ bias,
    scalar_t* __restrict__ output,
    int64_t count,
    int channels,
    int depth,
    int height,
    int width)
{
    const int output_depth = 2 * depth + 1;
    const int output_height = 2 * height + 1;
    const int output_width = 2 * width + 1;
    const int64_t compact_volume = int64_t(depth) * height * width;

    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += int64_t(blockDim.x) * gridDim.x) {
        int64_t coordinate = index;
        const int w = coordinate % output_width;
        coordinate /= output_width;
        const int h = coordinate % output_height;
        coordinate /= output_height;
        const int d = coordinate % output_depth;
        const int64_t batch_channel = coordinate / output_depth;

        scalar_t value = scalar_t(0);
        if ((d & h & w & 1) != 0) {
            const int64_t source_index =
                batch_channel * compact_volume
                + (int64_t(d >> 1) * height + (h >> 1)) * width
                + (w >> 1);
            value = compact[source_index];
        }

        if (HasBias) {
            value = static_cast<scalar_t>(
                value + bias[batch_channel % channels]);
        }
        output[index] = value;
    }
}

torch::Tensor reverse_weights_cuda(torch::Tensor weight)
{
    TORCH_CHECK(weight.is_cuda() && weight.is_contiguous(),
                "weight must be contiguous and on CUDA");
    TORCH_CHECK(weight.dim() == 5 &&
                weight.size(2) == 3 &&
                weight.size(3) == 3 &&
                weight.size(4) == 3,
                "expected a 3x3x3 kernel");

    const c10::cuda::CUDAGuard guard(weight.device());
    const int input_channels = weight.size(0);
    const int output_channels = weight.size(1);

    auto result = torch::empty(
        {output_channels, input_channels, 3, 3, 3}, weight.options());

    const int threads = 256;
    const int blocks = static_cast<int>(
        std::min<int64_t>((weight.numel() + threads - 1) / threads, 65535));
    const auto stream = c10::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        weight.scalar_type(), "reverse_weights_cuda", [&] {
            reverse_weights_kernel<scalar_t><<<blocks, threads, 0, stream.stream()>>>(
                weight.data_ptr<scalar_t>(),
                result.data_ptr<scalar_t>(),
                weight.numel(),
                input_channels,
                output_channels);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return result;
}

torch::Tensor expand_output_cuda(
    torch::Tensor compact,
    torch::Tensor bias)
{
    TORCH_CHECK(compact.is_cuda() && compact.is_contiguous(),
                "compact must be contiguous and on CUDA");
    TORCH_CHECK(compact.dim() == 5, "expected NCDHW input");

    const c10::cuda::CUDAGuard guard(compact.device());
    const int64_t batch = compact.size(0);
    const int channels = compact.size(1);
    const int depth = compact.size(2);
    const int height = compact.size(3);
    const int width = compact.size(4);
    const bool has_bias = bias.numel() != 0;

    if (has_bias) {
        TORCH_CHECK(bias.device() == compact.device() &&
                    bias.scalar_type() == compact.scalar_type() &&
                    bias.is_contiguous() &&
                    bias.numel() == channels,
                    "incompatible bias");
    }

    auto output = torch::empty(
        {batch, channels, 2LL * depth + 1,
         2LL * height + 1, 2LL * width + 1},
        compact.options());

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

    const int threads = 256;
    const int blocks = static_cast<int>(
        std::min<int64_t>((output.numel() + threads - 1) / threads, 65535));
    const auto stream = c10::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        compact.scalar_type(), "expand_output_cuda", [&] {
            const scalar_t* bias_pointer =
                has_bias ? bias.data_ptr<scalar_t>() : nullptr;

            if (has_bias) {
                expand_output_kernel<scalar_t, true>
                    <<<blocks, threads, 0, stream.stream()>>>(
                        compact.data_ptr<scalar_t>(), bias_pointer,
                        output.data_ptr<scalar_t>(), output.numel(),
                        channels, depth, height, width);
            } else {
                expand_output_kernel<scalar_t, false>
                    <<<blocks, threads, 0, stream.stream()>>>(
                        compact.data_ptr<scalar_t>(), nullptr,
                        output.data_ptr<scalar_t>(), output.numel(),
                        channels, depth, height, width);
            }
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


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


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: int,
        stride: int = 1,
        padding: int = 0,
        dilation: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv_transpose3d = nn.ConvTranspose3d(
            in_channels,
            out_channels,
            kernel_size=(kernel_size, kernel_size, kernel_size),
            stride=stride,
            padding=padding,
            dilation=dilation,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose3d
        supported = (
            x.is_cuda
            and x.dim() == 5
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
            and x.dtype == conv.weight.dtype
            and conv.kernel_size == (3, 3, 3)
            and conv.stride == (2, 2, 2)
            and conv.padding == (1, 1, 1)
            and conv.dilation == (2, 2, 2)
            and conv.output_padding == (0, 0, 0)
            and conv.groups == 1
        )
        if not supported:
            return conv(x)

        extension = _get_extension()
        weight = extension.reverse_weights_cuda(conv.weight.contiguous())

        # Only odd coordinates along all three output axes receive
        # convolution contributions. Compute those values densely.
        compact = F.conv3d(
            x, weight, bias=None, stride=1, padding=1
        ).contiguous()

        bias = (
            conv.bias.to(dtype=compact.dtype).contiguous()
            if conv.bias is not None
            else compact.new_empty(0)
        )
        return extension.expand_output_cuda(compact, bias)
```