From 88ff1e4eb7bbdd5a0d1c9025548e1451d8128a93 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 27 Jun 2023 11:58:21 -0700 Subject: [PATCH] Clean up `cuco::pair` (#319) Contributes to #110 Depends on #314 This PR: - deprecates `cuco::pair_type` alias - fixes issues with `cuco::make_pair` - separates `pair` declarations and implementation details --- .../hash_table/dynamic_map/contains_bench.cu | 2 +- .../hash_table/dynamic_map/erase_bench.cu | 2 +- .../hash_table/dynamic_map/find_bench.cu | 2 +- .../hash_table/dynamic_map/insert_bench.cu | 2 +- .../hash_table/static_map/contains_bench.cu | 2 +- .../hash_table/static_map/erase_bench.cu | 2 +- .../hash_table/static_map/find_bench.cu | 2 +- .../hash_table/static_map/insert_bench.cu | 2 +- .../hash_table/static_multimap/count_bench.cu | 2 +- .../static_multimap/insert_bench.cu | 2 +- .../hash_table/static_multimap/query_bench.cu | 4 +- .../static_multimap/retrieve_bench.cu | 4 +- examples/static_multimap/host_bulk_example.cu | 2 +- include/cuco/detail/dynamic_map.inl | 2 +- .../cuco/detail/open_addressing_ref_impl.cuh | 2 +- include/cuco/detail/pair.inl | 51 ++++++ include/cuco/detail/probe_sequence_impl.cuh | 9 +- .../cuco/detail/static_multimap/kernels.cuh | 4 +- include/cuco/detail/traits.hpp | 59 +++++++ include/cuco/detail/utils.cuh | 92 +++++++++++ include/cuco/dynamic_map.cuh | 2 +- include/cuco/pair.cuh | 146 ++++++++++++++++++ include/cuco/static_map.cuh | 8 +- include/cuco/static_multimap.cuh | 6 +- include/cuco/utility/traits.hpp | 3 + tests/dynamic_map/unique_sequence_test.cu | 6 +- tests/static_map/custom_type_test.cu | 25 ++- tests/static_map/duplicate_keys_test.cu | 2 +- tests/static_map/heterogeneous_lookup_test.cu | 6 +- tests/static_map/key_sentinel_test.cu | 32 ++-- tests/static_map/stream_test.cu | 6 +- tests/static_map/unique_sequence_test.cu | 78 +++++----- .../custom_pair_retrieve_test.cu | 10 +- tests/static_multimap/custom_type_test.cu | 46 +++--- .../heterogeneous_lookup_test.cu | 6 +- tests/static_multimap/insert_if_test.cu | 4 +- tests/static_multimap/multiplicity_test.cu | 70 ++++----- tests/static_multimap/non_match_test.cu | 74 ++++----- tests/static_multimap/pair_function_test.cu | 10 +- tests/utility/storage_test.cu | 2 +- 40 files changed, 567 insertions(+), 224 deletions(-) create mode 100644 include/cuco/detail/pair.inl create mode 100644 include/cuco/detail/traits.hpp create mode 100644 include/cuco/pair.cuh diff --git a/benchmarks/hash_table/dynamic_map/contains_bench.cu b/benchmarks/hash_table/dynamic_map/contains_bench.cu index 8e41b8e2d..ff349bc53 100644 --- a/benchmarks/hash_table/dynamic_map/contains_bench.cu +++ b/benchmarks/hash_table/dynamic_map/contains_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_contains( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const initial_size = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE); diff --git a/benchmarks/hash_table/dynamic_map/erase_bench.cu b/benchmarks/hash_table/dynamic_map/erase_bench.cu index b815515e8..96f5ec7ec 100644 --- a/benchmarks/hash_table/dynamic_map/erase_bench.cu +++ b/benchmarks/hash_table/dynamic_map/erase_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_erase( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const initial_size = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE); diff --git a/benchmarks/hash_table/dynamic_map/find_bench.cu b/benchmarks/hash_table/dynamic_map/find_bench.cu index 12576ccc1..b06cfab4e 100644 --- a/benchmarks/hash_table/dynamic_map/find_bench.cu +++ b/benchmarks/hash_table/dynamic_map/find_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_find( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const initial_size = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE); diff --git a/benchmarks/hash_table/dynamic_map/insert_bench.cu b/benchmarks/hash_table/dynamic_map/insert_bench.cu index de2fa8a4a..8e8cc8a84 100644 --- a/benchmarks/hash_table/dynamic_map/insert_bench.cu +++ b/benchmarks/hash_table/dynamic_map/insert_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_insert( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const initial_size = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE); diff --git a/benchmarks/hash_table/static_map/contains_bench.cu b/benchmarks/hash_table/static_map/contains_bench.cu index 09737a136..0b5d482a1 100644 --- a/benchmarks/hash_table/static_map/contains_bench.cu +++ b/benchmarks/hash_table/static_map/contains_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_contains( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_map/erase_bench.cu b/benchmarks/hash_table/static_map/erase_bench.cu index 3f26504a7..c6e56eb07 100644 --- a/benchmarks/hash_table/static_map/erase_bench.cu +++ b/benchmarks/hash_table/static_map/erase_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_erase( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_map/find_bench.cu b/benchmarks/hash_table/static_map/find_bench.cu index 4a1ccca11..276a35e0b 100644 --- a/benchmarks/hash_table/static_map/find_bench.cu +++ b/benchmarks/hash_table/static_map/find_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_find( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_map/insert_bench.cu b/benchmarks/hash_table/static_map/insert_bench.cu index b6fadc057..ef997bef8 100644 --- a/benchmarks/hash_table/static_map/insert_bench.cu +++ b/benchmarks/hash_table/static_map/insert_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_insert( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_multimap/count_bench.cu b/benchmarks/hash_table/static_multimap/count_bench.cu index e087e3243..fa71c8d0c 100644 --- a/benchmarks/hash_table/static_multimap/count_bench.cu +++ b/benchmarks/hash_table/static_multimap/count_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_count( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_multimap/insert_bench.cu b/benchmarks/hash_table/static_multimap/insert_bench.cu index c045f3a91..aa41044bb 100644 --- a/benchmarks/hash_table/static_multimap/insert_bench.cu +++ b/benchmarks/hash_table/static_multimap/insert_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_insert( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); diff --git a/benchmarks/hash_table/static_multimap/query_bench.cu b/benchmarks/hash_table/static_multimap/query_bench.cu index 783c83556..7d6202297 100644 --- a/benchmarks/hash_table/static_multimap/query_bench.cu +++ b/benchmarks/hash_table/static_multimap/query_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_query( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); @@ -99,4 +99,4 @@ NVBENCH_BENCH_TYPES(static_multimap_query, .set_name("static_multimap_query_uniform_multiplicity") .set_type_axes_names({"Key", "Value", "Distribution"}) .set_max_noise(defaults::MAX_NOISE) - .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE); \ No newline at end of file + .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE); diff --git a/benchmarks/hash_table/static_multimap/retrieve_bench.cu b/benchmarks/hash_table/static_multimap/retrieve_bench.cu index 432bd3485..e30fbe547 100644 --- a/benchmarks/hash_table/static_multimap/retrieve_bench.cu +++ b/benchmarks/hash_table/static_multimap/retrieve_bench.cu @@ -35,7 +35,7 @@ template std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_retrieve( nvbench::state& state, nvbench::type_list) { - using pair_type = cuco::pair_type; + using pair_type = cuco::pair; auto const num_keys = state.get_int64_or_default("NumInputs", defaults::N); auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY); @@ -98,4 +98,4 @@ NVBENCH_BENCH_TYPES(static_multimap_retrieve, .set_name("static_multimap_retrieve_uniform_multiplicity") .set_type_axes_names({"Key", "Value", "Distribution"}) .set_max_noise(defaults::MAX_NOISE) - .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE); \ No newline at end of file + .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE); diff --git a/examples/static_multimap/host_bulk_example.cu b/examples/static_multimap/host_bulk_example.cu index a7d5a95a7..d1fe5589a 100644 --- a/examples/static_multimap/host_bulk_example.cu +++ b/examples/static_multimap/host_bulk_example.cu @@ -60,7 +60,7 @@ int main(void) // The `_outer` suffix indicates that the occurrence of a non-match is 1. auto const output_size = map.count_outer(keys_to_find.begin(), keys_to_find.end()); - thrust::device_vector> d_results(output_size); + thrust::device_vector> d_results(output_size); // Finds all keys {0, 1, 2, ...} and stores associated key/value pairs into `d_results` // If a key `keys_to_find[i]` doesn't exist, `d_results[i].second == empty_value_sentinel` diff --git a/include/cuco/detail/dynamic_map.inl b/include/cuco/detail/dynamic_map.inl index e329eefeb..7b5145190 100644 --- a/include/cuco/detail/dynamic_map.inl +++ b/include/cuco/detail/dynamic_map.inl @@ -146,7 +146,7 @@ void dynamic_map::insert( auto const n = std::min(capacity_remaining, num_to_insert); auto const grid_size = (tile_size * n + stride * block_size - 1) / (stride * block_size); - detail::insert> + detail::insert> <<>>(first, first + n, submap_views_.data().get(), diff --git a/include/cuco/detail/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing_ref_impl.cuh index dd3a84434..56b699104 100644 --- a/include/cuco/detail/open_addressing_ref_impl.cuh +++ b/include/cuco/detail/open_addressing_ref_impl.cuh @@ -17,7 +17,7 @@ #pragma once #include -#include +#include #include #include diff --git a/include/cuco/detail/pair.inl b/include/cuco/detail/pair.inl new file mode 100644 index 000000000..56d16e4fb --- /dev/null +++ b/include/cuco/detail/pair.inl @@ -0,0 +1,51 @@ +/* + * Copyright (c) 2023, 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. + */ + +#pragma once + +#include +#include + +namespace cuco { + +template +__host__ __device__ constexpr pair::pair(First const& f, Second const& s) + : first{f}, second{s} +{ +} + +template +template +__host__ __device__ constexpr pair::pair(pair const& p) + : first{p.first}, second{p.second} +{ +} + +template +__host__ __device__ constexpr pair, std::decay_t> make_pair(F&& f, + S&& s) noexcept +{ + return pair, std::decay_t>(std::forward(f), std::forward(s)); +} + +template +__host__ __device__ constexpr bool operator==(cuco::pair const& lhs, + cuco::pair const& rhs) noexcept +{ + return lhs.first == rhs.first and lhs.second == rhs.second; +} + +} // namespace cuco diff --git a/include/cuco/detail/probe_sequence_impl.cuh b/include/cuco/detail/probe_sequence_impl.cuh index 14124b639..c108840b2 100644 --- a/include/cuco/detail/probe_sequence_impl.cuh +++ b/include/cuco/detail/probe_sequence_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -16,7 +16,8 @@ #pragma once -#include +#include +#include #include @@ -71,13 +72,13 @@ template class probe_sequence_impl_base { protected: - using value_type = cuco::pair_type; ///< Type of key/value pairs + using value_type = cuco::pair; ///< Type of key/value pairs using key_type = Key; ///< Key type using mapped_type = Value; ///< Type of mapped values using atomic_key_type = cuda::atomic; ///< Type of atomic keys using atomic_mapped_type = cuda::atomic; ///< Type of atomic mapped values /// Pair type of atomic key and atomic mapped value - using pair_atomic_type = cuco::pair_type; + using pair_atomic_type = cuco::pair; /// Type of the forward iterator to `pair_atomic_type` using iterator = pair_atomic_type*; /// Type of the forward iterator to `const pair_atomic_type` diff --git a/include/cuco/detail/static_multimap/kernels.cuh b/include/cuco/detail/static_multimap/kernels.cuh index c010fa8f3..ca5f898a5 100644 --- a/include/cuco/detail/static_multimap/kernels.cuh +++ b/include/cuco/detail/static_multimap/kernels.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -15,7 +15,7 @@ */ #pragma once -#include +#include #include diff --git a/include/cuco/detail/traits.hpp b/include/cuco/detail/traits.hpp new file mode 100644 index 000000000..602a93251 --- /dev/null +++ b/include/cuco/detail/traits.hpp @@ -0,0 +1,59 @@ +/* + * Copyright (c) 2023, 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 + */ + +#pragma once + +#include +#include + +#include +#include + +namespace cuco::detail { + +template +struct is_std_pair_like : cuda::std::false_type { +}; + +template +struct is_std_pair_like(cuda::std::declval())), + decltype(cuda::std::get<1>(cuda::std::declval()))>> + : cuda::std::conditional_t::value == 2, + cuda::std::true_type, + cuda::std::false_type> { +}; + +template +struct is_thrust_pair_like_impl : cuda::std::false_type { +}; + +template +struct is_thrust_pair_like_impl< + T, + cuda::std::void_t(cuda::std::declval())), + decltype(thrust::get<1>(cuda::std::declval()))>> + : cuda::std::conditional_t::value == 2, + cuda::std::true_type, + cuda::std::false_type> { +}; + +template +struct is_thrust_pair_like + : is_thrust_pair_like_impl()))>> { +}; + +} // namespace cuco::detail diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh index 0a337d4f5..fdded70f5 100644 --- a/include/cuco/detail/utils.cuh +++ b/include/cuco/detail/utils.cuh @@ -19,6 +19,8 @@ #include +#include + namespace cuco { namespace detail { @@ -102,5 +104,95 @@ struct strong_type { T value; ///< Underlying value }; +/** + * @brief Gives value to use as alignment for a pair type that is at least the + * size of the sum of the size of the first type and second type, or 16, + * whichever is smaller. + */ +template +constexpr std::size_t pair_alignment() +{ + return std::min(std::size_t{16}, cuda::std::bit_ceil(sizeof(First) + sizeof(Second))); +} + +/** + * @brief Denotes the equivalent packed type based on the size of the object. + * + * @tparam N The size of the object + */ +template +struct packed { + using type = void; ///< `void` type by default +}; +/** + * @brief Denotes the packed type when the size of the object is 8. + */ +template <> +struct packed { + using type = uint64_t; ///< Packed type as `uint64_t` if the size of the object is 8 +}; +/** + * @brief Denotes the packed type when the size of the object is 4. + */ +template <> +struct packed { + using type = uint32_t; ///< Packed type as `uint32_t` if the size of the object is 4 +}; +template +using packed_t = typename packed::type; + +/** + * @brief Indicates if a pair type can be packed. + * + * When the size of the key,value pair being inserted into the hash table is + * equal in size to a type where atomicCAS is natively supported, it is more + * efficient to "pack" the pair and insert it with a single atomicCAS. + * + * Pair types whose key and value have the same object representation may be + * packed. Also, the `Pair` must not contain any padding bits otherwise + * accessing the packed value would be undefined. + * + * @tparam Pair The pair type that will be packed + * + * @return true If the pair type can be packed + * @return false If the pair type cannot be packed + */ +template +constexpr bool is_packable() +{ + return not std::is_void>::value and std::has_unique_object_representations_v; +} +/** + * @brief Allows viewing a pair in a packed representation. + * + * Used as an optimization for inserting when a pair can be inserted with a + * single atomicCAS + */ +template +union pair_converter { + using packed_type = packed_t; ///< The packed pair type + packed_type packed; ///< The pair in the packed representation + Pair pair; ///< The pair in the pair representation + + /** + * @brief Constructs a pair converter by copying from `p` + * + * @tparam T Type that is convertible to `Pair` + * + * @param p The pair to copy from + */ + template + __device__ pair_converter(T&& p) : pair{p} + { + } + + /** + * @brief Constructs a pair converter by copying from `p` + * + * @param p The packed data to copy from + */ + __device__ pair_converter(packed_type p) : packed{p} {} +}; + } // namespace detail } // namespace cuco diff --git a/include/cuco/dynamic_map.cuh b/include/cuco/dynamic_map.cuh index b9dcf9f22..8aa87163f 100644 --- a/include/cuco/dynamic_map.cuh +++ b/include/cuco/dynamic_map.cuh @@ -101,7 +101,7 @@ class dynamic_map { static_assert(std::is_arithmetic::value, "Unsupported, non-arithmetic key type."); public: - using value_type = cuco::pair_type; ///< Type of key/value pairs + using value_type = cuco::pair; ///< Type of key/value pairs using key_type = Key; ///< Key type using mapped_type = Value; ///< Type of mapped values using atomic_ctr_type = cuda::atomic; ///< Atomic counter type diff --git a/include/cuco/pair.cuh b/include/cuco/pair.cuh new file mode 100644 index 000000000..0a804cc04 --- /dev/null +++ b/include/cuco/pair.cuh @@ -0,0 +1,146 @@ +/* + * Copyright (c) 2020-2023, 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. + */ + +#pragma once + +#include +#include + +#include +#include + +#include +#include + +namespace cuco { + +/** + * @brief Custom pair type + * + * @note This is necessary because `thrust::pair` is under aligned. + * + * @tparam First Type of the first value in the pair + * @tparam Second Type of the second value in the pair + */ +template +struct alignas(detail::pair_alignment()) pair { + using first_type = First; ///< Type of the first value in the pair + using second_type = Second; ///< Type of the second value in the pair + + pair() = default; + ~pair() = default; + pair(pair const&) = default; ///< Copy constructor + pair(pair&&) = default; ///< Move constructor + + /** + * @brief Replaces the contents of the pair with another pair. + * + * @return Reference of the current pair object + */ + pair& operator=(pair const&) = default; + + /** + * @brief Replaces the contents of the pair with another pair. + * + * @return Reference of the current pair object + */ + pair& operator=(pair&&) = default; + + /** + * @brief Constructs a pair from objects `f` and `s`. + * + * @param f The object to copy into `first` + * @param s The object to copy into `second` + */ + __host__ __device__ constexpr pair(First const& f, Second const& s); + + /** + * @brief Constructs a pair by copying from the given pair `p`. + * + * @tparam F Type of the first value of `p` + * @tparam S Type of the second value of `p` + * + * @param p The pair to copy from + */ + template + __host__ __device__ constexpr pair(pair const& p); + + /** + * @brief Constructs a pair from the given std::pair-like `p`. + * + * @tparam T Type of the pair to copy from + * + * @param p The input pair to copy from + */ + template ::value>* = nullptr> + __host__ __device__ constexpr pair(T const& p) + : pair{std::get<0>(thrust::raw_reference_cast(p)), std::get<1>(thrust::raw_reference_cast(p))} + { + } + + /** + * @brief Constructs a pair from the given thrust::pair-like `p`. + * + * @tparam T Type of the pair to copy from + * + * @param p The input pair to copy from + */ + template ::value>* = nullptr> + __host__ __device__ constexpr pair(T const& p) + : pair{thrust::get<0>(thrust::raw_reference_cast(p)), + thrust::get<1>(thrust::raw_reference_cast(p))} + { + } + + First first; ///< The first value in the pair + Second second; ///< The second value in the pair +}; + +/** + * @brief Creates a pair with the given first and second elements + * + * @tparam F Type of first element + * @tparam S Type of second element + * + * @param f First element + * @param s Second element + * + * @return A pair with first element `f` and second element `s`. + */ +template +__host__ __device__ constexpr pair, std::decay_t> make_pair(F&& f, + S&& s) noexcept; + +/** + * @brief Tests if both elements of lhs and rhs are equal + * + * @tparam T1 Type of the first element of the left-hand side pair + * @tparam T2 Type of the second element of the left-hand side pair + * @tparam U1 Type of the first element of the right-hand side pair + * @tparam U2 Type of the second element of the right-hand side pair + * + * @param lhs Left-hand side pair + * @param rhs Right-hand side pair + * + * @return True if two pairs are equal. False otherwise + */ +template +__host__ __device__ constexpr bool operator==(cuco::pair const& lhs, + cuco::pair const& rhs) noexcept; + +} // namespace cuco + +#include diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh index c1aeb6965..f296e9ed1 100644 --- a/include/cuco/static_map.cuh +++ b/include/cuco/static_map.cuh @@ -18,9 +18,9 @@ #include #include -#include #include #include +#include #include #include #include @@ -577,14 +577,14 @@ class static_map { friend class dynamic_map; ///< Dynamic map as friend class public: - using value_type = cuco::pair_type; ///< Type of key/value pairs + using value_type = cuco::pair; ///< Type of key/value pairs using key_type = Key; ///< Key type using mapped_type = Value; ///< Type of mapped values using atomic_key_type = cuda::atomic; ///< Type of atomic keys using atomic_mapped_type = cuda::atomic; ///< Type of atomic mapped values using pair_atomic_type = - cuco::pair_type; ///< Pair type of atomic key and atomic mapped value + cuco::pair; ///< Pair type of atomic key and atomic mapped value using slot_type = pair_atomic_type; ///< Type of hash map slots using atomic_ctr_type = cuda::atomic; ///< Atomic counter type using allocator_type = Allocator; ///< Allocator type diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh index fe68da32b..075848dd2 100644 --- a/include/cuco/static_multimap.cuh +++ b/include/cuco/static_multimap.cuh @@ -149,14 +149,14 @@ class static_multimap { "cuco::linear_probing."); public: - using value_type = cuco::pair_type; ///< Type of key/value pairs + using value_type = cuco::pair; ///< Type of key/value pairs using key_type = Key; ///< Key type using mapped_type = Value; ///< Type of mapped values using atomic_key_type = cuda::atomic; ///< Type of atomic keys using atomic_mapped_type = cuda::atomic; ///< Type of atomic mapped values using pair_atomic_type = - cuco::pair_type; ///< Pair type of atomic key and atomic mapped value + cuco::pair; ///< Pair type of atomic key and atomic mapped value using atomic_ctr_type = cuda::atomic; ///< Atomic counter type using allocator_type = Allocator; ///< Allocator type using slot_allocator_type = typename std::allocator_traits::rebind_alloc< diff --git a/include/cuco/utility/traits.hpp b/include/cuco/utility/traits.hpp index 78e8dabcb..1a6252dcb 100644 --- a/include/cuco/utility/traits.hpp +++ b/include/cuco/utility/traits.hpp @@ -16,6 +16,9 @@ #pragma once +#include +#include + #include namespace cuco { diff --git a/tests/dynamic_map/unique_sequence_test.cu b/tests/dynamic_map/unique_sequence_test.cu index 820fb95f8..aa01ca51a 100644 --- a/tests/dynamic_map/unique_sequence_test.cu +++ b/tests/dynamic_map/unique_sequence_test.cu @@ -48,9 +48,9 @@ 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(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto pairs_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); diff --git a/tests/static_map/custom_type_test.cu b/tests/static_map/custom_type_test.cu index ac743037a..e23216ca3 100644 --- a/tests/static_map/custom_type_test.cu +++ b/tests/static_map/custom_type_test.cu @@ -131,9 +131,9 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", insert_values.begin(), [] __device__(auto i) { return Value{i}; }); - auto insert_pairs = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto insert_pairs = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); SECTION("All inserted keys-value pairs should be correctly recovered during find") { @@ -212,7 +212,7 @@ 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_type const& pair) { + insert_pairs, insert_pairs + num, [view] __device__(cuco::pair const& pair) { return view.contains(pair.first, hash_custom_key{}, custom_key_equals{}); })); } @@ -220,12 +220,11 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", 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_type const& pair) mutable { - return m_view.insert(pair, hash_custom_key{}, custom_key_equals{}); - })); + REQUIRE(cuco::test::all_of(insert_pairs, + insert_pairs + num, + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert(pair, hash_custom_key{}, custom_key_equals{}); + })); } SECTION("Cannot find any key in an empty hash map") @@ -236,7 +235,7 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", REQUIRE(cuco::test::all_of( insert_pairs, insert_pairs + num, - [view] __device__(cuco::pair_type const& pair) mutable { + [view] __device__(cuco::pair const& pair) mutable { return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end(); })); } @@ -245,9 +244,7 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", { auto const view = map.get_device_view(); REQUIRE(cuco::test::all_of( - insert_pairs, - insert_pairs + num, - [view] __device__(cuco::pair_type const& pair) { + insert_pairs, insert_pairs + num, [view] __device__(cuco::pair const& pair) { return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end(); })); } diff --git a/tests/static_map/duplicate_keys_test.cu b/tests/static_map/duplicate_keys_test.cu index 54d1c42f1..5620fa4e9 100644 --- a/tests/static_map/duplicate_keys_test.cu +++ b/tests/static_map/duplicate_keys_test.cu @@ -49,7 +49,7 @@ TEMPLATE_TEST_CASE_SIG("Duplicate keys", auto pairs_begin = thrust::make_transform_iterator( thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair_type(i / 2, i / 2); }); + [] __device__(auto i) { return cuco::pair(i / 2, i / 2); }); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); diff --git a/tests/static_map/heterogeneous_lookup_test.cu b/tests/static_map/heterogeneous_lookup_test.cu index 17b7d5662..e842612b1 100644 --- a/tests/static_map/heterogeneous_lookup_test.cu +++ b/tests/static_map/heterogeneous_lookup_test.cu @@ -99,9 +99,9 @@ TEMPLATE_TEST_CASE("Heterogeneous lookup", cuco::static_map map{ capacity, cuco::empty_key{sentinel_key}, cuco::empty_value{sentinel_value}}; - auto insert_pairs = thrust::make_transform_iterator( - thrust::counting_iterator(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto insert_pairs = + thrust::make_transform_iterator(thrust::counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator(0), [] __device__(auto i) { return ProbeKey(i); }); diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu index dcf88c99c..d8441e9cf 100644 --- a/tests/static_map/key_sentinel_test.cu +++ b/tests/static_map/key_sentinel_test.cu @@ -51,19 +51,19 @@ 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(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto pairs_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(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_type const& pair) mutable { - return m_view.insert(pair, cuco::murmurhash3_32{}, custom_equals{}); - })); + REQUIRE(cuco::test::all_of(pairs_begin, + pairs_begin + num_keys, + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert( + pair, cuco::murmurhash3_32{}, custom_equals{}); + })); } SECTION( @@ -72,13 +72,11 @@ TEMPLATE_TEST_CASE_SIG( map.insert( pairs_begin, pairs_begin + num_keys, cuco::murmurhash3_32{}, custom_equals{}); // 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_type 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); - })); + REQUIRE(cuco::test::all_of( + pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair 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); + })); } } diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu index 35cd7e821..2bf71e2e6 100644 --- a/tests/static_map/stream_test.cu +++ b/tests/static_map/stream_test.cu @@ -53,9 +53,9 @@ 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(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto pairs_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); auto hash_fn = cuco::murmurhash3_32{}; auto equal_fn = thrust::equal_to{}; diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu index 54ba4e6f1..3fa4ef219 100644 --- a/tests/static_map/unique_sequence_test.cu +++ b/tests/static_map/unique_sequence_test.cu @@ -51,9 +51,9 @@ 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(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto pairs_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); @@ -87,69 +87,63 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", SECTION("Inserting unique keys should return insert success.") { - REQUIRE( - cuco::test::all_of(pairs_begin, - pairs_begin + num_keys, - [m_view] __device__(cuco::pair_type const& pair) mutable { - return m_view.insert(pair); - })); + REQUIRE(cuco::test::all_of(pairs_begin, + pairs_begin + num_keys, + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert(pair); + })); } SECTION("Cannot find any key in an empty hash map with non-const view") { SECTION("non-const view") - { - REQUIRE( - cuco::test::all_of(pairs_begin, - pairs_begin + num_keys, - [view] __device__(cuco::pair_type const& pair) mutable { - return view.find(pair.first) == view.end(); - })); - } - SECTION("const view") { REQUIRE(cuco::test::all_of(pairs_begin, pairs_begin + num_keys, - [view] __device__(cuco::pair_type const& pair) { + [view] __device__(cuco::pair const& pair) mutable { return view.find(pair.first) == view.end(); })); } + SECTION("const view") + { + REQUIRE(cuco::test::all_of( + pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair const& pair) { + return view.find(pair.first) == view.end(); + })); + } } SECTION("Keys are all found after inserting many keys.") { // Bulk insert keys - thrust::for_each(thrust::device, - pairs_begin, - pairs_begin + num_keys, - [m_view] __device__(cuco::pair_type const& pair) mutable { - m_view.insert(pair); - }); + thrust::for_each( + thrust::device, + pairs_begin, + pairs_begin + num_keys, + [m_view] __device__(cuco::pair const& pair) mutable { m_view.insert(pair); }); SECTION("non-const view") - { - // All keys should be found - REQUIRE(cuco::test::all_of( - pairs_begin, - pairs_begin + num_keys, - [view] __device__(cuco::pair_type const& pair) mutable { - auto const found = view.find(pair.first); - return (found != view.end()) and - (found->first.load() == pair.first and found->second.load() == pair.second); - })); - } - SECTION("const view") { // All keys should be found REQUIRE(cuco::test::all_of(pairs_begin, pairs_begin + num_keys, - [view] __device__(cuco::pair_type const& pair) { + [view] __device__(cuco::pair const& pair) mutable { auto const found = view.find(pair.first); return (found != view.end()) and (found->first.load() == pair.first and found->second.load() == pair.second); })); } + SECTION("const view") + { + // All keys should be found + REQUIRE(cuco::test::all_of( + pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair 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); + })); + } } } @@ -165,10 +159,10 @@ __inline__ void test_unique_sequence(Map& map, size_type num_keys) thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); - auto keys_begin = d_keys.begin(); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto keys_begin = d_keys.begin(); + auto pairs_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); thrust::device_vector d_contained(num_keys); auto zip_equal = [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }; diff --git a/tests/static_multimap/custom_pair_retrieve_test.cu b/tests/static_multimap/custom_pair_retrieve_test.cu index 563abd835..b9aba7e76 100644 --- a/tests/static_multimap/custom_pair_retrieve_test.cu +++ b/tests/static_multimap/custom_pair_retrieve_test.cu @@ -35,8 +35,8 @@ // Custom pair equal template struct pair_equal { - __device__ bool operator()(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) const + __device__ bool operator()(const cuco::pair& lhs, + const cuco::pair& rhs) const { return lhs.first == rhs.first; } @@ -86,7 +86,7 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - thrust::device_vector> d_pairs(num_pairs); + thrust::device_vector> d_pairs(num_pairs); // pair multiplicity = 2 thrust::transform(thrust::device, @@ -94,7 +94,7 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) thrust::counting_iterator(num_pairs), d_pairs.begin(), [] __device__(auto i) { - return cuco::pair_type{i / 2, i}; + return cuco::pair{i / 2, i}; }); auto pair_begin = d_pairs.begin(); @@ -107,7 +107,7 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) thrust::counting_iterator(num_pairs), pair_begin, [] __device__(auto i) { - return cuco::pair_type{i, i}; + return cuco::pair{i, i}; }); // create an array of prefix sum diff --git a/tests/static_multimap/custom_type_test.cu b/tests/static_multimap/custom_type_test.cu index d76404b18..f53719205 100644 --- a/tests/static_multimap/custom_type_test.cu +++ b/tests/static_multimap/custom_type_test.cu @@ -98,7 +98,7 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) auto count = map.count(key_begin, key_begin + num_pairs, stream, key_pair_equals{}); REQUIRE(count == num_pairs); - thrust::device_vector> found_pairs(num_pairs); + thrust::device_vector> found_pairs(num_pairs); auto output_end = map.retrieve( key_begin, key_begin + num_pairs, found_pairs.begin(), stream, key_pair_equals{}); std::size_t const size = std::distance(found_pairs.begin(), output_end); @@ -110,16 +110,17 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) thrust::device, found_pairs.begin(), found_pairs.end(), - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { return lhs.first.a < rhs.first.a; }); - - REQUIRE(cuco::test::equal( - pair_begin, - pair_begin + num_pairs, - found_pairs.begin(), - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first.a == rhs.first.a; - })); + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + return lhs.first.a < rhs.first.a; + }); + + REQUIRE( + cuco::test::equal(pair_begin, + pair_begin + num_pairs, + found_pairs.begin(), + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first.a == rhs.first.a; + })); } SECTION("Non-matches are not included in the output") @@ -141,7 +142,7 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) auto count = map.count(query_key_begin, query_key_begin + num, stream, key_pair_equals{}); REQUIRE(count == num_pairs); - thrust::device_vector> found_pairs(num_pairs); + thrust::device_vector> found_pairs(num_pairs); auto output_end = map.retrieve( query_key_begin, query_key_begin + num, found_pairs.begin(), stream, key_pair_equals{}); std::size_t const size = std::distance(found_pairs.begin(), output_end); @@ -153,15 +154,16 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) thrust::device, found_pairs.begin(), found_pairs.end(), - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { return lhs.first.a < rhs.first.a; }); - REQUIRE(cuco::test::equal( - pair_begin, - pair_begin + num_pairs, - found_pairs.begin(), - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first.a == rhs.first.a; - })); + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + return lhs.first.a < rhs.first.a; + }); + REQUIRE( + cuco::test::equal(pair_begin, + pair_begin + num_pairs, + found_pairs.begin(), + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first.a == rhs.first.a; + })); } SECTION("Outer functions include non-matches in the output") @@ -183,7 +185,7 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) map.count_outer(query_key_begin, query_key_begin + num, stream, key_pair_equals{}); REQUIRE(count_outer == num); - thrust::device_vector> found_pairs(num); + thrust::device_vector> found_pairs(num); auto output_end = map.retrieve_outer( query_key_begin, query_key_begin + num, found_pairs.begin(), stream, key_pair_equals{}); std::size_t const size_outer = std::distance(found_pairs.begin(), output_end); diff --git a/tests/static_multimap/heterogeneous_lookup_test.cu b/tests/static_multimap/heterogeneous_lookup_test.cu index 9b724d43c..5a5b8b242 100644 --- a/tests/static_multimap/heterogeneous_lookup_test.cu +++ b/tests/static_multimap/heterogeneous_lookup_test.cu @@ -103,9 +103,9 @@ TEMPLATE_TEST_CASE("Heterogeneous lookup", cuco::linear_probing<1, custom_hasher>> map{capacity, cuco::empty_key{sentinel_key}, cuco::empty_value{sentinel_value}}; - auto insert_pairs = thrust::make_transform_iterator( - thrust::counting_iterator(0), - [] __device__(auto i) { return cuco::pair_type(i, i); }); + auto insert_pairs = + thrust::make_transform_iterator(thrust::counting_iterator(0), + [] __device__(auto i) { return cuco::pair(i, i); }); auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator(0), [] __device__(auto i) { return ProbeKey(i); }); diff --git a/tests/static_multimap/insert_if_test.cu b/tests/static_multimap/insert_if_test.cu index 0d560ff6e..33f0b2ce3 100644 --- a/tests/static_multimap/insert_if_test.cu +++ b/tests/static_multimap/insert_if_test.cu @@ -55,7 +55,7 @@ TEMPLATE_TEST_CASE_SIG( constexpr std::size_t num_keys{1'000}; thrust::device_vector d_keys(num_keys); - thrust::device_vector> d_pairs(num_keys); + thrust::device_vector> d_pairs(num_keys); thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); // multiplicity = 1 @@ -64,7 +64,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(num_keys), d_pairs.begin(), [] __device__(auto i) { - return cuco::pair_type{i, i}; + return cuco::pair{i, i}; }); using probe = std::conditional_t< diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu index f21d52c3d..650145a41 100644 --- a/tests/static_multimap/multiplicity_test.cu +++ b/tests/static_multimap/multiplicity_test.cu @@ -36,7 +36,7 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) using Value = typename Map::mapped_type; thrust::device_vector d_keys(num_items / 2); - thrust::device_vector> d_pairs(num_items); + thrust::device_vector> d_pairs(num_items); thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); // multiplicity = 2 @@ -45,10 +45,10 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) thrust::counting_iterator(num_items), d_pairs.begin(), [] __device__(auto i) { - return cuco::pair_type{i / 2, i}; + return cuco::pair{i / 2, i}; }); - thrust::device_vector> d_results(num_items); + thrust::device_vector> d_results(num_items); auto key_begin = d_keys.begin(); auto pair_begin = d_pairs.begin(); @@ -91,22 +91,22 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) REQUIRE(size == num_items); // sort before compare - thrust::sort(thrust::device, - d_results.begin(), - d_results.end(), - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); - - REQUIRE(cuco::test::equal( - pair_begin, - pair_begin + num_items, - output_begin, - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + thrust::sort( + thrust::device, + d_results.begin(), + d_results.end(), + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + }); + + REQUIRE( + cuco::test::equal(pair_begin, + pair_begin + num_items, + output_begin, + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + })); } SECTION("count and count_outer should return the same value.") @@ -129,22 +129,22 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) REQUIRE(size == size_outer); // sort before compare - thrust::sort(thrust::device, - d_results.begin(), - d_results.end(), - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); - - REQUIRE(cuco::test::equal( - pair_begin, - pair_begin + num_items, - output_begin, - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + thrust::sort( + thrust::device, + d_results.begin(), + d_results.end(), + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + }); + + REQUIRE( + cuco::test::equal(pair_begin, + pair_begin + num_items, + output_begin, + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + })); } } diff --git a/tests/static_multimap/non_match_test.cu b/tests/static_multimap/non_match_test.cu index be76a38ce..afa6a938c 100644 --- a/tests/static_multimap/non_match_test.cu +++ b/tests/static_multimap/non_match_test.cu @@ -39,7 +39,7 @@ __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, s SECTION("Output of count and retrieve should be coherent.") { auto num = map.count(key_begin, key_begin + num_keys); - thrust::device_vector> d_results(num); + thrust::device_vector> d_results(num); REQUIRE(num == num_keys); @@ -50,28 +50,28 @@ __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, s REQUIRE(size == num_keys); // sort before compare - thrust::sort(thrust::device, - output_begin, - output_end, - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); - - REQUIRE(cuco::test::equal( - pair_begin, - pair_begin + num_keys, + thrust::sort( + thrust::device, output_begin, - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + output_end, + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + }); + + REQUIRE( + cuco::test::equal(pair_begin, + pair_begin + num_keys, + output_begin, + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + })); } SECTION("Output of count_outer and retrieve_outer should be coherent.") { auto num = map.count_outer(key_begin, key_begin + num_keys); - thrust::device_vector> d_results(num); + thrust::device_vector> d_results(num); REQUIRE(num == (num_keys + num_keys / 2)); @@ -82,34 +82,34 @@ __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, s REQUIRE(size == (num_keys + num_keys / 2)); // sort before compare - thrust::sort(thrust::device, - output_begin, - output_end, - [] __device__(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); + thrust::sort( + thrust::device, + output_begin, + output_end, + [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + }); // create gold reference - thrust::device_vector> gold(size); + thrust::device_vector> gold(size); auto gold_begin = gold.begin(); thrust::transform(thrust::device, thrust::counting_iterator(0), thrust::counting_iterator(size), gold_begin, [num_keys] __device__(auto i) { - if (i < num_keys) { return cuco::pair_type{i / 2, i}; } - return cuco::pair_type{i - num_keys / 2, -1}; + if (i < num_keys) { return cuco::pair{i / 2, i}; } + return cuco::pair{i - num_keys / 2, -1}; }); - REQUIRE(cuco::test::equal( - gold_begin, - gold_begin + size, - output_begin, - [] __device__(cuco::pair_type lhs, cuco::pair_type rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + REQUIRE( + cuco::test::equal(gold_begin, + gold_begin + size, + output_begin, + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + })); } } @@ -127,7 +127,7 @@ TEMPLATE_TEST_CASE_SIG( constexpr std::size_t num_keys{1'000}; thrust::device_vector d_keys(num_keys); - thrust::device_vector> d_pairs(num_keys); + thrust::device_vector> d_pairs(num_keys); thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); // multiplicity = 2 @@ -136,7 +136,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(num_keys), d_pairs.begin(), [] __device__(auto i) { - return cuco::pair_type{i / 2, i}; + return cuco::pair{i / 2, i}; }); using probe = std::conditional_t< diff --git a/tests/static_multimap/pair_function_test.cu b/tests/static_multimap/pair_function_test.cu index 8edecd6f1..2eae61a66 100644 --- a/tests/static_multimap/pair_function_test.cu +++ b/tests/static_multimap/pair_function_test.cu @@ -32,8 +32,8 @@ // Custom pair equal template struct pair_equal { - __device__ bool operator()(const cuco::pair_type& lhs, - const cuco::pair_type& rhs) const + __device__ bool operator()(const cuco::pair& lhs, + const cuco::pair& rhs) const { return lhs.first == rhs.first; } @@ -54,7 +54,7 @@ __inline__ void test_pair_functions(Map& map, PairIt pair_begin, std::size_t num thrust::counting_iterator(num_pairs), pair_begin, [] __device__(auto i) { - return cuco::pair_type{i, i}; + return cuco::pair{i, i}; }); SECTION("pair_contains returns true for all inserted pairs and false for non-inserted ones.") @@ -121,7 +121,7 @@ TEMPLATE_TEST_CASE_SIG( (int64_t, int64_t, cuco::test::probe_sequence::double_hashing)) { constexpr std::size_t num_pairs{4}; - thrust::device_vector> d_pairs(num_pairs); + thrust::device_vector> d_pairs(num_pairs); // pair multiplicity = 2 thrust::transform(thrust::device, @@ -129,7 +129,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(num_pairs), d_pairs.begin(), [] __device__(auto i) { - return cuco::pair_type{i / 2, i}; + return cuco::pair{i / 2, i}; }); using probe = std::conditional_t< diff --git a/tests/utility/storage_test.cu b/tests/utility/storage_test.cu index c82b5ab44..afb9848d3 100644 --- a/tests/utility/storage_test.cu +++ b/tests/utility/storage_test.cu @@ -16,9 +16,9 @@ #include -#include #include #include +#include #include #include