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

__global__ void maxpool3d_128_kernel(
    const float* __restrict__ x,
    float* __restrict__ out
) {
    constexpr int INPUT_SIZE = 128;
    constexpr int OUTPUT_SIZE = 62;

    const int ow = threadIdx.x & 63;
    const int oh = blockIdx.x * 4 + (threadIdx.x >> 6);
    const int od = blockIdx.y;

    if (ow >= OUTPUT_SIZE || oh >= OUTPUT_SIZE)
        return;

    const int iw = ow * 2 - 1;
    const int ih = oh * 2 - 1;
    const int id = od * 2 - 1;
    const int64_t plane = static_cast<int64_t>(blockIdx.z);

    const float* input = x + plane * INPUT_SIZE * INPUT_SIZE * INPUT_SIZE;
    float maximum = -CUDART_INF_F;

    #pragma unroll
    for (int kz = 0; kz < 3; ++kz) {
        const int z = id + kz * 3;
        if (z < 0)
            continue;

        #pragma unroll
        for (int ky = 0; ky < 3; ++ky) {
            const int y = ih + ky * 3;
            if (y < 0)
                continue;

            const int row = (z * INPUT_SIZE + y) * INPUT_SIZE;

            #pragma unroll
            for (int kx = 0; kx < 3; ++kx) {
                const int w = iw + kx * 3;
                if (w >= 0) {
                    const float value = input[row + w];
                    if (value > maximum || isnan(value))
                        maximum = value;
                }
            }
        }
    }

    const int64_t offset =
        ((plane * OUTPUT_SIZE + od) * OUTPUT_SIZE + oh) * OUTPUT_SIZE + ow;
    out[offset] = maximum;
}

torch::Tensor maxpool3d_128_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.is_contiguous(), "Expected contiguous NCDHW storage");
    TORCH_CHECK(
        x.dim() == 5 &&
        x.size(2) == 128 && x.size(3) == 128 && x.size(4) == 128,
        "Expected spatial dimensions (128, 128, 128)"
    );

    const c10::cuda::CUDAGuard device_guard(x.device());
    const int64_t planes = x.size(0) * x.size(1);
    TORCH_CHECK(planes <= 65535, "Too many batch-channel planes");

    auto out = torch::empty(
        {x.size(0), x.size(1), 62, 62, 62}, x.options()
    );
    if (planes == 0)
        return out;

    const dim3 grid(16, 62, static_cast<unsigned int>(planes));
    maxpool3d_128_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(x.data_ptr<float>(), out.data_ptr<float>());

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_extension = load_inline(
    name="maxpool3d_dilated_128_custom",
    cpp_sources=(
        "torch::Tensor maxpool3d_128_cuda(torch::Tensor x);"
    ),
    cuda_sources=_cuda_source,
    functions=["maxpool3d_128_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,
        )
        self._use_custom = (
            _triple(kernel_size) == (3, 3, 3)
            and _triple(kernel_size if stride is None else 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):
        if (
            self._use_custom
            and x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and tuple(x.shape[2:]) == (128, 128, 128)
            and x.shape[0] * x.shape[1] <= 65535
            and x.shape[1] > 0
            and x.is_contiguous()
            and not (torch.is_grad_enabled() and x.requires_grad)
        ):
            return _extension.maxpool3d_128_cuda(x)
        return self.maxpool(x)
```