Skip to content
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 4 additions & 16 deletions include/cuco/detail/static_multimap/static_multimap.inl
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@
* limitations under the License.
*/

#include <cuco/detail/utils.cuh>
#include <cuco/detail/utils.hpp>

#include <thrust/count.h>
Expand All @@ -22,18 +23,6 @@

#include <iterator>

namespace {
/**
* @brief Device functor used to determine if a slot is filled.
*/
template <typename Key>
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 <typename Key,
Expand Down Expand Up @@ -864,11 +853,10 @@ template <typename Key,
std::size_t static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::get_size(
cudaStream_t stream) const noexcept
{
auto begin = thrust::make_transform_iterator(
raw_slots(), [] __device__(cuco::pair_type<Key, Value> const& pair) { return pair.first; });
slot_is_filled<Key> filled(empty_key_sentinel_);
auto begin = thrust::make_transform_iterator(raw_slots(), detail::slot_to_tuple<Key, Value>{});
auto filled = cuco::detail::slot_is_filled<Key>{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 <typename Key,
Expand Down