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


_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>

__global__ 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,
    int spatial_size)
{
    constexpr int TILE = 128;
    constexpr int REDUCTION = 27;

    __shared__ float input_tile[REDUCTION][TILE];
    __shared__ float weight_tile[64][REDUCTION];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int channel_group = tid >> 5;
    const int spatial_start = blockIdx.x * TILE;
    const int batch = blockIdx.y;

    const int64_t input_channel_stride =
        static_cast<int64_t>(height) * width * depth;
    const int64_t input_batch_offset =
        static_cast<int64_t>(batch) * 3 * input_channel_stride;
    const int output_row_stride = output_width * depth;

    for (int index = tid; index < REDUCTION * TILE; index += 256) {
        const int k = index / TILE;
        const int column = index % TILE;
        const int position = spatial_start + column;

        float value = 0.0f;
        if (position < spatial_size) {
            const int channel = k / 9;
            const int kh = (k % 9) / 3;
            const int kw = k % 3;

            const int oh = position / output_row_stride;
            const int remainder = position % output_row_stride;
            const int ow = remainder / depth;
            const int od = remainder % depth;

            const int64_t offset =
                input_batch_offset +
                channel * input_channel_stride +
                (static_cast<int64_t>(oh + kh) * width + ow + kw)
                    * depth + od;
            value = input[offset];
        }
        input_tile[k][column] = value;
    }

    for (int index = tid; index < 64 * REDUCTION; index += 256) {
        weight_tile[index / REDUCTION][index % REDUCTION] = 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 1
    for (int k = 0; k < REDUCTION; ++k) {
        const float x0 = input_tile[k][lane];
        const float x1 = input_tile[k][lane + 32];
        const float x2 = input_tile[k][lane + 64];
        const float x3 = input_tile[k][lane + 96];

        #pragma unroll
        for (int c = 0; c < 8; ++c) {
            const float w = weight_tile[channel_group + c * 8][k];
            accumulators[c][0] = fmaf(x0, w, accumulators[c][0]);
            accumulators[c][1] = fmaf(x1, w, accumulators[c][1]);
            accumulators[c][2] = fmaf(x2, w, accumulators[c][2]);
            accumulators[c][3] = fmaf(x3, w, accumulators[c][3]);
        }
    }

    #pragma unroll
    for (int c = 0; c < 8; ++c) {
        const int channel = channel_group + c * 8;
        const float b = bias == nullptr ? 0.0f : bias[channel];
        const int64_t output_offset =
            (static_cast<int64_t>(batch) * 64 + channel) * spatial_size;

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

torch::Tensor conv3d_specialized(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias)
{
    const c10::cuda::CUDAGuard device_guard(input.device());

    const int 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 int spatial_size = output_height * output_width * depth;

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

    if (output.numel() == 0) {
        return output;
    }

    const dim3 blocks((spatial_size + 127) / 128, batch);
    const cudaStream_t stream =
        c10::cuda::getCurrentCUDAStream(input.get_device()).stream();

    conv3d_3x3x1_kernel<<<blocks, 256, 0, stream>>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        height,
        width,
        depth,
        output_width,
        spatial_size);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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

torch::Tensor conv3d_specialized(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias);
"""

_conv3d_extension = load_inline(
    name="asymmetric_conv3d_tiled_fp32_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv3d_specialized"],
    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

        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 weight.device == x.device
            and weight.dtype == torch.float32
            and weight.is_contiguous()
            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 not torch.is_autocast_enabled()
            and (
                bias is None
                or (
                    bias.device == x.device
                    and bias.dtype == torch.float32
                    and bias.is_contiguous()
                )
            )
        ):
            return _conv3d_extension.conv3d_specialized(
                x,
                weight,
                weight if bias is None else bias,
                bias is not None,
            )

        return conv(x)
```