Support purging non-empty null elements from LIST/STRING columns (#10701

)

Fixes #10291.

With certain operations in `libcudf`, it is possible to produce `LIST` columns with `NULL` rows that are not also empty. 
For instance, consider a `STRUCT` column is constructed with an explicit validity buffer and a `LIST` child column:
```c++
auto const lists   = lists_column_wrapper<int32_t>{ {0,1}, {2,3}, {4,5} };
auto const structs = structs_column_wrapper{ {lists}, null_at(1) };
```
Since `structs[1] == NULL`, its `LIST` member is also deemed null. However, for efficiency, the null-ness is recorded in the `LIST`'s validity buffer, without purging the unnecessary values from its child. The `LIST` columns appears as follows:
```
Validity: 101
Offsets:  [0, 2, 4, 6]
Child:    [0, 1, 2, 3, 4, 5]
```
Even though Row#1 is null, its size is `4-2 = 2`, and not `0`. (Row#1 is thus a non-empty null row.)

This commit adds a `cudf::purge_nonempty_nulls()` function that purges such rows, and reduces such columns to a more space-efficient representation, i.e.:
```
Validity: 101
Offsets:  [0, 2, 2, 4]
Child:    [0, 1, 4, 5]
```

This commit also modifies `cudf::gather()` not to produce `STRING`/`LIST` columns with "dirty" rows. Further, it adds two new functions to determine if a specified column needs such purging:
1. `cudf::may_have_nonempty_nulls()`: A fast check to check a column for the *possibility* of having non-empty nulls. This only checks whether the column or its descendants have null rows at all. If there are no nulls anywhere in the hierarchy, it does not need purging.
2. `cudf::has_nonempty_nulls()`: A deeper, more expensive check that categorically confirms whether non-empty null rows exist in any column in the hierarchy.

Authors:
  - MithunR (https://github.com/mythrocks)

Approvers:
  - Jake Hemstad (https://github.com/jrhemstad)
  - https://github.com/nvdbaranec
  - Jordan Jacobelli (https://github.com/Ethyling)

URL: #10701

conda/recipes/libcudf/meta.yaml

@@ -79,6 +79,7 @@ outputs:
         - 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.hpp
+        - test -f $PREFIX/include/cudf/detail/copy.cuh
         - test -f $PREFIX/include/cudf/detail/datetime.hpp
         - test -f $PREFIX/include/cudf/detail/fill.hpp
         - test -f $PREFIX/include/cudf/detail/gather.hpp

cpp/CMakeLists.txt

@@ -238,6 +238,7 @@ add_library(
   src/copying/gather.cu
   src/copying/get_element.cu
   src/copying/pack.cpp
+  src/copying/purge_nonempty_nulls.cu
   src/copying/reverse.cu
   src/copying/sample.cu
   src/copying/scatter.cu

cpp/include/cudf/copying.hpp

@@ -17,7 +17,10 @@
 #pragma once
 
 #include <cudf/column/column_view.hpp>
+#include <cudf/lists/lists_column_view.hpp>
 #include <cudf/scalar/scalar.hpp>
+#include <cudf/strings/strings_column_view.hpp>
+#include <cudf/structs/structs_column_view.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/types.hpp>
 
@@ -939,5 +942,155 @@ std::unique_ptr<table> sample(
   int64_t const seed                  = 0,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Checks if a column or its descendants have non-empty null rows
+ *
+ * @note This function is exact. If it returns `true`, there exists one or more
+ * non-empty null elements.
+ *
+ * A LIST or STRING column might have non-empty rows that are marked as null.
+ * A STRUCT OR LIST column might have child columns that have non-empty null rows.
+ * Other types of columns are deemed incapable of having non-empty null rows.
+ * E.g. Fixed width columns have no concept of an "empty" row.
+ *
+ * @param input The column which is (and whose descendants are) to be checked for
+ * non-empty null rows.
+ * @return true If either the column or its descendants have non-empty null rows.
+ * @return false If neither the column or its descendants have non-empty null rows.
+ */
+bool has_nonempty_nulls(column_view const& input);
+
+/**
+ * @brief Approximates if a column or its descendants *may* have non-empty null elements
+ *
+ * @note This function is approximate.
+ * - `true`: Non-empty null elements could exist
+ * - `false`: Non-empty null elements definitely do not exist
+ *
+ * False positives are possible, but false negatives are not.
+ *
+ * Compared to the exact `has_nonempty_nulls()` function, this function is typically
+ * more efficient.
+ *
+ * Complexity:
+ * - Best case: `O(count_descendants(input))`
+ * - Worst case: `O(count_descendants(input)) * m`, where `m` is the number of rows in the largest
+ * descendant
+ *
+ * @param input The column which is (and whose descendants are) to be checked for
+ * non-empty null rows
+ * @return true If either the column or its decendants have null rows
+ * @return false If neither the column nor its descendants have null rows
+ */
+bool may_have_nonempty_nulls(column_view const& input);
+
+/**
+ * @brief Copies `input`, purging any non-empty null rows in the column or its descendants
+ *
+ * LIST columns may have non-empty null rows.
+ * For example:
+ * @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);
+ *
+ * lists[1] is now null, but the lists child column still stores `{2,3}`.
+ * The lists column contents will be:
+ *   Validity: 101
+ *   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:
+ *   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/decendant 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);
+ *
+ * strings[1] is now null, but the strings column still stores `"CD"`.
+ * The lists column contents will be:
+ *   Validity: 101
+ *   Offsets:  [0, 2, 4, 6]
+ *   Child:    [A, B, C, D, E, F]
+ *
+ * After purging the contents of the list's null rows, the column's contents
+ * will be:
+ *   Validity: 101
+ *   Offsets:  [0, 2, 2, 4]
+ *   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/decendant 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) };
+ *
+ * structs[1].child is now null, but the lists column still stores `{2,3}`.
+ * The lists column contents will be:
+ *   Validity: 101
+ *   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:
+ *   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/decendant 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(
+  structs_column_view const& input,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /** @} */
 }  // namespace cudf

cpp/include/cudf/detail/copy.cuh

@@ -0,0 +1,47 @@
+/*
+ * Copyright (c) 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.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cudf/detail/copy.hpp>
+#include <cudf/detail/gather.cuh>
+
+namespace cudf::detail {
+
+/**
+ * @copydoc cudf::purge_nonempty_nulls(structs_column_view const&, rmm::mr::device_memory_resource*)
+ *
+ * @tparam ColumnViewT View type (lists_column_view, strings_column_view, or strings_column_view)
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ */
+template <typename ColumnViewT>
+std::unique_ptr<cudf::column> purge_nonempty_nulls(ColumnViewT const& input,
+                                                   rmm::cuda_stream_view stream,
+                                                   rmm::mr::device_memory_resource* mr)
+{
+  // Implement via identity gather.
+  auto const input_column = input.parent();
+  auto const gather_begin = thrust::counting_iterator<cudf::size_type>(0);
+  auto const gather_end   = gather_begin + input_column.size();
+
+  auto gathered_table = cudf::detail::gather(table_view{{input_column}},
+                                             gather_begin,
+                                             gather_end,
+                                             out_of_bounds_policy::DONT_CHECK,
+                                             stream,
+                                             mr);
+  return std::move(gathered_table->release()[0]);
+}
+
+}  // namespace cudf::detail

cpp/include/cudf/detail/copy.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2018-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.
@@ -299,5 +299,22 @@ std::unique_ptr<scalar> get_element(
   size_type index,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @copydoc cudf::has_nonempty_nulls
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+bool has_nonempty_nulls(column_view const& input,
+                        rmm::cuda_stream_view stream = rmm::cuda_stream_default);
+
+/**
+ * @copydoc cudf::may_have_nonempty_nulls
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+bool may_have_nonempty_nulls(column_view const& input,
+                             rmm::cuda_stream_view stream = rmm::cuda_stream_default);
+
 }  // namespace detail
 }  // namespace cudf

cpp/include/cudf/lists/detail/gather.cuh

@@ -18,6 +18,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/get_value.cuh>
 #include <cudf/lists/lists_column_view.hpp>
+#include <cudf/utilities/bit.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
@@ -82,6 +83,7 @@ gather_data make_gather_data(cudf::lists_column_view const& source_column,
   auto dst_offsets_c = cudf::make_fixed_width_column(
     data_type{type_id::INT32}, offset_count, mask_state::UNALLOCATED, stream, mr);
   mutable_column_view dst_offsets_v = dst_offsets_c->mutable_view();
+  auto const source_column_nullmask = source_column.null_mask();
 
   // generate the compacted outgoing offsets.
   auto count_iter = thrust::make_counting_iterator<int32_t>(0);
@@ -90,12 +92,23 @@ gather_data make_gather_data(cudf::lists_column_view const& source_column,
     count_iter,
     count_iter + offset_count,
     dst_offsets_v.begin<int32_t>(),
-    [gather_map, output_count, src_offsets, src_size] __device__(int32_t index) -> int32_t {
+    [source_column_nullmask,
+     source_column_offset = source_column.offset(),
+     gather_map,
+     output_count,
+     src_offsets,
+     src_size] __device__(int32_t index) -> int32_t {
       int32_t offset_index = index < output_count ? gather_map[index] : 0;
 
       // if this is an invalid index, this will be a NULL list
       if (NullifyOutOfBounds && ((offset_index < 0) || (offset_index >= src_size))) { return 0; }
 
+      // If the source row is null, the output row size must be 0.
+      if (source_column_nullmask != nullptr &&
+          not cudf::bit_is_set(source_column_nullmask, source_column_offset + offset_index)) {
+        return 0;
+      }
+
       // the length of this list
       return src_offsets[offset_index + 1] - src_offsets[offset_index];
     },
@@ -110,15 +123,27 @@ gather_data make_gather_data(cudf::lists_column_view const& source_column,
 
   // generate the base offsets
   rmm::device_uvector<int32_t> base_offsets = rmm::device_uvector<int32_t>(output_count, stream);
-  thrust::transform(rmm::exec_policy(stream),
-                    gather_map,
-                    gather_map + output_count,
-                    base_offsets.data(),
-                    [src_offsets, src_size, shift] __device__(int32_t index) {
-                      // if this is an invalid index, this will be a NULL list
-                      if (NullifyOutOfBounds && ((index < 0) || (index >= src_size))) { return 0; }
-                      return src_offsets[index] - shift;
-                    });
+  thrust::transform(
+    rmm::exec_policy(stream),
+    gather_map,
+    gather_map + output_count,
+    base_offsets.data(),
+    [source_column_nullmask,
+     source_column_offset = source_column.offset(),
+     src_offsets,
+     src_size,
+     shift] __device__(int32_t index) {
+      // if this is an invalid index, this will be a NULL list
+      if (NullifyOutOfBounds && ((index < 0) || (index >= src_size))) { return 0; }
+
+      // If the source row is null, the output row size must be 0.
+      if (source_column_nullmask != nullptr &&
+          not cudf::bit_is_set(source_column_nullmask, source_column_offset + index)) {
+        return 0;
+      }
+
+      return src_offsets[index] - shift;
+    });
 
   // Retrieve size of the resulting gather map for level N+1 (the last offset)
   size_type child_gather_map_size =

cpp/include/cudf/strings/detail/gather.cuh

@@ -303,14 +303,17 @@ std::unique_ptr<cudf::column> gather(
     data_type{type_id::INT32}, output_count + 1, mask_state::UNALLOCATED, stream, mr);
   auto const d_out_offsets = out_offsets_column->mutable_view().template data<int32_t>();
   auto const d_in_offsets  = (strings_count > 0) ? strings.offsets_begin() : nullptr;
-  thrust::transform(rmm::exec_policy(stream),
-                    begin,
-                    end,
-                    d_out_offsets,
-                    [d_in_offsets, strings_count] __device__(size_type in_idx) {
-                      if (NullifyOutOfBounds && (in_idx < 0 || in_idx >= strings_count)) return 0;
-                      return d_in_offsets[in_idx + 1] - d_in_offsets[in_idx];
-                    });
+  auto const d_strings     = column_device_view::create(strings.parent(), stream);
+  thrust::transform(
+    rmm::exec_policy(stream),
+    begin,
+    end,
+    d_out_offsets,
+    [d_strings = *d_strings, d_in_offsets, strings_count] __device__(size_type in_idx) {
+      if (NullifyOutOfBounds && (in_idx < 0 || in_idx >= strings_count)) return 0;
+      if (not d_strings.is_valid(in_idx)) return 0;
+      return d_in_offsets[in_idx + 1] - d_in_offsets[in_idx];
+    });
 
   // check total size is not too large
   size_t const total_bytes = thrust::transform_reduce(
@@ -329,7 +332,6 @@ std::unique_ptr<cudf::column> gather(
 
   // build chars column
   cudf::device_span<int32_t const> const d_out_offsets_span(d_out_offsets, output_count + 1);
-  auto const d_strings  = column_device_view::create(strings.parent(), stream);
   auto out_chars_column = gather_chars(d_strings->begin<string_view>(),
                                        begin,
                                        end,

cpp/include/cudf/structs/structs_column_view.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 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.
@@ -41,6 +41,11 @@ class structs_column_view : public column_view {
 
   explicit structs_column_view(column_view const& rhs);
 
+  /**
+   * @brief Returns the parent column.
+   */
+  [[nodiscard]] column_view parent() const;
+
   using column_view::child_begin;
   using column_view::child_end;
   using column_view::has_nulls;