```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 conv3d_3x3x1_cuda(
    torch::Tensor input,
    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>

__global__ __launch_bounds__(256)
void conv3d_3x3x1_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int height,
    int width,
    int depth,
    int output_width,
    int64_t output_spatial)
{
    __shared__ float input_tile[27][128];
    __shared__ float weight_tile[27][64];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int channel_base = (tid >> 5) * 8;
    const int64_t tile_start = static_cast<int64_t>(blockIdx.x) * 128;
    const int64_t batch = blockIdx.y;

    const int64_t load_position = tile_start + (tid & 127);
    const int z = load_position % depth;
    const int64_t row = load_position / depth;
    const int y = row % output_width;
    const int x = row / output_width;

    #pragma unroll
    for (int k = tid >> 7; k < 27; k += 2) {
        float value = 0.0f;
        if (load_position < output_spatial) {
            const int channel = k / 9;
            const int kernel_x = (k % 9) / 3;
            const int kernel_y = k % 3;

            const int64_t offset =
                (((batch * 3 + channel) * height + x + kernel_x)
                    * width + y + kernel_y) * depth + z;

            value = input[offset];
        }
        input_tile[k][tid & 127] = value;
    }

    for (int index = tid; index < 64 * 27; index += 256) {
        weight_tile[index % 27][index / 27] = weight[index];
    }

    __syncthreads();

    float accumulators[8][4];

    #pragma unroll
    for (int c = 0; c < 8; ++c) {
        #pragma unroll
        for (int p = 0; p < 4; ++p) {
            accumulators[c][p] = 0.0f;
        }
    }

    #pragma unroll
    for (int k = 0; k < 27; ++k) {
        const float v0 = input_tile[k][lane];
        const float v1 = input_tile[k][lane + 32];
        const float v2 = input_tile[k][lane + 64];
        const float v3 = input_tile[k][lane + 96];

        #pragma unroll
        for (int c = 0; c < 8; ++c) {
            const float w = weight_tile[k][channel_base + c];
            accumulators[c][0] = fmaf(w, v0, accumulators[c][0]);
            accumulators[c][1] = fmaf(w, v1, accumulators[c][1]);
            accumulators[c][2] = fmaf(w, v2, accumulators[c][2]);
            accumulators[c][3] = fmaf(w, v3, accumulators[c][3]);
        }
    }

    #pragma unroll
    for (int c = 0; c < 8; ++c) {
        const int channel = channel_base + c;
        const float b = bias == nullptr ? 0.0f : bias[channel];
        const int64_t output_base =
            (batch * 64 + channel) * output_spatial;

        #pragma unroll
        for (int p = 0; p < 4; ++p) {
            const int64_t position = tile_start + lane + p * 32;
            if (position < output_spatial) {
                output[output_base + position] = accumulators[c][p] + b;
            }
        }
    }
}

torch::Tensor conv3d_3x3x1_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Input and weight must be float32");
    TORCH_CHECK(input.is_contiguous() && weight.is_contiguous(),
                "Input and weight must be contiguous");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must be on the same device");
    TORCH_CHECK(input.dim() == 5 && input.size(1) == 3,
                "Expected input shape (N, 3, H, W, D)");
    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,
        "Expected weight shape (64, 3, 3, 3, 1)");
    TORCH_CHECK(input.size(2) >= 3 && input.size(3) >= 3 &&
                input.size(4) > 0 && input.size(0) <= 65535,
                "Unsupported input dimensions");

    const c10::cuda::CUDAGuard device_guard(input.device());

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

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

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(
            b.device() == input.device() &&
            b.scalar_type() == at::kFloat &&
            b.is_contiguous() && b.dim() == 1 && b.numel() == 64,
            "Expected a contiguous float32 bias of shape (64,)");
        bias_ptr = b.data_ptr<float>();
    }

    if (batch == 0) {
        return output;
    }

    const dim3 grid(
        static_cast<unsigned int>((spatial + 127) / 128),
        static_cast<unsigned int>(batch));

    conv3d_3x3x1_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        height,
        width,
        depth,
        output_width,
        spatial);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_conv3d_extension = load_inline(
    name="conv3d_asymmetric_tiled_fp32",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv3d_3x3x1_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 x.shape[0] <= 65535
            and x.shape[1] == 3
            and x.shape[2] >= 3
            and x.shape[3] >= 3
            and x.shape[4] > 0
            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 _conv3d_extension.conv3d_3x3x1_cuda(
                x, conv.weight, conv.bias
            )
        return conv(x)
```