Skip to content

Commit e4a8d75

Browse files
PointKernelPointKernel
committed
Merge remote-tracking branch 'upstream/dev' into preallocate-example
2 parents b19c3ac + 162b99e commit e4a8d75

File tree

5 files changed

+12
-12
lines changed

5 files changed

+12
-12
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -696,14 +696,14 @@ CUCO_KERNEL void retrieve(InputProbeIt input_probe,
696696
input_probe + block_end_offset,
697697
output_probe,
698698
output_match,
699-
atomic_counter);
699+
*atomic_counter);
700700
} else {
701701
ref.template retrieve<BlockSize>(block,
702702
input_probe + block_begin_offset,
703703
input_probe + block_end_offset,
704704
output_probe,
705705
output_match,
706-
atomic_counter);
706+
*atomic_counter);
707707
}
708708
}
709709
}

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

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1062,7 +1062,7 @@ class open_addressing_ref_impl {
10621062
* @param output_probe Beginning of the sequence of keys corresponding to matching elements in
10631063
* `output_match`
10641064
* @param output_match Beginning of the sequence of matching elements
1065-
* @param atomic_counter Pointer to an atomic object of integral type that is used to count the
1065+
* @param atomic_counter Atomic object of integral type that is used to count the
10661066
* number of output elements
10671067
*/
10681068
template <int32_t BlockSize,
@@ -1075,7 +1075,7 @@ class open_addressing_ref_impl {
10751075
InputProbeIt input_probe_end,
10761076
OutputProbeIt output_probe,
10771077
OutputMatchIt output_match,
1078-
AtomicCounter* atomic_counter) const
1078+
AtomicCounter& atomic_counter) const
10791079
{
10801080
auto constexpr is_outer = false;
10811081
auto const n = cuco::detail::distance(input_probe_begin, input_probe_end); // TODO include
@@ -1111,7 +1111,7 @@ class open_addressing_ref_impl {
11111111
* @param output_probe Beginning of the sequence of keys corresponding to matching elements in
11121112
* `output_match`
11131113
* @param output_match Beginning of the sequence of matching elements
1114-
* @param atomic_counter Pointer to an atomic object of integral type that is used to count the
1114+
* @param atomic_counter Atomic object of integral type that is used to count the
11151115
* number of output elements
11161116
*/
11171117
template <int32_t BlockSize,
@@ -1124,7 +1124,7 @@ class open_addressing_ref_impl {
11241124
InputProbeIt input_probe_end,
11251125
OutputProbeIt output_probe,
11261126
OutputMatchIt output_match,
1127-
AtomicCounter* atomic_counter) const
1127+
AtomicCounter& atomic_counter) const
11281128
{
11291129
auto constexpr is_outer = true;
11301130
auto const n = cuco::detail::distance(input_probe_begin, input_probe_end); // TODO include
@@ -1161,7 +1161,7 @@ class open_addressing_ref_impl {
11611161
* @param output_probe Beginning of the sequence of keys corresponding to matching elements in
11621162
* `output_match`
11631163
* @param output_match Beginning of the sequence of matching elements
1164-
* @param atomic_counter Pointer to an atomic object of integral type that is used to count the
1164+
* @param atomic_counter Atomic object of integral type that is used to count the
11651165
* number of output elements
11661166
*/
11671167
template <bool IsOuter,
@@ -1175,7 +1175,7 @@ class open_addressing_ref_impl {
11751175
cuco::detail::index_type n,
11761176
OutputProbeIt output_probe,
11771177
OutputMatchIt output_match,
1178-
AtomicCounter* atomic_counter) const
1178+
AtomicCounter& atomic_counter) const
11791179
{
11801180
namespace cg = cooperative_groups;
11811181

@@ -1212,7 +1212,7 @@ class open_addressing_ref_impl {
12121212
size_type offset = 0;
12131213
auto const count = counters[flushing_tile_id];
12141214
auto const rank = tile.thread_rank();
1215-
if (rank == 0) { offset = atomic_counter->fetch_add(count, cuda::memory_order_relaxed); }
1215+
if (rank == 0) { offset = atomic_counter.fetch_add(count, cuda::memory_order_relaxed); }
12161216
offset = tile.shfl(offset, 0);
12171217

12181218
// flush_buffers

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1521,7 +1521,7 @@ class operator_impl<
15211521
InputProbeIt input_probe_end,
15221522
OutputProbeIt output_probe,
15231523
OutputMatchIt output_match,
1524-
AtomicCounter* atomic_counter) const
1524+
AtomicCounter& atomic_counter) const
15251525
{
15261526
auto const& ref_ = static_cast<ref_type const&>(*this);
15271527
ref_.impl_.template retrieve<BlockSize>(

‎include/cuco/detail/static_multimap/static_multimap_ref.inl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -835,7 +835,7 @@ class operator_impl<
835835
InputProbeIt input_probe_end,
836836
OutputProbeIt output_probe,
837837
OutputMatchIt output_match,
838-
AtomicCounter* atomic_counter) const
838+
AtomicCounter& atomic_counter) const
839839
{
840840
auto const& ref_ = static_cast<ref_type const&>(*this);
841841
ref_.impl_.template retrieve<BlockSize>(

‎include/cuco/detail/static_set/static_set_ref.inl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -831,7 +831,7 @@ class operator_impl<op::retrieve_tag,
831831
InputProbeIt input_probe_end,
832832
OutputProbeIt output_probe,
833833
OutputMatchIt output_match,
834-
AtomicCounter* atomic_counter) const
834+
AtomicCounter& atomic_counter) const
835835
{
836836
auto const& ref_ = static_cast<ref_type const&>(*this);
837837
ref_.impl_.template retrieve<BlockSize>(

0 commit comments

Comments
 (0)