diff --git a/.clang-tidy b/.clang-tidy
index 37f90618f..97f5bdcc0 100644
--- a/.clang-tidy
+++ b/.clang-tidy
@@ -23,6 +23,7 @@ hicpp-*,\
 misc-*,\
 -misc-const-correctness,\
 -misc-include-cleaner,\
+-misc-use-internal-linkage,\
 modernize-*,\
 -modernize-avoid-c-arrays,\
 -modernize-use-auto,\
diff --git a/examples/createAndDestroyDeviceObject.cpp b/examples/createAndDestroyDeviceObject.cpp
index eebcb2604..26a57471b 100644
--- a/examples/createAndDestroyDeviceObject.cpp
+++ b/examples/createAndDestroyDeviceObject.cpp
@@ -69,7 +69,7 @@ class Image
     STDGPU_HOST_DEVICE std::uint8_t&
     operator()(const stdgpu::index_t x, const stdgpu::index_t y)
     {
-        return _values[y * _width + x];
+        return _values[(y * _width) + x];
     }
 
     STDGPU_HOST_DEVICE stdgpu::index_t
diff --git a/src/stdgpu/impl/algorithm_detail.h b/src/stdgpu/impl/algorithm_detail.h
index 9f275fd81..f589e49ea 100644
--- a/src/stdgpu/impl/algorithm_detail.h
+++ b/src/stdgpu/impl/algorithm_detail.h
@@ -29,14 +29,14 @@ template <class T>
 constexpr STDGPU_HOST_DEVICE const T&
 min(const T& a, const T& b)
 {
-    return (b < a) ? b : a;
+    return (b < a) ? b : a; // NOLINT(bugprone-return-const-ref-from-parameter)
 }
 
 template <class T>
 constexpr STDGPU_HOST_DEVICE const T&
 max(const T& a, const T& b)
 {
-    return (a < b) ? b : a;
+    return (a < b) ? b : a; // NOLINT(bugprone-return-const-ref-from-parameter)
 }
 
 template <class T>
@@ -45,7 +45,8 @@ clamp(const T& v, const T& lower, const T& upper)
 {
     STDGPU_EXPECTS(!(upper < lower));
 
-    return v < lower ? lower : upper < v ? upper : v; // NOLINT(readability-avoid-nested-conditional-operator)
+    // NOLINTNEXTLINE(bugprone-return-const-ref-from-parameter,readability-avoid-nested-conditional-operator)
+    return v < lower ? lower : upper < v ? upper : v;
 }
 
 template <typename IndexType,
@@ -58,7 +59,7 @@ for_each_index(ExecutionPolicy&& policy, IndexType size, UnaryFunction f)
     thrust::for_each(std::forward<ExecutionPolicy>(policy),
                      thrust::counting_iterator<IndexType>(0),
                      thrust::counting_iterator<IndexType>(size),
-                     f);
+                     std::move(f));
 }
 
 namespace detail
diff --git a/src/stdgpu/impl/bitset_detail.cuh b/src/stdgpu/impl/bitset_detail.cuh
index c5c8fb73a..46f6b79ce 100644
--- a/src/stdgpu/impl/bitset_detail.cuh
+++ b/src/stdgpu/impl/bitset_detail.cuh
@@ -17,6 +17,7 @@
 #define STDGPU_BITSET_DETAIL_H
 
 #include <limits>
+#include <type_traits>
 
 #include <stdgpu/algorithm.h>
 #include <stdgpu/atomic.cuh>
@@ -119,7 +120,7 @@ div_up(const index_t a, const index_t b) noexcept
     STDGPU_EXPECTS(a >= 0);
     STDGPU_EXPECTS(b > 0);
 
-    index_t result = (a % b != 0) ? (a / b + 1) : (a / b);
+    index_t result = (a % b != 0) ? ((a / b) + 1) : (a / b);
 
     STDGPU_ENSURES(result * b >= a);
 
@@ -146,7 +147,7 @@ private:
     STDGPU_HOST_DEVICE Block
     block_mask(const index_t i) const
     {
-        index_t remaining_bits = _size - i * _bits_per_block;
+        index_t remaining_bits = _size - (i * _bits_per_block);
         return (remaining_bits >= _bits_per_block)
                        ? ~static_cast<Block>(0)
                        : (static_cast<Block>(1) << static_cast<Block>(remaining_bits)) - static_cast<Block>(1);
@@ -256,7 +257,7 @@ template <typename ExecutionPolicy,
 inline void
 bitset<Block, Allocator>::set(ExecutionPolicy&& policy)
 {
-    fill(std::forward<ExecutionPolicy>(policy), device_begin(_bit_blocks), device_end(_bit_blocks), ~block_type(0));
+    fill(std::decay_t<ExecutionPolicy>{ policy }, device_begin(_bit_blocks), device_end(_bit_blocks), ~block_type(0));
 
     STDGPU_ENSURES(count(std::forward<ExecutionPolicy>(policy)) == size());
 }
@@ -284,7 +285,7 @@ template <typename ExecutionPolicy,
 inline void
 bitset<Block, Allocator>::reset(ExecutionPolicy&& policy)
 {
-    fill(std::forward<ExecutionPolicy>(policy), device_begin(_bit_blocks), device_end(_bit_blocks), block_type(0));
+    fill(std::decay_t<ExecutionPolicy>{ policy }, device_begin(_bit_blocks), device_end(_bit_blocks), block_type(0));
 
     STDGPU_ENSURES(count(std::forward<ExecutionPolicy>(policy)) == 0);
 }
diff --git a/src/stdgpu/impl/deque_detail.cuh b/src/stdgpu/impl/deque_detail.cuh
index fbda4d320..637953aa5 100644
--- a/src/stdgpu/impl/deque_detail.cuh
+++ b/src/stdgpu/impl/deque_detail.cuh
@@ -16,6 +16,8 @@
 #ifndef STDGPU_DEQUE_DETAIL_H
 #define STDGPU_DEQUE_DETAIL_H
 
+#include <type_traits>
+
 #include <stdgpu/contract.h>
 #include <stdgpu/iterator.h>
 #include <stdgpu/memory.h>
@@ -42,19 +44,21 @@ deque<T, Allocator>::createDeviceObject(ExecutionPolicy&& policy, const index_t&
 
     deque<T, Allocator> result(
             mutex_array<mutex_default_type, mutex_array_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     capacity,
                     mutex_array_allocator_type(allocator)),
             bitset<bitset_default_type, bitset_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     capacity,
                     bitset_allocator_type(allocator)),
-            atomic<int, atomic_int_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
+            atomic<int, atomic_int_allocator_type>::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                        atomic_int_allocator_type(allocator)),
-            atomic<unsigned int, atomic_uint_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
-                                                                                 atomic_uint_allocator_type(allocator)),
-            atomic<unsigned int, atomic_uint_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
-                                                                                 atomic_uint_allocator_type(allocator)),
+            atomic<unsigned int, atomic_uint_allocator_type>::createDeviceObject(
+                    std::decay_t<ExecutionPolicy>{ policy },
+                    atomic_uint_allocator_type(allocator)),
+            atomic<unsigned int, atomic_uint_allocator_type>::createDeviceObject(
+                    std::decay_t<ExecutionPolicy>{ policy },
+                    atomic_uint_allocator_type(allocator)),
             allocator);
     result._data = allocator_traits<allocator_type>::allocate(result._allocator, capacity);
     result._range_indices = allocator_traits<index_allocator_type>::allocate(result._index_allocator, capacity);
@@ -77,7 +81,7 @@ deque<T, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, deque<T, Allo
 {
     if (!detail::is_destroy_optimizable<value_type>())
     {
-        device_object.clear(std::forward<ExecutionPolicy>(policy));
+        device_object.clear(std::decay_t<ExecutionPolicy>{ policy });
     }
 
     allocator_traits<allocator_type>::deallocate(device_object._allocator,
@@ -89,13 +93,13 @@ deque<T, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, deque<T, Allo
     device_object._data = nullptr;
     device_object._range_indices = nullptr;
     mutex_array<mutex_default_type, mutex_array_allocator_type>::destroyDeviceObject(
-            std::forward<ExecutionPolicy>(policy),
+            std::decay_t<ExecutionPolicy>{ policy },
             device_object._locks);
-    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                             device_object._occupied);
-    atomic<int, atomic_int_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    atomic<int, atomic_int_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                 device_object._size);
-    atomic<unsigned int, atomic_uint_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    atomic<unsigned int, atomic_uint_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                           device_object._begin);
     atomic<unsigned int, atomic_uint_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
                                                                           device_object._end);
@@ -550,48 +554,50 @@ template <typename ExecutionPolicy,
 inline void
 deque<T, Allocator>::clear(ExecutionPolicy&& policy)
 {
-    if (empty(std::forward<ExecutionPolicy>(policy)))
+    if (empty(std::decay_t<ExecutionPolicy>{ policy }))
     {
         return;
     }
 
     if (!detail::is_destroy_optimizable<value_type>())
     {
-        const index_t begin = static_cast<index_t>(_begin.load(std::forward<ExecutionPolicy>(policy)));
-        const index_t end = static_cast<index_t>(_end.load(std::forward<ExecutionPolicy>(policy)));
+        const index_t begin = static_cast<index_t>(_begin.load(std::decay_t<ExecutionPolicy>{ policy }));
+        const index_t end = static_cast<index_t>(_end.load(std::decay_t<ExecutionPolicy>{ policy }));
 
         // Full, i.e. one large block and begin == end
-        if (full(std::forward<ExecutionPolicy>(policy)))
+        if (full(std::decay_t<ExecutionPolicy>{ policy }))
         {
-            detail::unoptimized_destroy(std::forward<ExecutionPolicy>(policy), device_begin(_data), device_end(_data));
+            detail::unoptimized_destroy(std::decay_t<ExecutionPolicy>{ policy },
+                                        device_begin(_data),
+                                        device_end(_data));
         }
         // One large block
         else if (begin <= end)
         {
-            detail::unoptimized_destroy(std::forward<ExecutionPolicy>(policy),
+            detail::unoptimized_destroy(std::decay_t<ExecutionPolicy>{ policy },
                                         make_device(_data + begin),
                                         make_device(_data + end));
         }
         // Two disconnected blocks
         else
         {
-            detail::unoptimized_destroy(std::forward<ExecutionPolicy>(policy),
+            detail::unoptimized_destroy(std::decay_t<ExecutionPolicy>{ policy },
                                         device_begin(_data),
                                         make_device(_data + end));
-            detail::unoptimized_destroy(std::forward<ExecutionPolicy>(policy),
+            detail::unoptimized_destroy(std::decay_t<ExecutionPolicy>{ policy },
                                         make_device(_data + begin),
                                         device_end(_data));
         }
     }
 
-    _occupied.reset(std::forward<ExecutionPolicy>(policy));
+    _occupied.reset(std::decay_t<ExecutionPolicy>{ policy });
 
-    _size.store(std::forward<ExecutionPolicy>(policy), 0);
+    _size.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
 
-    _begin.store(std::forward<ExecutionPolicy>(policy), 0);
-    _end.store(std::forward<ExecutionPolicy>(policy), 0);
+    _begin.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
+    _end.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
 
-    STDGPU_ENSURES(empty(std::forward<ExecutionPolicy>(policy)));
+    STDGPU_ENSURES(empty(std::decay_t<ExecutionPolicy>{ policy }));
     STDGPU_ENSURES(valid(std::forward<ExecutionPolicy>(policy)));
 }
 
@@ -614,9 +620,9 @@ deque<T, Allocator>::valid(ExecutionPolicy&& policy) const
         return true;
     }
 
-    return (size_valid(std::forward<ExecutionPolicy>(policy)) &&
-            occupied_count_valid(std::forward<ExecutionPolicy>(policy)) &&
-            _locks.valid(std::forward<ExecutionPolicy>(policy)));
+    return (size_valid(std::decay_t<ExecutionPolicy>{ policy }) &&
+            occupied_count_valid(std::decay_t<ExecutionPolicy>{ policy }) &&
+            _locks.valid(std::decay_t<ExecutionPolicy>{ policy }));
 }
 
 template <typename T, typename Allocator>
@@ -632,18 +638,18 @@ template <typename ExecutionPolicy,
 stdgpu::device_indexed_range<T>
 deque<T, Allocator>::device_range(ExecutionPolicy&& policy)
 {
-    const index_t begin = static_cast<index_t>(_begin.load(std::forward<ExecutionPolicy>(policy)));
-    const index_t end = static_cast<index_t>(_end.load(std::forward<ExecutionPolicy>(policy)));
+    const index_t begin = static_cast<index_t>(_begin.load(std::decay_t<ExecutionPolicy>{ policy }));
+    const index_t end = static_cast<index_t>(_end.load(std::decay_t<ExecutionPolicy>{ policy }));
 
     // Full, i.e. one large block and begin == end
-    if (full(std::forward<ExecutionPolicy>(policy)))
+    if (full(std::decay_t<ExecutionPolicy>{ policy }))
     {
-        iota(std::forward<ExecutionPolicy>(policy), device_begin(_range_indices), device_end(_range_indices), 0);
+        iota(std::decay_t<ExecutionPolicy>{ policy }, device_begin(_range_indices), device_end(_range_indices), 0);
     }
     // One large block, including empty block
     else if (begin <= end)
     {
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices),
              device_begin(_range_indices) + (end - begin),
              begin);
@@ -651,11 +657,11 @@ deque<T, Allocator>::device_range(ExecutionPolicy&& policy)
     // Two disconnected blocks
     else
     {
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices),
              device_begin(_range_indices) + end,
              0);
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices) + end,
              device_begin(_range_indices) + (end + capacity() - begin),
              begin);
@@ -679,18 +685,18 @@ template <typename ExecutionPolicy,
 stdgpu::device_indexed_range<const T>
 deque<T, Allocator>::device_range(ExecutionPolicy&& policy) const
 {
-    const index_t begin = static_cast<index_t>(_begin.load(std::forward<ExecutionPolicy>(policy)));
-    const index_t end = static_cast<index_t>(_end.load(std::forward<ExecutionPolicy>(policy)));
+    const index_t begin = static_cast<index_t>(_begin.load(std::decay_t<ExecutionPolicy>{ policy }));
+    const index_t end = static_cast<index_t>(_end.load(std::decay_t<ExecutionPolicy>{ policy }));
 
     // Full, i.e. one large block and begin == end
-    if (full(std::forward<ExecutionPolicy>(policy)))
+    if (full(std::decay_t<ExecutionPolicy>{ policy }))
     {
-        iota(std::forward<ExecutionPolicy>(policy), device_begin(_range_indices), device_end(_range_indices), 0);
+        iota(std::decay_t<ExecutionPolicy>{ policy }, device_begin(_range_indices), device_end(_range_indices), 0);
     }
     // One large block, including empty block
     else if (begin <= end)
     {
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices),
              device_begin(_range_indices) + (end - begin),
              begin);
@@ -698,11 +704,11 @@ deque<T, Allocator>::device_range(ExecutionPolicy&& policy) const
     // Two disconnected blocks
     else
     {
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices),
              device_begin(_range_indices) + end,
              0);
-        iota(std::forward<ExecutionPolicy>(policy),
+        iota(std::decay_t<ExecutionPolicy>{ policy },
              device_begin(_range_indices) + end,
              device_begin(_range_indices) + (end + capacity() - begin),
              begin);
@@ -726,7 +732,7 @@ template <typename ExecutionPolicy,
 bool
 deque<T, Allocator>::occupied_count_valid(ExecutionPolicy&& policy) const
 {
-    index_t size_count = size(std::forward<ExecutionPolicy>(policy));
+    index_t size_count = size(std::decay_t<ExecutionPolicy>{ policy });
     index_t size_sum = _occupied.count(std::forward<ExecutionPolicy>(policy));
 
     return (size_count == size_sum);
diff --git a/src/stdgpu/impl/memory_detail.h b/src/stdgpu/impl/memory_detail.h
index 91fc1e47b..33238eb89 100644
--- a/src/stdgpu/impl/memory_detail.h
+++ b/src/stdgpu/impl/memory_detail.h
@@ -174,7 +174,7 @@ unoptimized_destroy(ExecutionPolicy&& policy, Iterator first, Iterator last)
 
 template <typename T>
 T*
-createDeviceArray(const stdgpu::index64_t count, const T default_value)
+createDeviceArray(const stdgpu::index64_t count, const T default_value) // NOLINT(performance-unnecessary-value-param)
 {
     T* device_array = nullptr;
 
@@ -441,8 +441,8 @@ template <typename ExecutionPolicy,
           typename... Args,
           STDGPU_DETAIL_OVERLOAD_DEFINITION_IF(is_execution_policy_v<remove_cvref_t<ExecutionPolicy>>)>
 device_unique_object<T>::device_unique_object(ExecutionPolicy&& policy, Args&&... args)
-  : _object(new T(T::createDeviceObject(std::forward<ExecutionPolicy>(policy), std::forward<Args>(args)...)),
-            [_policy = std::forward<ExecutionPolicy>(policy)](T* ptr)
+  : _object(new T(T::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy }, std::forward<Args>(args)...)),
+            [_policy = std::decay_t<ExecutionPolicy>{ policy }](T* ptr)
             {
                 T::destroyDeviceObject(_policy, *ptr);
                 delete ptr;
diff --git a/src/stdgpu/impl/unordered_base_detail.cuh b/src/stdgpu/impl/unordered_base_detail.cuh
index e5e350053..ad72e0fbb 100644
--- a/src/stdgpu/impl/unordered_base_detail.cuh
+++ b/src/stdgpu/impl/unordered_base_detail.cuh
@@ -17,6 +17,7 @@
 #define STDGPU_UNORDERED_BASE_DETAIL_H
 
 #include <cmath>
+#include <type_traits>
 
 #include <stdgpu/algorithm.h>
 #include <stdgpu/bit.h>
@@ -168,9 +169,9 @@ template <typename ExecutionPolicy,
 device_indexed_range<typename unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::value_type>
 unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::device_range(ExecutionPolicy&& policy)
 {
-    _range_indices_end.store(std::forward<ExecutionPolicy>(policy), 0);
+    _range_indices_end.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
 
-    for_each_index(std::forward<ExecutionPolicy>(policy),
+    for_each_index(std::decay_t<ExecutionPolicy>{ policy },
                    total_count(),
                    unordered_base_collect_positions<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>(*this));
 
@@ -193,9 +194,9 @@ template <typename ExecutionPolicy,
 device_indexed_range<const typename unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::value_type>
 unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::device_range(ExecutionPolicy&& policy) const
 {
-    _range_indices_end.store(std::forward<ExecutionPolicy>(policy), 0);
+    _range_indices_end.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
 
-    for_each_index(std::forward<ExecutionPolicy>(policy),
+    for_each_index(std::decay_t<ExecutionPolicy>{ policy },
                    total_count(),
                    unordered_base_collect_positions<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>(*this));
 
@@ -324,16 +325,16 @@ loop_free(ExecutionPolicy&& policy, const unordered_base<Key, Value, KeyFromValu
     using flags_allocator_type = typename stdgpu::allocator_traits<Allocator>::template rebind_alloc<int>;
     flags_allocator_type flags_allocator = flags_allocator_type(base.get_allocator());
 
-    int* flags = allocator_traits<flags_allocator_type>::allocate_filled(std::forward<ExecutionPolicy>(policy),
+    int* flags = allocator_traits<flags_allocator_type>::allocate_filled(std::decay_t<ExecutionPolicy>{ policy },
                                                                          flags_allocator,
                                                                          base.total_count(),
                                                                          0);
 
-    for_each_index(std::forward<ExecutionPolicy>(policy),
+    for_each_index(std::decay_t<ExecutionPolicy>{ policy },
                    base.bucket_count(),
                    count_visits<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>(base, flags));
 
-    bool result = transform_reduce_index(std::forward<ExecutionPolicy>(policy),
+    bool result = transform_reduce_index(std::decay_t<ExecutionPolicy>{ policy },
                                          base.total_count(),
                                          true,
                                          logical_and<>(),
@@ -459,7 +460,7 @@ inline bool
 occupied_count_valid(ExecutionPolicy&& policy,
                      const unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>& base)
 {
-    index_t size_count = base.size(std::forward<ExecutionPolicy>(policy));
+    index_t size_count = base.size(std::decay_t<ExecutionPolicy>{ policy });
     index_t size_sum = base._occupied.count(std::forward<ExecutionPolicy>(policy));
 
     return (size_count == size_sum);
@@ -1197,13 +1198,13 @@ unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::valid(Execu
         return true;
     }
 
-    return (offset_range_valid(std::forward<ExecutionPolicy>(policy), *this) &&
-            loop_free(std::forward<ExecutionPolicy>(policy), *this) &&
-            values_reachable(std::forward<ExecutionPolicy>(policy), *this) &&
-            unique(std::forward<ExecutionPolicy>(policy), *this) &&
-            occupied_count_valid(std::forward<ExecutionPolicy>(policy), *this) &&
-            _locks.valid(std::forward<ExecutionPolicy>(policy)) &&
-            _excess_list_positions.valid(std::forward<ExecutionPolicy>(policy)));
+    return (offset_range_valid(std::decay_t<ExecutionPolicy>{ policy }, *this) &&
+            loop_free(std::decay_t<ExecutionPolicy>{ policy }, *this) &&
+            values_reachable(std::decay_t<ExecutionPolicy>{ policy }, *this) &&
+            unique(std::decay_t<ExecutionPolicy>{ policy }, *this) &&
+            occupied_count_valid(std::decay_t<ExecutionPolicy>{ policy }, *this) &&
+            _locks.valid(std::decay_t<ExecutionPolicy>{ policy }) &&
+            _excess_list_positions.valid(std::decay_t<ExecutionPolicy>{ policy }));
 }
 
 template <typename Key, typename Value, typename KeyFromValue, typename Hash, typename KeyEqual, typename Allocator>
@@ -1219,23 +1220,23 @@ template <typename ExecutionPolicy,
 void
 unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::clear(ExecutionPolicy&& policy)
 {
-    if (empty(std::forward<ExecutionPolicy>(policy)))
+    if (empty(std::decay_t<ExecutionPolicy>{ policy }))
     {
         return;
     }
 
     if (!detail::is_destroy_optimizable<Value>())
     {
-        for_each_index(std::forward<ExecutionPolicy>(policy),
+        for_each_index(std::decay_t<ExecutionPolicy>{ policy },
                        total_count(),
                        destroy_values<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>(*this));
     }
 
-    fill(std::forward<ExecutionPolicy>(policy), device_begin(_offsets), device_end(_offsets), 0);
+    fill(std::decay_t<ExecutionPolicy>{ policy }, device_begin(_offsets), device_end(_offsets), 0);
 
-    _occupied.reset(std::forward<ExecutionPolicy>(policy));
+    _occupied.reset(std::decay_t<ExecutionPolicy>{ policy });
 
-    _occupied_count.store(std::forward<ExecutionPolicy>(policy), 0);
+    _occupied_count.store(std::decay_t<ExecutionPolicy>{ policy }, 0);
 
     detail::vector_clear_iota(std::forward<ExecutionPolicy>(policy), _excess_list_positions, bucket_count());
 }
@@ -1263,24 +1264,24 @@ unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::createDevic
 
     unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator> result(
             bitset<bitset_default_type, bitset_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     total_count,
                     bitset_allocator_type(allocator)),
-            atomic<int, atomic_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
+            atomic<int, atomic_allocator_type>::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                    atomic_allocator_type(allocator)),
-            vector<index_t, index_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
+            vector<index_t, index_allocator_type>::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                       excess_count,
                                                                       index_allocator_type(allocator)),
             mutex_array<mutex_default_type, mutex_array_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     total_count,
                     mutex_array_allocator_type(allocator)),
-            atomic<int, atomic_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
+            atomic<int, atomic_allocator_type>::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                    atomic_allocator_type(allocator)),
             allocator);
     result._bucket_count = bucket_count;
     result._values = allocator_traits<allocator_type>::allocate(result._allocator, total_count);
-    result._offsets = allocator_traits<index_allocator_type>::allocate_filled(std::forward<ExecutionPolicy>(policy),
+    result._offsets = allocator_traits<index_allocator_type>::allocate_filled(std::decay_t<ExecutionPolicy>{ policy },
                                                                               result._index_allocator,
                                                                               total_count,
                                                                               0);
@@ -1289,7 +1290,7 @@ unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::createDevic
     result._hash = hasher();
     result._key_equal = key_equal();
 
-    detail::vector_clear_iota(std::forward<ExecutionPolicy>(policy), result._excess_list_positions, bucket_count);
+    detail::vector_clear_iota(std::decay_t<ExecutionPolicy>{ policy }, result._excess_list_positions, bucket_count);
 
     STDGPU_ENSURES(result._excess_list_positions.full(std::forward<ExecutionPolicy>(policy)));
 
@@ -1306,28 +1307,28 @@ unordered_base<Key, Value, KeyFromValue, Hash, KeyEqual, Allocator>::destroyDevi
 {
     if (!detail::is_destroy_optimizable<value_type>())
     {
-        device_object.clear(std::forward<ExecutionPolicy>(policy));
+        device_object.clear(std::decay_t<ExecutionPolicy>{ policy });
     }
 
     device_object._bucket_count = 0;
     allocator_traits<allocator_type>::deallocate(device_object._allocator,
                                                  device_object._values,
                                                  device_object.total_count());
-    allocator_traits<index_allocator_type>::deallocate_filled(std::forward<ExecutionPolicy>(policy),
+    allocator_traits<index_allocator_type>::deallocate_filled(std::decay_t<ExecutionPolicy>{ policy },
                                                               device_object._index_allocator,
                                                               device_object._offsets,
                                                               device_object.total_count());
     allocator_traits<index_allocator_type>::deallocate(device_object._index_allocator,
                                                        device_object._range_indices,
                                                        device_object.total_count());
-    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                             device_object._occupied);
-    atomic<int, atomic_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    atomic<int, atomic_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                             device_object._occupied_count);
     mutex_array<mutex_default_type, mutex_array_allocator_type>::destroyDeviceObject(
-            std::forward<ExecutionPolicy>(policy),
+            std::decay_t<ExecutionPolicy>{ policy },
             device_object._locks);
-    vector<index_t, index_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    vector<index_t, index_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                device_object._excess_list_positions);
     atomic<int, atomic_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
                                                             device_object._range_indices_end);
diff --git a/src/stdgpu/impl/vector_detail.cuh b/src/stdgpu/impl/vector_detail.cuh
index 159645f2d..fd84cd4d4 100644
--- a/src/stdgpu/impl/vector_detail.cuh
+++ b/src/stdgpu/impl/vector_detail.cuh
@@ -16,6 +16,8 @@
 #ifndef STDGPU_VECTOR_DETAIL_H
 #define STDGPU_VECTOR_DETAIL_H
 
+#include <type_traits>
+
 #include <stdgpu/algorithm.h>
 #include <stdgpu/contract.h>
 #include <stdgpu/iterator.h>
@@ -43,14 +45,14 @@ vector<T, Allocator>::createDeviceObject(ExecutionPolicy&& policy, const index_t
 
     vector<T, Allocator> result(
             mutex_array<mutex_default_type, mutex_array_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     capacity,
                     mutex_array_allocator_type(allocator)),
             bitset<bitset_default_type, bitset_allocator_type>::createDeviceObject(
-                    std::forward<ExecutionPolicy>(policy),
+                    std::decay_t<ExecutionPolicy>{ policy },
                     capacity,
                     bitset_allocator_type(allocator)),
-            atomic<int, atomic_allocator_type>::createDeviceObject(std::forward<ExecutionPolicy>(policy),
+            atomic<int, atomic_allocator_type>::createDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                    atomic_allocator_type(allocator)),
             allocator);
     result._data = allocator_traits<allocator_type>::allocate(result._allocator, capacity);
@@ -73,7 +75,7 @@ vector<T, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, vector<T, Al
 {
     if (!detail::is_destroy_optimizable<value_type>())
     {
-        device_object.clear(std::forward<ExecutionPolicy>(policy));
+        device_object.clear(std::decay_t<ExecutionPolicy>{ policy });
     }
 
     allocator_traits<allocator_type>::deallocate(device_object._allocator,
@@ -81,9 +83,9 @@ vector<T, Allocator>::destroyDeviceObject(ExecutionPolicy&& policy, vector<T, Al
                                                  device_object.capacity());
     device_object._data = nullptr;
     mutex_array<mutex_default_type, mutex_array_allocator_type>::destroyDeviceObject(
-            std::forward<ExecutionPolicy>(policy),
+            std::decay_t<ExecutionPolicy>{ policy },
             device_object._locks);
-    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy),
+    bitset<bitset_default_type, bitset_allocator_type>::destroyDeviceObject(std::decay_t<ExecutionPolicy>{ policy },
                                                                             device_object._occupied);
     atomic<int, atomic_allocator_type>::destroyDeviceObject(std::forward<ExecutionPolicy>(policy), device_object._size);
 }
@@ -344,8 +346,8 @@ template <typename ExecutionPolicy, typename T, typename Allocator>
 void
 vector_clear_iota(ExecutionPolicy&& policy, vector<T, Allocator>& v, const T& value)
 {
-    iota(std::forward<ExecutionPolicy>(policy), device_begin(v.data()), device_end(v.data()), value);
-    v._occupied.set(std::forward<ExecutionPolicy>(policy));
+    iota(std::decay_t<ExecutionPolicy>{ policy }, device_begin(v.data()), device_end(v.data()), value);
+    v._occupied.set(std::decay_t<ExecutionPolicy>{ policy });
     v._size.store(std::forward<ExecutionPolicy>(policy), static_cast<int>(v.capacity()));
 }
 
@@ -370,14 +372,14 @@ vector<T, Allocator>::insert(ExecutionPolicy&& policy,
                              ValueIterator begin,
                              ValueIterator end)
 {
-    if (position != device_end(std::forward<ExecutionPolicy>(policy)))
+    if (position != device_end(std::decay_t<ExecutionPolicy>{ policy }))
     {
         printf("stdgpu::vector::insert : Position not equal to device_end()\n");
         return;
     }
 
     index_t N = static_cast<index_t>(end - begin);
-    index_t new_size = size(std::forward<ExecutionPolicy>(policy)) + N;
+    index_t new_size = size(std::decay_t<ExecutionPolicy>{ policy }) + N;
 
     if (new_size > capacity())
     {
@@ -388,7 +390,7 @@ vector<T, Allocator>::insert(ExecutionPolicy&& policy,
         return;
     }
 
-    for_each_index(std::forward<ExecutionPolicy>(policy),
+    for_each_index(std::decay_t<ExecutionPolicy>{ policy },
                    N,
                    detail::vector_insert<T, Allocator, ValueIterator, true>(*this, size(), begin));
 
@@ -408,14 +410,14 @@ template <typename ExecutionPolicy,
 inline void
 vector<T, Allocator>::erase(ExecutionPolicy&& policy, device_ptr<const T> begin, device_ptr<const T> end)
 {
-    if (end != device_end(std::forward<ExecutionPolicy>(policy)))
+    if (end != device_end(std::decay_t<ExecutionPolicy>{ policy }))
     {
         printf("stdgpu::vector::erase : End iterator not equal to device_end()\n");
         return;
     }
 
     index_t N = static_cast<index_t>(end - begin);
-    index_t new_size = size(std::forward<ExecutionPolicy>(policy)) - N;
+    index_t new_size = size(std::decay_t<ExecutionPolicy>{ policy }) - N;
 
     if (new_size < 0)
     {
@@ -424,7 +426,9 @@ vector<T, Allocator>::erase(ExecutionPolicy&& policy, device_ptr<const T> begin,
         return;
     }
 
-    for_each_index(std::forward<ExecutionPolicy>(policy), N, detail::vector_erase<T, Allocator, true>(*this, new_size));
+    for_each_index(std::decay_t<ExecutionPolicy>{ policy },
+                   N,
+                   detail::vector_erase<T, Allocator, true>(*this, new_size));
 
     _size.store(std::forward<ExecutionPolicy>(policy), static_cast<int>(new_size));
 }
@@ -569,25 +573,25 @@ template <typename ExecutionPolicy,
 inline void
 vector<T, Allocator>::clear(ExecutionPolicy&& policy)
 {
-    if (empty(std::forward<ExecutionPolicy>(policy)))
+    if (empty(std::decay_t<ExecutionPolicy>{ policy }))
     {
         return;
     }
 
     if (!detail::is_destroy_optimizable<value_type>())
     {
-        const index_t current_size = size(std::forward<ExecutionPolicy>(policy));
+        const index_t current_size = size(std::decay_t<ExecutionPolicy>{ policy });
 
-        detail::unoptimized_destroy(std::forward<ExecutionPolicy>(policy),
+        detail::unoptimized_destroy(std::decay_t<ExecutionPolicy>{ policy },
                                     stdgpu::device_begin(_data),
                                     stdgpu::device_begin(_data) + current_size);
     }
 
-    _occupied.reset(std::forward<ExecutionPolicy>(policy));
+    _occupied.reset(std::decay_t<ExecutionPolicy>{ policy });
 
-    _size.store(std::forward<ExecutionPolicy>(policy), static_cast<int>(0));
+    _size.store(std::decay_t<ExecutionPolicy>{ policy }, static_cast<int>(0));
 
-    STDGPU_ENSURES(empty(std::forward<ExecutionPolicy>(policy)));
+    STDGPU_ENSURES(empty(std::decay_t<ExecutionPolicy>{ policy }));
     STDGPU_ENSURES(valid(std::forward<ExecutionPolicy>(policy)));
 }
 
@@ -610,9 +614,9 @@ vector<T, Allocator>::valid(ExecutionPolicy&& policy) const
         return true;
     }
 
-    return (size_valid(std::forward<ExecutionPolicy>(policy)) &&
-            occupied_count_valid(std::forward<ExecutionPolicy>(policy)) &&
-            _locks.valid(std::forward<ExecutionPolicy>(policy)));
+    return (size_valid(std::decay_t<ExecutionPolicy>{ policy }) &&
+            occupied_count_valid(std::decay_t<ExecutionPolicy>{ policy }) &&
+            _locks.valid(std::decay_t<ExecutionPolicy>{ policy }));
 }
 
 template <typename T, typename Allocator>
@@ -758,7 +762,7 @@ template <typename ExecutionPolicy>
 bool
 vector<T, Allocator>::occupied_count_valid(ExecutionPolicy&& policy) const
 {
-    index_t size_count = size(std::forward<ExecutionPolicy>(policy));
+    index_t size_count = size(std::decay_t<ExecutionPolicy>{ policy });
     index_t size_sum = _occupied.count(std::forward<ExecutionPolicy>(policy));
 
     return (size_count == size_sum);