     CudaSetDeviceScope scope{device.index()};
 
     Array actual_gx = EmptyLike(x, device);
-    Array y_cont = AsContiguousArray(y);
+    Array out_cont = AsContiguousArray(out);
     Array gout_cont = AsContiguousArray(gout);
     Array x_cont = AsContiguousArray(x);
 
-    cuda_internal::CudnnTensorDescriptor y_desc{y_cont};
+    cuda_internal::CudnnTensorDescriptor out_desc{out_cont};
     cuda_internal::CudnnTensorDescriptor gout_desc{gout_cont};
     cuda_internal::CudnnTensorDescriptor x_desc{x_cont};
     cuda_internal::CudnnTensorDescriptor gx_desc{actual_gx};

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

    // TODO(hvy): Rename to `actual_ggout`.
    Array ggy = EmptyLike(y, device);

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

        IndexableArray<const T> ggx_iarray{ggx};
        IndexableArray<const T> x_iarray{x};
        IndexableArray<const T> y_iarray{y};
        IndexableArray<T> ggy_iarray{ggy};

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

        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;
        int64_t block_size = std::min<int64_t>(total_size, kMaxBlockSize);

        MaxPoolDoubleBackwardKernel<<<grid_size, block_size>>>(
                ggx_iarray,
                x_iarray,
                y_iarray,
                ggy_iarray,
                x_indexer,
                y_indexer,
                kernel_indexer,
                CudaStackVector{stride},
                CudaStackVector{pad});
    });

    return ggy;
}

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


Array MaxPoolGradGrad(
        const Array& x,
        const Array& out,
        const Array& ggx,
        StackVector<int64_t, kMaxNdim> kernel_size,
        StackVector<int64_t, kMaxNdim> stride,
