memory: Add and use destroy* functions

src/stdgpu/impl/deque_detail.cuh

@@ -16,7 +16,6 @@
 #ifndef STDGPU_DEQUE_DETAIL_H
 #define STDGPU_DEQUE_DETAIL_H
 
-#include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 
@@ -469,17 +468,13 @@ deque<T>::clear()
     // One large block
     if (begin <= end)
     {
-        thrust::for_each(stdgpu::make_device(_data + begin), stdgpu::make_device(_data + end),
-                         stdgpu::detail::destroy_value<T>());
+        stdgpu::destroy(stdgpu::make_device(_data + begin), stdgpu::make_device(_data + end));
     }
     // Two disconnected blocks
     else
     {
-        thrust::for_each(stdgpu::device_begin(_data), stdgpu::make_device(_data + end),
-                         stdgpu::detail::destroy_value<T>());
-
-        thrust::for_each(stdgpu::make_device(_data + begin), stdgpu::device_end(_data),
-                         stdgpu::detail::destroy_value<T>());
+        stdgpu::destroy(stdgpu::device_begin(_data), stdgpu::make_device(_data + end));
+        stdgpu::destroy(stdgpu::make_device(_data + begin), stdgpu::device_end(_data));
     }
 
 

src/stdgpu/impl/memory_detail.h

@@ -13,8 +13,8 @@
  *  limitations under the License.
  */
 
-#ifndef STDGPU_MEMORYDETAIL_H
-#define STDGPU_MEMORYDETAIL_H
+#ifndef STDGPU_MEMORY_DETAIL_H
+#define STDGPU_MEMORY_DETAIL_H
 
 #include <cstdio>
 #include <type_traits>
@@ -226,8 +226,7 @@ destroyDeviceArray(T*& device_array)
 {
     #if !STDGPU_USE_FAST_DESTROY
         #if STDGPU_BACKEND != STDGPU_BACKEND_CUDA || STDGPU_DEVICE_COMPILER == STDGPU_DEVICE_COMPILER_NVCC
-            thrust::for_each(stdgpu::device_begin(device_array), stdgpu::device_end(device_array),
-                             stdgpu::detail::destroy_value<T>());
+            stdgpu::destroy(stdgpu::device_begin(device_array), stdgpu::device_end(device_array));
 
             stdgpu::detail::workaround_synchronize_device_thrust();
         #else
@@ -254,8 +253,7 @@ void
 destroyHostArray(T*& host_array)
 {
     #if !STDGPU_USE_FAST_DESTROY
-        thrust::for_each(stdgpu::host_begin(host_array), stdgpu::host_end(host_array),
-                         stdgpu::detail::destroy_value<T>());
+        stdgpu::destroy(stdgpu::host_begin(host_array), stdgpu::host_end(host_array));
     #endif
 
     stdgpu::safe_host_allocator<T> host_allocator;
@@ -271,8 +269,7 @@ destroyManagedArray(T*& managed_array)
 {
     #if !STDGPU_USE_FAST_DESTROY
         // Call on host since the initialization place is not known
-        thrust::for_each(stdgpu::host_begin(managed_array), stdgpu::host_end(managed_array),
-                         stdgpu::detail::destroy_value<T>());
+        stdgpu::destroy(stdgpu::host_begin(managed_array), stdgpu::host_end(managed_array));
     #endif
 
     stdgpu::safe_managed_allocator<T> managed_allocator;
@@ -506,11 +503,42 @@ default_allocator_traits::construct(T* p,
 template <typename T>
 STDGPU_HOST_DEVICE void
 default_allocator_traits::destroy(T* p)
+{
+    destroy_at(p);
+}
+
+
+template <typename T>
+STDGPU_HOST_DEVICE void
+destroy_at(T* p)
 {
     p->~T();
 }
 
 
+template <typename Iterator>
+void
+destroy(Iterator first,
+        Iterator last)
+{
+    thrust::for_each(first, last,
+                     detail::destroy_value<typename std::iterator_traits<Iterator>::value_type>());
+}
+
+
+template <typename Iterator, typename Size>
+Iterator
+destroy_n(Iterator first,
+          Size n)
+{
+    Iterator last = first + n;
+
+    destroy(first, last);
+
+    return last;
+}
+
+
 template <>
 dynamic_memory_type
 get_dynamic_memory_type(void* array);
@@ -563,4 +591,4 @@ struct [[deprecated("Replaced by stdgpu::safe_host_allocator<T>")]] safe_pinned_
 
 
 
-#endif // STDGPU_MEMORYDETAIL_H
+#endif // STDGPU_MEMORY_DETAIL_H

src/stdgpu/impl/vector_detail.cuh

@@ -16,8 +16,6 @@
 #ifndef STDGPU_VECTOR_DETAIL_H
 #define STDGPU_VECTOR_DETAIL_H
 
-#include <thrust/fill.h>
-
 #include <stdgpu/contract.h>
 #include <stdgpu/iterator.h>
 #include <stdgpu/memory.h>
@@ -331,8 +329,7 @@ vector<T>::clear()
 
     const index_t current_size = size();
 
-    thrust::for_each(stdgpu::device_begin(_data), stdgpu::device_begin(_data) + current_size,
-                     stdgpu::detail::destroy_value<T>());
+    stdgpu::destroy(stdgpu::device_begin(_data), stdgpu::device_begin(_data) + current_size);
 
     _occupied.reset();
 

src/stdgpu/memory.h

@@ -400,6 +400,42 @@ struct default_allocator_traits
 };
 
 
+/**
+ * \brief Destroys the value at the given pointer
+ * \tparam T The value type
+ * \param[in] p A pointer to the value to destroy
+ */
+template <typename T>
+STDGPU_HOST_DEVICE void
+destroy_at(T* p);
+
+
+/**
+ * \brief Destroys the range of values
+ * \tparam Iterator The iterator type of the values
+ * \param[in] first An iterator to the begin of the value range
+ * \param[in] last An iterator to the end of the value range
+ */
+template <typename Iterator>
+void
+destroy(Iterator first,
+        Iterator last);
+
+
+/**
+ * \brief Destroys the range of values
+ * \tparam Iterator The iterator type of the values
+ * \tparam Size The size type
+ * \param[in] first An iterator to the begin of the value range
+ * \param[in] n The number of elements in the value range
+ * \return An iterator to the end of the value range
+ */
+template <typename Iterator, typename Size>
+Iterator
+destroy_n(Iterator first,
+          Size n);
+
+
 /**
  * \brief Returns the total number of allocations of a specific memory type
  * \param[in] memory_type A dynamic memory type