Merge branch 'dev' into fea/support-for-span-like-arrow-filter-policy

benchmarks/benchmark_utils.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -29,10 +29,10 @@ auto dist_from_state(nvbench::state const& state)
   if constexpr (std::is_same_v<Dist, cuco::utility::distribution::unique>) {
     return Dist{};
   } else if constexpr (std::is_same_v<Dist, cuco::utility::distribution::uniform>) {
-    auto const multiplicity = state.get_int64_or_default("Multiplicity", defaults::MULTIPLICITY);
+    auto const multiplicity = state.get_int64("Multiplicity");
     return Dist{multiplicity};
   } else if constexpr (std::is_same_v<Dist, cuco::utility::distribution::gaussian>) {
-    auto const skew = state.get_float64_or_default("Skew", defaults::SKEW);
+    auto const skew = state.get_float64("Skew");
     return Dist{skew};
   } else {
     CUCO_FAIL("Unexpected distribution type");

benchmarks/bloom_filter/add_bench.cu

@@ -49,7 +49,7 @@ void bloom_filter_add(nvbench::state& state,
 
   auto const num_keys       = state.get_int64("NumInputs");
   auto const filter_size_mb = state.get_int64("FilterSizeMB");
-  auto const pattern_bits   = state.get_int64_or_default("PatternBits", WordsPerBlock);
+  auto const pattern_bits   = WordsPerBlock;
 
   try {
     auto const policy = policy_type{static_cast<uint32_t>(pattern_bits)};
@@ -171,4 +171,4 @@ NVBENCH_BENCH_TYPES(arrow_bloom_filter_add,
   .set_type_axes_names({"Key", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
   .add_int64_axis("NumInputs", {defaults::BF_N})
-  .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE);
+  .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE);

benchmarks/bloom_filter/contains_bench.cu

@@ -51,7 +51,7 @@ void bloom_filter_contains(
 
   auto const num_keys       = state.get_int64("NumInputs");
   auto const filter_size_mb = state.get_int64("FilterSizeMB");
-  auto const pattern_bits   = state.get_int64_or_default("PatternBits", WordsPerBlock);
+  auto const pattern_bits   = WordsPerBlock;
 
   try {
     auto const policy = policy_type{static_cast<uint32_t>(pattern_bits)};
@@ -181,4 +181,4 @@ NVBENCH_BENCH_TYPES(arrow_bloom_filter_contains,
   .set_type_axes_names({"Key", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
   .add_int64_axis("NumInputs", {defaults::BF_N})
-  .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE);
+  .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE);

benchmarks/static_multiset/retrieve_bench.cu

@@ -34,9 +34,9 @@ using namespace cuco::utility;
 template <typename Key, typename Dist>
 void static_multiset_retrieve(nvbench::state& state, nvbench::type_list<Key, Dist>)
 {
-  auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
-  auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);
-  auto const matching_rate = state.get_float64_or_default("MatchingRate", defaults::MATCHING_RATE);
+  auto const num_keys      = state.get_int64("NumInputs");
+  auto const occupancy     = state.get_float64("Occupancy");
+  auto const matching_rate = state.get_float64("MatchingRate");
 
   std::size_t const size = num_keys / occupancy;
 
@@ -68,20 +68,29 @@ NVBENCH_BENCH_TYPES(static_multiset_retrieve,
   .set_name("static_multiset_retrieve_uniform_occupancy")
   .set_type_axes_names({"Key", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
-  .add_float64_axis("Occupancy", defaults::OCCUPANCY_RANGE);
+  .add_int64_axis("NumInputs", {defaults::N})
+  .add_float64_axis("Occupancy", defaults::OCCUPANCY_RANGE)
+  .add_float64_axis("MatchingRate", {defaults::MATCHING_RATE})
+  .add_int64_axis("Multiplicity", {defaults::MULTIPLICITY});
 
 NVBENCH_BENCH_TYPES(static_multiset_retrieve,
                     NVBENCH_TYPE_AXES(defaults::KEY_TYPE_RANGE,
                                       nvbench::type_list<distribution::uniform>))
   .set_name("static_multiset_retrieve_uniform_matching_rate")
   .set_type_axes_names({"Key", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
-  .add_float64_axis("MatchingRate", defaults::MATCHING_RATE_RANGE);
+  .add_int64_axis("NumInputs", {defaults::N})
+  .add_float64_axis("Occupancy", {defaults::OCCUPANCY})
+  .add_float64_axis("MatchingRate", defaults::MATCHING_RATE_RANGE)
+  .add_int64_axis("Multiplicity", {defaults::MULTIPLICITY});
 
 NVBENCH_BENCH_TYPES(static_multiset_retrieve,
                     NVBENCH_TYPE_AXES(defaults::KEY_TYPE_RANGE,
                                       nvbench::type_list<distribution::uniform>))
   .set_name("static_multiset_retrieve_uniform_multiplicity")
   .set_type_axes_names({"Key", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
+  .add_int64_axis("NumInputs", {defaults::N})
+  .add_float64_axis("Occupancy", {defaults::OCCUPANCY})
+  .add_float64_axis("MatchingRate", {defaults::MATCHING_RATE})
   .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE);

tests/CMakeLists.txt

@@ -75,9 +75,11 @@ ConfigureTest(STATIC_SET_TEST
 # - static_map tests ------------------------------------------------------------------------------
 ConfigureTest(STATIC_MAP_TEST
     static_map/capacity_test.cu
+    static_map/contains_test.cu
     static_map/custom_type_test.cu
     static_map/duplicate_keys_test.cu
     static_map/erase_test.cu
+    static_map/find_test.cu
     static_map/for_each_test.cu
     static_map/hash_test.cu
     static_map/heterogeneous_lookup_test.cu
@@ -87,7 +89,6 @@ ConfigureTest(STATIC_MAP_TEST
     static_map/key_sentinel_test.cu
     static_map/shared_memory_test.cu
     static_map/stream_test.cu
-    static_map/unique_sequence_test.cu
     static_map/rehash_test.cu)
 
 ###################################################################################################

tests/bloom_filter/arrow_policy_test.cu

@@ -150,8 +150,12 @@ void test_filter_bitset(Filter& filter, size_t num_keys)
     })));
 }
 
-TEMPLATE_TEST_CASE_SIG(
-  "Arrow filter policy bitset validation", "", (class Key), (int32_t), (int64_t), (float))
+TEMPLATE_TEST_CASE_SIG("bloom_filter arrow filter policy bitset validation",
+                       "",
+                       (class Key),
+                       (int32_t),
+                       (int64_t),
+                       (float))
 {
   // Get test settings
   auto const [sub_filters, num_keys] = get_arrow_filter_test_settings<Key>();

tests/bloom_filter/unique_sequence_test.cu

@@ -85,7 +85,7 @@ void test_unique_sequence(Filter& filter, size_type num_keys)
 }
 
 TEMPLATE_TEST_CASE_SIG(
-  "Unique sequence with default policy",
+  "bloom_filter default policy tests",
   "",
   ((class Key, class Policy), Key, Policy),
   (int32_t, cuco::default_filter_policy<cuco::xxhash_64<int32_t>, uint32_t, 1>),
@@ -105,7 +105,7 @@ TEMPLATE_TEST_CASE_SIG(
   test_unique_sequence(filter, num_keys);
 }
 
-TEMPLATE_TEST_CASE_SIG("Unique sequence with arrow policy",
+TEMPLATE_TEST_CASE_SIG("bloom_filter arrow policy tests",
                        "",
                        ((class Key, class Policy), Key, Policy),
                        (int32_t, cuco::arrow_filter_policy<int32_t>),

tests/dynamic_bitset/find_next_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -37,7 +37,7 @@ __global__ void find_next_kernel(BitsetRef ref, size_type num_elements, OutputIt
 
 extern bool modulo_bitgen(uint64_t i);  // Defined in get_test.cu
 
-TEST_CASE("Find next set test", "")
+TEST_CASE("dynamic_bitset find next set test", "")
 {
   cuco::experimental::detail::dynamic_bitset bv;
 

tests/dynamic_bitset/get_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -37,7 +37,7 @@ __global__ void test_kernel(BitsetRef ref, size_type num_elements, OutputIt outp
 
 bool modulo_bitgen(uint64_t i) { return i % 7 == 0; }
 
-TEST_CASE("Get test", "")
+TEST_CASE("dynamic_bitset get test", "")
 {
   cuco::experimental::detail::dynamic_bitset bv;
 

tests/dynamic_bitset/rank_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,7 +26,7 @@
 
 extern bool modulo_bitgen(uint64_t i);  // Defined in get_test.cu
 
-TEST_CASE("Rank test", "")
+TEST_CASE("dynamic_bitset rank test", "")
 {
   cuco::experimental::detail::dynamic_bitset bv;
 

tests/dynamic_bitset/select_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -37,7 +37,7 @@ __global__ void select_false_kernel(BitsetRef ref, size_type num_elements, Outpu
 
 extern bool modulo_bitgen(uint64_t i);  // Defined in get_test.cu
 
-TEST_CASE("Select test", "")
+TEST_CASE("dynamic_bitset select test", "")
 {
   cuco::experimental::detail::dynamic_bitset bv;
 

tests/dynamic_bitset/size_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -18,7 +18,7 @@
 
 #include <catch2/catch_test_macros.hpp>
 
-TEST_CASE("Size computation", "")
+TEST_CASE("dynamic_bitset size computation test", "")
 {
   cuco::experimental::detail::dynamic_bitset bv;
   using size_type = std::size_t;

tests/dynamic_map/erase_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,7 +25,7 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
-TEMPLATE_TEST_CASE_SIG("erase key",
+TEMPLATE_TEST_CASE_SIG("dynamic_map erase tests",
                        "",
                        ((typename Key, typename Value), Key, Value),
                        (int32_t, int32_t),

tests/dynamic_map/unique_sequence_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,7 +30,7 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
-TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
+TEMPLATE_TEST_CASE_SIG("dynamic_map unique sequence tests",
                        "",
                        ((typename Key, typename Value), Key, Value),
                        (int32_t, int32_t),

tests/static_map/capacity_test.cu

@@ -18,7 +18,7 @@
 
 #include <catch2/catch_test_macros.hpp>
 
-TEST_CASE("Static map capacity", "")
+TEST_CASE("static_map capacity test", "")
 {
   using Key        = int32_t;
   using T          = int32_t;

tests/static_map/unique_sequence_test.cu → tests/static_map/contains_test.cu

@@ -33,8 +33,6 @@
 
 using size_type = int32_t;
 
-int32_t constexpr SENTINEL = -1;
-
 template <typename Map>
 void test_unique_sequence(Map& map, size_type num_keys)
 {
@@ -63,17 +61,6 @@ void test_unique_sequence(Map& map, size_type num_keys)
     REQUIRE(cuco::test::none_of(d_contained.begin(), d_contained.end(), thrust::identity{}));
   }
 
-  SECTION("Non-inserted keys have no matches")
-  {
-    thrust::device_vector<Value> d_results(num_keys);
-
-    map.find(keys_begin, keys_begin + num_keys, d_results.begin());
-    auto zip = thrust::make_zip_iterator(thrust::make_tuple(
-      d_results.begin(), thrust::constant_iterator<Key>{map.empty_key_sentinel()}));
-
-    REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal));
-  }
-
   SECTION("All conditionally inserted keys should be contained")
   {
     auto const inserted = map.insert_if(
@@ -92,7 +79,6 @@ void test_unique_sequence(Map& map, size_type num_keys)
   }
 
   map.insert(pairs_begin, pairs_begin + num_keys);
-  REQUIRE(map.size() == num_keys);
 
   SECTION("All inserted keys should be contained.")
   {
@@ -115,37 +101,6 @@ void test_unique_sequence(Map& map, size_type num_keys)
     REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal));
   }
 
-  SECTION("All inserted keys should be correctly recovered during find")
-  {
-    thrust::device_vector<Value> d_results(num_keys);
-
-    map.find(keys_begin, keys_begin + num_keys, d_results.begin());
-    auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), keys_begin));
-
-    REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal));
-  }
-
-  SECTION("Conditional find should return valid values on even inputs.")
-  {
-    auto found_results = thrust::device_vector<Key>(num_keys);
-    auto gold_fn       = cuda::proclaim_return_type<Value>([] __device__(auto const& i) {
-      return i % 2 == 0 ? static_cast<Value>(i) : Value{SENTINEL};
-    });
-
-    map.find_if(keys_begin,
-                keys_begin + num_keys,
-                thrust::counting_iterator<std::size_t>{0},
-                is_even,
-                found_results.begin());
-
-    REQUIRE(cuco::test::equal(
-      found_results.begin(),
-      found_results.end(),
-      thrust::make_transform_iterator(thrust::counting_iterator<Key>{0}, gold_fn),
-      cuda::proclaim_return_type<bool>(
-        [] __device__(auto const& found, auto const& gold) { return found == gold; })));
-  }
-
   SECTION("All inserted key-values should be properly retrieved")
   {
     thrust::device_vector<Value> d_values(num_keys);
@@ -163,33 +118,23 @@ void test_unique_sequence(Map& map, size_type num_keys)
 }
 
 TEMPLATE_TEST_CASE_SIG(
-  "static_map: unique sequence",
+  "static_map: contains + retrieve_all tests",
   "",
   ((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 size_type num_keys{400};
-  constexpr size_type gold_capacity = CGSize == 1 ? 422   // 211 x 1 x 2
-                                                  : 412;  // 103 x 2 x 2
 
   // XXX: testing static extent is intended, DO NOT CHANGE
   using extent_type = cuco::extent<size_type, num_keys>;
@@ -206,9 +151,7 @@ TEMPLATE_TEST_CASE_SIG(
                               probe,
                               cuco::cuda_allocator<cuda::std::byte>,
                               cuco::storage<2>>{
-    extent_type{}, cuco::empty_key<Key>{SENTINEL}, cuco::empty_value<Value>{SENTINEL}};
-
-  REQUIRE(map.capacity() == gold_capacity);
+    extent_type{}, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
   test_unique_sequence(map, num_keys);
 }

tests/static_map/custom_type_test.cu

@@ -98,7 +98,7 @@ struct custom_key_equals {
   }
 };
 
-TEMPLATE_TEST_CASE_SIG("User defined key and value type",
+TEMPLATE_TEST_CASE_SIG("static_map custom key and value type tests",
                        "",
                        ((typename Key, typename Value), Key, Value),
 #if defined(CUCO_HAS_INDEPENDENT_THREADS)  // Key type larger than 8B only supported for sm_70 and