```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 transform_weight(torch::Tensor weight);
torch::Tensor expand_output(
    torch::Tensor input, torch::Tensor bias, bool has_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 transform_weight_kernel(
    const scalar_t* __restrict__ weight,
    scalar_t* __restrict__ output,
    int in_channels,
    int out_channels,
    int64_t count)
{
    // Physical output layout: [out_channels, 3, 3, 3, in_channels].
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        const int ic = i % in_channels;
        const int64_t q = i / in_channels;
        const int k = q % 27;
        const int oc = q / 27;

        output[i] =
            weight[(int64_t(ic) * out_channels + oc) * 27 + 26 - k];
    }
}

template <bool HAS_BIAS>
__global__ void expand_channels_last_64_kernel(
    const float4* __restrict__ input,
    const float4* __restrict__ bias,
    float4* __restrict__ output)
{
    constexpr int D = 16;
    constexpr int H = 32;
    constexpr int W = 32;
    constexpr int CV = 16;
    constexpr int OD = 33;
    constexpr int OH = 65;
    constexpr int OW = 65;
    constexpr int PLANE_VECTORS = OH * OW * CV;

    const int slab = blockIdx.y;
    const int n = slab / OD;
    const int d = slab - n * OD;
    const int first = blockIdx.x * blockDim.x + threadIdx.x;
    const int step = blockDim.x * gridDim.x;

    float4 b = make_float4(0.f, 0.f, 0.f, 0.f);
    if (HAS_BIAS)
        b = bias[threadIdx.x & (CV - 1)];

    float4* dst = output + int64_t(slab) * PLANE_VECTORS;

    // Entire even depth planes contain only zero or bias.
    if ((d & 1) == 0) {
        for (int i = first; i < PLANE_VECTORS; i += step)
            dst[i] = b;
        return;
    }

    const float4* src =
        input + (int64_t(n) * D + d / 2) * H * W * CV;

    for (int i = first; i < PLANE_VECTORS; i += step) {
        const int pixel = i / CV;
        const int h = pixel / OW;
        const int w = pixel - h * OW;

        float4 value = b;
        if ((h & w & 1) != 0) {
            value = src[((h / 2) * W + w / 2) * CV + (i & (CV - 1))];
            if (HAS_BIAS) {
                value.x += b.x;
                value.y += b.y;
                value.z += b.z;
                value.w += b.w;
            }
        }
        dst[i] = value;
    }
}

template <typename scalar_t, bool HAS_BIAS>
__global__ void expand_generic_kernel(
    const scalar_t* __restrict__ input,
    const scalar_t* __restrict__ bias,
    scalar_t* __restrict__ output,
    int channels,
    int depth,
    int height,
    int width,
    int64_t count)
{
    const int od = 2 * depth + 1;
    const int oh = 2 * height + 1;
    const int ow = 2 * width + 1;

    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        int64_t q = i;
        const int c = q % channels;
        q /= channels;
        const int w = q % ow;
        q /= ow;
        const int h = q % oh;
        q /= oh;
        const int d = q % od;
        const int64_t n = q / od;

        scalar_t value = scalar_t(0);
        if ((d & h & w & 1) != 0) {
            const int64_t source =
                ((((n * depth + d / 2) * height + h / 2) * width
                  + w / 2) * channels + c);
            value = input[source];
        }
        if (HAS_BIAS)
            value = value + bias[c];

        output[i] = value;
    }
}

torch::Tensor transform_weight(torch::Tensor weight)
{
    const c10::cuda::CUDAGuard guard(weight.device());
    const int64_t ic = weight.size(0);
    const int64_t oc = weight.size(1);

    auto output = torch::empty_strided(
        {oc, ic, 3, 3, 3},
        {27 * ic, 1, 9 * ic, 3 * ic, ic},
        weight.options());

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

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        weight.scalar_type(), "transform_weight_channels_last", [&] {
            transform_weight_kernel<scalar_t>
                <<<blocks, threads, 0, stream.stream()>>>(
                    weight.data_ptr<scalar_t>(),
                    output.data_ptr<scalar_t>(),
                    static_cast<int>(ic),
                    static_cast<int>(oc),
                    output.numel());
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

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

    const int64_t batch = input.size(0);
    const int64_t channels = input.size(1);
    const int64_t depth = input.size(2);
    const int64_t height = input.size(3);
    const int64_t width = input.size(4);
    const int64_t od = 2 * depth + 1;
    const int64_t oh = 2 * height + 1;
    const int64_t ow = 2 * width + 1;

    auto output = torch::empty_strided(
        {batch, channels, od, oh, ow},
        {od * oh * ow * channels, 1,
         oh * ow * channels, ow * channels, channels},
        input.options());

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

    const int threads = 256;
    const auto stream = c10::cuda::getCurrentCUDAStream();

    if (input.scalar_type() == at::ScalarType::Float &&
        channels == 64 && depth == 16 && height == 32 && width == 32 &&
        batch * od <= 65535) {
        const dim3 grid(8, static_cast<unsigned int>(batch * od));

        const float4* src =
            reinterpret_cast<const float4*>(input.data_ptr<float>());
        float4* dst =
            reinterpret_cast<float4*>(output.data_ptr<float>());

        if (has_bias) {
            expand_channels_last_64_kernel<true>
                <<<grid, threads, 0, stream.stream()>>>(
                    src,
                    reinterpret_cast<const float4*>(bias.data_ptr<float>()),
                    dst);
        } else {
            expand_channels_last_64_kernel<false>
                <<<grid, threads, 0, stream.stream()>>>(
                    src, nullptr, dst);
        }
    } else {
        const int blocks = static_cast<int>(
            std::min<int64_t>((count + threads - 1) / threads, 4096));

        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half, at::ScalarType::BFloat16,
            input.scalar_type(), "expand_channels_last_generic", [&] {
                if (has_bias) {
                    expand_generic_kernel<scalar_t, true>
                        <<<blocks, threads, 0, stream.stream()>>>(
                            input.data_ptr<scalar_t>(),
                            bias.data_ptr<scalar_t>(),
                            output.data_ptr<scalar_t>(),
                            channels, depth, height, width, count);
                } else {
                    expand_generic_kernel<scalar_t, false>
                        <<<blocks, threads, 0, stream.stream()>>>(
                            input.data_ptr<scalar_t>(),
                            nullptr,
                            output.data_ptr<scalar_t>(),
                            channels, depth, height, width, count);
                }
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_ops = load_inline(
    name="sparse_transpose3d_channels_last_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["transform_weight", "expand_output"],
    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 == conv.weight.dtype
            and x.dtype in (
                torch.float16,
                torch.bfloat16,
                torch.float32,
                torch.float64,
            )
            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)

        weight = _ops.transform_weight(conv.weight.contiguous())
        compact = F.conv3d(
            x.contiguous(memory_format=torch.channels_last_3d),
            weight,
            bias=None,
            stride=1,
            padding=1,
            dilation=1,
        )
        compact = compact.contiguous(memory_format=torch.channels_last_3d)

        has_bias = conv.bias is not None
        bias_tensor = (
            conv.bias.to(dtype=compact.dtype).contiguous()
            if has_bias
            else compact
        )
        return _ops.expand_output(compact, bias_tensor, has_bias)
```