             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;

        Indexer<> x_indexer,
        Indexer<> y_indexer,
        Indexer<> kernel_indexer,
        int64_t* stride,
        int64_t* pad,
        NdimIndex x_index) {
    for (auto it_y = y_indexer.It(blockIdx.x * blockDim.x + threadIdx.x, blockDim.x * gridDim.x); it_y; ++it_y) {
        x_index.index()[0] = it_y.index()[0];  // batch.

                    x_indexer,
                    y_indexer,
                    kernel_indexer,
                    CudaStackVector{stride_},
                    CudaStackVector{pad_},
                    NdimIndex{x_iarray.ndim()});
        });

