             *x_desc,
             internal::GetRawOffsetData(x_cont),
             cuda_internal::GetCudnnCoefficientPtr<0>(dtype),
-            *y_desc,
-            internal::GetRawOffsetData(y));
-
-    return y;
+            *out_desc,
+            internal::GetRawOffsetData(actual_out));
+
+    return actual_out;
 }
 
 Array PoolGrad(
         cudnnPoolingMode_t cudnn_pooling_mode,
         const Array& x,
-        const Array& y,
+        const Array& out,
         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(out.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));


    CudaSetDeviceScope scope{device.index()};

    Array y = Empty(out_shape, dtype, device);
    Array x_cont = AsContiguousArray(x);

    cuda_internal::CudnnTensorDescriptor x_desc{x_cont};
    cuda_internal::CudnnTensorDescriptor y_desc{y};

    cuda_internal::CudnnPoolingDescriptor pool_desc{cudnn_pooling_mode, CUDNN_NOT_PROPAGATE_NAN, kernel_size, pad, stride};


    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.
