From 202c2f29e76ee37b96ad0cc817738b6ce35933bd Mon Sep 17 00:00:00 2001
From: Karthikeyan Natarajan <karthikeyann@users.noreply.github.com>
Date: Wed, 25 May 2022 17:34:23 +0530
Subject: [PATCH 1/7] add missing docs in table headers

---
 .../cudf/table/experimental/row_operators.cuh | 127 ++++++++++++++----
 cpp/include/cudf/table/row_operators.cuh      | 116 ++++++++++++++--
 cpp/include/cudf/table/table.hpp              |  16 ++-
 cpp/include/cudf/table/table_device_view.cuh  | 120 ++++++++++++++++-
 cpp/include/cudf/table/table_view.hpp         |  65 +++++++--
 5 files changed, 398 insertions(+), 46 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index f01bc41a2bc..ee5fbdc9b09 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -67,6 +67,7 @@ namespace experimental {
  */
 template <cudf::type_id t>
 struct dispatch_void_if_nested {
+  /// The type to dispatch to if the type is nested
   using type = std::conditional_t<cudf::is_nested(data_type(t)), void, id_to_type<t>>;
 };
 
@@ -95,11 +96,17 @@ struct strong_index_iterator : public thrust::iterator_facade<strong_index_itera
                                                               thrust::random_access_traversal_tag,
                                                               Index,
                                                               Underlying> {
-  using super_t = thrust::iterator_adaptor<strong_index_iterator<Index>, Index>;
+  using super_t =
+    thrust::iterator_adaptor<strong_index_iterator<Index>, Index>;  ///< The base class
 
+  /**
+   * @brief Constructs a strong index iterator
+   *
+   * @param n The beginning index
+   */
   explicit constexpr strong_index_iterator(Underlying n) : begin{n} {}
 
-  friend class thrust::iterator_core_access;
+  friend class thrust::iterator_core_access;  ///< Allow access to the base class
 
  private:
   __device__ constexpr void increment() { ++begin; }
@@ -154,8 +161,8 @@ namespace lexicographic {
  */
 template <typename Nullate>
 class device_row_comparator {
-  friend class self_comparator;
-  friend class two_table_comparator;
+  friend class self_comparator;       ///< Allow self_comparator to access private members
+  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
 
   /**
    * @brief Construct a function object for performing a lexicographic
@@ -378,7 +385,12 @@ template <typename Comparator>
 using less_equivalent_comparator =
   weak_ordering_comparator_impl<Comparator, weak_ordering::LESS, weak_ordering::EQUIVALENT>;
 
+/**
+ * @brief Preprocessed table for use with lexicographical comparison
+ *
+ */
 struct preprocessed_table {
+  /// Type of table device view owner for the preprocessed table.
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
 
@@ -396,6 +408,7 @@ struct preprocessed_table {
    * values compare to all other for every column. If it is nullptr, then null precedence would be
    * `null_order::BEFORE` for all columns.
    * @param stream The stream to launch kernels and h->d copies on while preprocessing.
+   * @return A preprocessed table as a shared pointer
    */
   static std::shared_ptr<preprocessed_table> create(table_view const& table,
                                                     host_span<order const> column_order,
@@ -403,8 +416,8 @@ struct preprocessed_table {
                                                     rmm::cuda_stream_view stream);
 
  private:
-  friend class self_comparator;
-  friend class two_table_comparator;
+  friend class self_comparator;       ///< Allow self_comparator to access private members
+  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
 
   preprocessed_table(table_device_view_owner&& table,
                      rmm::device_uvector<order>&& column_order,
@@ -525,6 +538,8 @@ class self_comparator {
    * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`.
    *
    * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+   * @param nullate Indicates if either input column contains nulls
+   * @return A binary callable object
    */
   template <typename Nullate>
   less_comparator<device_row_comparator<Nullate>> device_comparator(Nullate nullate = {}) const
@@ -537,6 +552,7 @@ class self_comparator {
   std::shared_ptr<preprocessed_table> d_t;
 };
 
+// @cond
 template <typename Comparator>
 struct strong_index_comparator_adapter {
   __device__ constexpr weak_ordering operator()(lhs_index_type const lhs_index,
@@ -563,6 +579,7 @@ struct strong_index_comparator_adapter {
 
   Comparator const comparator;
 };
+// @endcond
 
 /**
  * @brief An owning object that can be used to lexicographically compare rows of two different
@@ -636,6 +653,8 @@ class two_table_comparator {
    * `j` of the left table.
    *
    * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+   * @param nullate Indicates if either input column contains nulls.
+   * @return A binary callable object
    */
   template <typename Nullate>
   less_comparator<strong_index_comparator_adapter<device_row_comparator<Nullate>>>
@@ -662,11 +681,15 @@ class row_hasher;
 }
 
 namespace equality {
-
+/**
+ * @brief Comparator for performing equality comparison between the rows of two tables.
+ *
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+ */
 template <typename Nullate>
 class device_row_comparator {
-  friend class self_comparator;
-  friend class two_table_comparator;
+  friend class self_comparator;       ///< Allow self_comparator to access private members
+  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
 
  public:
   /**
@@ -860,24 +883,30 @@ class device_row_comparator {
   null_equality const nulls_are_equal;
 };
 
+/**
+ * @brief Preprocess table for use with row equality comparison or row hashing
+ *
+ * Sets up the table for use with row equality comparison or row hashing. The resulting
+ * preprocessed table can be passed to the constructor of `equality::self_comparator` to
+ * avoid preprocessing again.
+ *
+ */
 struct preprocessed_table {
   /**
-   * @brief Preprocess table for use with row equality comparison or row hashing
-   *
-   * Sets up the table for use with row equality comparison or row hashing. The resulting
-   * preprocessed table can be passed to the constructor of `equality::self_comparator` to
-   * avoid preprocessing again.
+   * @brief Factory to construct preprocessed_table for use with
+   * row equality comparison or row hashing
    *
    * @param table The table to preprocess
    * @param stream The cuda stream to use while preprocessing.
+   * @return A preprocessed table as shared pointer
    */
   static std::shared_ptr<preprocessed_table> create(table_view const& table,
                                                     rmm::cuda_stream_view stream);
 
  private:
-  friend class self_comparator;
-  friend class two_table_comparator;
-  friend class hash::row_hasher;
+  friend class self_comparator;       ///< Allow self_comparator to access private members
+  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
+  friend class hash::row_hasher;      ///< Allow row_hasher to access private members
 
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
@@ -899,6 +928,10 @@ struct preprocessed_table {
   std::vector<rmm::device_buffer> _null_buffers;
 };
 
+/**
+ * @brief Comparator for performing equality comparisons between two rows of the same table.
+ *
+ */
 class self_comparator {
  public:
   /**
@@ -932,11 +965,14 @@ class self_comparator {
    *
    * `F(i,j)` returns true if and only if row `i` compares equal to row `j`.
    *
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+   * @param nullate Indicates if either input column contains nulls
+   * @param nulls_are_equal Indicates if nulls are equal
+   * @return A binary callable object
    */
   template <typename Nullate>
   device_row_comparator<Nullate> device_comparator(
-    Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const
+    Nullate nullate = {}, null_equality nulls_are_equal = n`ull_equality::EQUAL) const
   {
     return device_row_comparator(nullate, *d_t, *d_t, nulls_are_equal);
   }
@@ -945,6 +981,7 @@ class self_comparator {
   std::shared_ptr<preprocessed_table> d_t;
 };
 
+// @cond
 template <typename Comparator>
 struct strong_index_comparator_adapter {
   __device__ constexpr bool operator()(lhs_index_type const lhs_index,
@@ -962,6 +999,7 @@ struct strong_index_comparator_adapter {
 
   Comparator const comparator;
 };
+// @endcond
 
 /**
  * @brief An owning object that can be used to equality compare rows of two different tables.
@@ -1022,7 +1060,10 @@ class two_table_comparator {
    * Similarly, `F(rhs_index_type i, lhs_index_type j)` returns true if and only if row `i` of the
    * right table compares equal to row `j` of the left table.
    *
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+   * @param nullate Indicates if either input column contains nulls
+   * @param nulls_are_equal Indicates if nulls are equal
+   * @return A binary callable object
    */
   template <typename Nullate>
   auto device_comparator(Nullate nullate               = {},
@@ -1050,6 +1091,13 @@ namespace hash {
 template <template <typename> class hash_function, typename Nullate>
 class element_hasher {
  public:
+  /**
+   * @brief Constructs an element_hasher object.
+   *
+   * @param nulls Indicates whether to check for nulls
+   * @param seed  The seed to use for the hash function
+   * @param null_hash The hash value to use for nulls
+   */
   __device__ element_hasher(
     Nullate nulls,
     uint32_t seed             = DEFAULT_HASH_SEED,
@@ -1058,6 +1106,14 @@ class element_hasher {
   {
   }
 
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view const& col,
                                         size_type row_index) const noexcept
@@ -1066,6 +1122,14 @@ class element_hasher {
     return hash_function<T>{_seed}(col.element<T>(row_index));
   }
 
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view const& col,
                                         size_type row_index) const noexcept
@@ -1073,9 +1137,9 @@ class element_hasher {
     CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
-  uint32_t _seed;
-  hash_value_type _null_hash;
-  Nullate _check_nulls;
+  uint32_t _seed;              ///< The seed to use for hashing
+  hash_value_type _null_hash;  ///< Hash value to use for null elements
+  Nullate _check_nulls;        ///< Whether to check for nulls
 };
 
 /**
@@ -1086,11 +1150,17 @@ class element_hasher {
  */
 template <template <typename> class hash_function, typename Nullate>
 class device_row_hasher {
-  friend class row_hasher;
+  friend class row_hasher;  ///< Allow row_hasher to access private members.
 
  public:
   device_row_hasher() = delete;
 
+  /**
+   * @brief Return the hash value of a row in the given table.
+   *
+   * @param row_index The row index to compute the hash value of
+   * @return The hash value of the row
+   */
   __device__ auto operator()(size_type row_index) const noexcept
   {
     auto it = thrust::make_transform_iterator(_table.begin(), [=](auto const& column) {
@@ -1187,6 +1257,10 @@ class device_row_hasher {
 // type and are interchangeable.
 using preprocessed_table = row::equality::preprocessed_table;
 
+/**
+ * @brief Computes the hash value of a row in the given table.
+ *
+ */
 class row_hasher {
  public:
   /**
@@ -1219,7 +1293,10 @@ class row_hasher {
    *
    * `F(i)` returns the hash of row i.
    *
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+   * @param nullate Indicates if either input column contains nulls
+   * @param seed The seed to use for the hash function
+   * @return A hash operator to use on the device
    */
   template <template <typename> class hash_function = detail::default_hash, typename Nullate>
   device_row_hasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index a181e9bae63..d8a662a72be 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/cpp/include/cudf/table/row_operators.cuh
@@ -78,10 +78,6 @@ __device__ weak_ordering compare_elements(Element lhs, Element rhs)
  * `[-Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN, null] (for null_order::AFTER)`
  * `[null, -Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN] (for null_order::BEFORE)`
  *
- * @param lhs first element
- * @param rhs second element
- * @return Indicates the relationship between the elements in
- * the `lhs` and `rhs` columns.
  */
 template <typename Element, std::enable_if_t<std::is_floating_point_v<Element>>* = nullptr>
 __device__ weak_ordering relational_compare(Element lhs, Element rhs)
@@ -121,10 +117,9 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord
  * @brief A specialization for non-floating-point `Element` type relational
  * comparison to derive the order of the elements with respect to `lhs`.
  *
- * @param[in] lhs first element
- * @param[in] rhs second element
- * @return Indicates the relationship between the elements in
- * the `lhs` and `rhs` columns.
+ * @param lhs first element
+ * @param rhs second element
+ * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns
  */
 template <typename Element, std::enable_if_t<not std::is_floating_point_v<Element>>* = nullptr>
 __device__ weak_ordering relational_compare(Element lhs, Element rhs)
@@ -215,12 +210,14 @@ class element_equality_comparator {
                             rhs.element<Element>(rhs_element_index));
   }
 
+  // @cond
   template <typename Element,
             std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* = nullptr>
   __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index)
   {
     CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
   }
+  // @endcond
 
  private:
   column_device_view lhs;
@@ -229,9 +226,22 @@ class element_equality_comparator {
   null_equality nulls_are_equal;
 };
 
+/**
+ * @brief Performs a relational comparison between two elements in two columns.
+ *
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls
+ */
 template <typename Nullate>
 class row_equality_comparator {
  public:
+  /**
+   * @brief Construct a new row equality comparator object
+   *
+   * @param has_nulls Indicates if either input column contains nulls
+   * @param lhs The column containing the first element
+   * @param rhs The column containing the second element (may be the same as lhs)
+   * @param nulls_are_equal Indicates if two null elements are treated as equivalent
+   */
   row_equality_comparator(Nullate has_nulls,
                           table_device_view lhs,
                           table_device_view rhs,
@@ -241,6 +251,13 @@ class row_equality_comparator {
     CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns.");
   }
 
+  /**
+   * @brief Compares the specified rows for equality.
+   *
+   * @param lhs_row_index The index of the first row to compare (in the lhs table)
+   * @param rhs_row_index The index of the second row to compare (in the rhs table)
+   * @return true if both rows are equal
+   */
   __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept
   {
     auto equal_elements = [=](column_device_view l, column_device_view r) {
@@ -287,6 +304,14 @@ class element_relational_comparator {
   {
   }
 
+  /**
+   * @brief Construct type-dispatched function object for performing a relational comparison between
+   * two elements in two columns.
+   *
+   * @param has_nulls Indicates if either input column contains nulls
+   * @param lhs The column containing the first element
+   * @param rhs The column containing the second element (may be the same as lhs)
+   */
   __host__ __device__ element_relational_comparator(Nullate has_nulls,
                                                     column_device_view lhs,
                                                     column_device_view rhs)
@@ -320,12 +345,14 @@ class element_relational_comparator {
                               rhs.element<Element>(rhs_element_index));
   }
 
+  // @cond
   template <typename Element,
             std::enable_if_t<not cudf::is_relationally_comparable<Element, Element>()>* = nullptr>
   __device__ weak_ordering operator()(size_type lhs_element_index, size_type rhs_element_index)
   {
     CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
   }
+  // @endcond
 
  private:
   column_device_view lhs;
@@ -432,6 +459,14 @@ class row_lexicographic_comparator {
 template <template <typename> class hash_function, typename Nullate>
 class element_hasher {
  public:
+  /**
+   * @brief Returns the hash value of the given element in the given column.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
@@ -439,28 +474,63 @@ class element_hasher {
     return hash_function<T>{}(col.element<T>(row_index));
   }
 
+  /**
+   * @brief Returns the hash value of the given element in the given column.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
     CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
-  Nullate has_nulls;
+  Nullate has_nulls;  ///< A cudf::nullate type describing how to check for nulls.
 };
 
+/**
+ * @brief Function object for computing the hash value of a row in a column.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls
+ */
 template <template <typename> class hash_function, typename Nullate>
 class element_hasher_with_seed {
  public:
+  /**
+   * @brief Constructs a function object for hashing an element in the given column
+   *
+   * @param has_nulls Indicates if either input column contains nulls
+   * @param seed The seed to use for the hash function
+   */
   __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed)
     : _seed{seed}, _has_nulls{has_nulls}
   {
   }
 
+  /**
+   * @brief Constructs a function object for hashing an element in the given column
+   *
+   * @param has_nulls Indicates if either input column contains nulls
+   * @param seed The seed to use for the hash function
+   * @param null_hash The hash value to use for null elements
+   */
   __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed, hash_value_type null_hash)
     : _seed{seed}, _null_hash{null_hash}, _has_nulls{has_nulls}
   {
   }
 
+  /**
+   * @brief Returns the hash value of the given element in the given column.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
@@ -468,6 +538,14 @@ class element_hasher_with_seed {
     return hash_function<T>{_seed}(col.element<T>(row_index));
   }
 
+  /**
+   * @brief Returns the hash value of the given element in the given column.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
@@ -490,15 +568,35 @@ template <template <typename> class hash_function, typename Nullate>
 class row_hasher {
  public:
   row_hasher() = delete;
+
+  /**
+   * @brief Constructs a row_hasher object.
+   *
+   * @param has_nulls Indicates if either input table contains nulls
+   * @param t A table_device_view to hash
+   */
   CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t)
     : _table{t}, _has_nulls{has_nulls}
   {
   }
+  /**
+   * @brief Constructs a row_hasher object with a seed value.
+   *
+   * @param has_nulls Indicates if either input table contains nulls
+   * @param t A table_device_view to hash
+   * @param seed A seed value to use for hashing
+   */
   CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t, uint32_t seed)
     : _table{t}, _seed(seed), _has_nulls{has_nulls}
   {
   }
 
+  /**
+   * @brief Computes the hash value of the row at `row_index` in the `table`
+   *
+   * @param row_index The index of the row in the `table` to hash
+   * @return The hash value of the row at `row_index` in the `table`
+   */
   __device__ auto operator()(size_type row_index) const
   {
     // Hash the first column w/ the seed
diff --git a/cpp/include/cudf/table/table.hpp b/cpp/include/cudf/table/table.hpp
index 4a3c31d08e9..d2c3226b213 100644
--- a/cpp/include/cudf/table/table.hpp
+++ b/cpp/include/cudf/table/table.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -39,12 +39,14 @@ class table {
  public:
   table()        = default;
   ~table()       = default;
-  table(table&&) = default;
+  table(table&&) = default;  ///< Move constructor
   table& operator=(table const&) = delete;
   table& operator=(table&&) = delete;
 
   /**
    * @brief Construct a new table by copying the contents of another table.
+   *
+   * @param other The table to copy
    */
   table(table const& other);
 
@@ -70,17 +72,23 @@ class table {
 
   /**
    * @brief Returns the number of columns in the table
+   *
+   * @return The number of columns in the table
    */
   [[nodiscard]] size_type num_columns() const noexcept { return _columns.size(); }
 
   /**
    * @brief Returns the number of rows
+   *
+   * @return  The number of rows
    */
   [[nodiscard]] size_type num_rows() const noexcept { return _num_rows; }
 
   /**
    * @brief Returns an immutable, non-owning `table_view` of the contents of
    *this `table`.
+   *
+   * @return An immutable, non-owning `table_view` of the contents of this `table`
    */
   [[nodiscard]] table_view view() const;
 
@@ -93,6 +101,8 @@ class table {
   /**
    * @brief Returns a mutable, non-owning `mutable_table_view` of the contents
    * of this `table`.
+   *
+   * @return A mutable, non-owning `mutable_table_view` of the contents of this `table`
    */
   mutable_table_view mutable_view();
 
@@ -107,6 +117,8 @@ class table {
    * `unique_ptr`s to the constituent columns.
    *
    * After `release()`, `num_columns() == 0` and `num_rows() == 0`
+   *
+   * @returns A vector of `unique_ptr`s to the constituent columns
    */
   std::vector<std::unique_ptr<column>> release();
 
diff --git a/cpp/include/cudf/table/table_device_view.cuh b/cpp/include/cudf/table/table_device_view.cuh
index 8d08a3fd0b0..52701cb5899 100644
--- a/cpp/include/cudf/table/table_device_view.cuh
+++ b/cpp/include/cudf/table/table_device_view.cuh
@@ -33,20 +33,63 @@
 
 namespace cudf {
 namespace detail {
+
+/**
+ * @brief Base class for a device table of `ColumnDeviceView`s
+ *
+ * This class should not be used directly. Instead:
+ * -`table_device_view` should be used for a table of `column_device_view`s
+ * -`mutable_table_device_view` should be used for a table of `mutable_column__device_view`s
+ *
+ * All public constructors and member functions of `table_device_view_base` are
+ * available in both `table_device_view` and `mutable_table_device_view`.
+ *
+ * @tparam ColumnDeviceView The type of column device view the table contains
+ * @tparam HostTableView The type of table view used to create the table device view
+ */
 template <typename ColumnDeviceView, typename HostTableView>
 class table_device_view_base {
  public:
   table_device_view_base()                              = delete;
   ~table_device_view_base()                             = default;
-  table_device_view_base(table_device_view_base const&) = default;
-  table_device_view_base(table_device_view_base&&)      = default;
+  table_device_view_base(table_device_view_base const&) = default;  ///< Copy constructor
+  table_device_view_base(table_device_view_base&&)      = default;  ///< Move constructor
+  /**
+   * @brief Copy assignment operator
+   *
+   * @return Reference to this object
+   */
   table_device_view_base& operator=(table_device_view_base const&) = default;
+  /**
+   * @brief Move assignment operator
+   *
+   * @return Reference to this object (after transferring ownership)
+   */
   table_device_view_base& operator=(table_device_view_base&&) = default;
 
+  /**
+   * @brief Returns an iterator to the first view in the `table`.
+   *
+   * @return An iterator to the first view in the `table`
+   */
   __device__ ColumnDeviceView* begin() const noexcept { return _columns; }
 
+  /**
+   * @brief Returns an iterator one past the last column view in the `table`.
+   *
+   * `end()` acts as a place holder. Attempting to dereference it results in
+   * undefined behavior.
+   *
+   * @return An iterator to one past the last column view in the `table`
+   */
   __device__ ColumnDeviceView* end() const noexcept { return _columns + _num_columns; }
 
+  /**
+   * @brief Returns a reference to the view of the specified column
+   *
+   * @param column_index The index of the desired column
+   * @return A reference to the desired column
+   */
   __device__ ColumnDeviceView const& column(size_type column_index) const noexcept
   {
     assert(column_index >= 0);
@@ -54,6 +97,12 @@ class table_device_view_base {
     return _columns[column_index];
   }
 
+  /**
+   * @brief Returns a reference to the view of the specified column
+   *
+   * @param column_index The index of the desired column
+   * @return A reference to the desired column
+   */
   __device__ ColumnDeviceView& column(size_type column_index) noexcept
   {
     assert(column_index >= 0);
@@ -61,10 +110,26 @@ class table_device_view_base {
     return _columns[column_index];
   }
 
+  /**
+   * @brief Returns the number of columns
+   *
+   * @return The number of columns
+   */
   [[nodiscard]] __host__ __device__ size_type num_columns() const noexcept { return _num_columns; }
 
+  /**
+   * @brief Returns the number of rows
+   *
+   * @return The number of rows
+   */
   [[nodiscard]] __host__ __device__ size_type num_rows() const noexcept { return _num_rows; }
 
+  /**
+   * @brief Destroy the `table_device_view` object.
+   *
+   * @note Does not free the table data, simply frees the device memory
+   * allocated to hold the constituent column views.
+   */
   void destroy();
 
  private:
@@ -73,14 +138,38 @@ class table_device_view_base {
   size_type _num_columns{};
 
  protected:
+  /**
+   * @brief Construct a new table device view base object from host table_view
+   *
+   * @param source_view The host table_view to create table device view from
+   * @param stream The CUDA stream to use for device memory allocation
+   */
   table_device_view_base(HostTableView source_view, rmm::cuda_stream_view stream);
 
+  /// Pointer to device memory holding the descendant storage
   rmm::device_buffer* _descendant_storage{};
 };
 }  // namespace detail
 
+/**
+ * @brief Table device view that is usable in device memory
+ */
 class table_device_view : public detail::table_device_view_base<column_device_view, table_view> {
  public:
+  /**
+   * @brief Factory to construct a table device view that is usable in device memory.
+   *
+   * Allocates and copies views of `source_view`'s children to device memory to
+   * make them accessible in device code.
+   *
+   * Returns a `std::unique_ptr<table_device_view>` with a custom deleter to
+   * free the device memory allocated for the children.
+   *
+   * @param source_view The table view whose contents will be copied to create a new table
+   * @param stream CUDA stream used for device memory operations
+   * @return A `unique_ptr` to a `table_device_view` that makes the data from `source_view`
+   * available in device memory
+   */
   static auto create(table_view source_view,
                      rmm::cuda_stream_view stream = rmm::cuda_stream_default)
   {
@@ -96,9 +185,27 @@ class table_device_view : public detail::table_device_view_base<column_device_vi
   }
 };
 
+/**
+ * @brief Mutable table device view that is usable in device memory
+ * Elements of the table can be modified in device memory.
+ */
 class mutable_table_device_view
   : public detail::table_device_view_base<mutable_column_device_view, mutable_table_view> {
  public:
+  /**
+   * @brief Factory to construct a mutable table device view that is usable in device memory.
+   *
+   * Allocates and copies views of `source_view`'s children to device memory to
+   * make them accessible in device code.
+   *
+   * Returns a `std::unique_ptr<mutable_table_device_view>` with a custom deleter to
+   * free the device memory allocated for the children.
+   *
+   * @param source_view The table view whose contents will be copied to create a new table
+   * @param stream CUDA stream used for device memory operations
+   * @return A `unique_ptr` to a `mutable_table_device_view` that makes the data from `source_view`
+   * available in device memory
+   */
   static auto create(mutable_table_view source_view,
                      rmm::cuda_stream_view stream = rmm::cuda_stream_default)
   {
@@ -115,6 +222,15 @@ class mutable_table_device_view
   }
 };
 
+/**
+ * @brief Copies the contents of a table_view to a column device view in contiguous device memory
+ *
+ * @tparam ColumnDeviceView The column device view type to copy to
+ * @tparam HostTableView The type of the table_view to copy from
+ * @param source_view The table_view to copy from
+ * @param stream The stream to use for device memory allocation
+ * @return tuple of device_buffer and @p ColumnDeviceView device pointer
+ */
 template <typename ColumnDeviceView, typename HostTableView>
 auto contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream)
 {
diff --git a/cpp/include/cudf/table/table_view.hpp b/cpp/include/cudf/table/table_view.hpp
index 77b9e539506..ef8b9803d3e 100644
--- a/cpp/include/cudf/table/table_view.hpp
+++ b/cpp/include/cudf/table/table_view.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -58,8 +58,8 @@ class table_view_base {
   size_type _num_rows{};               ///< The number of elements in every column
 
  public:
-  using iterator       = decltype(std::begin(_columns));
-  using const_iterator = decltype(std::cbegin(_columns));
+  using iterator       = decltype(std::begin(_columns));   ///< Iterator type for the table
+  using const_iterator = decltype(std::cbegin(_columns));  ///< const iterator type for the table
 
   /**
    * @brief Construct a table from a vector of column views
@@ -81,11 +81,15 @@ class table_view_base {
 
   /**
    * @brief Returns an iterator to the first view in the `table`.
+   *
+   * @return An iterator to the first view in the `table`
    */
   iterator begin() noexcept { return std::begin(_columns); }
 
   /**
    * @brief Returns an iterator to the first view in the `table`.
+   *
+   * @return An iterator to the first view in the `table`
    */
   [[nodiscard]] const_iterator begin() const noexcept { return std::begin(_columns); }
 
@@ -94,6 +98,8 @@ class table_view_base {
    *
    * `end()` acts as a place holder. Attempting to dereference it results in
    * undefined behavior.
+   *
+   * @return An iterator to one past the last column view in the `table`
    */
   iterator end() noexcept { return std::end(_columns); }
 
@@ -102,6 +108,8 @@ class table_view_base {
    *
    * `end()` acts as a place holder. Attempting to dereference it results in
    * undefined behavior.
+   *
+   * @return An iterator to one past the last column view in the `table`
    */
   [[nodiscard]] const_iterator end() const noexcept { return std::end(_columns); }
 
@@ -118,16 +126,22 @@ class table_view_base {
 
   /**
    * @brief Returns the number of columns
+   *
+   * @return The number of columns
    */
   [[nodiscard]] size_type num_columns() const noexcept { return _columns.size(); }
 
   /**
    * @brief Returns the number of rows
+   *
+   * @return The number of rows
    */
   [[nodiscard]] size_type num_rows() const noexcept { return _num_rows; }
 
   /**
    * @brief Returns true if `num_columns()` returns zero, or false otherwise
+   *
+   * @return True if `num_columns()` returns zero, or false otherwise
    */
   [[nodiscard]] size_type is_empty() const noexcept { return num_columns() == 0; }
 
@@ -135,10 +149,20 @@ class table_view_base {
 
   ~table_view_base() = default;
 
-  table_view_base(table_view_base const&) = default;
+  table_view_base(table_view_base const&) = default;  ///< Copy constructor
 
-  table_view_base(table_view_base&&) = default;
+  table_view_base(table_view_base&&) = default;  ///< Move constructor
+  /**
+   * @brief Copy assignment operator
+   *
+   * @return Reference to this object
+   */
   table_view_base& operator=(table_view_base const&) = default;
+  /**
+   * @brief Move assignment operator
+   *
+   * @return Reference to this object (after transferring ownership)
+   */
   table_view_base& operator=(table_view_base&&) = default;
 };
 }  // namespace detail
@@ -155,7 +179,7 @@ class table_view : public detail::table_view_base<column_view> {
   using detail::table_view_base<column_view>::table_view_base;
 
  public:
-  using ColumnView = column_view;
+  using ColumnView = column_view;  ///< The type of column view the table contains
 
   table_view() = default;
 
@@ -223,10 +247,16 @@ class mutable_table_view : public detail::table_view_base<mutable_column_view> {
   using detail::table_view_base<mutable_column_view>::table_view_base;
 
  public:
-  using ColumnView = mutable_column_view;
+  using ColumnView = mutable_column_view;  ///< The type of column views in the table
 
   mutable_table_view() = default;
 
+  /**
+   * @brief Returns column at specified index
+   *
+   * @param column_index The index of the desired column
+   * @return A mutable column view reference to the desired column
+   */
   [[nodiscard]] mutable_column_view& column(size_type column_index) const
   {
     return const_cast<mutable_column_view&>(table_view_base::column(column_index));
@@ -257,16 +287,34 @@ class mutable_table_view : public detail::table_view_base<mutable_column_view> {
   mutable_table_view(std::vector<mutable_table_view> const& views);
 };
 
+/**
+ * @brief Returns True if any of the columns in the table is nullable. (not entire hierarchy)
+ *
+ * @param view The table to check for nullability
+ * @return True if any of the columns in the table is nullable, false otherwise
+ */
 inline bool nullable(table_view const& view)
 {
   return std::any_of(view.begin(), view.end(), [](auto const& col) { return col.nullable(); });
 }
 
+/**
+ * @brief Returns True if the table has nulls in any of its columns. (not entire hierarchy)
+ *
+ * @param view The table to check for nulls
+ * @return True if the table has nulls in any of its columns, false otherwise
+ */
 inline bool has_nulls(table_view const& view)
 {
   return std::any_of(view.begin(), view.end(), [](auto const& col) { return col.has_nulls(); });
 }
 
+/**
+ * @brief Returns True if the table has nulls in any of its columns hierarchy
+ *
+ * @param input The table to check for nulls
+ * @return True if the table has nulls in any of its columns hierarchy, false otherwise
+ */
 inline bool has_nested_nulls(table_view const& input)
 {
   return std::any_of(input.begin(), input.end(), [](auto const& col) {
@@ -322,10 +370,11 @@ namespace detail {
  */
 template <typename TableView>
 bool is_relationally_comparable(TableView const& lhs, TableView const& rhs);
-
+// @cond
 extern template bool is_relationally_comparable<table_view>(table_view const& lhs,
                                                             table_view const& rhs);
 extern template bool is_relationally_comparable<mutable_table_view>(mutable_table_view const& lhs,
                                                                     mutable_table_view const& rhs);
+// @endcond
 }  // namespace detail
 }  // namespace cudf

From e8abed04e3cd203705b4b72b8e7afbb4336faa9d Mon Sep 17 00:00:00 2001
From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com>
Date: Sat, 28 May 2022 05:13:04 +0530
Subject: [PATCH 2/7] Apply suggestions from code review

Co-authored-by: Vyas Ramasubramani <vyas.ramasubramani@gmail.com>
Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/table/experimental/row_operators.cuh | 4 ++--
 cpp/include/cudf/table/row_operators.cuh              | 6 +++---
 2 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index ee5fbdc9b09..0275b66ddd3 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -408,7 +408,7 @@ struct preprocessed_table {
    * values compare to all other for every column. If it is nullptr, then null precedence would be
    * `null_order::BEFORE` for all columns.
    * @param stream The stream to launch kernels and h->d copies on while preprocessing.
-   * @return A preprocessed table as a shared pointer
+   * @return A shared pointer to a preprocessed table
    */
   static std::shared_ptr<preprocessed_table> create(table_view const& table,
                                                     host_span<order const> column_order,
@@ -972,7 +972,7 @@ class self_comparator {
    */
   template <typename Nullate>
   device_row_comparator<Nullate> device_comparator(
-    Nullate nullate = {}, null_equality nulls_are_equal = n`ull_equality::EQUAL) const
+    Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const
   {
     return device_row_comparator(nullate, *d_t, *d_t, nulls_are_equal);
   }
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index d8a662a72be..cdf7f0a7ef6 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/cpp/include/cudf/table/row_operators.cuh
@@ -117,8 +117,8 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord
  * @brief A specialization for non-floating-point `Element` type relational
  * comparison to derive the order of the elements with respect to `lhs`.
  *
- * @param lhs first element
- * @param rhs second element
+ * @param lhs The first element
+ * @param rhs The second element
  * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns
  */
 template <typename Element, std::enable_if_t<not std::is_floating_point_v<Element>>* = nullptr>
@@ -256,7 +256,7 @@ class row_equality_comparator {
    *
    * @param lhs_row_index The index of the first row to compare (in the lhs table)
    * @param rhs_row_index The index of the second row to compare (in the rhs table)
-   * @return true if both rows are equal
+   * @return true if both rows are equal, otherwise false
    */
   __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept
   {

From 5840fa9fd66b94ac41496fe01d06ab5f07f886b7 Mon Sep 17 00:00:00 2001
From: Karthikeyan Natarajan <karthikeyann@users.noreply.github.com>
Date: Sat, 28 May 2022 05:39:59 +0530
Subject: [PATCH 3/7] address review comments

---
 cpp/include/cudf/table/experimental/row_operators.cuh | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index 0275b66ddd3..3ab7e7546d8 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -884,11 +884,7 @@ class device_row_comparator {
 };
 
 /**
- * @brief Preprocess table for use with row equality comparison or row hashing
- *
- * Sets up the table for use with row equality comparison or row hashing. The resulting
- * preprocessed table can be passed to the constructor of `equality::self_comparator` to
- * avoid preprocessing again.
+ * @brief Preprocessed table for use with row equality comparison or row hashing
  *
  */
 struct preprocessed_table {
@@ -896,6 +892,10 @@ struct preprocessed_table {
    * @brief Factory to construct preprocessed_table for use with
    * row equality comparison or row hashing
    *
+   * Sets up the table for use with row equality comparison or row hashing. The resulting
+   * preprocessed table can be passed to the constructor of `equality::self_comparator` to
+   * avoid preprocessing again.
+   *
    * @param table The table to preprocess
    * @param stream The cuda stream to use while preprocessing.
    * @return A preprocessed table as shared pointer

From 41f1e8c96a8e37cdbbf93b60772156be76b5b86f Mon Sep 17 00:00:00 2001
From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com>
Date: Wed, 1 Jun 2022 18:14:26 +0530
Subject: [PATCH 4/7] Apply suggestions from code review

Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/table/table_device_view.cuh | 1 +
 cpp/include/cudf/table/table_view.hpp        | 2 +-
 2 files changed, 2 insertions(+), 1 deletion(-)

diff --git a/cpp/include/cudf/table/table_device_view.cuh b/cpp/include/cudf/table/table_device_view.cuh
index 52701cb5899..3391fe461d8 100644
--- a/cpp/include/cudf/table/table_device_view.cuh
+++ b/cpp/include/cudf/table/table_device_view.cuh
@@ -187,6 +187,7 @@ class table_device_view : public detail::table_device_view_base<column_device_vi
 
 /**
  * @brief Mutable table device view that is usable in device memory
+ 
  * Elements of the table can be modified in device memory.
  */
 class mutable_table_device_view
diff --git a/cpp/include/cudf/table/table_view.hpp b/cpp/include/cudf/table/table_view.hpp
index ef8b9803d3e..b8fa3585306 100644
--- a/cpp/include/cudf/table/table_view.hpp
+++ b/cpp/include/cudf/table/table_view.hpp
@@ -82,7 +82,7 @@ class table_view_base {
   /**
    * @brief Returns an iterator to the first view in the `table`.
    *
-   * @return An iterator to the first view in the `table`
+   * @return An iterator to the first column_view
    */
   iterator begin() noexcept { return std::begin(_columns); }
 

From d7b715aa79daaf63ea4ac31be1e457467cdd6240 Mon Sep 17 00:00:00 2001
From: Karthikeyan Natarajan <karthikeyann@users.noreply.github.com>
Date: Wed, 1 Jun 2022 18:29:04 +0530
Subject: [PATCH 5/7] newline style fix

---
 cpp/include/cudf/table/table_device_view.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/table/table_device_view.cuh b/cpp/include/cudf/table/table_device_view.cuh
index 3391fe461d8..13c09719be8 100644
--- a/cpp/include/cudf/table/table_device_view.cuh
+++ b/cpp/include/cudf/table/table_device_view.cuh
@@ -187,7 +187,7 @@ class table_device_view : public detail::table_device_view_base<column_device_vi
 
 /**
  * @brief Mutable table device view that is usable in device memory
- 
+ *
  * Elements of the table can be modified in device memory.
  */
 class mutable_table_device_view

From ffbb1812fb0287eda0d81e8354f6afbc5bec36f7 Mon Sep 17 00:00:00 2001
From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com>
Date: Wed, 1 Jun 2022 20:11:51 +0530
Subject: [PATCH 6/7] Update cpp/include/cudf/table/table_view.hpp

Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/table/table_view.hpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/cpp/include/cudf/table/table_view.hpp b/cpp/include/cudf/table/table_view.hpp
index b8fa3585306..2e35ef2ca3a 100644
--- a/cpp/include/cudf/table/table_view.hpp
+++ b/cpp/include/cudf/table/table_view.hpp
@@ -299,8 +299,9 @@ inline bool nullable(table_view const& view)
 }
 
 /**
- * @brief Returns True if the table has nulls in any of its columns. (not entire hierarchy)
+ * @brief Returns True if the table has nulls in any of its columns.
  *
+ * This checks for nulls in the columns and but not in any of the columns' children.
  * @param view The table to check for nulls
  * @return True if the table has nulls in any of its columns, false otherwise
  */

From 7e2fe0005ff4e53e66e71d2b718c59dc71834269 Mon Sep 17 00:00:00 2001
From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com>
Date: Wed, 1 Jun 2022 20:27:58 +0530
Subject: [PATCH 7/7] Update cpp/include/cudf/table/table_view.hpp

Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/table/table_view.hpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/cpp/include/cudf/table/table_view.hpp b/cpp/include/cudf/table/table_view.hpp
index 2e35ef2ca3a..2df37887d9a 100644
--- a/cpp/include/cudf/table/table_view.hpp
+++ b/cpp/include/cudf/table/table_view.hpp
@@ -302,6 +302,7 @@ inline bool nullable(table_view const& view)
  * @brief Returns True if the table has nulls in any of its columns.
  *
  * This checks for nulls in the columns and but not in any of the columns' children.
+ *
  * @param view The table to check for nulls
  * @return True if the table has nulls in any of its columns, false otherwise
  */