Update docs

Author: sleeepyjack

include/cuco/detail/open_addressing/kernels.cuh

@@ -399,7 +399,34 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void find(InputIt first,
   }
 }
 
-// TODO docs
+/**
+ * @brief Retrieves the equivalent container elements of all keys in the range `[input_probe,
+ * input_probe + n)`.
+ *
+ * If key `k = *(input_probe + i)` has one or more matches in the container, copies `k` to
+ * `output_probe` and associated slot contents to `output_match`, respectively. The output order is
+ * unspecified.
+ *
+ * @tparam IsOuter Flag indicating whether it's an outer count or not
+ * @tparam block_size The size of the thread block
+ * @tparam InputProbeIt Device accessible input iterator
+ * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
+ * convertible to the `InputProbeIt`'s `value_type`
+ * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
+ * convertible to the container's `value_type`
+ * @tparam AtomicCounter Integral atomic type that follows the same semantics as
+ * `cuda::(std::)atomic(_ref)`
+ * @tparam Ref Type of non-owning device ref allowing access to storage
+ *
+ * @param input_probe Beginning of the sequence of input keys
+ * @param n Number of the keys to query
+ * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
+ * `output_match`
+ * @param output_match Beginning of the sequence of matching elements
+ * @param atomic_counter Pointer to an atomic object of integral type that is used to count the
+ * number of output elements
+ * @param ref Non-owning container device ref used to access the slot storage
+ */
 template <bool IsOuter,
           int32_t BlockSize,
           class InputProbeIt,

include/cuco/detail/open_addressing/open_addressing_impl.cuh

@@ -585,10 +585,9 @@ class open_addressing_impl {
    *
    * This function synchronizes the given CUDA stream.
    *
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
    * @tparam Ref Type of non-owning device container ref allowing access to storage
@@ -630,10 +629,9 @@ class open_addressing_impl {
    *
    * This function synchronizes the given CUDA stream.
    *
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
    * @tparam Ref Type of non-owning device container ref allowing access to storage

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -989,23 +989,22 @@ class open_addressing_ref_impl {
    * Use `count()` to determine the size of the output range.
    *
    * @tparam BlockSize Size of the thread block this operation is executed in
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
-   * @tparam AtomicCounter Atomic counter type that follows the same semantics as
-   * `cuda::atomic(_ref)`
+   * @tparam AtomicCounter Integral atomic counter type that follows the same semantics as
+   * `cuda::(std::)atomic(_ref)`
    *
    * @param block Thread block this operation is executed in
    * @param input_probe_begin Beginning of the input sequence of keys
    * @param input_probe_end End of the input sequence of keys
    * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
    * `output_match`
    * @param output_match Beginning of the sequence of matching elements
-   * @param atomic_counter Counter that is used to determine the next free position in the output
-   * sequences
+   * @param atomic_counter Pointer to an atomic object of integral type that is used to count the
+   * number of output elements
    */
   template <int32_t BlockSize,
             class InputProbeIt,
@@ -1039,23 +1038,22 @@ class open_addressing_ref_impl {
    * to the output sequence.
    *
    * @tparam BlockSize Size of the thread block this operation is executed in
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
-   * @tparam AtomicCounter Atomic counter type that follows the same semantics as
-   * `cuda::atomic(_ref)`
+   * @tparam AtomicCounter Integral atomic counter type that follows the same semantics as
+   * `cuda::(std::)atomic(_ref)`
    *
    * @param block Thread block this operation is executed in
    * @param input_probe_begin Beginning of the input sequence of keys
    * @param input_probe_end End of the input sequence of keys
    * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
    * `output_match`
    * @param output_match Beginning of the sequence of matching elements
-   * @param atomic_counter Counter that is used to determine the next free position in the output
-   * sequences
+   * @param atomic_counter Pointer to an atomic object of integral type that is used to count the
+   * number of output elements
    */
   template <int32_t BlockSize,
             class InputProbeIt,
@@ -1090,23 +1088,22 @@ class open_addressing_ref_impl {
    *
    * @tparam IsOuter Flag indicating if an inner or outer retrieve operation should be performed
    * @tparam BlockSize Size of the thread block this operation is executed in
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
-   * @tparam AtomicCounter Atomic counter type that follows the same semantics as
-   * `cuda::atomic(_ref)`
+   * @tparam AtomicCounter Integral atomic type that follows the same semantics as
+   * `cuda::(std::)atomic(_ref)`
    *
    * @param block Thread block this operation is executed in
    * @param input_probe_begin Beginning of the input sequence of keys
    * @param input_probe_end End of the input sequence of keys
    * @param output_probe Beginning of the sequence of keys corresponding to matching elements in
    * `output_match`
    * @param output_match Beginning of the sequence of matching elements
-   * @param atomic_counter Counter that is used to determine the next free position in the output
-   * sequences
+   * @param atomic_counter Pointer to an atomic object of integral type that is used to count the
+   * number of output elements
    */
   template <bool IsOuter,
             int32_t BlockSize,

include/cuco/static_multiset.cuh

@@ -487,10 +487,9 @@ class static_multiset {
    *
    * This function synchronizes the given CUDA stream.
    *
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
    *
@@ -524,10 +523,9 @@ class static_multiset {
    *
    * This function synchronizes the given CUDA stream.
    *
-   * @tparam InputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * @tparam InputProbeIt Device accessible input iterator
    * @tparam OutputProbeIt Device accessible input iterator whose `value_type` is
-   * convertible to the container's `key_type`
+   * convertible to the `InputProbeIt`'s `value_type`
    * @tparam OutputMatchIt Device accessible input iterator whose `value_type` is
    * convertible to the container's `value_type`
    *