```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_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];
    }
}

__device__ __forceinline__ void store_streaming(
    float4* address, float4 value)
{
    asm volatile(
        "st.global.cs.v4.f32 [%0], {%1, %2, %3, %4};"
        :
        : "l"(address), "f"(value.x), "f"(value.y),
          "f"(value.z), "f"(value.w)
        : "memory");
}

// Each block expands a complete compact plane into two output planes.
// All threads load useful input, and each issues eight coalesced stores.
__global__ __launch_bounds__(256) void expand_planes_kernel(
    const float2* __restrict__ input,
    float4* __restrict__ output,
    uint32_t compact_planes)
{
    const uint32_t tid = threadIdx.x;
    const float4 zero = make_float4(0.f, 0.f, 0.f, 0.f);

    if (blockIdx.x == 0) {
        const uint32_t tail_base = compact_planes * 2048u;
        for (uint32_t i = tid; i < 1041u; i += 256u) {
            store_streaming(output + tail_base + i, zero);
        }
    }

    const uint32_t offset = (tid & 15u) | ((tid & 240u) << 1);

    for (uint32_t plane = blockIdx.x;
         plane < compact_planes;
         plane += gridDim.x) {
        const uint32_t source = plane * 512u + tid;
        const float2 first = __ldg(input + source);
        const float2 second = __ldg(input + source + 256u);

        float4* destination = output + plane * 2048u + offset;

        store_streaming(destination, zero);
        store_streaming(destination + 16u, zero);
        store_streaming(destination + 512u, zero);
        store_streaming(destination + 528u, zero);

        store_streaming(destination + 1024u, zero);
        store_streaming(
            destination + 1040u,
            make_float4(0.f, first.x, 0.f, first.y));

        store_streaming(destination + 1536u, zero);
        store_streaming(
            destination + 1552u,
            make_float4(0.f, second.x, 0.f, second.y));
    }
}

torch::Tensor reverse_transpose_weight(torch::Tensor weight)
{
    TORCH_CHECK(weight.is_cuda());
    const c10::cuda::CUDAGuard guard(weight.device());
    TORCH_CHECK(weight.scalar_type() == at::kFloat);
    TORCH_CHECK(weight.is_contiguous());
    TORCH_CHECK(
        weight.dim() == 5 &&
        weight.size(2) == 3 &&
        weight.size(3) == 3 &&
        weight.size(4) == 3);

    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_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)
{
    TORCH_CHECK(input.is_cuda());
    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 = 4096;
    constexpr int64_t channel_stride = 131072;
    constexpr int64_t tail_count = 4164;

    const int64_t main_count = batch * channels * channel_stride;
    const int64_t total_count = main_count + tail_count;
    TORCH_CHECK(batch > 0 && channels > 0);
    TORCH_CHECK(total_count / 4 < UINT32_MAX);

    auto storage = torch::empty({total_count}, input.options());
    const uint32_t compact_planes =
        static_cast<uint32_t>(batch * channels * 16);
    const int blocks =
        static_cast<int>(std::min<uint32_t>(compact_planes, 4096u));
    const auto stream = c10::cuda::getCurrentCUDAStream();

    expand_planes_kernel<<<blocks, 256, 0, stream.stream()>>>(
        reinterpret_cast<const float2*>(input.data_ptr<float>()),
        reinterpret_cast<float4*>(storage.data_ptr<float>()),
        compact_planes);

    C10_CUDA_KERNEL_LAUNCH_CHECK();

    // Overlapping boundary rows and planes contain only zeros.
    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_plane_expansion_v6",
    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)
```