diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index aa1d08054..097a403ed 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -132,20 +132,20 @@ template template [[nodiscard]] auto constexpr make_bucket_extent(extent ext) { - return make_bucket_extent(ext); - /* - // TODO -if constexpr (cuco::is_double_hashing::value) { -} else { - auto const size = cuco::detail::int_div_ceil( - cuda::std::max(static_cast(ext), static_cast(1)), - ProbingScheme::cg_size * Storage::bucket_size) + cuda::std::min(static_cast(ext), -static_cast(1)); if constexpr (N == dynamic_extent) { return bucket_extent{size -* ProbingScheme::cg_size}; } else { return bucket_extent{}; + if constexpr (cuco::is_double_hashing::value) { + return make_bucket_extent(ext); + } else { + auto const size = cuco::detail::int_div_ceil( + cuda::std::max(static_cast(ext), static_cast(1)), + ProbingScheme::cg_size * Storage::bucket_size) + + cuda::std::min(static_cast(ext), static_cast(1)); + if constexpr (N == dynamic_extent) { + return bucket_extent{size * ProbingScheme::cg_size}; + } else { + return bucket_extent{}; + } } } -*/ -} template [[nodiscard]] auto constexpr make_bucket_extent(SizeType size) diff --git a/tests/static_map/capacity_test.cu b/tests/static_map/capacity_test.cu index db93006d7..e7bae86f5 100644 --- a/tests/static_map/capacity_test.cu +++ b/tests/static_map/capacity_test.cu @@ -159,7 +159,7 @@ TEST_CASE("Static map capacity", "") SECTION("Dynamic extent is evaluated at run time.") { - auto constexpr gold_capacity = 412; // 103 x 2 x 2 + auto constexpr gold_capacity = 404; // 101 x 2 x 2 using probe = cuco::linear_probing<2, cuco::default_hash_function>; auto map = diff --git a/tests/static_map/erase_test.cu b/tests/static_map/erase_test.cu index cd0f26d1b..d5bdffd67 100644 --- a/tests/static_map/erase_test.cu +++ b/tests/static_map/erase_test.cu @@ -119,8 +119,10 @@ TEMPLATE_TEST_CASE_SIG( thrust::equal_to, probe, cuco::cuda_allocator, - cuco::storage<2>>{ - num_keys, cuco::empty_key{-1}, cuco::empty_value{-1}, cuco::erased_key{-2}}; + cuco::storage<2>>{num_keys * 2, + cuco::empty_key{-1}, + cuco::empty_value{-1}, + cuco::erased_key{-2}}; test_erase(map, num_keys); } diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu index 22cfd2d4a..4c432b2fd 100644 --- a/tests/static_map/unique_sequence_test.cu +++ b/tests/static_map/unique_sequence_test.cu @@ -166,24 +166,32 @@ TEMPLATE_TEST_CASE_SIG( (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 + + using probe = std::conditional_t< + Probe == cuco::test::probe_sequence::linear_probing, + cuco::linear_probing>, + cuco::double_hashing, cuco::murmurhash3_32>>; + + auto constexpr gold_capacity = [&]() { + if constexpr (cuco::is_double_hashing::value) { + return CGSize == 1 ? 422 // 211 x 1 x 2 + : 412; // 103 x 2 x 2 + } else { + return CGSize == 1 ? 402 // 201 x 1 x 2 + : 404; // 101 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>>{ + auto map = cuco::static_map, + probe, + cuco::cuda_allocator, + cuco::storage<2>>{ extent_type{}, cuco::empty_key{-1}, cuco::empty_value{-1}}; REQUIRE(map.capacity() == gold_capacity); diff --git a/tests/static_multiset/insert_test.cu b/tests/static_multiset/insert_test.cu index b1754f44b..0c7d27dbb 100644 --- a/tests/static_multiset/insert_test.cu +++ b/tests/static_multiset/insert_test.cu @@ -89,13 +89,21 @@ TEMPLATE_TEST_CASE_SIG( (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 using probe = std::conditional_t>, cuco::double_hashing>>; + auto constexpr gold_capacity = [&]() { + if constexpr (cuco::is_double_hashing::value) { + return CGSize == 1 ? 422 // 211 x 1 x 2 + : 412; // 103 x 2 x 2 + } else { + return CGSize == 1 ? 402 // 201 x 1 x 2 + : 404; // 101 x 2 x 2 + } + }(); + auto set = cuco::static_multiset{num_keys, cuco::empty_key{-1}, {}, probe{}, {}, cuco::storage<2>{}}; diff --git a/tests/static_set/capacity_test.cu b/tests/static_set/capacity_test.cu index 433e25287..2f6441799 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 tests", "") { using Key = int32_t; using ProbeT = cuco::double_hashing<1, cuco::default_hash_function>; @@ -120,9 +120,9 @@ TEST_CASE("Static set capacity", "") REQUIRE(ref_capacity == gold_capacity); } - SECTION("Dynamic extent is evaluated at run time.") + SECTION("Dynamic extent for linear porbing.") { - auto constexpr gold_capacity = 412; // 103 x 2 x 2 + auto constexpr gold_capacity = 404; // 101 x 2 x 2 using probe = cuco::linear_probing<2, cuco::default_hash_function>; auto set = cuco::static_set>, cuco::double_hashing>>; + auto constexpr gold_capacity = [&]() { + if constexpr (cuco::is_double_hashing::value) { + return CGSize == 1 ? 409 // 409 x 1 x 1 + : 422; // 211 x 2 x 1 + } else { + return CGSize == 1 ? 401 // 401 x 1 x 1 + : 402; // 201 x 2 x 1 + } + }(); + auto set = cuco::static_set{num_keys, desired_load_factor, cuco::empty_key{-1}, {}, probe{}}; REQUIRE(set.capacity() == gold_capacity); diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu index 199c7cbff..453ce5fbe 100644 --- a/tests/static_set/unique_sequence_test.cu +++ b/tests/static_set/unique_sequence_test.cu @@ -132,14 +132,21 @@ TEMPLATE_TEST_CASE_SIG( (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 - ; using probe = std::conditional_t>, cuco::double_hashing>>; + auto constexpr gold_capacity = [&]() { + if constexpr (cuco::is_double_hashing::value) { + return CGSize == 1 ? 422 // 211 x 1 x 2 + : 412; // 103 x 2 x 2 + } else { + return CGSize == 1 ? 402 // 201 x 1 x 2 + : 404; // 101 x 2 x 2 + } + }(); + auto set = cuco::static_set{num_keys, cuco::empty_key{-1}, {}, probe{}, {}, cuco::storage<2>{}};