```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_tiled_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__ void conv3d_tiled_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels_out,
    int depth,
    int height,
    int width,
    int out_height,
    int out_width,
    int spatial_out)
{
    // Each block computes 64 output channels at 64 spatial positions.
    __shared__ float input_tile[27][64];
    __shared__ float weight_tile[64][28];

    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int warp = tid >> 5;
    const int spatial_base = blockIdx.x * 64;
    const int channel_base = blockIdx.y * 64;
    const int batch = blockIdx.z;

    const int p0 = spatial_base + lane;
    const int p1 = p0 + 32;

    float accum0[8];
    float accum1[8];

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        accum0[i] = 0.0f;
        accum1[i] = 0.0f;
    }

    const int64_t input_volume =
        static_cast<int64_t>(depth) * height * width;
    const float* batch_input = input + batch * 3LL * input_volume;

    #pragma unroll
    for (int ic = 0; ic < 3; ++ic) {
        for (int index = tid; index < 27 * 64; index += 256) {
            const int k = index / 64;
            const int position = spatial_base + index % 64;

            float value = 0.0f;
            if (position < spatial_out) {
                const int ow = position % out_width;
                const int oh = (position / out_width) % out_height;
                const int od = position / (out_width * out_height);

                const int kd = k / 9;
                const int kh = (k / 3) % 3;
                const int kw = k % 3;

                const int64_t offset =
                    ic * input_volume +
                    (static_cast<int64_t>(od + kd) * height + oh + kh)
                        * width + ow + kw;
                value = batch_input[offset];
            }
            input_tile[k][index % 64] = value;
        }

        for (int index = tid; index < 64 * 27; index += 256) {
            const int local_channel = index / 27;
            const int k = index % 27;
            const int channel = channel_base + local_channel;

            weight_tile[local_channel][k] =
                channel < channels_out
                    ? weight[static_cast<int64_t>(channel) * 81 + ic * 27 + k]
                    : 0.0f;
        }

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < 27; ++k) {
            const float x0 = input_tile[k][lane];
            const float x1 = input_tile[k][lane + 32];

            #pragma unroll
            for (int i = 0; i < 8; ++i) {
                const float w = weight_tile[warp * 8 + i][k];
                accum0[i] = fmaf(x0, w, accum0[i]);
                accum1[i] = fmaf(x1, w, accum1[i]);
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        const int channel = channel_base + warp * 8 + i;
        if (channel < channels_out) {
            const float b = bias == nullptr ? 0.0f : bias[channel];
            const int64_t offset =
                (static_cast<int64_t>(batch) * channels_out + channel)
                * spatial_out;

            if (p0 < spatial_out)
                output[offset + p0] = accum0[i] + b;
            if (p1 < spatial_out)
                output[offset + p1] = accum1[i] + b;
        }
    }
}

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

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

    const int out_depth = depth - 2;
    const int out_height = height - 2;
    const int out_width = width - 2;
    const int spatial_out = out_depth * out_height * out_width;

    auto output = torch::empty(
        {batch, channels_out, out_depth, out_height, out_width},
        input.options());

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

    const float* bias_ptr =
        bias.has_value() ? bias.value().data_ptr<float>() : nullptr;

    const dim3 grid(
        (spatial_out + 63) / 64,
        (channels_out + 63) / 64,
        batch);

    conv3d_tiled_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        channels_out,
        depth,
        height,
        width,
        out_height,
        out_width,
        spatial_out);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_conv3d_extension = load_inline(
    name="conv3d_fp32_tiled_64x64",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv3d_tiled_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, kernel_size),
            stride=stride,
            padding=padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )

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