Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

Use CUB's ScanByKey implementation #1555

Merged
merged 5 commits into from
May 5, 2022
Merged

Use CUB's ScanByKey implementation #1555

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
9911b8d
Add missing header.
alliepiper Oct 26, 2021
699f55b
Add utilities to convert contiguous iterators to pointers.
alliepiper Apr 25, 2022
f218be1
Update the CUDA scan_by_key impl to use cub's ScanByKey.
alliepiper Oct 26, 2021
5f794f6
Split the scan_by_key test into inclusive/exclusive tests.
alliepiper May 2, 2022
b9c1328
Address review suggestions.
alliepiper May 5, 2022
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
92 changes: 92 additions & 0 deletions testing/is_contiguous_iterator.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -134,3 +134,95 @@ void test_is_contiguous_iterator_vectors()
}
DECLARE_VECTOR_UNITTEST(test_is_contiguous_iterator_vectors);


struct expect_pointer{};
struct expect_passthrough{};

template <typename IteratorT,
typename PointerT,
typename expected_unwrapped_type /* = expect_[pointer|passthrough] */>
struct check_unwrapped_iterator
{
using unwrapped_t = typename std::remove_reference<
decltype(thrust::detail::try_unwrap_contiguous_iterator(
std::declval<IteratorT>()))>::type;

static constexpr bool value =
std::is_same<expected_unwrapped_type, expect_pointer>::value
? std::is_same<unwrapped_t, PointerT>::value
: std::is_same<unwrapped_t, IteratorT>::value;
};

template <typename T>
void test_try_unwrap_contiguous_iterator()
{
// Raw pointers should pass whether expecting pointers or passthrough.
THRUST_STATIC_ASSERT((check_unwrapped_iterator<T *,
T *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<T *,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<T const *,
T const *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<T const *,
T const *,
expect_passthrough>::value));

THRUST_STATIC_ASSERT((check_unwrapped_iterator<thrust::device_ptr<T>,
T *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<thrust::device_ptr<T const>,
T const *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::vector<T>::iterator,
T *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::vector<T>::reverse_iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::array<T, 1>::iterator,
T *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::array<T const, 1>::iterator,
T const *,
expect_pointer>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::list<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::deque<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::set<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::multiset<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::map<T, T>::iterator,
std::pair<T const, T> *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::multimap<T, T>::iterator,
std::pair<T const, T> *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::unordered_set<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::unordered_multiset<T>::iterator,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::unordered_map<T, T>::iterator,
std::pair<T const, T> *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<typename std::unordered_multimap<T, T>::iterator,
std::pair<T const, T> *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<std::istream_iterator<T>,
T *,
expect_passthrough>::value));
THRUST_STATIC_ASSERT((check_unwrapped_iterator<std::ostream_iterator<T>,
void,
expect_passthrough>::value));
}
DECLARE_GENERIC_UNITTEST(test_try_unwrap_contiguous_iterator);
Loading