@@ -449,7 +449,7 @@ class operator_impl<
449449 * @return True if the given element is successfully inserted
450450 */
451451 template <typename Value>
452- __device__ bool insert (Value const & value) noexcept
452+ __device__ bool insert (Value value) noexcept
453453 {
454454 ref_type& ref_ = static_cast <ref_type&>(*this );
455455 return ref_.impl_ .insert (value);
@@ -467,7 +467,7 @@ class operator_impl<
467467 */
468468 template <typename Value>
469469 __device__ bool insert (cooperative_groups::thread_block_tile<cg_size> const & group,
470- Value const & value) noexcept
470+ Value value) noexcept
471471 {
472472 auto & ref_ = static_cast <ref_type&>(*this );
473473 if (ref_.erased_key_sentinel () != ref_.empty_key_sentinel ()) {
@@ -507,7 +507,7 @@ class operator_impl<
507507 * @param value The element to insert
508508 */
509509 template <typename Value>
510- __device__ void insert_or_assign (Value const & value) noexcept
510+ __device__ void insert_or_assign (Value value) noexcept
511511 {
512512 static_assert (cg_size == 1 , " Non-CG operation is incompatible with the current probing scheme" );
513513
@@ -558,7 +558,7 @@ class operator_impl<
558558 */
559559 template <typename Value>
560560 __device__ void insert_or_assign (cooperative_groups::thread_block_tile<cg_size> const & group,
561- Value const & value) noexcept
561+ Value value) noexcept
562562 {
563563 ref_type& ref_ = static_cast <ref_type&>(*this );
564564
@@ -630,7 +630,7 @@ class operator_impl<
630630 * @return Returns `true` if the given `value` is inserted or `value` has a match in the map.
631631 */
632632 template <typename Value>
633- __device__ constexpr bool attempt_insert_or_assign (value_type* slot, Value const & value) noexcept
633+ __device__ constexpr bool attempt_insert_or_assign (value_type* slot, Value value) noexcept
634634 {
635635 ref_type& ref_ = static_cast <ref_type&>(*this );
636636 auto expected_key = ref_.impl_ .empty_slot_sentinel ().first ;
@@ -688,7 +688,7 @@ class operator_impl<
688688 */
689689
690690 template <typename Value, typename Op>
691- __device__ bool insert_or_apply (Value const & value, Op op)
691+ __device__ bool insert_or_apply (Value value, Op op)
692692 {
693693 static_assert (cg_size == 1 , " Non-CG operation is incompatible with the current probing scheme" );
694694
@@ -724,7 +724,7 @@ class operator_impl<
724724 typename Init,
725725 typename Op,
726726 typename = cuda::std::enable_if_t <std::is_convertible_v<Value, value_type>>>
727- __device__ bool insert_or_apply (Value const & value, Init init, Op op)
727+ __device__ bool insert_or_apply (Value value, Init init, Op op)
728728 {
729729 static_assert (cg_size == 1 , " Non-CG operation is incompatible with the current probing scheme" );
730730
@@ -755,7 +755,7 @@ class operator_impl<
755755
756756 template <typename Value, typename Op>
757757 __device__ bool insert_or_apply (cooperative_groups::thread_block_tile<cg_size> const & group,
758- Value const & value,
758+ Value value,
759759 Op op)
760760 {
761761 static_assert (
@@ -787,7 +787,7 @@ class operator_impl<
787787 */
788788 template <typename Value, typename Init, typename Op>
789789 __device__ bool insert_or_apply (cooperative_groups::thread_block_tile<cg_size> const & group,
790- Value const & value,
790+ Value value,
791791 Init init,
792792 Op op)
793793 {
@@ -817,7 +817,7 @@ class operator_impl<
817817 * @return Returns `true` if the given `value` is inserted successfully.
818818 */
819819 template <typename Value, typename Init, typename Op>
820- __device__ bool dispatch_insert_or_apply (Value const & value, Init init, Op op)
820+ __device__ bool dispatch_insert_or_apply (Value value, Init init, Op op)
821821 {
822822 ref_type& ref_ = static_cast <ref_type&>(*this );
823823 // if init equals sentinel value, then we can just `apply` op instead of write
@@ -845,10 +845,7 @@ class operator_impl<
845845 */
846846 template <typename Value, typename Init, typename Op>
847847 __device__ bool dispatch_insert_or_apply (
848- cooperative_groups::thread_block_tile<cg_size> const & group,
849- Value const & value,
850- Init init,
851- Op op)
848+ cooperative_groups::thread_block_tile<cg_size> const & group, Value value, Init init, Op op)
852849 {
853850 ref_type& ref_ = static_cast <ref_type&>(*this );
854851 // if init equals sentinel value, then we can just `apply` op instead of write
@@ -878,7 +875,7 @@ class operator_impl<
878875 * @return Returns `true` if the given `value` is inserted successfully.
879876 */
880877 template <bool UseDirectApply, typename Value, typename Op>
881- __device__ bool insert_or_apply_impl (Value const & value, Op op)
878+ __device__ bool insert_or_apply_impl (Value value, Op op)
882879 {
883880 ref_type& ref_ = static_cast <ref_type&>(*this );
884881
@@ -955,7 +952,7 @@ class operator_impl<
955952 */
956953 template <bool UseDirectApply, typename Value, typename Op>
957954 __device__ bool insert_or_apply_impl (cooperative_groups::thread_block_tile<cg_size> const & group,
958- Value const & value,
955+ Value value,
959956 Op op)
960957 {
961958 ref_type& ref_ = static_cast <ref_type&>(*this );
@@ -1050,10 +1047,8 @@ class operator_impl<
10501047 * and the element to insert.
10511048 */
10521049 template <bool UseDirectApply, typename Value, typename Op>
1053- [[nodiscard]] __device__ insert_result attempt_insert_or_apply (value_type* address,
1054- value_type const & expected,
1055- Value const & desired,
1056- Op op) noexcept
1050+ [[nodiscard]] __device__ insert_result
1051+ attempt_insert_or_apply (value_type* address, value_type expected, Value desired, Op op) noexcept
10571052 {
10581053 ref_type& ref_ = static_cast <ref_type&>(*this );
10591054
@@ -1128,7 +1123,7 @@ class operator_impl<
11281123 * insertion is successful or not.
11291124 */
11301125 template <typename Value>
1131- __device__ cuda::std::pair<iterator, bool > insert_and_find (Value const & value) noexcept
1126+ __device__ cuda::std::pair<iterator, bool > insert_and_find (Value value) noexcept
11321127 {
11331128 ref_type& ref_ = static_cast <ref_type&>(*this );
11341129 return ref_.impl_ .insert_and_find (value);
@@ -1151,7 +1146,7 @@ class operator_impl<
11511146 */
11521147 template <typename Value>
11531148 __device__ cuda::std::pair<iterator, bool > insert_and_find (
1154- cooperative_groups::thread_block_tile<cg_size> const & group, Value const & value) noexcept
1149+ cooperative_groups::thread_block_tile<cg_size> const & group, Value value) noexcept
11551150 {
11561151 ref_type& ref_ = static_cast <ref_type&>(*this );
11571152 return ref_.impl_ .insert_and_find (group, value);
@@ -1187,7 +1182,7 @@ class operator_impl<
11871182 * @return True if the given element is successfully erased
11881183 */
11891184 template <typename ProbeKey>
1190- __device__ bool erase (ProbeKey const & key) noexcept
1185+ __device__ bool erase (ProbeKey key) noexcept
11911186 {
11921187 ref_type& ref_ = static_cast <ref_type&>(*this );
11931188 return ref_.impl_ .erase (key);
@@ -1205,7 +1200,7 @@ class operator_impl<
12051200 */
12061201 template <typename ProbeKey>
12071202 __device__ bool erase (cooperative_groups::thread_block_tile<cg_size> const & group,
1208- ProbeKey const & key) noexcept
1203+ ProbeKey key) noexcept
12091204 {
12101205 auto & ref_ = static_cast <ref_type&>(*this );
12111206 return ref_.impl_ .erase (group, key);
@@ -1244,7 +1239,7 @@ class operator_impl<
12441239 * @return A boolean indicating whether the probe key is present
12451240 */
12461241 template <typename ProbeKey>
1247- [[nodiscard]] __device__ bool contains (ProbeKey const & key) const noexcept
1242+ [[nodiscard]] __device__ bool contains (ProbeKey key) const noexcept
12481243 {
12491244 // CRTP: cast `this` to the actual ref type
12501245 auto const & ref_ = static_cast <ref_type const &>(*this );
@@ -1266,7 +1261,7 @@ class operator_impl<
12661261 */
12671262 template <typename ProbeKey>
12681263 [[nodiscard]] __device__ bool contains (
1269- cooperative_groups::thread_block_tile<cg_size> const & group, ProbeKey const & key) const noexcept
1264+ cooperative_groups::thread_block_tile<cg_size> const & group, ProbeKey key) const noexcept
12701265 {
12711266 auto const & ref_ = static_cast <ref_type const &>(*this );
12721267 return ref_.impl_ .contains (group, key);
@@ -1307,7 +1302,7 @@ class operator_impl<
13071302 * @return An iterator to the position at which the equivalent key is stored
13081303 */
13091304 template <typename ProbeKey>
1310- [[nodiscard]] __device__ iterator find (ProbeKey const & key) const noexcept
1305+ [[nodiscard]] __device__ iterator find (ProbeKey key) const noexcept
13111306 {
13121307 // CRTP: cast `this` to the actual ref type
13131308 auto const & ref_ = static_cast <ref_type const &>(*this );
@@ -1328,8 +1323,8 @@ class operator_impl<
13281323 * @return An iterator to the position at which the equivalent key is stored
13291324 */
13301325 template <typename ProbeKey>
1331- [[nodiscard]] __device__ iterator find (
1332- cooperative_groups::thread_block_tile<cg_size> const & group, ProbeKey const & key) const noexcept
1326+ [[nodiscard]] __device__ iterator
1327+ find ( cooperative_groups::thread_block_tile<cg_size> const & group, ProbeKey key) const noexcept
13331328 {
13341329 auto const & ref_ = static_cast <ref_type const &>(*this );
13351330 return ref_.impl_ .find (group, key);
@@ -1370,7 +1365,7 @@ class operator_impl<
13701365 * @param callback_op Function to apply to the copy of the matched key-value pair
13711366 */
13721367 template <class ProbeKey , class CallbackOp >
1373- __device__ void for_each (ProbeKey const & key, CallbackOp&& callback_op) const noexcept
1368+ __device__ void for_each (ProbeKey key, CallbackOp&& callback_op) const noexcept
13741369 {
13751370 // CRTP: cast `this` to the actual ref type
13761371 auto const & ref_ = static_cast <ref_type const &>(*this );
@@ -1397,7 +1392,7 @@ class operator_impl<
13971392 */
13981393 template <class ProbeKey , class CallbackOp >
13991394 __device__ void for_each (cooperative_groups::thread_block_tile<cg_size> const & group,
1400- ProbeKey const & key,
1395+ ProbeKey key,
14011396 CallbackOp&& callback_op) const noexcept
14021397 {
14031398 // CRTP: cast `this` to the actual ref type
@@ -1436,7 +1431,7 @@ class operator_impl<
14361431 * @return Number of occurrences found by the current thread
14371432 */
14381433 template <typename ProbeKey>
1439- __device__ size_type count (ProbeKey const & key) const noexcept
1434+ __device__ size_type count (ProbeKey key) const noexcept
14401435 {
14411436 auto const & ref_ = static_cast <ref_type const &>(*this );
14421437 return ref_.impl_ .count (key);
@@ -1454,7 +1449,7 @@ class operator_impl<
14541449 */
14551450 template <typename ProbeKey>
14561451 __device__ size_type count (cooperative_groups::thread_block_tile<cg_size> const & group,
1457- ProbeKey const & key) const noexcept
1452+ ProbeKey key) const noexcept
14581453 {
14591454 auto const & ref_ = static_cast <ref_type const &>(*this );
14601455 return ref_.impl_ .count (group, key);
@@ -1486,8 +1481,9 @@ class operator_impl<
14861481 * @brief Retrieves all the slots corresponding to all keys in the range `[input_probe_begin,
14871482 * input_probe_end)`.
14881483 *
1489- * If key `k = *(first + i)` exists in the container, copies `k` to `output_probe` and associated
1490- * slot content to `output_match`, respectively. The output order is unspecified.
1484+ * If key `k = *(first + i)` exists in the container, copies `k` to `
1485+ * output_probe` and associated slot content to `output_match`, respectively. The output order is
1486+ * unspecified.
14911487 *
14921488 * Behavior is undefined if the size of the output range exceeds the number of retrieved slots.
14931489 * Use `count()` to determine the size of the output range.
@@ -1529,4 +1525,4 @@ class operator_impl<
15291525 }
15301526};
15311527} // namespace detail
1532- } // namespace cuco
1528+ } // namespace cuco
0 commit comments