 #include "xchainer/cuda/cuda_device.h"
 
+#include <algorithm>
 #include <cassert>
 #include <cstdint>
 #include <memory>

            Indexer<> y_indexer{y_.shape()};
            Indexer<> kernel_indexer{Shape{kernel_size_.begin(), kernel_size_.end()}};

            std::shared_ptr<void> stride = device.Allocate(stride_.size() * sizeof(int64_t));
            CheckCudaError(cudaMemcpy(stride.get(), stride_.data(), stride_.size() * sizeof(int64_t), cudaMemcpyHostToDevice));

            std::shared_ptr<void> pad = device.Allocate(pad_.size() * sizeof(int64_t));
            CheckCudaError(cudaMemcpy(pad.get(), pad_.data(), pad_.size() * sizeof(int64_t), cudaMemcpyHostToDevice));

            static const int kMaxBlockSize = CudaOccupancyMaxPotentialBlockSize(&MaxPoolDoubleBackwardKernel<T>).block_size;
            int64_t total_size = y_indexer.total_size();
            int64_t grid_size = (total_size + kMaxBlockSize - 1) / kMaxBlockSize;

        // Iterate over the kernel in the reverse order, since the resulting index should the be first match.
        for (auto it_kernel = kernel_indexer.It(kernel_indexer.total_size() - 1); it_kernel.raw_index() >= 0; --it_kernel) {
            for (int8_t i = 2; i < x_indexer.ndim(); ++i) {
                int64_t idx = it_y.index()[i] * stride.data[i - 2] - pad.data[i - 2] + it_kernel.index()[i - 2];
                idx = max(idx, int64_t{0});
                idx = min(idx, x_indexer.shape()[i] - 1);
                x_index.index()[i] = idx;
