deque: Add at() and shrink_to_fit()

src/stdgpu/deque.cuh

@@ -97,6 +97,24 @@ class deque
          */
         deque() = default;
 
+        /**
+         * \brief Reads the value at position n
+         * \param[in] n The position
+         * \return The value at this position
+         * \pre 0 <= n < size()
+         */
+        STDGPU_DEVICE_ONLY reference
+        at(const index_type n);
+
+        /**
+         * \brief Reads the value at position n
+         * \param[in] n The position
+         * \return The value at this position
+         * \pre 0 <= n < size()
+         */
+        STDGPU_DEVICE_ONLY const_reference
+        at(const index_type n) const;
+
         /**
          * \brief Reads the value at position n
          * \param[in] n The position
@@ -226,6 +244,13 @@ class deque
         STDGPU_HOST_DEVICE index_t
         capacity() const;
 
+        /**
+         * \brief Requests to shrink the capacity to the current size
+         * \note This is non-binding and may not have any effect
+         */
+        void
+        shrink_to_fit();
+
         /**
          * \brief Returns a pointer to the underlying data
          * \return The underlying array

src/stdgpu/impl/deque_detail.cuh

@@ -68,15 +68,15 @@ deque<T>::destroyDeviceObject(deque<T>& device_object)
 
 template <typename T>
 inline STDGPU_DEVICE_ONLY typename deque<T>::reference
-deque<T>::operator[](const deque<T>::index_type n)
+deque<T>::at(const deque<T>::index_type n)
 {
-    return const_cast<reference>(static_cast<const deque<T>*>(this)->operator[](n));
+    return const_cast<reference>(static_cast<const deque<T>*>(this)->at(n));
 }
 
 
 template <typename T>
 inline STDGPU_DEVICE_ONLY typename deque<T>::const_reference
-deque<T>::operator[](const deque<T>::index_type n) const
+deque<T>::at(const deque<T>::index_type n) const
 {
     STDGPU_EXPECTS(0 <= n);
     STDGPU_EXPECTS(n < size());
@@ -90,6 +90,22 @@ deque<T>::operator[](const deque<T>::index_type n) const
 }
 
 
+template <typename T>
+inline STDGPU_DEVICE_ONLY typename deque<T>::reference
+deque<T>::operator[](const deque<T>::index_type n)
+{
+    return at(n);
+}
+
+
+template <typename T>
+inline STDGPU_DEVICE_ONLY typename deque<T>::const_reference
+deque<T>::operator[](const deque<T>::index_type n) const
+{
+    return at(n);
+}
+
+
 template <typename T>
 inline STDGPU_DEVICE_ONLY typename deque<T>::reference
 deque<T>::front()
@@ -407,6 +423,14 @@ deque<T>::capacity() const
 }
 
 
+template <typename T>
+inline void
+deque<T>::shrink_to_fit()
+{
+    // Reject request for performance reasons
+}
+
+
 template <typename T>
 inline const T*
 deque<T>::data() const

test/stdgpu/deque.inc

@@ -1390,6 +1390,46 @@ TEST_F(stdgpu_deque, simultaneous_push_back_and_pop_front)
 }
 
 
+struct at_non_const_deque
+{
+    stdgpu::deque<int> pool;
+
+    at_non_const_deque(stdgpu::deque<int> pool)
+        : pool(pool)
+    {
+
+    }
+
+    STDGPU_DEVICE_ONLY void
+    operator()(const int x)
+    {
+        pool.at(x) = x * x;
+    }
+};
+
+
+TEST_F(stdgpu_deque, at_non_const)
+{
+    const stdgpu::index_t N = 100000;
+
+    stdgpu::deque<int> pool = stdgpu::deque<int>::createDeviceObject(N);
+
+    fill_deque(pool);
+
+    thrust::for_each(thrust::counting_iterator<int>(0), thrust::counting_iterator<int>(N),
+                     at_non_const_deque(pool));
+
+    int* host_numbers = copyCreateDevice2HostArray(pool.data(), N);
+    for (stdgpu::index_t i = 0; i < N; ++i)
+    {
+        EXPECT_EQ(host_numbers[i], i * i);
+    }
+
+    stdgpu::deque<int>::destroyDeviceObject(pool);
+    destroyHostArray<int>(host_numbers);
+}
+
+
 struct access_operator_non_const_deque
 {
     stdgpu::deque<int> pool;
@@ -1430,3 +1470,72 @@ TEST_F(stdgpu_deque, access_operator_non_const)
 }
 
 
+TEST_F(stdgpu_deque, shrink_to_fit_empty)
+{
+    const stdgpu::index_t N = 10000;
+
+    stdgpu::deque<int> pool = stdgpu::deque<int>::createDeviceObject(N);
+
+    ASSERT_EQ(pool.size(), 0);
+    ASSERT_EQ(pool.capacity(), N);
+    ASSERT_TRUE(pool.valid());
+
+    pool.shrink_to_fit();
+
+    EXPECT_EQ(pool.size(), 0);
+    EXPECT_TRUE(pool.capacity() == 0 || pool.capacity() == N); // Capacity may have changed or not
+    EXPECT_TRUE(pool.valid());
+
+    stdgpu::deque<int>::destroyDeviceObject(pool);
+}
+
+
+TEST_F(stdgpu_deque, shrink_to_fit_full)
+{
+    const stdgpu::index_t N = 10000;
+
+    stdgpu::deque<int> pool = stdgpu::deque<int>::createDeviceObject(N);
+
+    fill_deque(pool);
+
+    ASSERT_EQ(pool.size(), N);
+    ASSERT_EQ(pool.capacity(), N);
+    ASSERT_TRUE(pool.valid());
+
+    pool.shrink_to_fit();
+
+    EXPECT_EQ(pool.size(), N);
+    EXPECT_EQ(pool.capacity(), N);
+    EXPECT_TRUE(pool.valid());
+
+    stdgpu::deque<int>::destroyDeviceObject(pool);
+}
+
+
+TEST_F(stdgpu_deque, shrink_to_fit)
+{
+    const stdgpu::index_t N            = 10000;
+    const stdgpu::index_t N_pop        = 100;
+    const stdgpu::index_t N_remaining  = N - N_pop;
+
+    stdgpu::deque<int> pool = stdgpu::deque<int>::createDeviceObject(N);
+
+    fill_deque(pool);
+
+    thrust::for_each(thrust::counting_iterator<int>(0), thrust::counting_iterator<int>(N_pop),
+                     pop_back_deque<int>(pool));
+
+    ASSERT_EQ(pool.size(), N_remaining);
+    ASSERT_EQ(pool.capacity(), N);
+    ASSERT_TRUE(pool.valid());
+
+    pool.shrink_to_fit();
+
+    EXPECT_EQ(pool.size(), N_remaining);
+    EXPECT_TRUE(pool.capacity() == N_remaining || pool.capacity() == N); // Capacity may have changed or not
+    EXPECT_TRUE(pool.valid());
+
+    stdgpu::deque<int>::destroyDeviceObject(pool);
+}
+
+