```python
import torch
import torch.nn as nn
from torch.nn.modules.utils import _triple
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>

__global__ void maxpool3d_128_kernel(
    const float* __restrict__ input,
    float* __restrict__ output)
{
    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;
    const int plane = blockIdx.z;

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

    const int iw = 2 * ow - 1;
    const int ih = 2 * oh - 1;
    const int id = 2 * od - 1;

    const int64_t input_base =
        static_cast<int64_t>(plane) * INPUT_SIZE * INPUT_SIZE * INPUT_SIZE;

    float maximum = -CUDART_INF_F;

    #pragma unroll
    for (int kd = 0; kd < 3; ++kd) {
        const int z = id + 3 * kd;
        if (z >= 0) {
            #pragma unroll
            for (int kh = 0; kh < 3; ++kh) {
                const int y = ih + 3 * kh;
                if (y >= 0) {
                    const int64_t row =
                        input_base + z * INPUT_SIZE * INPUT_SIZE
                        + y * INPUT_SIZE;

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

    const int64_t output_offset =
        ((static_cast<int64_t>(plane) * OUTPUT_SIZE + od)
          * OUTPUT_SIZE + oh) * OUTPUT_SIZE + ow;

    output[output_offset] = maximum;
}

torch::Tensor maxpool3d_128_cuda(torch::Tensor input)
{
    const c10::cuda::CUDAGuard device_guard(input.device());

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

    const unsigned int planes =
        static_cast<unsigned int>(input.size(0) * input.size(1));

    const dim3 grid(16, 62, planes);
    const auto stream =
        at::cuda::getCurrentCUDAStream(input.get_device());

    maxpool3d_128_kernel<<<grid, 256, 0, stream>>>(
        input.data_ptr<float>(),
        output.data_ptr<float>());

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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


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_specialized = (
            tuple(_triple(kernel_size)) == (3, 3, 3)
            and tuple(_triple(effective_stride)) == (2, 2, 2)
            and tuple(_triple(padding)) == (1, 1, 1)
            and tuple(_triple(dilation)) == (3, 3, 3)
            and not return_indices
            and not ceil_mode
        )

    def forward(self, x: torch.Tensor):
        if (
            self._use_specialized
            and x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and tuple(x.shape[2:]) == (128, 128, 128)
            and x.is_contiguous()
            and 0 < x.shape[0] * x.shape[1] <= 65535
            and not (torch.is_grad_enabled() and x.requires_grad)
        ):
            return _extension.maxpool3d_128_cuda(x)
        return self.maxpool(x)
```