diff --git a/benchmarks/benchmark_utils.hpp b/benchmarks/benchmark_utils.hpp index 680345ebe..2f0c50c4f 100644 --- a/benchmarks/benchmark_utils.hpp +++ b/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) { return Dist{}; } else if constexpr (std::is_same_v) { - 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) { - 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"); diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu index 8e4d0b762..8b502d0d5 100644 --- a/benchmarks/bloom_filter/add_bench.cu +++ b/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(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); \ No newline at end of file + .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index bf81f5d83..3d2ed1e54 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/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(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); \ No newline at end of file + .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); diff --git a/benchmarks/static_multiset/retrieve_bench.cu b/benchmarks/static_multiset/retrieve_bench.cu index efd694946..01ce7599b 100644 --- a/benchmarks/static_multiset/retrieve_bench.cu +++ b/benchmarks/static_multiset/retrieve_bench.cu @@ -34,9 +34,9 @@ using namespace cuco::utility; template void static_multiset_retrieve(nvbench::state& state, nvbench::type_list) { - 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,7 +68,10 @@ 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, @@ -76,7 +79,10 @@ NVBENCH_BENCH_TYPES(static_multiset_retrieve, .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, @@ -84,4 +90,7 @@ NVBENCH_BENCH_TYPES(static_multiset_retrieve, .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); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index dfe478251..4a7855f72 100644 --- a/tests/CMakeLists.txt +++ b/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) ################################################################################################### diff --git a/tests/bloom_filter/arrow_policy_test.cu b/tests/bloom_filter/arrow_policy_test.cu index 1b7f5384f..82ba25021 100644 --- a/tests/bloom_filter/arrow_policy_test.cu +++ b/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(); diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 28379f425..b23564270 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/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, 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), diff --git a/tests/dynamic_bitset/find_next_test.cu b/tests/dynamic_bitset/find_next_test.cu index efc80f271..07cca1f52 100644 --- a/tests/dynamic_bitset/find_next_test.cu +++ b/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; diff --git a/tests/dynamic_bitset/get_test.cu b/tests/dynamic_bitset/get_test.cu index 3e25af15b..39a8f9d00 100644 --- a/tests/dynamic_bitset/get_test.cu +++ b/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; diff --git a/tests/dynamic_bitset/rank_test.cu b/tests/dynamic_bitset/rank_test.cu index ed462963d..cf8cc23cb 100644 --- a/tests/dynamic_bitset/rank_test.cu +++ b/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; diff --git a/tests/dynamic_bitset/select_test.cu b/tests/dynamic_bitset/select_test.cu index ceead7a23..5557fa7b2 100644 --- a/tests/dynamic_bitset/select_test.cu +++ b/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; diff --git a/tests/dynamic_bitset/size_test.cu b/tests/dynamic_bitset/size_test.cu index 611159dc3..6585b8f6f 100644 --- a/tests/dynamic_bitset/size_test.cu +++ b/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 -TEST_CASE("Size computation", "") +TEST_CASE("dynamic_bitset size computation test", "") { cuco::experimental::detail::dynamic_bitset bv; using size_type = std::size_t; diff --git a/tests/dynamic_map/erase_test.cu b/tests/dynamic_map/erase_test.cu index 0701c9495..fc74e8e69 100644 --- a/tests/dynamic_map/erase_test.cu +++ b/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 -TEMPLATE_TEST_CASE_SIG("erase key", +TEMPLATE_TEST_CASE_SIG("dynamic_map erase tests", "", ((typename Key, typename Value), Key, Value), (int32_t, int32_t), diff --git a/tests/dynamic_map/unique_sequence_test.cu b/tests/dynamic_map/unique_sequence_test.cu index ac3c1d8b6..7f17cb7b8 100644 --- a/tests/dynamic_map/unique_sequence_test.cu +++ b/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 -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), diff --git a/tests/static_map/capacity_test.cu b/tests/static_map/capacity_test.cu index db93006d7..441415eb3 100644 --- a/tests/static_map/capacity_test.cu +++ b/tests/static_map/capacity_test.cu @@ -18,7 +18,7 @@ #include -TEST_CASE("Static map capacity", "") +TEST_CASE("static_map capacity test", "") { using Key = int32_t; using T = int32_t; diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/contains_test.cu similarity index 70% rename from tests/static_map/unique_sequence_test.cu rename to tests/static_map/contains_test.cu index 8e8d6c9ad..3f3a6312c 100644 --- a/tests/static_map/unique_sequence_test.cu +++ b/tests/static_map/contains_test.cu @@ -33,8 +33,6 @@ using size_type = int32_t; -int32_t constexpr SENTINEL = -1; - template 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 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{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 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(num_keys); - auto gold_fn = cuda::proclaim_return_type([] __device__(auto const& i) { - return i % 2 == 0 ? static_cast(i) : Value{SENTINEL}; - }); - - map.find_if(keys_begin, - keys_begin + num_keys, - thrust::counting_iterator{0}, - is_even, - found_results.begin()); - - REQUIRE(cuco::test::equal( - found_results.begin(), - found_results.end(), - thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), - cuda::proclaim_return_type( - [] __device__(auto const& found, auto const& gold) { return found == gold; }))); - } - SECTION("All inserted key-values should be properly retrieved") { thrust::device_vector d_values(num_keys); @@ -163,7 +118,7 @@ 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, @@ -171,25 +126,15 @@ TEMPLATE_TEST_CASE_SIG( 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; @@ -206,9 +151,7 @@ TEMPLATE_TEST_CASE_SIG( probe, cuco::cuda_allocator, cuco::storage<2>>{ - extent_type{}, cuco::empty_key{SENTINEL}, cuco::empty_value{SENTINEL}}; - - REQUIRE(map.capacity() == gold_capacity); + extent_type{}, cuco::empty_key{-1}, cuco::empty_value{-1}}; test_unique_sequence(map, num_keys); } diff --git a/tests/static_map/custom_type_test.cu b/tests/static_map/custom_type_test.cu index d2d5138aa..9b8361b4d 100644 --- a/tests/static_map/custom_type_test.cu +++ b/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 diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu new file mode 100644 index 000000000..7406f2326 --- /dev/null +++ b/tests/static_map/find_test.cu @@ -0,0 +1,149 @@ +/* + * 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +using size_type = int32_t; + +int32_t constexpr SENTINEL = -1; + +template +void test_unique_sequence(Map& map, size_type num_keys) +{ + using Key = typename Map::key_type; + using Value = typename Map::mapped_type; + + auto keys_begin = thrust::counting_iterator{0}; + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair{i, i}; })); + + auto zip_equal = cuda::proclaim_return_type( + [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }); + auto is_even = + cuda::proclaim_return_type([] __device__(auto const& i) { return i % 2 == 0; }); + + SECTION("Non-inserted keys have no matches") + { + thrust::device_vector 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{map.empty_key_sentinel()})); + + REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); + } + + map.insert(pairs_begin, pairs_begin + num_keys); + + SECTION("All inserted keys should be correctly recovered during find") + { + thrust::device_vector 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(num_keys); + auto gold_fn = cuda::proclaim_return_type([] __device__(auto const& i) { + return i % 2 == 0 ? static_cast(i) : Value{SENTINEL}; + }); + + map.find_if(keys_begin, + keys_begin + num_keys, + thrust::counting_iterator{0}, + is_even, + found_results.begin()); + + REQUIRE(cuco::test::equal( + found_results.begin(), + found_results.end(), + thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), + cuda::proclaim_return_type( + [] __device__(auto const& found, auto const& gold) { return found == gold; }))); + } +} + +TEMPLATE_TEST_CASE_SIG( + "static_map: find 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; + using probe = std::conditional_t< + Probe == cuco::test::probe_sequence::linear_probing, + cuco::linear_probing>, + cuco::double_hashing, cuco::murmurhash3_32>>; + + auto map = cuco::static_map, + probe, + cuco::cuda_allocator, + cuco::storage<2>>{ + extent_type{}, cuco::empty_key{SENTINEL}, cuco::empty_value{SENTINEL}}; + + REQUIRE(map.capacity() == gold_capacity); + + test_unique_sequence(map, num_keys); +} diff --git a/tests/static_map/heterogeneous_lookup_test.cu b/tests/static_map/heterogeneous_lookup_test.cu index 08d31f0ed..8ba2d22ba 100644 --- a/tests/static_map/heterogeneous_lookup_test.cu +++ b/tests/static_map/heterogeneous_lookup_test.cu @@ -82,7 +82,7 @@ struct custom_key_equal { } }; -TEMPLATE_TEST_CASE_SIG("Heterogeneous lookup", +TEMPLATE_TEST_CASE_SIG("static_map heterogeneous lookup tests", "", ((typename T, int CGSize), T, CGSize), #if defined(CUCO_HAS_INDEPENDENT_THREADS) // Key type larger than 8B only supported for sm_70 and diff --git a/tests/static_map/insert_or_assign_test.cu b/tests/static_map/insert_or_assign_test.cu index a6ad767d6..7c82f0817 100644 --- a/tests/static_map/insert_or_assign_test.cu +++ b/tests/static_map/insert_or_assign_test.cu @@ -71,7 +71,7 @@ void test_insert_or_assign(Map& map, size_type num_keys) } TEMPLATE_TEST_CASE_SIG( - "Insert or assign", + "static_map insert_or_assign tests", "", ((typename Key, typename Value, cuco::test::probe_sequence Probe, int CGSize), Key, diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu index 8ef4a0b69..3f502b6bd 100644 --- a/tests/static_map/key_sentinel_test.cu +++ b/tests/static_map/key_sentinel_test.cu @@ -33,8 +33,7 @@ struct custom_equals { __device__ bool operator()(T lhs, T rhs) const { return A[lhs] == A[rhs]; } }; -TEMPLATE_TEST_CASE_SIG( - "Key comparison against sentinel", "", ((typename T), T), (int32_t), (int64_t)) +TEMPLATE_TEST_CASE_SIG("static_map key sentinel tests", "", ((typename T), T), (int32_t), (int64_t)) { using Key = T; using Value = T; diff --git a/tests/static_map/rehash_test.cu b/tests/static_map/rehash_test.cu index e9b8df498..62e3355e5 100644 --- a/tests/static_map/rehash_test.cu +++ b/tests/static_map/rehash_test.cu @@ -22,7 +22,7 @@ #include -TEST_CASE("Rehash", "") +TEST_CASE("static_map rehash test", "") { using key_type = int; using mapped_type = long; diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu index 775d6b6b3..0059b61dd 100644 --- a/tests/static_map/shared_memory_test.cu +++ b/tests/static_map/shared_memory_test.cu @@ -66,7 +66,7 @@ __global__ void shared_memory_test_kernel(Ref* maps, } } -TEMPLATE_TEST_CASE_SIG("Shared memory static map", +TEMPLATE_TEST_CASE_SIG("static_map shared memory tests", "", ((typename Key, typename Value), Key, Value), (int32_t, int32_t), diff --git a/tests/static_multimap/custom_pair_retrieve_test.cu b/tests/static_multimap/custom_pair_retrieve_test.cu index f8a13b1c0..ebfb35c61 100644 --- a/tests/static_multimap/custom_pair_retrieve_test.cu +++ b/tests/static_multimap/custom_pair_retrieve_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -188,7 +188,7 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) } TEMPLATE_TEST_CASE_SIG( - "Tests of non-shared-memory pair_retrieve", + "static_multimap non-shared-memory pair_retrieve tests", "", ((typename Key, typename Value, cuco::test::probe_sequence Probe), Key, Value, Probe), (int32_t, int32_t, cuco::test::probe_sequence::linear_probing), diff --git a/tests/static_multimap/custom_type_test.cu b/tests/static_multimap/custom_type_test.cu index 74210f5bd..9d1a7b98f 100644 --- a/tests/static_multimap/custom_type_test.cu +++ b/tests/static_multimap/custom_type_test.cu @@ -208,7 +208,7 @@ void test_custom_key_value_type(Map& map, std::size_t num_pairs) } } -TEMPLATE_TEST_CASE_SIG("User defined key and value type", +TEMPLATE_TEST_CASE_SIG("static_multimap user defined key and value type tests", "", ((cuco::test::probe_sequence Probe), Probe), (cuco::test::probe_sequence::linear_probing), diff --git a/tests/static_multimap/heterogeneous_lookup_test.cu b/tests/static_multimap/heterogeneous_lookup_test.cu index db80af597..ead1bae59 100644 --- a/tests/static_multimap/heterogeneous_lookup_test.cu +++ b/tests/static_multimap/heterogeneous_lookup_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. @@ -80,7 +80,7 @@ struct custom_key_equal { } }; -TEMPLATE_TEST_CASE("Heterogeneous lookup", +TEMPLATE_TEST_CASE("static_multimap heterogeneous lookup tests", "", #if defined(CUCO_HAS_INDEPENDENT_THREADS) // Key type larger than 8B only supported for sm_70 and // up diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu index dfbba5128..699345237 100644 --- a/tests/static_multimap/multiplicity_test.cu +++ b/tests/static_multimap/multiplicity_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -147,7 +147,7 @@ void test_multiplicity_two(Map& map, std::size_t num_items) } TEMPLATE_TEST_CASE_SIG( - "Multiplicity equals two", + "static_multimap multiplicity tests", "", ((typename Key, typename Value, cuco::test::probe_sequence Probe), Key, Value, Probe), (int32_t, int32_t, cuco::test::probe_sequence::linear_probing), diff --git a/tests/static_multimap/non_match_test.cu b/tests/static_multimap/non_match_test.cu index 729116d50..69642d39d 100644 --- a/tests/static_multimap/non_match_test.cu +++ b/tests/static_multimap/non_match_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -114,7 +114,7 @@ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, std::size_t } TEMPLATE_TEST_CASE_SIG( - "Tests of non-matches", + "static_multimap non-match tests", "", ((typename Key, typename Value, cuco::test::probe_sequence Probe), Key, Value, Probe), (int32_t, int32_t, cuco::test::probe_sequence::linear_probing), diff --git a/tests/static_multimap/pair_function_test.cu b/tests/static_multimap/pair_function_test.cu index 448449806..5e86f2052 100644 --- a/tests/static_multimap/pair_function_test.cu +++ b/tests/static_multimap/pair_function_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -108,7 +108,7 @@ void test_pair_functions(Map& map, PairIt pair_begin, std::size_t num_pairs) } TEMPLATE_TEST_CASE_SIG( - "Tests of pair functions", + "static_multimap pair functions tests", "", ((typename Key, typename Value, cuco::test::probe_sequence Probe), Key, Value, Probe), (int32_t, int32_t, cuco::test::probe_sequence::linear_probing), diff --git a/tests/static_set/capacity_test.cu b/tests/static_set/capacity_test.cu index 433e25287..f1f2e5ba1 100644 --- a/tests/static_set/capacity_test.cu +++ b/tests/static_set/capacity_test.cu @@ -18,7 +18,7 @@ #include -TEST_CASE("Static set capacity", "") +TEST_CASE("static_set capacity test", "") { using Key = int32_t; using ProbeT = cuco::double_hashing<1, cuco::default_hash_function>; diff --git a/tests/static_set/heterogeneous_lookup_test.cu b/tests/static_set/heterogeneous_lookup_test.cu index 9adf35371..463184cf1 100644 --- a/tests/static_set/heterogeneous_lookup_test.cu +++ b/tests/static_set/heterogeneous_lookup_test.cu @@ -82,8 +82,11 @@ struct custom_key_equal { } }; -TEMPLATE_TEST_CASE_SIG( - "Heterogeneous lookup", "", ((typename T, int CGSize), T, CGSize), (int32_t, 1), (int32_t, 2)) +TEMPLATE_TEST_CASE_SIG("static_set heterogeneous lookup tests", + "", + ((typename T, int CGSize), T, CGSize), + (int32_t, 1), + (int32_t, 2)) { using Key = T; using InsertKey = key_pair; diff --git a/tests/static_set/rehash_test.cu b/tests/static_set/rehash_test.cu index c36a1dc4a..e2078d183 100644 --- a/tests/static_set/rehash_test.cu +++ b/tests/static_set/rehash_test.cu @@ -21,7 +21,7 @@ #include -TEST_CASE("Rehash", "") +TEST_CASE("static_set rehash test", "") { using key_type = int; diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu index a11c3e8a8..f28a7c209 100644 --- a/tests/static_set/retrieve_all_test.cu +++ b/tests/static_set/retrieve_all_test.cu @@ -58,7 +58,7 @@ void test_unique_sequence(Set& set, std::size_t num_keys) } TEMPLATE_TEST_CASE_SIG( - "Retrieve all", + "static_set::retrieve_all tests", "", ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize), (int32_t, cuco::test::probe_sequence::double_hashing, 1), diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu index 7ea88d418..3c054ba6f 100644 --- a/tests/static_set/retrieve_test.cu +++ b/tests/static_set/retrieve_test.cu @@ -73,7 +73,7 @@ void test_unique_sequence(Set& set, std::size_t num_keys) } TEMPLATE_TEST_CASE_SIG( - "SetRetrieve", + "static_set retrieve tests", "", ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize), (int32_t, cuco::test::probe_sequence::double_hashing, 1), diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu index f8cc9b4f1..b970e735f 100644 --- a/tests/static_set/shared_memory_test.cu +++ b/tests/static_set/shared_memory_test.cu @@ -64,7 +64,8 @@ __global__ void shared_memory_test_kernel(Ref* sets, } } -TEMPLATE_TEST_CASE_SIG("Shared memory static set", "", ((typename Key), Key), (int32_t), (int64_t)) +TEMPLATE_TEST_CASE_SIG( + "static_set shared memory tests", "", ((typename Key), Key), (int32_t), (int64_t)) { constexpr std::size_t number_of_sets = 1000; constexpr std::size_t elements_in_set = 500; @@ -187,7 +188,7 @@ __global__ void shared_memory_hash_set_kernel(bool* key_found) if (retrieved_it != find_ref.end() && *retrieved_it == rank) { key_found[index] = true; } } -TEST_CASE("static set shared memory slots.", "") +TEST_CASE("static_set shared memory slots test", "") { constexpr std::size_t N = 256; [[maybe_unused]] auto constexpr num_windows = diff --git a/tests/static_set/size_test.cu b/tests/static_set/size_test.cu index 4a858eb70..1b087a566 100644 --- a/tests/static_set/size_test.cu +++ b/tests/static_set/size_test.cu @@ -22,7 +22,7 @@ #include -TEST_CASE("Size computation", "") +TEST_CASE("static_set size test", "") { constexpr std::size_t num_keys{400}; diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu index 0cd2924a9..ba9161169 100644 --- a/tests/static_set/unique_sequence_test.cu +++ b/tests/static_set/unique_sequence_test.cu @@ -136,7 +136,7 @@ void test_unique_sequence(Set& set, size_type num_keys) } TEMPLATE_TEST_CASE_SIG( - "Unique sequence", + "static_set unique sequence tests", "", ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize), (int32_t, cuco::test::probe_sequence::double_hashing, 1), diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index 0f8d53c16..96e97e4b9 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -21,7 +21,7 @@ #include TEMPLATE_TEST_CASE_SIG( - "Extent tests", "", ((typename SizeType), SizeType), (int32_t), (int64_t), (std::size_t)) + "utility extent tests", "", ((typename SizeType), SizeType), (int32_t), (int64_t), (std::size_t)) { SizeType constexpr num = 1234; SizeType constexpr gold_reference = 314; // 157 x 2 diff --git a/tests/utility/hash_test.cu b/tests/utility/hash_test.cu index 5e0edda74..665ba1d85 100644 --- a/tests/utility/hash_test.cu +++ b/tests/utility/hash_test.cu @@ -67,7 +67,7 @@ __global__ void check_identity_hash_result_kernel(OutputIter result) cuda::std::numeric_limits::max(), cuda::std::numeric_limits::max()); } -TEST_CASE("Test cuco::identity_hash", "") +TEST_CASE("utility cuco::identity_hash test", "") { SECTION("Check if host-generated hash values match the identity function.") { @@ -121,7 +121,7 @@ __global__ void check_hash_result_kernel_64(OutputIter result) check_hash_result>>(123456789, 2031761887105658523, 0); } -TEST_CASE("Test cuco::xxhash_64", "") +TEST_CASE("utility cuco::xxhash_64 test", "") { // Reference hash values were computed using https://github.com/Cyan4973/xxHash SECTION("Check if host-generated hash values match the reference implementation.") @@ -184,7 +184,7 @@ __global__ void check_hash_result_kernel_32(OutputIter result) result[i++] = check_hash_result>>(123456789, 3715432378, 0); } -TEST_CASE("Test cuco::xxhash_32", "") +TEST_CASE("utility cuco::xxhash_32 test", "") { // Reference hash values were computed using https://github.com/Cyan4973/xxHash SECTION("Check if host-generated hash values match the reference implementation.") @@ -221,7 +221,7 @@ TEST_CASE("Test cuco::xxhash_32", "") } } -TEMPLATE_TEST_CASE_SIG("Static vs. dynamic key hash test", +TEMPLATE_TEST_CASE_SIG("utility hasher compute_hash tests", "", ((typename Hash), Hash), (cuco::murmurhash3_32), @@ -323,7 +323,7 @@ __global__ void check_murmurhash3_128_result_kernel(OutputIter result) 1024); } -TEST_CASE("Test cuco::murmurhash3_x64_128", "") +TEST_CASE("utility cuco::murmurhash3_x64_128 test", "") { // Reference hash values were computed using // https://github.com/aappleby/smhasher/blob/master/src/MurmurHash3.cpp diff --git a/tests/utility/probing_scheme_test.cu b/tests/utility/probing_scheme_test.cu index 8f282108d..994d5e443 100644 --- a/tests/utility/probing_scheme_test.cu +++ b/tests/utility/probing_scheme_test.cu @@ -77,7 +77,7 @@ __global__ void generate_cg_probing_sequence(Key key, } TEMPLATE_TEST_CASE_SIG( - "probing_scheme scalar vs CGSize 1 test", + "utility probing_scheme tests", "", ((typename Key, cuco::test::probe_sequence Probe, int32_t WindowSize), Key, Probe, WindowSize), (int32_t, cuco::test::probe_sequence::double_hashing, 1), @@ -105,4 +105,4 @@ TEMPLATE_TEST_CASE_SIG( REQUIRE(cuco::test::equal( scalar_seq.begin(), scalar_seq.end(), cg_seq.begin(), thrust::equal_to{})); -} \ No newline at end of file +} diff --git a/tests/utility/storage_test.cu b/tests/utility/storage_test.cu index 74da3ffd7..f07da8c3a 100644 --- a/tests/utility/storage_test.cu +++ b/tests/utility/storage_test.cu @@ -23,7 +23,7 @@ #include -TEMPLATE_TEST_CASE_SIG("Storage tests", +TEMPLATE_TEST_CASE_SIG("utility storage tests", "", ((typename Key, typename Value), Key, Value), (int32_t, int32_t),