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


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

template<bool Fixed128>
__global__ void maxpool3d_kernel(
    const float* __restrict__ x,
    float* __restrict__ out,
    int depth,
    int height,
    int width,
    int out_depth,
    int out_height,
    int out_width)
{
    const int H = Fixed128 ? 128 : height;
    const int W = Fixed128 ? 128 : width;
    const int OH = Fixed128 ? 62 : out_height;
    const int OW = Fixed128 ? 62 : out_width;

    const int spatial = blockIdx.x * blockDim.x + threadIdx.x;
    if (spatial >= OH * OW)
        return;

    const int od = blockIdx.y;
    const int nc = blockIdx.z;
    const int oh = spatial / OW;
    const int ow = spatial - oh * OW;

    const int d0 = od * 2 - 1;
    const int h0 = oh * 2 - 1;
    const int w0 = ow * 2 - 1;

    const float* input = x + static_cast<int64_t>(nc) * depth * H * W;
    float maximum = -CUDART_INF_F;

    const bool interior =
        d0 >= 0 && d0 + 6 < depth &&
        h0 >= 0 && h0 + 6 < H &&
        w0 >= 0 && w0 + 6 < W;

    if (interior) {
        const int base = (d0 * H + h0) * W + w0;
        #pragma unroll
        for (int kd = 0; kd < 3; ++kd) {
            #pragma unroll
            for (int kh = 0; kh < 3; ++kh) {
                #pragma unroll
                for (int kw = 0; kw < 3; ++kw) {
                    const float value =
                        input[base + kd * 3 * H * W + kh * 3 * W + kw * 3];
                    if (value > maximum || isnan(value))
                        maximum = value;
                }
            }
        }
    } else {
        #pragma unroll
        for (int kd = 0; kd < 3; ++kd) {
            const int d = d0 + kd * 3;
            if (static_cast<unsigned>(d) < static_cast<unsigned>(depth)) {
                #pragma unroll
                for (int kh = 0; kh < 3; ++kh) {
                    const int h = h0 + kh * 3;
                    if (static_cast<unsigned>(h) < static_cast<unsigned>(H)) {
                        #pragma unroll
                        for (int kw = 0; kw < 3; ++kw) {
                            const int w = w0 + kw * 3;
                            if (static_cast<unsigned>(w) < static_cast<unsigned>(W)) {
                                const float value = input[(d * H + h) * W + w];
                                if (value > maximum || isnan(value))
                                    maximum = value;
                            }
                        }
                    }
                }
            }
        }
    }

    out[(static_cast<int64_t>(nc) * out_depth + od) * OH * OW + spatial]
        = maximum;
}

torch::Tensor maxpool3d_cuda(torch::Tensor x) {
    TORCH_CHECK(x.is_cuda(), "Expected a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.dim() == 5 && x.is_contiguous(),
                "Expected contiguous NCDHW input");

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

    const int64_t N = x.size(0);
    const int64_t C = x.size(1);
    const int D = static_cast<int>(x.size(2));
    const int H = static_cast<int>(x.size(3));
    const int W = static_cast<int>(x.size(4));

    TORCH_CHECK(D >= 5 && H >= 5 && W >= 5, "Input is too small");
    TORCH_CHECK(x.numel() <= 2147483647LL, "Input is too large");

    const int OD = (D - 5) / 2 + 1;
    const int OH = (H - 5) / 2 + 1;
    const int OW = (W - 5) / 2 + 1;

    TORCH_CHECK(OD <= 65535 && N * C <= 65535,
                "Input exceeds supported launch dimensions");

    auto out = torch::empty({N, C, OD, OH, OW}, x.options());
    if (out.numel() == 0)
        return out;

    constexpr int threads = 256;
    const dim3 grid((OH * OW + threads - 1) / threads, OD, N * C);
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (H == 128 && W == 128) {
        maxpool3d_kernel<true><<<grid, threads, 0, stream>>>(
            x.data_ptr<float>(), out.data_ptr<float>(),
            D, H, W, OD, OH, OW);
    } else {
        maxpool3d_kernel<false><<<grid, threads, 0, stream>>>(
            x.data_ptr<float>(), out.data_ptr<float>(),
            D, H, W, OD, OH, OW);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_extension = load_inline(
    name="maxpool3d_k3_s2_p1_d3_ext",
    cpp_sources="torch::Tensor maxpool3d_cuda(torch::Tensor x);",
    cuda_sources=_cuda_source,
    functions=["maxpool3d_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


def _triple(value):
    if isinstance(value, int):
        return (value, value, value)
    return tuple(value)


class ModelNew(nn.Module):
    def __init__(
        self,
        kernel_size: int,
        stride: int = None,
        padding: int = 0,
        dilation: int = 1,
        return_indices: bool = False,
        ceil_mode: bool = False,
    ):
        super().__init__()
        self.maxpool = nn.MaxPool3d(
            kernel_size=kernel_size,
            stride=stride,
            padding=padding,
            dilation=dilation,
            return_indices=return_indices,
            ceil_mode=ceil_mode,
        )
        effective_stride = kernel_size if stride is None else stride
        self._use_custom = (
            _triple(kernel_size) == (3, 3, 3)
            and _triple(effective_stride) == (2, 2, 2)
            and _triple(padding) == (1, 1, 1)
            and _triple(dilation) == (3, 3, 3)
            and not return_indices
            and not ceil_mode
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if (
            self._use_custom
            and x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and x.is_contiguous()
            and not (torch.is_grad_enabled() and x.requires_grad)
            and x.shape[1] > 0
            and min(x.shape[2:]) >= 5
            and x.numel() <= 2147483647
            and x.shape[0] * x.shape[1] <= 65535
            and (x.shape[2] - 5) // 2 + 1 <= 65535
        ):
            return _extension.maxpool3d_cuda(x)
        return self.maxpool(x)
```