```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 transpose_flip_weight_cuda(torch::Tensor weight);

torch::Tensor expand_sparse_cuda(
    torch::Tensor compact,
    torch::Tensor bias,
    bool has_bias);
"""

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

template <typename scalar_t>
__global__ void transpose_flip_weight_kernel(
    const scalar_t* __restrict__ weight,
    scalar_t* __restrict__ transformed,
    int64_t count,
    int in_channels,
    int out_channels)
{
    for (int64_t idx = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         idx < count;
         idx += int64_t(blockDim.x) * gridDim.x) {
        const int kernel_idx = idx % 27;
        const int64_t channel_idx = idx / 27;
        const int ic = channel_idx % in_channels;
        const int oc = channel_idx / in_channels;

        // Reversing the flattened cubic kernel reverses all three axes.
        const int64_t source =
            (int64_t(ic) * out_channels + oc) * 27 + 26 - kernel_idx;
        transformed[idx] = weight[source];
    }
}

torch::Tensor transpose_flip_weight_cuda(torch::Tensor weight)
{
    const c10::cuda::CUDAGuard guard(weight.device());

    const int in_channels = weight.size(0);
    const int out_channels = weight.size(1);

    auto transformed = torch::empty(
        {out_channels, in_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, 4096));

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half,
        at::ScalarType::BFloat16,
        weight.scalar_type(),
        "transpose_flip_weight_cuda",
        [&] {
            transpose_flip_weight_kernel<scalar_t>
                <<<blocks, threads, 0,
                   at::cuda::getCurrentCUDAStream()>>>(
                    weight.data_ptr<scalar_t>(),
                    transformed.data_ptr<scalar_t>(),
                    weight.numel(),
                    in_channels,
                    out_channels);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return transformed;
}

template <typename scalar_t, int FixedD, int FixedH, int FixedW>
__global__ void expand_sparse_kernel(
    const scalar_t* __restrict__ compact,
    const scalar_t* __restrict__ bias,
    scalar_t* __restrict__ output,
    int64_t count,
    int channels,
    int runtime_d,
    int runtime_h,
    int runtime_w)
{
    const int d = FixedD != 0 ? FixedD : runtime_d;
    const int h = FixedH != 0 ? FixedH : runtime_h;
    const int w = FixedW != 0 ? FixedW : runtime_w;

    const int od_size = 2 * d + 1;
    const int oh_size = 2 * h + 1;
    const int ow_size = 2 * w + 1;
    const int64_t out_volume =
        int64_t(od_size) * oh_size * ow_size;
    const int64_t compact_volume = int64_t(d) * h * w;

    for (int64_t idx = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         idx < count;
         idx += int64_t(blockDim.x) * gridDim.x) {
        const int64_t nc = idx / out_volume;
        const int64_t spatial = idx - nc * out_volume;

        const int ow = spatial % ow_size;
        const int64_t dh = spatial / ow_size;
        const int oh = dh % oh_size;
        const int od = dh / oh_size;

        scalar_t value = scalar_t(0);
        if ((ow & 1) && (oh & 1) && (od & 1)) {
            const int64_t compact_idx =
                nc * compact_volume +
                (int64_t(od >> 1) * h + (oh >> 1)) * w +
                (ow >> 1);
            value = compact[compact_idx];
        }

        if (bias != nullptr) {
            value = value + bias[nc % channels];
        }
        output[idx] = value;
    }
}

torch::Tensor expand_sparse_cuda(
    torch::Tensor compact,
    torch::Tensor bias,
    bool has_bias)
{
    const c10::cuda::CUDAGuard guard(compact.device());

    const int64_t batch = compact.size(0);
    const int channels = compact.size(1);
    const int d = compact.size(2);
    const int h = compact.size(3);
    const int w = compact.size(4);

    auto output = torch::empty(
        {batch, channels, 2 * d + 1, 2 * h + 1, 2 * w + 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, 8192));

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

            if (d == 16 && h == 32 && w == 32) {
                expand_sparse_kernel<scalar_t, 16, 32, 32>
                    <<<blocks, threads, 0,
                       at::cuda::getCurrentCUDAStream()>>>(
                        compact.data_ptr<scalar_t>(),
                        bias_ptr,
                        output.data_ptr<scalar_t>(),
                        output.numel(),
                        channels, d, h, w);
            } else {
                expand_sparse_kernel<scalar_t, 0, 0, 0>
                    <<<blocks, threads, 0,
                       at::cuda::getCurrentCUDAStream()>>>(
                        compact.data_ptr<scalar_t>(),
                        bias_ptr,
                        output.data_ptr<scalar_t>(),
                        output.numel(),
                        channels, d, h, w);
            }
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_sparse_deconv = load_inline(
    name="sparse_deconv3d_stride2_dilation2_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=[
        "transpose_flip_weight_cuda",
        "expand_sparse_cuda",
    ],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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.ndim == 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
            and not torch.is_autocast_enabled()
        )

        if not supported:
            return conv(x)

        # Nonzero convolution contributions occur only at odd coordinates.
        # Compute that compact lattice with a dense convolution, then expand.
        weight = _sparse_deconv.transpose_flip_weight_cuda(
            conv.weight.contiguous()
        )
        compact = F.conv3d(
            x.contiguous(),
            weight,
            bias=None,
            stride=1,
            padding=1,
            dilation=1,
        ).contiguous()

        has_bias = conv.bias is not None
        return _sparse_deconv.expand_sparse_cuda(
            compact,
            conv.bias.contiguous() if has_bias else compact,
            has_bias,
        )
```