```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 maxpool3d_128_cuda(torch::Tensor x);
"""

_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>
#include <cstdint>

template <bool CheckPadding>
__device__ __forceinline__ float pool_window(
    const float* __restrict__ input,
    const int depth,
    const int height,
    const int width
) {
    float best = -CUDART_INF_F;

    #pragma unroll
    for (int kd = 0; kd < 3; ++kd) {
        const int z = depth + 3 * kd;
        if (CheckPadding && z < 0) {
            continue;
        }

        #pragma unroll
        for (int kh = 0; kh < 3; ++kh) {
            const int y = height + 3 * kh;
            if (CheckPadding && y < 0) {
                continue;
            }

            #pragma unroll
            for (int kw = 0; kw < 3; ++kw) {
                const int w = width + 3 * kw;
                if (CheckPadding && w < 0) {
                    continue;
                }

                const float value = __ldg(input + z * 16384 + y * 128 + w);
                best = (value > best || isnan(value)) ? value : best;
            }
        }
    }
    return best;
}

__global__ void maxpool3d_128_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    const int64_t count
) {
    constexpr int64_t input_volume = 128LL * 128 * 128;
    constexpr int64_t output_volume = 62LL * 62 * 62;

    for (
        int64_t index = static_cast<int64_t>(blockIdx.x) * blockDim.x
                      + threadIdx.x;
        index < count;
        index += static_cast<int64_t>(blockDim.x) * gridDim.x
    ) {
        const int64_t plane = index / output_volume;
        const int position = static_cast<int>(index % output_volume);
        const int ow = position % 62;
        const int oh = (position / 62) % 62;
        const int od = position / (62 * 62);

        const float* plane_input = input + plane * input_volume;
        const int d = 2 * od - 1;
        const int h = 2 * oh - 1;
        const int w = 2 * ow - 1;

        // With this geometry, only the first output along each axis
        // intersects padding; the final window ends at input index 127.
        if (od > 0 && oh > 0 && ow > 0) {
            output[index] = pool_window<false>(plane_input, d, h, w);
        } else {
            output[index] = pool_window<true>(plane_input, d, h, w);
        }
    }
}

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 input");
    TORCH_CHECK(
        x.dim() == 5 &&
        x.size(2) == 128 &&
        x.size(3) == 128 &&
        x.size(4) == 128,
        "Expected input shape (N, C, 128, 128, 128)"
    );

    const c10::cuda::CUDAGuard device_guard(x.device());
    auto output = torch::empty(
        {x.size(0), x.size(1), 62, 62, 62}, x.options()
    );

    const int64_t count = output.numel();
    if (count == 0) {
        return output;
    }

    constexpr int threads = 256;
    const int64_t needed_blocks = (count + threads - 1) / threads;
    const int blocks = static_cast<int>(
        needed_blocks < 65535 ? needed_blocks : 65535
    );

    maxpool3d_128_kernel<<<
        blocks, threads, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        output.data_ptr<float>(),
        count
    );
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="maxpool3d_dilated_128_ext",
    cpp_sources=_cpp_source,
    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,
        )
        effective_stride = kernel_size if stride is None else stride
        self._specialized = (
            _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):
        if (
            self._specialized
            and x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 5
            and tuple(x.shape[2:]) == (128, 128, 128)
            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)
```