From fbbee7d49ae0ca2336b457afaa5d83cde5593358 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Daniel=20J=C3=BCnger?= <djuenger@nvidia.com>
Date: Thu, 17 Oct 2024 15:53:55 +0000
Subject: [PATCH] Update docs

---
 .../cuco/detail/open_addressing/kernels.cuh   | 29 +++++++++++++-
 .../open_addressing/open_addressing_impl.cuh  | 10 ++---
 .../open_addressing_ref_impl.cuh              | 39 +++++++++----------
 include/cuco/static_multiset.cuh              | 10 ++---
 4 files changed, 54 insertions(+), 34 deletions(-)

diff --git a/include/cuco/detail/open_addressing/kernels.cuh b/include/cuco/detail/open_addressing/kernels.cuh
index 5e0d49944..4726683a3 100644
--- a/include/cuco/detail/open_addressing/kernels.cuh
+++ b/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,
diff --git a/include/cuco/detail/open_addressing/open_addressing_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_impl.cuh
index e2f84054a..9b638ab76 100644
--- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh
+++ b/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
diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
index 9c17eb90c..fb0ca0720 100644
--- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
+++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
@@ -989,14 +989,13 @@ 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
@@ -1004,8 +1003,8 @@ class open_addressing_ref_impl {
    * @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,14 +1038,13 @@ 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
@@ -1054,8 +1052,8 @@ class open_addressing_ref_impl {
    * @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,14 +1088,13 @@ 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
@@ -1105,8 +1102,8 @@ class open_addressing_ref_impl {
    * @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,
diff --git a/include/cuco/static_multiset.cuh b/include/cuco/static_multiset.cuh
index b75c25b68..a2a65de8c 100644
--- a/include/cuco/static_multiset.cuh
+++ b/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`
    *