```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 conv3d_small_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 CHANNEL_TILE = 8;
constexpr int THREADS = 256;

__global__ void conv3d_small_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ out,
    int channels_out,
    int height,
    int width,
    int depth,
    int64_t spatial_out)
{
    __shared__ float weights[27][CHANNEL_TILE];

    const int oc0 = blockIdx.y * CHANNEL_TILE;

    for (int i = threadIdx.x; i < 27 * CHANNEL_TILE; i += THREADS) {
        const int k = i / CHANNEL_TILE;
        const int oc = oc0 + i % CHANNEL_TILE;
        weights[k][i % CHANNEL_TILE] =
            oc < channels_out ? weight[oc * 27 + k] : 0.0f;
    }
    __syncthreads();

    const int64_t p =
        static_cast<int64_t>(blockIdx.x) * THREADS + threadIdx.x;

    if (p >= spatial_out)
        return;

    const int output_width = width - 2;
    const int d = p % depth;
    const int64_t hw = p / depth;
    const int w = hw % output_width;
    const int h = hw / output_width;

    const int64_t row_stride = static_cast<int64_t>(width) * depth;
    const int64_t channel_stride =
        static_cast<int64_t>(height) * row_stride;

    const float* input = x
        + static_cast<int64_t>(blockIdx.z) * 3 * channel_stride
        + static_cast<int64_t>(h) * row_stride
        + static_cast<int64_t>(w) * depth + d;

    float accum[CHANNEL_TILE];

    #pragma unroll
    for (int o = 0; o < CHANNEL_TILE; ++o)
        accum[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 float value = __ldg(
                    input + c * channel_stride
                    + kh * row_stride + kw * depth);
                const int k = c * 9 + kh * 3 + kw;

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

    #pragma unroll
    for (int o = 0; o < CHANNEL_TILE; ++o) {
        const int oc = oc0 + o;
        if (oc < channels_out) {
            const int64_t offset =
                (static_cast<int64_t>(blockIdx.z) * channels_out + oc)
                * spatial_out + p;
            out[offset] = accum[o] + (bias ? bias[oc] : 0.0f);
        }
    }
}

torch::Tensor conv3d_small_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.dim() == 5 && x.size(1) == 3,
                "Expected NCDHW input with three channels");
    TORCH_CHECK(weight.dim() == 5 &&
                weight.size(1) == 3 &&
                weight.size(2) == 3 &&
                weight.size(3) == 3 &&
                weight.size(4) == 1,
                "Expected weight shape [Cout, 3, 3, 3, 1]");
    TORCH_CHECK(x.device() == weight.device(),
                "Input and weight devices must match");
    TORCH_CHECK(x.size(2) >= 3 && x.size(3) >= 3 && x.size(4) > 0,
                "Invalid input dimensions");

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t batch = x.size(0);
    const int channels_out = weight.size(0);
    const int height = x.size(2);
    const int width = x.size(3);
    const int depth = x.size(4);

    const float* bias_ptr = 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() == channels_out,
                    "Invalid bias tensor");
        bias_ptr = b.data_ptr<float>();
    }

    auto out = torch::empty(
        {batch, channels_out, height - 2, width - 2, depth},
        x.options());

    if (out.numel() == 0)
        return out;

    const int64_t spatial_out =
        static_cast<int64_t>(height - 2) * (width - 2) * depth;

    const dim3 grid(
        static_cast<unsigned>((spatial_out + THREADS - 1) / THREADS),
        (channels_out + CHANNEL_TILE - 1) / CHANNEL_TILE,
        static_cast<unsigned>(batch));

    conv3d_small_kernel<<<
        grid, THREADS, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        out.data_ptr<float>(),
        channels_out,
        height,
        width,
        depth,
        spatial_out);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_extension = load_inline(
    name="conv3d_asymmetric_small_channels_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv3d_small_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
        weight = conv.weight
        bias = conv.bias

        use_custom = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.is_contiguous()
            and weight.is_contiguous()
            and weight.dtype == x.dtype
            and weight.device == x.device
            and conv.in_channels == 3
            and x.shape[1] == 3
            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[2] >= 3
            and x.shape[3] >= 3
            and x.shape[4] > 0
            and x.shape[0] <= 65535
            and conv.out_channels <= 65535 * 8
            and not torch.is_autocast_enabled()
            and (
                bias is None
                or (
                    bias.dtype == x.dtype
                    and bias.device == x.device
                    and bias.is_contiguous()
                )
            )
        )

        if use_custom:
            return _extension.conv3d_small_cuda(x, weight, bias)

        return F.conv3d(
            x,
            weight,
            bias,
            conv.stride,
            conv.padding,
            conv.dilation,
            conv.groups,
        )
```