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

_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 = 256;
constexpr int CHANNEL_TILE = 8;
constexpr int PIXELS = 2;

__global__ void conv3x3x1_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ out,
    int H, int W, int D, int O)
{
    __shared__ float weights[27][CHANNEL_TILE];

    const int tid = threadIdx.x;
    const int oc0 = blockIdx.y * CHANNEL_TILE;
    const int n = blockIdx.z;

    if (tid < 27 * CHANNEL_TILE) {
        const int c = tid / 27;
        const int k = tid % 27;
        weights[k][c] = weight[(oc0 + c) * 27 + k];
    }
    __syncthreads();

    const int output_row = (W - 2) * D;
    const int output_plane = (H - 2) * output_row;
    const int input_plane = H * W * D;
    const int tile_start = blockIdx.x * THREADS * PIXELS + tid;

    int positions[PIXELS];
    int input_offsets[PIXELS];
    float sums[PIXELS][CHANNEL_TILE];

    #pragma unroll
    for (int p = 0; p < PIXELS; ++p) {
        positions[p] = tile_start + p * THREADS;
        const int h = positions[p] / output_row;
        const int remainder = positions[p] - h * output_row;
        input_offsets[p] = n * 3 * input_plane + h * W * D + remainder;

        #pragma unroll
        for (int c = 0; c < CHANNEL_TILE; ++c)
            sums[p][c] = 0.0f;
    }

    #pragma unroll
    for (int ic = 0; ic < 3; ++ic) {
        #pragma unroll
        for (int kh = 0; kh < 3; ++kh) {
            #pragma unroll
            for (int kw = 0; kw < 3; ++kw) {
                const int k = ic * 9 + kh * 3 + kw;
                const int offset = ic * input_plane + kh * W * D + kw * D;

                float values[PIXELS];
                #pragma unroll
                for (int p = 0; p < PIXELS; ++p) {
                    values[p] = positions[p] < output_plane
                        ? x[input_offsets[p] + offset] : 0.0f;
                }

                #pragma unroll
                for (int c = 0; c < CHANNEL_TILE; ++c) {
                    const float w = weights[k][c];
                    #pragma unroll
                    for (int p = 0; p < PIXELS; ++p)
                        sums[p][c] = fmaf(values[p], w, sums[p][c]);
                }
            }
        }
    }

    #pragma unroll
    for (int c = 0; c < CHANNEL_TILE; ++c) {
        const float b = bias ? bias[oc0 + c] : 0.0f;
        const int64_t base =
            (static_cast<int64_t>(n) * O + oc0 + c) * output_plane;

        #pragma unroll
        for (int p = 0; p < PIXELS; ++p) {
            if (positions[p] < output_plane)
                out[base + positions[p]] = sums[p][c] + b;
        }
    }
}

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

    const int N = x.size(0);
    const int H = x.size(2);
    const int W = x.size(3);
    const int D = x.size(4);
    const int O = weight.size(0);

    auto out = torch::empty({N, O, H - 2, W - 2, D}, x.options());
    if (out.numel() == 0)
        return out;

    const int plane = (H - 2) * (W - 2) * D;
    const dim3 grid(
        (plane + THREADS * PIXELS - 1) / (THREADS * PIXELS),
        O / CHANNEL_TILE,
        N);

    conv3x3x1_kernel<<<
        grid, THREADS, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias.has_value() ? bias->data_ptr<float>() : nullptr,
        out.data_ptr<float>(),
        H, W, D, O);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_extension = load_inline(
    name="asymmetric_conv3d_custom_v1",
    cpp_sources=r"""
#include <torch/extension.h>
torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias);
""",
    cuda_sources=_cuda_source,
    functions=["asymmetric_conv_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,
        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 conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            and (
                conv.bias is None
                or (
                    conv.bias.dtype == torch.float32
                    and conv.bias.device == x.device
                    and conv.bias.is_contiguous()
                )
            )
            and not torch.is_autocast_enabled()
            and x.shape[1] == conv.in_channels == 3
            and conv.out_channels % 8 == 0
            and conv.out_channels // 8 <= 65535
            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 x.shape[0] <= 65535
            and x.shape[2] >= 3
            and x.shape[3] >= 3
            and x.shape[4] > 0
            and x.numel() < 2**31 - 4096
        ):
            return _extension.asymmetric_conv_cuda(
                x, conv.weight, conv.bias
            )
        return conv(x)
```