Skip to content

Commit

Permalink
Fix a bug: expose public device APIs
Browse files Browse the repository at this point in the history
  • Loading branch information
@PointKernel
PointKernel committed Dec 14, 2021
1 parent 57986f4 commit 79e2a18
Show file tree
Hide file tree
Showing 3 changed files with 143 additions and 116 deletions.
2 changes: 2 additions & 0 deletions include/cuco/detail/probe_sequence_impl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -108,6 +108,7 @@ class probe_sequence_impl_base {
{
}

public:
/**
* @brief Returns the capacity of the hash map.
*/
Expand All @@ -126,6 +127,7 @@ class probe_sequence_impl_base {
*/
__device__ __forceinline__ const_iterator get_slots() const noexcept { return slots_; }

protected:
iterator slots_; ///< Pointer to beginning of the hash map slots
const std::size_t capacity_; ///< Total number of slots
}; // class probe_sequence_impl_base
Expand Down
102 changes: 51 additions & 51 deletions include/cuco/detail/static_multimap/device_view_impl.inl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,27 +24,6 @@ template <typename Key,
typename Allocator,
class ProbeSequence>
class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_impl_base {
public:
/**
* @brief Gets the sentinel value used to represent an empty key slot.
*
* @return The sentinel value used to represent an empty key slot
*/
__host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
{
return empty_key_sentinel_;
}

/**
* @brief Gets the sentinel value used to represent an empty value slot.
*
* @return The sentinel value used to represent an empty value slot
*/
__host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
{
return empty_value_sentinel_;
}

protected:
// Import member type definitions from `static_multimap`
using value_type = value_type;
Expand Down Expand Up @@ -81,26 +60,6 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
{
}

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type* get_slots() noexcept
{
return probe_sequence_.get_slots();
}

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
{
return probe_sequence_.get_slots();
}

/**
* @brief Returns the initial slot for a given key `k`
*
Expand Down Expand Up @@ -161,16 +120,6 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
return probe_sequence_.next_slot(s);
}

/**
* @brief Gets the maximum number of elements the hash map can hold.
*
* @return The maximum number of elements the hash map can hold
*/
__host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
{
return probe_sequence_.get_capacity();
}

/**
* @brief Load two key/value pairs from the given slot to the target pair array.
*
Expand All @@ -189,6 +138,57 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
}
}

public:
/**
* @brief Gets the sentinel value used to represent an empty key slot.
*
* @return The sentinel value used to represent an empty key slot
*/
__host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
{
return empty_key_sentinel_;
}

/**
* @brief Gets the sentinel value used to represent an empty value slot.
*
* @return The sentinel value used to represent an empty value slot
*/
__host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
{
return empty_value_sentinel_;
}

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type* get_slots() noexcept
{
return probe_sequence_.get_slots();
}

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
{
return probe_sequence_.get_slots();
}

/**
* @brief Gets the maximum number of elements the hash map can hold.
*
* @return The maximum number of elements the hash map can hold
*/
__host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
{
return probe_sequence_.get_capacity();
}

private:
probe_sequence_type probe_sequence_; ///< Probe sequence used to probe the hash map
Key empty_key_sentinel_{}; ///< Key value that represents an empty slot
Expand Down
155 changes: 90 additions & 65 deletions include/cuco/static_multimap.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -559,6 +559,78 @@ class static_multimap {
class device_mutable_view_impl;
class device_view_impl;

template <typename ViewImpl>
class device_view_base {
protected:
// Import member type definitions from `static_multimap`
using value_type = value_type;
using key_type = Key;
using mapped_type = Value;
using pair_atomic_type = pair_atomic_type;
using iterator = pair_atomic_type*;
using const_iterator = pair_atomic_type const*;
using probe_sequence_type = probe_sequence_type;

__host__ __device__ device_view_base(pair_atomic_type* slots,
std::size_t capacity,
Key empty_key_sentinel,
Value empty_value_sentinel) noexcept
: impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
}

public:
/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type* get_slots() noexcept { return impl_.get_slots(); }

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
{
return impl_.get_slots();
}

/**
* @brief Gets the maximum number of elements the hash map can hold.
*
* @return The maximum number of elements the hash map can hold
*/
__host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
{
return impl_.get_capacity();
}

/**
* @brief Gets the sentinel value used to represent an empty key slot.
*
* @return The sentinel value used to represent an empty key slot
*/
__host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
{
return impl_.get_empty_key_sentinel();
}

/**
* @brief Gets the sentinel value used to represent an empty value slot.
*
* @return The sentinel value used to represent an empty value slot
*/
__host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
{
return impl_.get_empty_value_sentinel();
}

protected:
ViewImpl impl_;
}; // class device_view_base

public:
/**
* @brief Mutable, non-owning view-type that may be used in device code to
Expand All @@ -580,14 +652,14 @@ class static_multimap {
* });
* \endcode
*/
class device_mutable_view {
class device_mutable_view : public device_view_base<device_mutable_view_impl> {
public:
// Import member type definitions from `static_multimap`
using value_type = value_type;
using key_type = Key;
using mapped_type = Value;
using iterator = pair_atomic_type*;
using const_iterator = pair_atomic_type const*;
using view_base_type = device_view_base<device_mutable_view_impl>;
using value_type = typename view_base_type::value_type;
using key_type = typename view_base_type::key_type;
using mapped_type = typename view_base_type::mapped_type;
using iterator = typename view_base_type::iterator;
using const_iterator = typename view_base_type::const_iterator;

/**
* @brief Construct a mutable view of the first `capacity` slots of the
Expand All @@ -604,7 +676,7 @@ class static_multimap {
std::size_t capacity,
Key empty_key_sentinel,
Value empty_value_sentinel) noexcept
: impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel}
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
}

Expand All @@ -620,7 +692,7 @@ class static_multimap {
value_type const& insert_pair) noexcept;

private:
device_mutable_view_impl impl_;
using device_view_base<device_mutable_view_impl>::impl_;
}; // class device mutable view

/**
Expand All @@ -631,14 +703,14 @@ class static_multimap {
* value.
*
*/
class device_view {
class device_view : public device_view_base<device_view_impl> {
public:
// Import member type definitions from `static_multimap`
using value_type = value_type;
using key_type = Key;
using mapped_type = Value;
using iterator = pair_atomic_type*;
using const_iterator = pair_atomic_type const*;
using view_base_type = device_view_base<device_view_impl>;
using value_type = typename view_base_type::value_type;
using key_type = typename view_base_type::key_type;
using mapped_type = typename view_base_type::mapped_type;
using iterator = typename view_base_type::iterator;
using const_iterator = typename view_base_type::const_iterator;

/**
* @brief Construct a view of the first `capacity` slots of the
Expand All @@ -655,55 +727,8 @@ class static_multimap {
std::size_t capacity,
Key empty_key_sentinel,
Value empty_value_sentinel) noexcept
: impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
}

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type* get_slots() noexcept { return impl_.get_slots(); }

/**
* @brief Gets slots array.
*
* @return Slots array
*/
__device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
{
return impl_.get_slots();
}

/**
* @brief Gets the maximum number of elements the hash map can hold.
*
* @return The maximum number of elements the hash map can hold
*/
__host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
{
return impl_.get_capacity();
}

/**
* @brief Gets the sentinel value used to represent an empty key slot.
*
* @return The sentinel value used to represent an empty key slot
*/
__host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
return impl_.get_empty_key_sentinel();
}

/**
* @brief Gets the sentinel value used to represent an empty value slot.
*
* @return The sentinel value used to represent an empty value slot
*/
__host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
{
return impl_.get_empty_value_sentinel();
}

/**
Expand Down Expand Up @@ -1053,7 +1078,7 @@ class static_multimap {
PairEqual pair_equal) noexcept;

private:
device_view_impl impl_;
using device_view_base<device_view_impl>::impl_;
}; // class device_view

/**
Expand Down

0 comments on commit 79e2a18

Please sign in to comment.