Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add specialization for is_bitwise_comparable<cuco::pair>> #335

Closed
wants to merge 2 commits into from
Closed

Add specialization for is_bitwise_comparable<cuco::pair>> #335

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4376e78
Add specialization for is_bitwise_comparable<cuco::pair>>
sleeepyjack Jul 18, 2023
bf39284
Check for padding
sleeepyjack Jul 18, 2023
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
9 changes: 9 additions & 0 deletions include/cuco/detail/traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,15 @@

namespace cuco::detail {

template <typename T, typename = void>
struct is_bitwise_comparable_impl : std::false_type {
};

template <typename T>
struct is_bitwise_comparable_impl<T, std::enable_if_t<std::has_unique_object_representations_v<T>>>
: std::true_type {
};

template <typename T, typename = void>
struct is_std_pair_like : cuda::std::false_type {
};
Expand Down
11 changes: 11 additions & 0 deletions include/cuco/pair.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@

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

#include <thrust/device_reference.h>
#include <thrust/tuple.h>
Expand Down Expand Up @@ -141,6 +142,16 @@ template <class T1, class T2, class U1, class U2>
__host__ __device__ constexpr bool operator==(cuco::pair<T1, T2> const& lhs,
cuco::pair<U1, U2> const& rhs) noexcept;

/**
* @brief A pair of bitwise comparable types is also bitwise comparable
*/
template <typename First, typename Second>
struct is_bitwise_comparable<pair<First, Second>>
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is dangerous. If someone used a pair<char, int> as their key type, then it would be considered bitwise comparable, even though it would have padding bits.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Right. would adding && sizeof(pair<First, Second>) == sizeof(First)+sizeof(Second) help?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Remind me where we need to perform a bitwise equality on a pair?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Here: https://github.com/NVIDIA/cuCollections/blob/a2833dbfb1b7d4915d530bd6adb45092a87a3b07/include/cuco/detail/open_addressing_ref_impl.cuh#L659C5-L659C76

For a static_map, the value_type of the underlying open_addressing_impl is cuco::pair<Key, Value>.

: std::integral_constant<bool,
is_bitwise_comparable_v<First> && is_bitwise_comparable_v<Second> &&
sizeof(pair<First, Second>) == sizeof(First) + sizeof(Second)> {
};

} // namespace cuco

#include <cuco/detail/pair.inl>
12 changes: 3 additions & 9 deletions include/cuco/utility/traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,8 +16,7 @@

#pragma once

#include <thrust/device_reference.h>
#include <thrust/tuple.h>
#include <cuco/detail/traits.hpp>

#include <type_traits>

Expand All @@ -36,15 +35,10 @@ namespace cuco {
* if a `NaN` bit pattern were used as the empty sentinel value, it may not compare bitwise equal to
* other `NaN` bit patterns.
*
* @note By default, only types with unique object representations are allowed
*/
template <typename T, typename = void>
struct is_bitwise_comparable : std::false_type {
};

/// By default, only types with unique object representations are allowed
template <typename T>
struct is_bitwise_comparable<T, std::enable_if_t<std::has_unique_object_representations_v<T>>>
: std::true_type {
struct is_bitwise_comparable : detail::is_bitwise_comparable_impl<T> {
};

template <typename T>
Expand Down