Merge remote-tracking branch 'upstream/dev' into cccl-utilities

Author: PointKernel

README.md

@@ -131,16 +131,16 @@ while [ "${#args[@]}" -ne 0 ]; do
     esac
 done
 
+if [ $VERBOSE ]; then
+    set -x
+fi
+
 # Convert to full paths:
 HOST_COMPILER=$(which ${HOST_COMPILER})
 CUDA_COMPILER=$(which ${CUDA_COMPILER})
 # Make CUDA arch list compatible with cmake
 CUDA_ARCHS=$(echo "$CUDA_ARCHS" | tr ' ,' ';;')
 
-if [ $VERBOSE ]; then
-    set -x
-fi
-
 # Begin processing unsets after option parsing
 set -u
 
@@ -217,4 +217,4 @@ if command -v sccache >/dev/null; then
     source "./sccache_stats.sh" end
 else
     echo "sccache stats: N/A"
-fi
+fi

ci/build.sh

@@ -18,8 +18,8 @@
 
 #include <cub/block/block_reduce.cuh>
 #include <cuda/std/atomic>
+#include <cuda/std/functional>
 #include <thrust/device_vector.h>
-#include <thrust/logical.h>
 #include <thrust/transform.h>
 
 #include <cmath>

examples/static_map/count_by_key_example.cu

@@ -16,10 +16,10 @@
 
 #include <cuco/static_map.cuh>
 
+#include <cuda/std/functional>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/iterator/zip_iterator.h>
-#include <thrust/logical.h>
 #include <thrust/sequence.h>
 #include <thrust/tuple.h>
 

examples/static_map/device_ref_example.cu

@@ -18,6 +18,7 @@
 #include <cuco/storage.cuh>
 
 #include <cuda/std/array>
+#include <cuda/std/functional>
 #include <thrust/device_vector.h>
 #include <thrust/reduce.h>
 #include <thrust/scan.h>

examples/static_set/device_subsets_example.cu

@@ -16,6 +16,8 @@
 
 #include <cuco/static_set.cuh>
 
+#include <cuda/std/functional>
+
 #include <cooperative_groups.h>
 
 /**

examples/static_set/shared_memory_example.cu

@@ -389,7 +389,7 @@ CUCO_KERNEL void find(InputIt first,
       auto submap_view = submap_views[i];
       auto found       = submap_view.find(key, hash, key_equal);
       if (found != submap_view.end()) {
-        found_value = found->second;
+        found_value = found->second.load(cuda::std::memory_order_relaxed);
         break;
       }
     }
@@ -466,7 +466,7 @@ CUCO_KERNEL void find(InputIt first,
       auto submap_view = submap_views[i];
       auto found       = submap_view.find(tile, key, hash, key_equal);
       if (found != submap_view.end()) {
-        found_value = found->second;
+        found_value = found->second.load(cuda::std::memory_order_relaxed);
         break;
       }
     }

include/cuco/detail/dynamic_map_kernels.cuh

@@ -27,7 +27,6 @@
 #include <cuda/std/functional>
 #include <cuda/std/iterator>
 #include <cuda/std/type_traits>
-#include <thrust/distance.h>
 #include <thrust/execution_policy.h>
 #include <thrust/logical.h>
 #include <thrust/reduce.h>
@@ -395,7 +394,7 @@ class open_addressing_ref_impl {
           if (eq_res == detail::equal_result::EQUAL) { return false; }
         }
         if (eq_res == detail::equal_result::AVAILABLE) {
-          auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
+          auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
           switch (attempt_insert((storage_ref_.data() + *probing_iter)->data() + intra_bucket_index,
                                  slot_content,
                                  val)) {
@@ -701,7 +700,7 @@ class open_addressing_ref_impl {
         if (eq_res == detail::equal_result::EMPTY) { return false; }
         // Key exists, return true if successfully deleted
         if (eq_res == detail::equal_result::EQUAL) {
-          auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
+          auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
           switch (attempt_insert_stable(
             (storage_ref_.data() + *probing_iter)->data() + intra_bucket_index,
             slot_content,

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -21,6 +21,7 @@
 #include <cuco/operator.hpp>
 
 #include <cuda/atomic>
+#include <cuda/std/iterator>
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
 
@@ -505,7 +506,7 @@ class operator_impl<
       for (auto& slot_content : bucket_slots) {
         auto const eq_res =
           ref_.impl_.predicate_.operator()<is_insert::YES>(key, slot_content.first);
-        auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
+        auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
         auto slot_ptr = (storage_ref.data() + *probing_iter)->data() + intra_bucket_index;
 
         // If the key is already in the container, update the payload and return
@@ -875,7 +876,7 @@ class operator_impl<
       for (auto& slot_content : bucket_slots) {
         auto const eq_res =
           ref_.impl_.predicate_.operator()<is_insert::YES>(key, slot_content.first);
-        auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
+        auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
         auto slot_ptr = (storage_ref.data() + *probing_iter)->data() + intra_bucket_index;
 
         // If the key is already in the container, update the payload and return

include/cuco/detail/static_map/static_map_ref.inl

@@ -22,6 +22,7 @@
 
 #include <cuda/std/cmath>
 #include <cuda/std/functional>  // TODO include <cuda/std/algorithm> instead once available
+#include <cuda/std/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/span>
 #include <thrust/device_vector.h>
@@ -284,7 +285,7 @@ class key_generator {
       thrust::sequence(exec_policy, out_begin, out_end, value_type{0});
       thrust::shuffle(exec_policy, out_begin, out_end, this->rng_);
     } else if constexpr (std::is_same_v<Dist, distribution::uniform>) {
-      size_t num_keys = thrust::distance(out_begin, out_end);
+      size_t num_keys = cuda::std::distance(out_begin, out_end);
       size_t seed     = this->rng_();
 
       thrust::transform(exec_policy,
@@ -293,7 +294,7 @@ class key_generator {
                         out_begin,
                         detail::generate_uniform_fn<value_type, Dist, RNG>{num_keys, dist, seed});
     } else if constexpr (std::is_same_v<Dist, distribution::gaussian>) {
-      size_t num_keys = thrust::distance(out_begin, out_end);
+      size_t num_keys = cuda::std::distance(out_begin, out_end);
 
       thrust::counting_iterator<size_t> seq(this->rng_());
 
@@ -369,7 +370,7 @@ class key_generator {
     CUCO_EXPECTS(keep_prob >= 0.0 and keep_prob <= 1.0, "Probability needs to be between 0 and 1");
 
     if (keep_prob < 1.0) {
-      size_t const num_keys = thrust::distance(begin, end);
+      size_t const num_keys = cuda::std::distance(begin, end);
 
       thrust::counting_iterator<size_t> seeds(rng_());