Update docs

include/cuco/detail/open_addressing/kernels.cuh

@@ -186,6 +186,8 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void erase(InputIt first,
  * @brief For each key in the range [first, first + n), applies the function object `callback_op` to
  * the copy of all corresponding matches found in the container.
  *
+ * @note The return value of `callback_op`, if any, is ignored.
+ *
  * @tparam CGSize Number of threads in each CG
  * @tparam BlockSize Number of threads in each block
  * @tparam InputIt Device accessible input iterator whose `value_type` is
@@ -195,7 +197,7 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void erase(InputIt first,
  *
  * @param first Beginning of the sequence of input elements
  * @param n Number of input elements
- * @param callback_op Function to call on every element found in the container
+ * @param callback_op Function to call on every matched slot found in the container
  * @param ref Non-owning container device ref used to access the slot storage
  */
 template <int32_t CGSize, int32_t BlockSize, typename InputIt, typename CallbackOp, typename Ref>

include/cuco/detail/open_addressing/open_addressing_impl.cuh

@@ -686,6 +686,8 @@ class open_addressing_impl {
    * @brief Asynchronously applies the given function object `callback_op` to the copy of every
    * filled slot in the container
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param callback_op Function to call on every filled slot in the container
@@ -712,6 +714,8 @@ class open_addressing_impl {
    * @brief For each key in the range [first, last), asynchronously applies the function object
    * `callback_op` to the copy of all corresponding matches found in the container.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam InputIt Device accessible random access input iterator
    * @tparam CallbackOp Type of unary callback function object
    * @tparam Ref Type of non-owning device container ref allowing access to storage

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -969,11 +969,13 @@ class open_addressing_ref_impl {
    * @brief For a given key, applies the function object `callback_op` to the copy of all
    * corresponding matches found in the container.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam ProbeKey Probe key type
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param key The key to search for
-   * @param callback_op Function to apply to every match
+   * @param callback_op Function to apply to every matched slot
    */
   template <class ProbeKey, class CallbackOp>
   __device__ void for_each(ProbeKey const& key, CallbackOp&& callback_op) const noexcept
@@ -1010,14 +1012,16 @@ class open_addressing_ref_impl {
    * callback if it finds a matching element. If multiple elements are found within the same group,
    * each thread with a match will call the callback with its associated element.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @note Synchronizing `group` within `callback_op` is undefined behavior.
    *
    * @tparam ProbeKey Probe key type
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param group The Cooperative Group used to perform this operation
    * @param key The key to search for
-   * @param callback_op Function to apply to every match
+   * @param callback_op Function to apply to every matched slot
    */
   template <class ProbeKey, class CallbackOp>
   __device__ void for_each(cooperative_groups::thread_block_tile<cg_size> const& group,
@@ -1064,6 +1068,8 @@ class open_addressing_ref_impl {
    *
    * @note Synchronizing `group` within `callback_op` is undefined behavior.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @note The `sync_op` function can be used to perform work that requires synchronizing threads in
    * `group` inbetween probing steps, where the number of probing steps performed between
    * synchronization points is capped by `window_size * cg_size`. The functor will be called right
@@ -1075,7 +1081,7 @@ class open_addressing_ref_impl {
    *
    * @param group The Cooperative Group used to perform this operation
    * @param key The key to search for
-   * @param callback_op Function to apply to every match
+   * @param callback_op Function to apply to every matched slot
    * @param sync_op Function that is allowed to synchronize `group` inbetween probing windows
    */
   template <class ProbeKey, class CallbackOp, class SyncOp>

include/cuco/detail/static_map/static_map_ref.inl

@@ -1274,11 +1274,13 @@ class operator_impl<
    * @brief For a given key, applies the function object `callback_op` to its match found in the
    * container.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam ProbeKey Probe key type
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param key The key to search for
-   * @param callback_op Function to apply to the match
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    */
   template <class ProbeKey, class CallbackOp>
   __device__ void for_each(ProbeKey const& key, CallbackOp&& callback_op) const noexcept
@@ -1293,8 +1295,9 @@ class operator_impl<
    * container.
    *
    * @note This function uses cooperative group semantics, meaning that any thread may call the
-   * callback if it finds a matching element. If multiple elements are found within the same group,
-   * each thread with a match will call the callback with its associated element.
+   * callback if it finds a matching key-value pair.
+   *
+   * @note The return value of `callback_op`, if any, is ignored.
    *
    * @note Synchronizing `group` within `callback_op` is undefined behavior.
    *
@@ -1303,7 +1306,7 @@ class operator_impl<
    *
    * @param group The Cooperative Group used to perform this operation
    * @param key The key to search for
-   * @param callback_op Function to apply to the match
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    */
   template <class ProbeKey, class CallbackOp>
   __device__ void for_each(cooperative_groups::thread_block_tile<cg_size> const& group,

include/cuco/static_map.cuh

@@ -766,9 +766,11 @@ class static_map {
    * @brief Applies the given function object `callback_op` to the copy of every filled slot in the
    * container
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam CallbackOp Type of unary callback function object
    *
-   * @param callback_op Function to call on every filled slot in the container
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    * @param stream CUDA stream used for this operation
    */
   template <typename CallbackOp>
@@ -778,9 +780,11 @@ class static_map {
    * @brief Asynchronously applies the given function object `callback_op` to the copy of every
    * filled slot in the container
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam CallbackOp Type of unary callback function object
    *
-   * @param callback_op Function to call on every filled slot in the container
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    * @param stream CUDA stream used for this operation
    */
   template <typename CallbackOp>
@@ -790,12 +794,14 @@ class static_map {
    * @brief For each key in the range [first, last), applies the function object `callback_op` to
    * the copy of all corresponding matches found in the container.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam InputIt Device accessible random access input iterator
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param first Beginning of the sequence of keys
    * @param last End of the sequence of keys
-   * @param callback_op Function to call on every match found in the container
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    * @param stream CUDA stream used for this operation
    */
   template <typename InputIt, typename CallbackOp>
@@ -808,12 +814,14 @@ class static_map {
    * @brief For each key in the range [first, last), asynchronously applies the function object
    * `callback_op` to the copy of all corresponding matches found in the container.
    *
+   * @note The return value of `callback_op`, if any, is ignored.
+   *
    * @tparam InputIt Device accessible random access input iterator
    * @tparam CallbackOp Type of unary callback function object
    *
    * @param first Beginning of the sequence of keys
    * @param last End of the sequence of keys
-   * @param callback_op Function to call on every match found in the container
+   * @param callback_op Function to apply to the copy of the matched key-value pair
    * @param stream CUDA stream used for this operation
    */
   template <typename InputIt, typename CallbackOp>