```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_transpose_weight(torch::Tensor weight);
torch::Tensor expand_sparse_volume(torch::Tensor input);
"""

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

__global__ void reverse_transpose_weight_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int in_channels,
    int out_channels,
    int64_t count)
{
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        const int k = i % 27;
        const int ic = (i / 27) % in_channels;
        const int oc = i / (27LL * in_channels);
        output[i] =
            input[(int64_t(ic) * out_channels + oc) * 27 + 26 - k];
    }
}

/*
Physical storage uses a 32 x 64 x 64 volume per channel.

The logical 33 x 65 x 65 volume also includes its zero boundary
planes. Those planes can share storage with subsequent zero rows
and planes. Extra initialized storage covers the final boundary.

Four output values are written per thread:
    [0, input[2*w], 0, input[2*w+1]]
on odd depth/height coordinates, and zeros everywhere else.
*/
__global__ void expand_sparse_volume_kernel(
    const float2* __restrict__ input,
    float4* __restrict__ output,
    uint32_t main_vectors,
    uint32_t total_vectors)
{
    const uint32_t step = blockDim.x * gridDim.x;

    for (uint32_t v = blockIdx.x * blockDim.x + threadIdx.x;
         v < total_vectors;
         v += step) {
        float4 value = make_float4(0.f, 0.f, 0.f, 0.f);

        // In vector coordinates, height parity is bit 4 and depth
        // parity is bit 10. Removing those bits gives the compact
        // input's float2 index.
        if (v < main_vectors && (v & 1040u) == 1040u) {
            const uint32_t source =
                (v & 15u) |
                ((v >> 1) & 496u) |
                ((v >> 2) & ~511u);

            const float2 pair = input[source];
            value.y = pair.x;
            value.w = pair.y;
        }

        output[v] = value;
    }
}

torch::Tensor reverse_transpose_weight(torch::Tensor weight)
{
    const c10::cuda::CUDAGuard guard(weight.device());
    TORCH_CHECK(weight.scalar_type() == at::kFloat);
    TORCH_CHECK(weight.is_contiguous());

    const int ic = weight.size(0);
    const int oc = weight.size(1);
    auto output = torch::empty({oc, ic, 3, 3, 3}, weight.options());

    constexpr 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();

    reverse_transpose_weight_kernel<<<blocks, threads, 0, stream.stream()>>>(
        weight.data_ptr<float>(),
        output.data_ptr<float>(),
        ic, oc, output.numel());

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

torch::Tensor expand_sparse_volume(torch::Tensor input)
{
    const c10::cuda::CUDAGuard guard(input.device());
    TORCH_CHECK(input.scalar_type() == at::kFloat);
    TORCH_CHECK(input.is_contiguous());
    TORCH_CHECK(
        input.dim() == 5 &&
        input.size(2) == 16 &&
        input.size(3) == 32 &&
        input.size(4) == 32);

    const int64_t batch = input.size(0);
    const int64_t channels = input.size(1);

    constexpr int64_t row_stride = 64;
    constexpr int64_t plane_stride = 64 * 64;
    constexpr int64_t channel_stride = 32 * 64 * 64;

    const int64_t main_count = batch * channels * channel_stride;

    // The largest logical offset is main_count + 4160.
    // Round the allocation up to a complete float4.
    constexpr int64_t tail_count = plane_stride + row_stride + 4;
    const int64_t total_count = main_count + tail_count;

    TORCH_CHECK(total_count / 4 <= UINT32_MAX);

    auto storage = torch::empty({total_count}, input.options());

    const uint32_t main_vectors = static_cast<uint32_t>(main_count / 4);
    const uint32_t total_vectors = static_cast<uint32_t>(total_count / 4);

    constexpr int threads = 256;
    const int blocks = static_cast<int>(
        std::min<uint32_t>(
            (total_vectors + threads - 1) / threads, 4096));
    const auto stream = c10::cuda::getCurrentCUDAStream();

    expand_sparse_volume_kernel<<<blocks, threads, 0, stream.stream()>>>(
        reinterpret_cast<const float2*>(input.data_ptr<float>()),
        reinterpret_cast<float4*>(storage.data_ptr<float>()),
        main_vectors,
        total_vectors);

    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return storage.as_strided(
        {batch, channels, 33, 65, 65},
        {
            channels * channel_stride,
            channel_stride,
            plane_stride,
            row_stride,
            1
        });
}
"""

_ops = load_inline(
    name="sparse_transpose3d_compact_boundaries_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["reverse_transpose_weight", "expand_sparse_volume"],
    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 == torch.float32
            and conv.weight.dtype == torch.float32
            and x.device == conv.weight.device
            and x.shape[0] > 0
            and x.shape[1] == conv.in_channels
            and x.shape[2:] == (16, 32, 32)
            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 conv.bias is None
            and not torch.is_autocast_enabled()
            and x.shape[0] * conv.out_channels < 131071
        )
        if not supported:
            return conv(x)

        weight = _ops.reverse_transpose_weight(conv.weight.contiguous())
        compact = F.conv3d(
            x.contiguous(),
            weight,
            bias=None,
            stride=1,
            padding=1,
            dilation=1,
        ).contiguous()

        return _ops.expand_sparse_volume(compact)
```