execution: Add is_execution_policy trait and conditionally enable pol…

…icy-based functions

src/stdgpu/atomic.cuh

@@ -30,6 +30,7 @@
 #include <cstddef>
 #include <type_traits>
 
+#include <stdgpu/execution.h>
 #include <stdgpu/impl/type_traits.h>
 #include <stdgpu/memory.h>
 #include <stdgpu/platform.h>
@@ -123,7 +124,7 @@ public:
      * \note The size is implicitly set to 1 (and not needed as a parameter) as the object only manages a single value
      */
     template <typename ExecutionPolicy,
-              STDGPU_DETAIL_OVERLOAD_IF(!std::is_same_v<std::remove_reference_t<ExecutionPolicy>, Allocator>)>
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static atomic
     createDeviceObject(ExecutionPolicy&& policy, const Allocator& allocator = Allocator());
 
@@ -140,7 +141,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \param[in] device_object The object allocated on the GPU (device)
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static void
     destroyDeviceObject(ExecutionPolicy&& policy, atomic& device_object);
 

src/stdgpu/bitset.cuh

@@ -31,6 +31,7 @@
 #include <type_traits>
 
 #include <stdgpu/cstddef.h>
+#include <stdgpu/execution.h>
 #include <stdgpu/memory.h>
 #include <stdgpu/platform.h>
 
@@ -171,7 +172,8 @@ public:
      * \param[in] allocator The allocator instance to use
      * \return A newly created object of this class allocated on the GPU (device)
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static bitset
     createDeviceObject(ExecutionPolicy&& policy, const index_t& size, const Allocator& allocator = Allocator());
 
@@ -188,7 +190,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \param[in] device_object The object allocated on the GPU (device)
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static void
     destroyDeviceObject(ExecutionPolicy&& policy, bitset& device_object);
 
@@ -210,7 +213,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \post count() == size()
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     void
     set(ExecutionPolicy&& policy);
 
@@ -244,7 +248,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \post count() == 0
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     void
     reset(ExecutionPolicy&& policy);
 
@@ -268,7 +273,8 @@ public:
      * \tparam ExecutionPolicy The type of the execution policy
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     void
     flip(ExecutionPolicy&& policy);
 
@@ -335,7 +341,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return The number of set bits
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     index_t
     count(ExecutionPolicy&& policy) const;
 
@@ -352,7 +359,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return True if all bits are set, false otherwise
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     bool
     all(ExecutionPolicy&& policy) const;
 
@@ -369,7 +377,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return True if any bits are set, false otherwise
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     bool
     any(ExecutionPolicy&& policy) const;
 
@@ -386,7 +395,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return True if none of the bits are set, false otherwise
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     bool
     none(ExecutionPolicy&& policy) const;
 

src/stdgpu/deque.cuh

@@ -30,6 +30,7 @@
 #include <stdgpu/atomic.cuh>
 #include <stdgpu/bitset.cuh>
 #include <stdgpu/cstddef.h>
+#include <stdgpu/execution.h>
 #include <stdgpu/iterator.h>
 #include <stdgpu/memory.h>
 #include <stdgpu/mutex.cuh>
@@ -90,7 +91,8 @@ public:
      * \param[in] allocator The allocator instance to use
      * \return A newly created object of this class allocated on the GPU (device)
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static deque<T, Allocator>
     createDeviceObject(ExecutionPolicy&& policy, const index_t& capacity, const Allocator& allocator = Allocator());
 
@@ -107,7 +109,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \param[in] device_object The object allocated on the GPU (device)
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     static void
     destroyDeviceObject(ExecutionPolicy&& policy, deque<T, Allocator>& device_object);
 
@@ -302,7 +305,8 @@ public:
      * \tparam ExecutionPolicy The type of the execution policy
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     void
     clear(ExecutionPolicy&& policy);
 
@@ -319,7 +323,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return True if the state is valid, false otherwise
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     bool
     valid(ExecutionPolicy&& policy) const;
 
@@ -336,7 +341,8 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return A range of the object
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     stdgpu::device_indexed_range<T>
     device_range(ExecutionPolicy&& policy);
 
@@ -353,15 +359,17 @@ public:
      * \param[in] policy The execution policy, e.g. host or device, corresponding to the allocator
      * \return A const range of the object
      */
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     stdgpu::device_indexed_range<const T>
     device_range(ExecutionPolicy&& policy) const;
 
 private:
     STDGPU_DEVICE_ONLY bool
     occupied(const index_t n) const;
 
-    template <typename ExecutionPolicy>
+    template <typename ExecutionPolicy,
+              STDGPU_DETAIL_OVERLOAD_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
     bool
     occupied_count_valid(ExecutionPolicy&& policy) const;
 

src/stdgpu/execution.h

@@ -23,27 +23,61 @@
 
 #include <type_traits>
 
+#include <stdgpu/impl/type_traits.h>
+
 /**
  * \file stdgpu/execution.h
  */
 
 #include <thrust/execution_policy.h>
 
+namespace stdgpu
+{
+
+/**
+ * \ingroup execution
+ * \brief Type traits to check whether the provided class is an execution policy
+ */
+template <typename T>
+struct is_execution_policy : std::is_base_of<thrust::execution_policy<T>, T>
+{
+};
+
+//! @cond Doxygen_Suppress
+template <typename T>
+inline constexpr bool is_execution_policy_v = is_execution_policy<T>::value;
+//! @endcond
+
+} // namespace stdgpu
+
 namespace stdgpu::execution
 {
 
+/**
+ * \ingroup execution
+ * \brief The base execution policy class from which all policies are derived and custom ones must be derived
+ */
+template <typename T>
+using execution_policy = thrust::execution_policy<T>;
+
 /**
  * \ingroup execution
  * \brief The device execution policy class
  */
 using device_policy = std::remove_const_t<decltype(thrust::device)>;
 
+static_assert(is_execution_policy_v<remove_cvref_t<device_policy>>,
+              "stdgpu::execution::device_policy: Should be an execution policy but is not");
+
 /**
  * \ingroup execution
  * \brief The host execution policy class
  */
 using host_policy = std::remove_const_t<decltype(thrust::host)>;
 
+static_assert(is_execution_policy_v<remove_cvref_t<host_policy>>,
+              "stdgpu::execution::host_policy: Should be an execution policy but is not");
+
 /**
  * \ingroup execution
  * \brief The device execution policy

src/stdgpu/impl/atomic_detail.cuh

@@ -138,7 +138,7 @@ atomic<T, Allocator>::createDeviceObject(const Allocator& allocator)
 
 template <typename T, typename Allocator>
 template <typename ExecutionPolicy,
-          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(!std::is_same_v<std::remove_reference_t<ExecutionPolicy>, Allocator>)>
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline atomic<T, Allocator>
 atomic<T, Allocator>::createDeviceObject(ExecutionPolicy&& policy, const Allocator& allocator)
 {
@@ -159,7 +159,8 @@ atomic<T, Allocator>::destroyDeviceObject(atomic<T, Allocator>& device_object)
 }
 
 template <typename T, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline void
 atomic<T, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, atomic<T, Allocator>& device_object)
 {

src/stdgpu/impl/bitset_detail.cuh

@@ -187,7 +187,8 @@ bitset<Block, Allocator>::createDeviceObject(const index_t& size, const Allocato
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline bitset<Block, Allocator>
 bitset<Block, Allocator>::createDeviceObject(ExecutionPolicy&& policy, const index_t& size, const Allocator& allocator)
 {
@@ -209,7 +210,8 @@ bitset<Block, Allocator>::destroyDeviceObject(bitset<Block, Allocator>& device_o
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline void
 bitset<Block, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, bitset<Block, Allocator>& device_object)
 {
@@ -249,7 +251,8 @@ bitset<Block, Allocator>::set()
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline void
 bitset<Block, Allocator>::set(ExecutionPolicy&& policy)
 {
@@ -276,7 +279,8 @@ bitset<Block, Allocator>::reset()
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline void
 bitset<Block, Allocator>::reset(ExecutionPolicy&& policy)
 {
@@ -303,7 +307,8 @@ bitset<Block, Allocator>::flip()
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline void
 bitset<Block, Allocator>::flip(ExecutionPolicy&& policy)
 {
@@ -380,7 +385,8 @@ bitset<Block, Allocator>::count() const
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline index_t
 bitset<Block, Allocator>::count(ExecutionPolicy&& policy) const
 {
@@ -404,7 +410,8 @@ bitset<Block, Allocator>::all() const
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline bool
 bitset<Block, Allocator>::all(ExecutionPolicy&& policy) const
 {
@@ -424,7 +431,8 @@ bitset<Block, Allocator>::any() const
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline bool
 bitset<Block, Allocator>::any(ExecutionPolicy&& policy) const
 {
@@ -444,7 +452,8 @@ bitset<Block, Allocator>::none() const
 }
 
 template <typename Block, typename Allocator>
-template <typename ExecutionPolicy>
+template <typename ExecutionPolicy,
+          STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 inline bool
 bitset<Block, Allocator>::none(ExecutionPolicy&& policy) const
 {