Test updates of CCCL (thrust, cub, libcudacxx) to 2.1.0.

CMakeLists.txt

@@ -15,6 +15,9 @@
 #=============================================================================
 cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR)
 
+set(rapids-cmake-repo bdice/rapids-cmake)
+set(rapids-cmake-branch cccl-update-2.1.0)
+
 if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)
     file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-23.10/RAPIDS.cmake
          ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)

examples/static_map/custom_type_example.cu

@@ -22,6 +22,8 @@
 #include <thrust/logical.h>
 #include <thrust/transform.h>
 
+#include <cuda/functional>
+
 // User-defined key type
 #if !defined(CUCO_HAS_INDEPENDENT_THREADS)
 struct custom_key_type {
@@ -88,7 +90,8 @@ int main(void)
   // Create an iterator of input key/value pairs
   auto pairs_begin = thrust::make_transform_iterator(
     thrust::make_counting_iterator<int32_t>(0),
-    [] __device__(auto i) { return cuco::make_pair(custom_key_type{i}, custom_value_type{i}); });
+    cuda::proclaim_return_type<cuco::pair<custom_key_type, custom_value_type>>(
+      [] __device__(auto i) { return cuco::make_pair(custom_key_type{i}, custom_value_type{i}); }));
 
   // Construct a map with 100,000 slots using the given empty key/value sentinels. Note the
   // capacity is chosen knowing we will insert 80,000 keys, for an load factor of 80%.
@@ -101,7 +104,8 @@ int main(void)
   // Reproduce inserted keys
   auto insert_keys =
     thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    [] __device__(auto i) { return custom_key_type{i}; });
+                                    cuda::proclaim_return_type<custom_key_type>(
+                                      [] __device__(auto i) { return custom_key_type{i}; }));
 
   thrust::device_vector<bool> contained(num_pairs);
 
@@ -114,7 +118,9 @@ int main(void)
 
   // All inserted keys are contained
   assert(
-    thrust::all_of(contained.begin(), contained.end(), [] __device__(auto const& b) { return b; }));
+    thrust::all_of(contained.begin(),
+                   contained.end(),
+                   cuda::proclaim_return_type<bool>([] __device__(auto const& b) { return b; })));
 
   return 0;
 }

tests/dynamic_map/unique_sequence_test.cu

@@ -27,6 +27,8 @@
 #include <thrust/sequence.h>
 #include <thrust/tuple.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
@@ -48,9 +50,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
   thrust::sequence(thrust::device, d_keys.begin(), d_keys.end());
   thrust::sequence(thrust::device, d_values.begin(), d_values.end());
 
-  auto pairs_begin =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
+  auto pairs_begin = thrust::make_transform_iterator(
+    thrust::make_counting_iterator<int>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); }));
 
   thrust::device_vector<Value> d_results(num_keys);
   thrust::device_vector<bool> d_contained(num_keys);
@@ -62,17 +65,20 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
     map.find(d_keys.begin(), d_keys.end(), d_results.begin());
     auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin()));
 
-    REQUIRE(cuco::test::all_of(zip, zip + num_keys, [] __device__(auto const& p) {
-      return thrust::get<0>(p) == thrust::get<1>(p);
-    }));
+    REQUIRE(cuco::test::all_of(
+      zip, zip + num_keys, cuda::proclaim_return_type<bool>([] __device__(auto const& p) {
+        return thrust::get<0>(p) == thrust::get<1>(p);
+      })));
   }
 
   SECTION("All non-inserted keys-value pairs should have the empty sentinel value recovered")
   {
     map.find(d_keys.begin(), d_keys.end(), d_results.begin());
 
     REQUIRE(cuco::test::all_of(
-      d_results.begin(), d_results.end(), [] __device__(auto const& p) { return p == -1; }));
+      d_results.begin(),
+      d_results.end(),
+      cuda::proclaim_return_type<bool>([] __device__(auto const& p) { return p == -1; })));
   }
 
   SECTION("All inserted keys-value pairs should be contained")

tests/static_map/custom_type_test.cu

@@ -25,6 +25,8 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/transform.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 #include <tuple>
@@ -123,17 +125,18 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
                     thrust::counting_iterator<int>(0),
                     thrust::counting_iterator<int>(num),
                     insert_keys.begin(),
-                    [] __device__(auto i) { return Key{i}; });
+                    cuda::proclaim_return_type<Key>([] __device__(auto i) { return Key{i}; }));
 
   thrust::transform(thrust::device,
                     thrust::counting_iterator<int>(0),
                     thrust::counting_iterator<int>(num),
                     insert_values.begin(),
-                    [] __device__(auto i) { return Value{i}; });
+                    cuda::proclaim_return_type<Value>([] __device__(auto i) { return Value{i}; }));
 
-  auto insert_pairs =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
+  auto insert_pairs = thrust::make_transform_iterator(
+    thrust::make_counting_iterator<int>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); }));
 
   SECTION("All inserted keys-value pairs should be correctly recovered during find")
   {
@@ -151,9 +154,9 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
     REQUIRE(cuco::test::equal(insert_values.begin(),
                               insert_values.end(),
                               found_values.begin(),
-                              [] __device__(Value lhs, Value rhs) {
+                              cuda::proclaim_return_type<bool>([] __device__(Value lhs, Value rhs) {
                                 return std::tie(lhs.f, lhs.s) == std::tie(rhs.f, rhs.s);
-                              }));
+                              })));
   }
 
   SECTION("All inserted keys-value pairs should be contained")
@@ -175,7 +178,7 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
       insert_pairs,
       insert_pairs + num,
       thrust::counting_iterator<int>(0),
-      [] __device__(auto const& key) { return (key % 2) == 0; },
+      cuda::proclaim_return_type<bool>([] __device__(auto const& key) { return (key % 2) == 0; }),
       hash_custom_key{},
       custom_key_equals{});
 
@@ -187,12 +190,13 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
                  hash_custom_key{},
                  custom_key_equals{});
 
-    REQUIRE(cuco::test::equal(contained.begin(),
-                              contained.end(),
-                              thrust::counting_iterator<int>(0),
-                              [] __device__(auto const& idx_contained, auto const& idx) {
-                                return ((idx % 2) == 0) == idx_contained;
-                              }));
+    REQUIRE(cuco::test::equal(
+      contained.begin(),
+      contained.end(),
+      thrust::counting_iterator<int>(0),
+      cuda::proclaim_return_type<bool>([] __device__(auto const& idx_contained, auto const& idx) {
+        return ((idx % 2) == 0) == idx_contained;
+      })));
   }
 
   SECTION("Non-inserted keys-value pairs should not be contained")
@@ -212,19 +216,23 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
     map.insert(insert_pairs, insert_pairs + num, hash_custom_key{}, custom_key_equals{});
     auto view = map.get_device_view();
     REQUIRE(cuco::test::all_of(
-      insert_pairs, insert_pairs + num, [view] __device__(cuco::pair<Key, Value> const& pair) {
+      insert_pairs,
+      insert_pairs + num,
+      cuda::proclaim_return_type<bool>([view] __device__(cuco::pair<Key, Value> const& pair) {
         return view.contains(pair.first, hash_custom_key{}, custom_key_equals{});
-      }));
+      })));
   }
 
   SECTION("Inserting unique keys should return insert success.")
   {
     auto m_view = map.get_device_mutable_view();
     REQUIRE(cuco::test::all_of(insert_pairs,
                                insert_pairs + num,
-                               [m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
-                                 return m_view.insert(pair, hash_custom_key{}, custom_key_equals{});
-                               }));
+                               cuda::proclaim_return_type<bool>(
+                                 [m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
+                                   return m_view.insert(
+                                     pair, hash_custom_key{}, custom_key_equals{});
+                                 })));
   }
 
   SECTION("Cannot find any key in an empty hash map")
@@ -235,18 +243,21 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type",
       REQUIRE(cuco::test::all_of(
         insert_pairs,
         insert_pairs + num,
-        [view] __device__(cuco::pair<Key, Value> const& pair) mutable {
-          return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end();
-        }));
+        cuda::proclaim_return_type<bool>(
+          [view] __device__(cuco::pair<Key, Value> const& pair) mutable {
+            return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end();
+          })));
     }
 
     SECTION("const view")
     {
       auto const view = map.get_device_view();
       REQUIRE(cuco::test::all_of(
-        insert_pairs, insert_pairs + num, [view] __device__(cuco::pair<Key, Value> const& pair) {
+        insert_pairs,
+        insert_pairs + num,
+        cuda::proclaim_return_type<bool>([view] __device__(cuco::pair<Key, Value> const& pair) {
           return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end();
-        }));
+        })));
     }
   }
 }

tests/static_map/duplicate_keys_test.cu

@@ -27,6 +27,8 @@
 #include <thrust/sequence.h>
 #include <thrust/sort.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Duplicate keys",
@@ -49,7 +51,8 @@ TEMPLATE_TEST_CASE_SIG("Duplicate keys",
 
   auto pairs_begin = thrust::make_transform_iterator(
     thrust::make_counting_iterator<int>(0),
-    [] __device__(auto i) { return cuco::pair<Key, Value>(i / 2, i / 2); });
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i / 2, i / 2); }));
 
   thrust::device_vector<Value> d_results(num_keys);
   thrust::device_vector<bool> d_contained(num_keys);

tests/static_map/heterogeneous_lookup_test.cu

@@ -25,6 +25,8 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/transform.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 #include <tuple>
@@ -99,11 +101,13 @@ TEMPLATE_TEST_CASE("Heterogeneous lookup",
   cuco::static_map<Key, Value> map{
     capacity, cuco::empty_key<Key>{sentinel_key}, cuco::empty_value<Value>{sentinel_value}};
 
-  auto insert_pairs =
-    thrust::make_transform_iterator(thrust::counting_iterator<int>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
-  auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator<int>(0),
-                                                    [] __device__(auto i) { return ProbeKey(i); });
+  auto insert_pairs = thrust::make_transform_iterator(
+    thrust::counting_iterator<int>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); }));
+  auto probe_keys = thrust::make_transform_iterator(
+    thrust::counting_iterator<int>(0),
+    cuda::proclaim_return_type<ProbeKey>([] __device__(auto i) { return ProbeKey(i); }));
 
   SECTION("All inserted keys-value pairs should be contained")
   {

tests/static_map/key_sentinel_test.cu

@@ -22,6 +22,8 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 #define SIZE 10
@@ -51,19 +53,21 @@ TEMPLATE_TEST_CASE_SIG(
   }
   CUCO_CUDA_TRY(cudaMemcpyToSymbol(A, h_A, SIZE * sizeof(int)));
 
-  auto pairs_begin =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<T>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
+  auto pairs_begin = thrust::make_transform_iterator(
+    thrust::make_counting_iterator<T>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); }));
 
   SECTION(
     "Tests of non-CG insert: The custom `key_equal` can never be used to compare against sentinel")
   {
-    REQUIRE(cuco::test::all_of(pairs_begin,
-                               pairs_begin + num_keys,
-                               [m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
-                                 return m_view.insert(
-                                   pair, cuco::default_hash_function<Key>{}, custom_equals<Key>{});
-                               }));
+    REQUIRE(cuco::test::all_of(
+      pairs_begin,
+      pairs_begin + num_keys,
+      cuda::proclaim_return_type<bool>(
+        [m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
+          return m_view.insert(pair, cuco::default_hash_function<Key>{}, custom_equals<Key>{});
+        })));
   }
 
   SECTION(
@@ -75,10 +79,12 @@ TEMPLATE_TEST_CASE_SIG(
                custom_equals<Key>{});
     // All keys inserted via custom `key_equal` should be found
     REQUIRE(cuco::test::all_of(
-      pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair<Key, Value> const& pair) {
+      pairs_begin,
+      pairs_begin + num_keys,
+      cuda::proclaim_return_type<bool>([view] __device__(cuco::pair<Key, Value> const& pair) {
         auto const found = view.find(pair.first);
         return (found != view.end()) and
                (found->first.load() == pair.first and found->second.load() == pair.second);
-      }));
+      })));
   }
 }

tests/static_map/shared_memory_test.cu

@@ -25,6 +25,8 @@
 #include <thrust/sequence.h>
 #include <thrust/tuple.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 #include <limits>
@@ -127,9 +129,11 @@ TEMPLATE_TEST_CASE_SIG("Shared memory static map",
     auto zip = thrust::make_zip_iterator(
       thrust::make_tuple(d_keys_exist.begin(), d_keys_and_values_correct.begin()));
 
-    REQUIRE(cuco::test::all_of(zip, zip + d_keys_exist.size(), [] __device__(auto const& z) {
-      return thrust::get<0>(z) and thrust::get<1>(z);
-    }));
+    REQUIRE(cuco::test::all_of(zip,
+                               zip + d_keys_exist.size(),
+                               cuda::proclaim_return_type<bool>([] __device__(auto const& z) {
+                                 return thrust::get<0>(z) and thrust::get<1>(z);
+                               })));
   }
 
   SECTION("No key is found before insertion.")

tests/static_map/stream_test.cu

@@ -27,6 +27,8 @@
 #include <thrust/sequence.h>
 #include <thrust/tuple.h>
 
+#include <cuda/functional>
+
 #include <catch2/catch_template_test_macros.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
@@ -53,9 +55,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
   thrust::sequence(thrust::device, d_keys.begin(), d_keys.end());
   thrust::sequence(thrust::device, d_values.begin(), d_values.end());
 
-  auto pairs_begin =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
+  auto pairs_begin = thrust::make_transform_iterator(
+    thrust::make_counting_iterator<int>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>(
+      [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); }));
 
   auto hash_fn  = cuco::default_hash_function<Key>{};
   auto equal_fn = thrust::equal_to<Value>{};
@@ -69,11 +72,12 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
     map.find(d_keys.begin(), d_keys.end(), d_results.begin(), hash_fn, equal_fn, stream);
     auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin()));
 
-    REQUIRE(cuco::test::all_of(
-      zip,
-      zip + num_keys,
-      [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); },
-      stream));
+    REQUIRE(cuco::test::all_of(zip,
+                               zip + num_keys,
+                               cuda::proclaim_return_type<bool>([] __device__(auto const& p) {
+                                 return thrust::get<0>(p) == thrust::get<1>(p);
+                               }),
+                               stream));
   }
 
   SECTION("All inserted keys-value pairs should be contained")