Migrate tests to excercise the new map (#393)

This PR migrates static map tests to test the new static map:

- Large type tests are not migrated since the new map doesn't support
keys larger than 8 bytes
- Shared memory tests require related functions to be added into ref
code thus would be in a separate PR

tests/static_map/custom_type_test.cu

@@ -27,6 +27,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cuda/functional>
+
 #include <tuple>
 
 // User-defined key type
@@ -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

@@ -29,16 +29,52 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
-TEMPLATE_TEST_CASE_SIG("Duplicate keys",
-                       "",
-                       ((typename Key, typename Value), Key, Value),
-                       (int32_t, int32_t),
-                       (int32_t, int64_t),
-                       (int64_t, int32_t),
-                       (int64_t, int64_t))
+#include <cuda/functional>
+
+using size_type = std::size_t;
+
+TEMPLATE_TEST_CASE_SIG(
+  "static_map duplicate keys",
+  "",
+  ((typename Key, typename Value, cuco::test::probe_sequence Probe, int CGSize),
+   Key,
+   Value,
+   Probe,
+   CGSize),
+  (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int64_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int64_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int64_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int64_t, int64_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int32_t, int32_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int32_t, int64_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int32_t, int32_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int32_t, int64_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int64_t, int32_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int64_t, int32_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2))
 {
-  constexpr std::size_t num_keys{500'000};
-  cuco::static_map<Key, Value> map{
+  constexpr size_type num_keys{500'000};
+
+  using probe =
+    std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                       cuco::experimental::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
+                       cuco::experimental::double_hashing<CGSize,
+                                                          cuco::murmurhash3_32<Key>,
+                                                          cuco::murmurhash3_32<Key>>>;
+
+  auto map = cuco::experimental::static_map<Key,
+                                            Value,
+                                            cuco::experimental::extent<size_type>,
+                                            cuda::thread_scope_device,
+                                            thrust::equal_to<Key>,
+                                            probe,
+                                            cuco::cuda_allocator<std::byte>,
+                                            cuco::experimental::storage<2>>{
     num_keys * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
   thrust::device_vector<Key> d_keys(num_keys);
@@ -49,7 +85,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);
@@ -68,7 +105,7 @@ TEMPLATE_TEST_CASE_SIG("Duplicate keys",
 
     map.insert(pairs_begin, pairs_begin + num_keys);
 
-    auto const num_entries = map.get_size();
+    auto const num_entries = map.size();
     REQUIRE(num_entries == gold);
 
     auto [key_out_end, value_out_end] =

tests/static_map/heterogeneous_lookup_test.cu

@@ -27,6 +27,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cuda/functional>
+
 #include <tuple>
 
 // insert key type
@@ -115,8 +117,9 @@ TEMPLATE_TEST_CASE_SIG("Heterogeneous lookup",
   auto insert_pairs = thrust::make_transform_iterator(
     thrust::counting_iterator<int>(0),
     [] __device__(auto i) { return cuco::pair<InsertKey, Value>(i, i); });
-  auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator<int>(0),
-                                                    [] __device__(auto i) { return ProbeKey(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/insert_and_find_test.cu

@@ -26,6 +26,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cuda/functional>
+
 static constexpr int Iters = 10'000;
 
 template <typename Ref>
@@ -129,7 +131,8 @@ TEMPLATE_TEST_CASE_SIG(
   thrust::sequence(thrust::device, d_keys.begin(), d_keys.end());
   map.find(d_keys.begin(), d_keys.end(), d_values.begin());
 
-  REQUIRE(cuco::test::all_of(d_values.begin(), d_values.end(), [] __device__(Value v) {
-    return v == (Blocks * Threads) / CGSize;
-  }));
+  REQUIRE(cuco::test::all_of(
+    d_values.begin(), d_values.end(), cuda::proclaim_return_type<bool>([] __device__(Value v) {
+      return v == (Blocks * Threads) / CGSize;
+    })));
 }

tests/static_map/insert_or_assign_test.cu

@@ -27,6 +27,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cuda/functional>
+
 using size_type = std::size_t;
 
 template <typename Map>
@@ -36,9 +38,11 @@ __inline__ void test_insert_or_assign(Map& map, size_type num_keys)
   using Value = typename Map::mapped_type;
 
   // Insert pairs
-  auto pairs_begin =
-    thrust::make_transform_iterator(thrust::counting_iterator<size_type>(0),
-                                    [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
+  auto pairs_begin = thrust::make_transform_iterator(
+    thrust::counting_iterator<size_type>(0),
+    cuda::proclaim_return_type<cuco::pair<Key, Value>>([] __device__(auto i) {
+      return cuco::pair<Key, Value>{i, i};
+    }));
 
   auto const initial_size = map.insert(pairs_begin, pairs_begin + num_keys);
   REQUIRE(initial_size == num_keys);  // all keys should be inserted
@@ -58,8 +62,9 @@ __inline__ void test_insert_or_assign(Map& map, size_type num_keys)
   thrust::device_vector<Key> d_values(num_keys);
   map.retrieve_all(d_keys.begin(), d_values.begin());
 
-  auto gold_values_begin = thrust::make_transform_iterator(thrust::counting_iterator<size_type>(0),
-                                                           [] __device__(auto i) { return i * 2; });
+  auto gold_values_begin = thrust::make_transform_iterator(
+    thrust::counting_iterator<size_type>(0),
+    cuda::proclaim_return_type<size_type>([] __device__(auto i) { return i * 2; }));
 
   thrust::sort(thrust::device, d_values.begin(), d_values.end());
   REQUIRE(cuco::test::equal(

tests/static_map/key_sentinel_test.cu

@@ -24,12 +24,14 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cuda/functional>
+
 #define SIZE 10
 __device__ int A[SIZE];
 
 template <typename T>
 struct custom_equals {
-  __device__ bool operator()(T lhs, T rhs) { return A[lhs] == A[rhs]; }
+  __device__ bool operator()(T lhs, T rhs) const { return A[lhs] == A[rhs]; }
 };
 
 TEMPLATE_TEST_CASE_SIG(
@@ -39,46 +41,51 @@ TEMPLATE_TEST_CASE_SIG(
   using Value = T;
 
   constexpr std::size_t num_keys{SIZE};
-  cuco::static_map<Key, Value> map{
-    SIZE * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
+  auto map = cuco::experimental::static_map{
+    SIZE * 2,
+    cuco::empty_key<Key>{-1},
+    cuco::empty_value<Value>{-1},
+    custom_equals<Key>{},
+    cuco::experimental::linear_probing<1, cuco::default_hash_function<Key>>{}};
 
-  auto m_view = map.get_device_mutable_view();
-  auto view   = map.get_device_view();
+  auto insert_ref = map.ref(cuco::experimental::op::insert);
+  auto find_ref   = map.ref(cuco::experimental::op::find);
 
   int h_A[SIZE];
   for (int i = 0; i < SIZE; i++) {
     h_A[i] = i;
   }
   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>(
+                           [insert_ref] __device__(cuco::pair<Key, Value> const& pair) mutable {
+                             return insert_ref.insert(pair);
+                           })));
   }
 
   SECTION(
     "Tests of CG insert: The custom `key_equal` can never be used to compare against sentinel")
   {
-    map.insert(pairs_begin,
-               pairs_begin + num_keys,
-               cuco::default_hash_function<Key>{},
-               custom_equals<Key>{});
+    map.insert(pairs_begin, pairs_begin + num_keys);
     // 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) {
-        auto const found = view.find(pair.first);
-        return (found != view.end()) and
-               (found->first.load() == pair.first and found->second.load() == pair.second);
-      }));
+      pairs_begin,
+      pairs_begin + num_keys,
+      cuda::proclaim_return_type<bool>([find_ref] __device__(cuco::pair<Key, Value> const& pair) {
+        auto const found = find_ref.find(pair.first);
+        return (found != find_ref.end()) and
+               (found->first == pair.first and found->second == pair.second);
+      })));
   }
 }