Skip to content

Commit fa7c7e1

Browse files
PointKernelPointKernel
committed
Merge remote-tracking branch 'upstream/dev' into add-flat-storage
2 parents da0f3db + 2303a7a commit fa7c7e1

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

43 files changed

+130
-118
lines changed

‎README.md‎

Lines changed: 5 additions & 5 deletions
Large diffs are not rendered by default.

‎examples/static_map/count_by_key_example.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,8 @@
1818

1919
#include <cub/block/block_reduce.cuh>
2020
#include <cuda/std/atomic>
21+
#include <cuda/std/functional>
2122
#include <thrust/device_vector.h>
22-
#include <thrust/logical.h>
2323
#include <thrust/transform.h>
2424

2525
#include <cmath>
@@ -132,7 +132,7 @@ int main(void)
132132
load_factor,
133133
cuco::empty_key{empty_key_sentinel},
134134
cuco::empty_value{empty_value_sentinel},
135-
thrust::equal_to<Key>{},
135+
cuda::std::equal_to<Key>{},
136136
cuco::linear_probing<1, cuco::default_hash_function<Key>>{}};
137137

138138
// Get a non-owning, mutable reference of the map that allows `insert_and_find` operation to pass

‎examples/static_map/device_ref_example.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,10 +16,10 @@
1616

1717
#include <cuco/static_map.cuh>
1818

19+
#include <cuda/std/functional>
1920
#include <thrust/device_vector.h>
2021
#include <thrust/execution_policy.h>
2122
#include <thrust/iterator/zip_iterator.h>
22-
#include <thrust/logical.h>
2323
#include <thrust/sequence.h>
2424
#include <thrust/tuple.h>
2525

@@ -138,7 +138,7 @@ int main(void)
138138
auto map = cuco::static_map{capacity,
139139
cuco::empty_key{empty_key_sentinel},
140140
cuco::empty_value{empty_value_sentinel},
141-
thrust::equal_to<Key>{},
141+
cuda::std::equal_to<Key>{},
142142
cuco::linear_probing<1, cuco::default_hash_function<Key>>{}};
143143

144144
// Get a non-owning, mutable reference of the map that allows inserts to pass by value into the

‎examples/static_multiset/host_bulk_example.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -79,4 +79,4 @@ int main(void)
7979
}
8080

8181
return 0;
82-
}
82+
}

‎examples/static_set/device_subsets_example.cu‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@
1818
#include <cuco/storage.cuh>
1919

2020
#include <cuda/std/array>
21+
#include <cuda/std/functional>
2122
#include <thrust/device_vector.h>
2223
#include <thrust/reduce.h>
2324
#include <thrust/scan.h>
@@ -59,7 +60,7 @@ using storage_type = cuco::bucket_storage<key_type, bucket_size>;
5960
using storage_ref_type = typename storage_type::ref_type;
6061
using ref_type = cuco::static_set_ref<key_type,
6162
cuda::thread_scope_device,
62-
thrust::equal_to<key_type>,
63+
cuda::std::equal_to<key_type>,
6364
probing_scheme_type,
6465
storage_ref_type>; ///< Set ref type
6566

‎examples/static_set/shared_memory_example.cu‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,8 @@
1616

1717
#include <cuco/static_set.cuh>
1818

19+
#include <cuda/std/functional>
20+
1921
#include <cooperative_groups.h>
2022

2123
/**
@@ -90,7 +92,7 @@ int main(void)
9092
using set_type = cuco::static_set<Key,
9193
extent_type,
9294
cuda::thread_scope_block,
93-
thrust::equal_to<Key>,
95+
cuda::std::equal_to<Key>,
9496
probing_scheme_type,
9597
cuco::cuda_allocator<Key>,
9698
cuco::storage<bucket_size>>;

‎include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -25,8 +25,8 @@
2525

2626
#include <cuda/atomic>
2727
#include <cuda/std/functional>
28+
#include <cuda/std/iterator>
2829
#include <cuda/std/type_traits>
29-
#include <thrust/distance.h>
3030
#include <thrust/execution_policy.h>
3131
#include <thrust/logical.h>
3232
#include <thrust/reduce.h>
@@ -407,7 +407,7 @@ class open_addressing_ref_impl {
407407
if (eq_res == detail::equal_result::EQUAL) { return false; }
408408
}
409409
if (eq_res == detail::equal_result::AVAILABLE) {
410-
auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
410+
auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
411411
switch (
412412
attempt_insert(get_slot_ptr(*probing_iter, intra_bucket_index), slot_content, val)) {
413413
case insert_result::DUPLICATE: {
@@ -709,7 +709,7 @@ class open_addressing_ref_impl {
709709
if (eq_res == detail::equal_result::EMPTY) { return false; }
710710
// Key exists, return true if successfully deleted
711711
if (eq_res == detail::equal_result::EQUAL) {
712-
auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
712+
auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
713713
switch (attempt_insert_stable(get_slot_ptr(*probing_iter, intra_bucket_index),
714714
slot_content,
715715
this->erased_slot_sentinel())) {

‎include/cuco/detail/static_map/static_map_ref.inl‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@
2121
#include <cuco/operator.hpp>
2222

2323
#include <cuda/atomic>
24+
#include <cuda/std/iterator>
2425
#include <cuda/std/type_traits>
2526
#include <cuda/std/utility>
2627

@@ -525,7 +526,7 @@ class operator_impl<
525526
for (auto& slot_content : bucket_slots) {
526527
auto const eq_res =
527528
ref_.impl_.predicate_.operator()<is_insert::YES>(key, slot_content.first);
528-
auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
529+
auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
529530
auto slot_ptr = ref_.impl_.get_slot_ptr(*probing_iter, intra_bucket_index);
530531

531532
// If the key is already in the container, update the payload and return
@@ -895,7 +896,7 @@ class operator_impl<
895896
for (auto& slot_content : bucket_slots) {
896897
auto const eq_res =
897898
ref_.impl_.predicate_.operator()<is_insert::YES>(key, slot_content.first);
898-
auto const intra_bucket_index = thrust::distance(bucket_slots.begin(), &slot_content);
899+
auto const intra_bucket_index = cuda::std::distance(bucket_slots.begin(), &slot_content);
899900
auto slot_ptr = ref_.impl_.get_slot_ptr(*probing_iter, intra_bucket_index);
900901

901902
// If the key is already in the container, update the payload and return

‎include/cuco/dynamic_map.cuh‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,8 @@
2222
#include <cuco/types.cuh>
2323

2424
#include <cuda/std/atomic>
25+
#include <cuda/std/functional>
2526
#include <thrust/device_vector.h>
26-
#include <thrust/functional.h>
2727

2828
#include <cstddef>
2929
#include <memory>
@@ -52,7 +52,7 @@ template <class Key,
5252
class T,
5353
class Extent = cuco::extent<std::size_t>,
5454
cuda::thread_scope Scope = cuda::thread_scope_device,
55-
class KeyEqual = thrust::equal_to<Key>,
55+
class KeyEqual = cuda::std::equal_to<Key>,
5656
class ProbingScheme = cuco::linear_probing<4, // CG size
5757
cuco::default_hash_function<Key>>,
5858
class Allocator = cuco::cuda_allocator<cuco::pair<Key, T>>,
@@ -347,7 +347,7 @@ class dynamic_map {
347347
*/
348348
template <typename InputIt,
349349
typename Hash = cuco::default_hash_function<key_type>,
350-
typename KeyEqual = thrust::equal_to<key_type>>
350+
typename KeyEqual = cuda::std::equal_to<key_type>>
351351
void insert(InputIt first,
352352
InputIt last,
353353
Hash hash = Hash{},
@@ -386,7 +386,7 @@ class dynamic_map {
386386
*/
387387
template <typename InputIt,
388388
typename Hash = cuco::default_hash_function<key_type>,
389-
typename KeyEqual = thrust::equal_to<key_type>>
389+
typename KeyEqual = cuda::std::equal_to<key_type>>
390390
void erase(InputIt first,
391391
InputIt last,
392392
Hash hash = Hash{},
@@ -416,7 +416,7 @@ class dynamic_map {
416416
template <typename InputIt,
417417
typename OutputIt,
418418
typename Hash = cuco::default_hash_function<key_type>,
419-
typename KeyEqual = thrust::equal_to<key_type>>
419+
typename KeyEqual = cuda::std::equal_to<key_type>>
420420
void find(InputIt first,
421421
InputIt last,
422422
OutputIt output_begin,
@@ -446,7 +446,7 @@ class dynamic_map {
446446
template <typename InputIt,
447447
typename OutputIt,
448448
typename Hash = cuco::default_hash_function<key_type>,
449-
typename KeyEqual = thrust::equal_to<key_type>>
449+
typename KeyEqual = cuda::std::equal_to<key_type>>
450450
void contains(InputIt first,
451451
InputIt last,
452452
OutputIt output_begin,

‎include/cuco/static_map.cuh‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -28,8 +28,8 @@
2828
#include <cuco/utility/traits.hpp>
2929

3030
#include <cuda/std/atomic>
31+
#include <cuda/std/functional>
3132
#include <cuda/stream_ref>
32-
#include <thrust/functional.h>
3333

3434
#if defined(CUCO_HAS_CUDA_BARRIER)
3535
#include <cuda/barrier>
@@ -89,7 +89,7 @@ template <class Key,
8989
class T,
9090
class Extent = cuco::extent<std::size_t>,
9191
cuda::thread_scope Scope = cuda::thread_scope_device,
92-
class KeyEqual = thrust::equal_to<Key>,
92+
class KeyEqual = cuda::std::equal_to<Key>,
9393
class ProbingScheme = cuco::linear_probing<4, // CG size
9494
cuco::default_hash_function<Key>>,
9595
class Allocator = cuco::cuda_allocator<cuco::pair<Key, T>>,
@@ -1392,7 +1392,7 @@ class static_map {
13921392
*/
13931393
template <typename InputIt,
13941394
typename Hash = cuco::default_hash_function<key_type>,
1395-
typename KeyEqual = thrust::equal_to<key_type>>
1395+
typename KeyEqual = cuda::std::equal_to<key_type>>
13961396
void insert(InputIt first,
13971397
InputIt last,
13981398
Hash hash = Hash{},
@@ -1426,7 +1426,7 @@ class static_map {
14261426
typename StencilIt,
14271427
typename Predicate,
14281428
typename Hash = cuco::default_hash_function<key_type>,
1429-
typename KeyEqual = thrust::equal_to<key_type>>
1429+
typename KeyEqual = cuda::std::equal_to<key_type>>
14301430
void insert_if(InputIt first,
14311431
InputIt last,
14321432
StencilIt stencil,
@@ -1464,7 +1464,7 @@ class static_map {
14641464
*/
14651465
template <typename InputIt,
14661466
typename Hash = cuco::default_hash_function<key_type>,
1467-
typename KeyEqual = thrust::equal_to<key_type>>
1467+
typename KeyEqual = cuda::std::equal_to<key_type>>
14681468
void erase(InputIt first,
14691469
InputIt last,
14701470
Hash hash = Hash{},
@@ -1493,7 +1493,7 @@ class static_map {
14931493
template <typename InputIt,
14941494
typename OutputIt,
14951495
typename Hash = cuco::default_hash_function<key_type>,
1496-
typename KeyEqual = thrust::equal_to<key_type>>
1496+
typename KeyEqual = cuda::std::equal_to<key_type>>
14971497
void find(InputIt first,
14981498
InputIt last,
14991499
OutputIt output_begin,
@@ -1548,7 +1548,7 @@ class static_map {
15481548
template <typename InputIt,
15491549
typename OutputIt,
15501550
typename Hash = cuco::default_hash_function<key_type>,
1551-
typename KeyEqual = thrust::equal_to<key_type>>
1551+
typename KeyEqual = cuda::std::equal_to<key_type>>
15521552
void contains(InputIt first,
15531553
InputIt last,
15541554
OutputIt output_begin,
@@ -2070,7 +2070,7 @@ class static_map {
20702070
* @return `true` if the insert was successful, `false` otherwise.
20712071
*/
20722072
template <typename Hash = cuco::default_hash_function<key_type>,
2073-
typename KeyEqual = thrust::equal_to<key_type>>
2073+
typename KeyEqual = cuda::std::equal_to<key_type>>
20742074
__device__ bool insert(value_type const& insert_pair,
20752075
Hash hash = Hash{},
20762076
KeyEqual key_equal = KeyEqual{}) noexcept;
@@ -2101,7 +2101,7 @@ class static_map {
21012101
* either `true` if the insert was successful, `false` otherwise.
21022102
*/
21032103
template <typename Hash = cuco::default_hash_function<key_type>,
2104-
typename KeyEqual = thrust::equal_to<key_type>>
2104+
typename KeyEqual = cuda::std::equal_to<key_type>>
21052105
__device__ thrust::pair<iterator, bool> insert_and_find(
21062106
value_type const& insert_pair, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) noexcept;
21072107

@@ -2126,7 +2126,7 @@ class static_map {
21262126
*/
21272127
template <typename CG,
21282128
typename Hash = cuco::default_hash_function<key_type>,
2129-
typename KeyEqual = thrust::equal_to<key_type>>
2129+
typename KeyEqual = cuda::std::equal_to<key_type>>
21302130
__device__ bool insert(CG const& g,
21312131
value_type const& insert_pair,
21322132
Hash hash = Hash{},
@@ -2147,7 +2147,7 @@ class static_map {
21472147
* @return `true` if the erasure was successful, `false` otherwise.
21482148
*/
21492149
template <typename Hash = cuco::default_hash_function<key_type>,
2150-
typename KeyEqual = thrust::equal_to<key_type>>
2150+
typename KeyEqual = cuda::std::equal_to<key_type>>
21512151
__device__ bool erase(key_type const& k,
21522152
Hash hash = Hash{},
21532153
KeyEqual key_equal = KeyEqual{}) noexcept;
@@ -2170,7 +2170,7 @@ class static_map {
21702170
*/
21712171
template <typename CG,
21722172
typename Hash = cuco::default_hash_function<key_type>,
2173-
typename KeyEqual = thrust::equal_to<key_type>>
2173+
typename KeyEqual = cuda::std::equal_to<key_type>>
21742174
__device__ bool erase(CG const& g,
21752175
key_type const& k,
21762176
Hash hash = Hash{},
@@ -2337,7 +2337,7 @@ class static_map {
23372337
* containing `k` was inserted
23382338
*/
23392339
template <typename Hash = cuco::default_hash_function<key_type>,
2340-
typename KeyEqual = thrust::equal_to<key_type>>
2340+
typename KeyEqual = cuda::std::equal_to<key_type>>
23412341
__device__ iterator find(Key const& k,
23422342
Hash hash = Hash{},
23432343
KeyEqual key_equal = KeyEqual{}) noexcept;
@@ -2357,7 +2357,7 @@ class static_map {
23572357
* containing `k` was inserted
23582358
*/
23592359
template <typename Hash = cuco::default_hash_function<key_type>,
2360-
typename KeyEqual = thrust::equal_to<key_type>>
2360+
typename KeyEqual = cuda::std::equal_to<key_type>>
23612361
__device__ const_iterator find(Key const& k,
23622362
Hash hash = Hash{},
23632363
KeyEqual key_equal = KeyEqual{}) const noexcept;
@@ -2384,7 +2384,7 @@ class static_map {
23842384
*/
23852385
template <typename CG,
23862386
typename Hash = cuco::default_hash_function<key_type>,
2387-
typename KeyEqual = thrust::equal_to<key_type>>
2387+
typename KeyEqual = cuda::std::equal_to<key_type>>
23882388
__device__ iterator
23892389
find(CG g, Key const& k, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) noexcept;
23902390

@@ -2410,7 +2410,7 @@ class static_map {
24102410
*/
24112411
template <typename CG,
24122412
typename Hash = cuco::default_hash_function<key_type>,
2413-
typename KeyEqual = thrust::equal_to<key_type>>
2413+
typename KeyEqual = cuda::std::equal_to<key_type>>
24142414
__device__ const_iterator
24152415
find(CG g, Key const& k, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) const noexcept;
24162416

@@ -2439,7 +2439,7 @@ class static_map {
24392439
*/
24402440
template <typename ProbeKey,
24412441
typename Hash = cuco::default_hash_function<key_type>,
2442-
typename KeyEqual = thrust::equal_to<key_type>>
2442+
typename KeyEqual = cuda::std::equal_to<key_type>>
24432443
__device__ bool contains(ProbeKey const& k,
24442444
Hash hash = Hash{},
24452445
KeyEqual key_equal = KeyEqual{}) const noexcept;
@@ -2474,7 +2474,7 @@ class static_map {
24742474
template <typename CG,
24752475
typename ProbeKey,
24762476
typename Hash = cuco::default_hash_function<key_type>,
2477-
typename KeyEqual = thrust::equal_to<key_type>>
2477+
typename KeyEqual = cuda::std::equal_to<key_type>>
24782478
__device__ std::enable_if_t<std::is_invocable_v<KeyEqual, ProbeKey, Key>, bool> contains(
24792479
CG const& g,
24802480
ProbeKey const& k,

0 commit comments

Comments
 (0)