```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_forward_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.h>
#include <cuda_runtime.h>

__global__ __launch_bounds__(256)
void conv3d_3x64_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int D, int H, int W,
    int OH, int OW, int spatial)
{
    __shared__ float weights[64][81];
    __shared__ float patch[9][128];

    const int tid = threadIdx.x;
    const int lane_x = tid & 15;
    const int lane_c = tid >> 4;
    const int batch = blockIdx.y;
    const int first = blockIdx.x * 128;

    for (int i = tid; i < 64 * 81; i += 256)
        weights[i / 81][i % 81] = weight[i];

    float accum[4][8];
    #pragma unroll
    for (int c = 0; c < 4; ++c)
        #pragma unroll
        for (int s = 0; s < 8; ++s)
            accum[c][s] = 0.0f;

    const int load_column = tid & 127;
    const int load_group = tid >> 7;
    const int position = first + load_column;
    const bool valid = position < spatial;

    const int od = position / (OH * OW);
    const int remainder = position - od * OH * OW;
    const int oh = remainder / OW;
    const int ow = remainder - oh * OW;

    const int64_t volume = int64_t(D) * H * W;
    const int64_t input_base =
        int64_t(batch) * 3 * volume + int64_t(od) * H * W + oh * W + ow;

    __syncthreads();

    #pragma unroll 1
    for (int stage = 0; stage < 9; ++stage) {
        const int channel = stage / 3;
        const int kd = stage % 3;
        const int64_t base =
            input_base + channel * volume + int64_t(kd) * H * W;

        #pragma unroll
        for (int k = load_group; k < 9; k += 2) {
            const int kh = k / 3;
            const int kw = k % 3;
            patch[k][load_column] =
                valid ? x[base + kh * W + kw] : 0.0f;
        }
        __syncthreads();

        #pragma unroll
        for (int k = 0; k < 9; ++k) {
            float values[8];
            #pragma unroll
            for (int s = 0; s < 8; ++s)
                values[s] = patch[k][lane_x + 16 * s];

            #pragma unroll
            for (int c = 0; c < 4; ++c) {
                const float w = weights[lane_c + 16 * c][stage * 9 + k];
                #pragma unroll
                for (int s = 0; s < 8; ++s)
                    accum[c][s] = fmaf(w, values[s], accum[c][s]);
            }
        }
        __syncthreads();
    }

    #pragma unroll
    for (int c = 0; c < 4; ++c) {
        const int channel = lane_c + 16 * c;
        const float b = bias ? bias[channel] : 0.0f;
        const int64_t output_base =
            (int64_t(batch) * 64 + channel) * spatial;

        #pragma unroll
        for (int s = 0; s < 8; ++s) {
            const int p = first + lane_x + 16 * s;
            if (p < spatial)
                output[output_base + p] = accum[c][s] + b;
        }
    }
}

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

    const int D = x.size(2);
    const int H = x.size(3);
    const int W = x.size(4);
    const int OD = D - 2;
    const int OH = H - 2;
    const int OW = W - 2;
    const int spatial = OD * OH * OW;

    auto output = torch::empty(
        {x.size(0), 64, OD, OH, OW}, x.options());

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

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

    const dim3 grid((spatial + 127) / 128, x.size(0));
    conv3d_3x64_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        D, H, W, OH, OW, spatial);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="conv3d_tiled_3x64_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["conv3d_forward_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _Conv3dFunction(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias):
        ctx.save_for_backward(x, weight)
        ctx.has_bias = bias is not None
        return _extension.conv3d_forward_cuda(x, weight, bias)

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        grad_x = grad_weight = grad_bias = None

        if ctx.needs_input_grad[0]:
            grad_x = torch.nn.grad.conv3d_input(
                x.shape, weight, grad_output,
                stride=1, padding=0, dilation=1, groups=1,
            )
        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv3d_weight(
                x, weight.shape, grad_output,
                stride=1, padding=0, dilation=1, groups=1,
            )
        if ctx.has_bias and ctx.needs_input_grad[2]:
            grad_bias = grad_output.sum(dim=(0, 2, 3, 4))

        return grad_x, grad_weight, grad_bias


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
        supported = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.is_contiguous()
            and x.shape[1] == 3
            and 0 < x.shape[0] <= 65535
            and min(x.shape[2:]) >= 3
            and x.shape[2] * x.shape[3] * x.shape[4] < 2**31
            and conv.in_channels == 3
            and conv.out_channels == 64
            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 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()
        )

        if supported:
            return _Conv3dFunction.apply(x, conv.weight, conv.bias)
        return conv(x)
```