```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_forward_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>

namespace {

constexpr int OUTPUT_SIDE = 62;
constexpr int OUTPUT_VOLUME = OUTPUT_SIDE * OUTPUT_SIDE * OUTPUT_SIDE;
constexpr int INPUT_VOLUME = 64 * 64 * 64;
constexpr int TILE = 128;

__global__ __launch_bounds__(256)
void conv3d_3x64_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output)
{
    __shared__ float input_tile[27][TILE];
    __shared__ float weight_tile[64][28];

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

    const int load_position = spatial_base + (tid & 127);
    const bool valid_load = load_position < OUTPUT_VOLUME;

    const int oz = load_position / (OUTPUT_SIDE * OUTPUT_SIDE);
    const int remainder = load_position % (OUTPUT_SIDE * OUTPUT_SIDE);
    const int oy = remainder / OUTPUT_SIDE;
    const int ox = remainder % OUTPUT_SIDE;

    const int input_base =
        batch * 3 * INPUT_VOLUME + oz * 64 * 64 + oy * 64 + ox;

    float accum[8][4];

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

    #pragma unroll 1
    for (int ic = 0; ic < 3; ++ic) {
        #pragma unroll
        for (int k = tid >> 7; k < 27; k += 2) {
            const int kz = k / 9;
            const int ky = (k / 3) % 3;
            const int kx = k % 3;

            float value = 0.0f;
            if (valid_load) {
                value = input[
                    input_base + ic * INPUT_VOLUME
                    + kz * 64 * 64 + ky * 64 + kx
                ];
            }
            input_tile[k][tid & 127] = value;
        }

        #pragma unroll
        for (int index = tid; index < 64 * 27; index += 256) {
            const int oc = index / 27;
            const int k = index % 27;
            weight_tile[oc][k] = weight[oc * 81 + ic * 27 + k];
        }

        __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];
            const float x2 = input_tile[k][lane + 64];
            const float x3 = input_tile[k][lane + 96];

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

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < 8; ++i) {
        const int oc = channel_group + i * 8;
        const float b = bias ? bias[oc] : 0.0f;
        const int64_t output_base =
            (static_cast<int64_t>(batch) * 64 + oc) * OUTPUT_VOLUME;

        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            const int position = spatial_base + lane + j * 32;
            if (position < OUTPUT_VOLUME) {
                output[output_base + position] = accum[i][j] + b;
            }
        }
    }
}

} // namespace

torch::Tensor conv3d_forward_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.dim() == 5 &&
                input.size(1) == 3 &&
                input.size(2) == 64 &&
                input.size(3) == 64 &&
                input.size(4) == 64,
                "Unsupported input shape");
    TORCH_CHECK(weight.dim() == 5 &&
                weight.size(0) == 64 &&
                weight.size(1) == 3 &&
                weight.size(2) == 3 &&
                weight.size(3) == 3 &&
                weight.size(4) == 3,
                "Unsupported weight shape");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must be on the same device");

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

    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.numel() == 64,
                    "Unsupported bias");
        bias_ptr = b.data_ptr<float>();
    }

    auto output = torch::empty(
        {input.size(0), 64, OUTPUT_SIDE, OUTPUT_SIDE, OUTPUT_SIDE},
        input.options());

    if (input.size(0) == 0) {
        return output;
    }

    const dim3 grid(
        (OUTPUT_VOLUME + TILE - 1) / TILE,
        static_cast<unsigned int>(input.size(0)));

    conv3d_3x64_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>());

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="custom_conv3d_c3_o64_k3_tile128",
    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_output = grad_output.contiguous()
        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
        use_custom = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and tuple(x.shape[1:]) == (3, 64, 64, 64)
            and x.shape[0] <= 65535
            and x.is_contiguous()
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            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 not torch.is_autocast_enabled()
        )

        if not use_custom:
            return conv(x)

        if torch.is_grad_enabled() and (
            x.requires_grad
            or conv.weight.requires_grad
            or (conv.bias is not None and conv.bias.requires_grad)
        ):
            return _Conv3dFunction.apply(x, conv.weight, conv.bias)

        return _extension.conv3d_forward_cuda(x, conv.weight, conv.bias)
```