             cudnnPoolingBackward,
             *pool_desc,
             cuda_internal::GetCudnnCoefficientPtr<1>(dtype),
-            *y_desc,
-            internal::GetRawOffsetData(y_cont),
+            *out_desc,
+            internal::GetRawOffsetData(out_cont),
             *gout_desc,
             internal::GetRawOffsetData(gout_cont),
             *x_desc,

            *x_desc,
            internal::GetRawOffsetData(x_cont),
            cuda_internal::GetCudnnCoefficientPtr<0>(dtype),
            *y_desc,
            internal::GetRawOffsetData(y));

    return y;
}

Array PoolGrad(
        cudnnPoolingMode_t cudnn_pooling_mode,
        const Array& x,
        const Array& y,
        const Array& gout,
        StackVector<int64_t, kMaxNdim> kernel_size,
        StackVector<int64_t, kMaxNdim> stride,
        StackVector<int64_t, kMaxNdim> pad,
        const nonstd::optional<Array>& gx) {
    CHAINERX_ASSERT(y.shape() == gout.shape());
    CHAINERX_ASSERT(kernel_size.size() == static_cast<size_t>(x.ndim() - 2));
    CHAINERX_ASSERT(stride.size() == static_cast<size_t>(x.ndim() - 2));
    CHAINERX_ASSERT(pad.size() == static_cast<size_t>(x.ndim() - 2));

    CudaDevice& device = dynamic_cast<CudaDevice&>(ggx.device());
    CudaSetDeviceScope scope{device.index()};

    Array actual_ggout = EmptyLike(out, device);

    VisitFloatingPointDtype(actual_ggout.dtype(), [&](auto pt) {
        using T = typename decltype(pt)::type;

        IndexableArray<const T> ggx_iarray{ggx};
        IndexableArray<const T> x_iarray{x};
        IndexableArray<const T> out_iarray{out};
        IndexableArray<T> ggout_iarray{actual_ggout};

        Indexer<> x_indexer{x.shape()};
        Indexer<> out_indexer{out.shape()};
        Indexer<> kernel_indexer{Shape{kernel_size.begin(), kernel_size.end()}};

        static const int kMaxBlockSize = CudaOccupancyMaxPotentialBlockSize(&MaxPoolDoubleBackwardKernel<T>).block_size;
        int64_t total_size = out_indexer.total_size();
        int64_t grid_size = (total_size + kMaxBlockSize - 1) / kMaxBlockSize;
        int64_t block_size = std::min<int64_t>(total_size, kMaxBlockSize);

        MaxPoolDoubleBackwardKernel<<<grid_size, block_size>>>(
                ggx_iarray,
                x_iarray,
                out_iarray,
                ggout_iarray,
                x_indexer,
                out_indexer,
                kernel_indexer,
                CudaStackVector{stride},
                CudaStackVector{pad});
    });

    return actual_ggout;
}

// Pooling states are identical for most CUDA pooling ops.
