From b7229465f0697aefa79e00f950a87e45d781b598 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Sun, 22 May 2022 11:47:17 -0400 Subject: [PATCH 1/2] Cleanups: remove redundant functor definition --- .../detail/static_multimap/static_multimap.inl | 18 +++--------------- 1 file changed, 3 insertions(+), 15 deletions(-) diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl index 175852acf..2c9bf4141 100644 --- a/include/cuco/detail/static_multimap/static_multimap.inl +++ b/include/cuco/detail/static_multimap/static_multimap.inl @@ -14,6 +14,7 @@ * limitations under the License. */ +#include #include #include @@ -22,18 +23,6 @@ #include -namespace { -/** - * @brief Device functor used to determine if a slot is filled. - */ -template -struct slot_is_filled { - slot_is_filled(Key s) : empty_key_sentinel{s} {} - __device__ __forceinline__ bool operator()(Key const& k) { return k != empty_key_sentinel; } - Key empty_key_sentinel; -}; -} // anonymous namespace - namespace cuco { template ::get_size( cudaStream_t stream) const noexcept { - auto begin = thrust::make_transform_iterator( - raw_slots(), [] __device__(cuco::pair_type const& pair) { return pair.first; }); - slot_is_filled filled(empty_key_sentinel_); + auto begin = thrust::make_transform_iterator(raw_slots(), detail::slot_to_tuple{}); + auto filled = cuco::detail::slot_is_filled{empty_key_sentinel_}; return thrust::count_if(thrust::cuda::par.on(stream), begin, begin + capacity_, filled); } From 4f851b3f3087f498bd9133b47601268f63379098 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Sun, 22 May 2022 14:49:41 -0400 Subject: [PATCH 2/2] Minor cleanups: use getters --- include/cuco/detail/static_multimap/static_multimap.inl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl index 2c9bf4141..d0be7c6c4 100644 --- a/include/cuco/detail/static_multimap/static_multimap.inl +++ b/include/cuco/detail/static_multimap/static_multimap.inl @@ -854,9 +854,9 @@ std::size_t static_multimap::get_si cudaStream_t stream) const noexcept { auto begin = thrust::make_transform_iterator(raw_slots(), detail::slot_to_tuple{}); - auto filled = cuco::detail::slot_is_filled{empty_key_sentinel_}; + auto filled = cuco::detail::slot_is_filled{get_empty_key_sentinel()}; - return thrust::count_if(thrust::cuda::par.on(stream), begin, begin + capacity_, filled); + return thrust::count_if(thrust::cuda::par.on(stream), begin, begin + get_capacity(), filled); } template