```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor asymmetric_conv3d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias);
"""

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

constexpr int THREADS = 128;
constexpr int CHANNEL_TILE = 8;

__global__ void asymmetric_conv3d_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int height,
    int width,
    int depth,
    int output_height,
    int output_width,
    int64_t spatial_size)
{
    __shared__ float coefficients[27][CHANNEL_TILE];

    const int channel_start = blockIdx.y * CHANNEL_TILE;
    const int batch = blockIdx.z;

    for (int i = threadIdx.x; i < 27 * CHANNEL_TILE; i += THREADS) {
        const int channel = i / 27;
        const int k = i % 27;
        coefficients[k][channel] =
            weight[(channel_start + channel) * 27 + k];
    }
    __syncthreads();

    const int64_t position =
        static_cast<int64_t>(blockIdx.x) * THREADS + threadIdx.x;
    if (position >= spatial_size)
        return;

    const int z = position % depth;
    const int64_t plane_position = position / depth;
    const int w = plane_position % output_width;
    const int h = plane_position / output_width;

    const int64_t input_plane =
        static_cast<int64_t>(height) * width * depth;
    const int64_t input_base =
        static_cast<int64_t>(batch) * 3 * input_plane +
        (static_cast<int64_t>(h) * width + w) * depth + z;

    float accumulator[CHANNEL_TILE];

    #pragma unroll
    for (int o = 0; o < CHANNEL_TILE; ++o)
        accumulator[o] = 0.0f;

    #pragma unroll
    for (int c = 0; c < 3; ++c) {
        #pragma unroll
        for (int kh = 0; kh < 3; ++kh) {
            #pragma unroll
            for (int kw = 0; kw < 3; ++kw) {
                const int k = c * 9 + kh * 3 + kw;
                const int64_t offset =
                    input_base + c * input_plane +
                    (static_cast<int64_t>(kh) * width + kw) * depth;
                const float value = __ldg(x + offset);

                #pragma unroll
                for (int o = 0; o < CHANNEL_TILE; ++o) {
                    accumulator[o] =
                        fmaf(value, coefficients[k][o], accumulator[o]);
                }
            }
        }
    }

    #pragma unroll
    for (int o = 0; o < CHANNEL_TILE; ++o) {
        float value = accumulator[o];
        if (bias != nullptr)
            value += bias[channel_start + o];

        const int64_t output_offset =
            (static_cast<int64_t>(batch) * 64 + channel_start + o) *
                spatial_size +
            position;
        output[output_offset] = value;
    }
}

torch::Tensor asymmetric_conv3d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    const c10::cuda::CUDAGuard device_guard(x.device());

    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat);
    TORCH_CHECK(x.is_contiguous() && x.dim() == 5 && x.size(1) == 3);
    TORCH_CHECK(weight.is_cuda() && weight.device() == x.device());
    TORCH_CHECK(weight.scalar_type() == at::kFloat && weight.is_contiguous());
    TORCH_CHECK(
        weight.dim() == 5 &&
        weight.size(0) == 64 && weight.size(1) == 3 &&
        weight.size(2) == 3 && weight.size(3) == 3 &&
        weight.size(4) == 1);
    TORCH_CHECK(x.size(2) >= 3 && x.size(3) >= 3 && x.size(4) > 0);

    const int64_t batch = x.size(0);
    const int height = x.size(2);
    const int width = x.size(3);
    const int depth = x.size(4);
    const int output_height = height - 2;
    const int output_width = width - 2;
    const int64_t spatial_size =
        static_cast<int64_t>(output_height) * output_width * depth;

    auto output = torch::empty(
        {batch, 64, output_height, output_width, depth}, x.options());

    if (batch == 0)
        return output;

    const float* bias_pointer = nullptr;
    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(
            b.device() == x.device() &&
            b.scalar_type() == at::kFloat &&
            b.is_contiguous() && b.numel() == 64);
        bias_pointer = b.data_ptr<float>();
    }

    const dim3 grid(
        static_cast<unsigned int>((spatial_size + THREADS - 1) / THREADS),
        64 / CHANNEL_TILE,
        static_cast<unsigned int>(batch));

    asymmetric_conv3d_kernel<<<
        grid, THREADS, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_pointer,
        output.data_ptr<float>(),
        height,
        width,
        depth,
        output_height,
        output_width,
        spatial_size);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="asymmetric_conv3d_direct_v1",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["asymmetric_conv3d_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,
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv3d = nn.Conv3d(
            in_channels,
            out_channels,
            (kernel_size, kernel_size, 1),
            stride=stride,
            padding=padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv3d
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.is_contiguous()
            and x.shape[1] == 3
            and x.shape[2] >= 3
            and x.shape[3] >= 3
            and x.shape[4] > 0
            and x.shape[0] <= 65535
            and conv.in_channels == 3
            and conv.out_channels == 64
            and conv.kernel_size == (3, 3, 1)
            and conv.stride == (1, 1, 1)
            and conv.padding == (0, 0, 0)
            and conv.dilation == (1, 1, 1)
            and conv.groups == 1
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            and not torch.is_autocast_enabled()
        ):
            return _get_extension().asymmetric_conv3d_cuda(
                x, conv.weight, conv.bias
            )
        return conv(x)
```