Refactor purge_nonempty_nulls

conda/recipes/libcudf/meta.yaml

@@ -78,7 +78,6 @@ outputs:
         - test -f $PREFIX/include/cudf/detail/binaryop.hpp
         - test -f $PREFIX/include/cudf/detail/calendrical_month_sequence.cuh
         - test -f $PREFIX/include/cudf/detail/concatenate.hpp
-        - test -f $PREFIX/include/cudf/detail/copy.cuh
         - test -f $PREFIX/include/cudf/detail/copy.hpp
         - test -f $PREFIX/include/cudf/detail/datetime.hpp
         - test -f $PREFIX/include/cudf/detail/fill.hpp

cpp/include/cudf/copying.hpp

@@ -1012,12 +1012,19 @@ bool has_nonempty_nulls(column_view const& input);
 bool may_have_nonempty_nulls(column_view const& input);
 
 /**
- * @brief Copies `input`, purging any non-empty null rows in the column or its descendants
+ * @brief Copy `input` into output while purging any non-empty null rows in the column or its
+ * descendants.
  *
- * LIST columns may have non-empty null rows.
- * For example:
- * @code{.pseudo}
+ * If the input column is not of compound type (LIST/STRING/STRUCT/DICTIONARY), the output will be
+ * the same as input.
+ *
+ * The purge operation only applies directly to LIST and STRING columns, but it applies indirectly
+ * to STRUCT/DICTIONARY columns as well, since these columns may have child columns that
+ * are LIST or STRING.
  *
+ * Examples:
+ *
+ * @code{.pseudo}
  * auto const lists   = lists_column_wrapper<int32_t>{ {0,1}, {2,3}, {4,5} }.release();
  * cudf::detail::set_null_mask(lists->null_mask(), 1, 2, false);
  *
@@ -1027,33 +1034,13 @@ bool may_have_nonempty_nulls(column_view const& input);
  *   Offsets:  [0, 2, 4, 6]
  *   Child:    [0, 1, 2, 3, 4, 5]
  *
- * After purging the contents of the list's null rows, the column's contents
- * will be:
+ * After purging the contents of the list's null rows, the column's contents will be:
  *   Validity: 101
  *   Offsets:  [0, 2, 2, 4]
  *   Child:    [0, 1, 4, 5]
  * @endcode
  *
- * The purge operation only applies directly to LIST and STRING columns, but it
- * applies indirectly to STRUCT columns as well, since LIST and STRUCT columns
- * may have child/descendant columns that are LIST or STRING.
- *
- * @param input The column whose null rows are to be checked and purged
- * @param mr Device memory resource used to allocate the returned column's device memory
- * @return std::unique_ptr<column> Column with equivalent contents to `input`, but with
- * the contents of null rows purged
- */
-std::unique_ptr<column> purge_nonempty_nulls(
-  lists_column_view const& input,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
-/**
- * @brief Copies `input`, purging any non-empty null rows in the column or its descendants
- *
- * STRING columns may have non-empty null rows.
- * For example:
  * @code{.pseudo}
- *
  * auto const strings = strings_column_wrapper{ "AB", "CD", "EF" }.release();
  * cudf::detail::set_null_mask(strings->null_mask(), 1, 2, false);
  *
@@ -1070,26 +1057,7 @@ std::unique_ptr<column> purge_nonempty_nulls(
  *   Child:    [A, B, E, F]
  * @endcode
  *
- * The purge operation only applies directly to LIST and STRING columns, but it
- * applies indirectly to STRUCT columns as well, since LIST and STRUCT columns
- * may have child/descendant columns that are LIST or STRING.
- *
- * @param input The column whose null rows are to be checked and purged
- * @param mr Device memory resource used to allocate the returned column's device memory
- * @return std::unique_ptr<column> Column with equivalent contents to `input`, but with
- * the contents of null rows purged
- */
-std::unique_ptr<column> purge_nonempty_nulls(
-  strings_column_view const& input,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
-/**
- * @brief Copies `input`, purging any non-empty null rows in the column or its descendants
- *
- * STRUCTS columns may have null rows, with non-empty child rows.
- * For example:
  * @code{.pseudo}
- *
  * auto const lists   = lists_column_wrapper<int32_t>{ {0,1}, {2,3}, {4,5} };
  * auto const structs = structs_column_wrapper{ {lists}, null_at(1) };
  *
@@ -1106,17 +1074,12 @@ std::unique_ptr<column> purge_nonempty_nulls(
  *   Child:    [0, 1, 4, 5]
  * @endcode
  *
- * The purge operation only applies directly to LIST and STRING columns, but it
- * applies indirectly to STRUCT columns as well, since LIST and STRUCT columns
- * may have child/descendant columns that are LIST or STRING.
- *
  * @param input The column whose null rows are to be checked and purged
  * @param mr Device memory resource used to allocate the returned column's device memory
- * @return std::unique_ptr<column> Column with equivalent contents to `input`, but with
- * the contents of null rows purged
+ * @return A new column with equivalent contents to `input`, but with null rows purged
  */
 std::unique_ptr<column> purge_nonempty_nulls(
-  structs_column_view const& input,
+  column_view const& input,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /** @} */

cpp/include/cudf/detail/copy.hpp

@@ -315,5 +315,15 @@ bool has_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream);
  */
 bool may_have_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream);
 
+/**
+ * @copydoc cudf::purge_nonempty_nulls
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+std::unique_ptr<column> purge_nonempty_nulls(
+  column_view const& input,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 }  // namespace detail
 }  // namespace cudf

cpp/src/copying/purge_nonempty_nulls.cu

@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 #include <cudf/copying.hpp>
-#include <cudf/detail/copy.cuh>
+#include <cudf/detail/gather.cuh>
 #include <cudf/utilities/default_stream.hpp>
 
 #include <thrust/count.h>
@@ -80,6 +80,24 @@ bool has_nonempty_nulls(cudf::column_view const& input, rmm::cuda_stream_view st
 
   return false;
 }
+
+std::unique_ptr<column> purge_nonempty_nulls(column_view const& input,
+                                             rmm::cuda_stream_view stream,
+                                             rmm::mr::device_memory_resource* mr)
+{
+  // If not compound types (LIST/STRING/STRUCT/DICTIONARY) then just copy the input into output.
+  if (!cudf::is_compound(input.type())) { return std::make_unique<column>(input, stream, mr); }
+
+  // Implement via identity gather.
+  auto gathered_table = cudf::detail::gather(table_view{{input}},
+                                             thrust::make_counting_iterator(0),
+                                             thrust::make_counting_iterator(input.size()),
+                                             out_of_bounds_policy::DONT_CHECK,
+                                             stream,
+                                             mr);
+  return std::move(gathered_table->release().front());
+}
+
 }  // namespace detail
 
 /**
@@ -110,27 +128,9 @@ bool has_nonempty_nulls(column_view const& input)
 }
 
 /**
- * @copydoc cudf::purge_nonempty_nulls(lists_column_view const&, rmm::mr::device_memory_resource*)
- */
-std::unique_ptr<cudf::column> purge_nonempty_nulls(lists_column_view const& input,
-                                                   rmm::mr::device_memory_resource* mr)
-{
-  return detail::purge_nonempty_nulls(input, cudf::get_default_stream(), mr);
-}
-
-/**
- * @copydoc cudf::purge_nonempty_nulls(structs_column_view const&, rmm::mr::device_memory_resource*)
- */
-std::unique_ptr<cudf::column> purge_nonempty_nulls(structs_column_view const& input,
-                                                   rmm::mr::device_memory_resource* mr)
-{
-  return detail::purge_nonempty_nulls(input, cudf::get_default_stream(), mr);
-}
-
-/**
- * @copydoc cudf::purge_nonempty_nulls(strings_column_view const&, rmm::mr::device_memory_resource*)
+ * @copydoc cudf::purge_nonempty_nulls(column_view const&, rmm::mr::device_memory_resource*)
  */
-std::unique_ptr<cudf::column> purge_nonempty_nulls(strings_column_view const& input,
+std::unique_ptr<cudf::column> purge_nonempty_nulls(column_view const& input,
                                                    rmm::mr::device_memory_resource* mr)
 {
   return detail::purge_nonempty_nulls(input, cudf::get_default_stream(), mr);

cpp/src/lists/set_operations.cu

@@ -17,7 +17,7 @@
 #include "utilities.hpp"
 
 #include <cudf/column/column_factories.hpp>
-#include <cudf/detail/copy.cuh>
+#include <cudf/detail/copy.hpp>
 #include <cudf/detail/copy_if.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
@@ -176,9 +176,8 @@ std::unique_ptr<column> intersect_distinct(lists_column_view const& lhs,
                                   stream,
                                   mr);
 
-  return null_count == 0
-           ? std::move(output)
-           : cudf::detail::purge_nonempty_nulls(lists_column_view{output->view()}, stream, mr);
+  return null_count == 0 ? std::move(output)
+                         : cudf::detail::purge_nonempty_nulls(output->view(), stream, mr);
 }
 
 std::unique_ptr<column> union_distinct(lists_column_view const& lhs,
@@ -253,9 +252,8 @@ std::unique_ptr<column> difference_distinct(lists_column_view const& lhs,
                                   stream,
                                   mr);
 
-  return null_count == 0
-           ? std::move(output)
-           : cudf::detail::purge_nonempty_nulls(lists_column_view{output->view()}, stream, mr);
+  return null_count == 0 ? std::move(output)
+                         : cudf::detail::purge_nonempty_nulls(output->view(), stream, mr);
 }
 
 }  // namespace detail