[REVIEW] Define and implement new stream compaction APIs copy_if, drop_nulls, apply_boolean_mask, drop_duplicate and unique_count.

CHANGELOG.md

@@ -79,6 +79,7 @@
 - PR #3239 Adding floating point specialization to comparators for NaNs
 - PR #3270 Move predicates files to legacy
 - PR #3282 Add `num_bitmask_words`
+- PR #3303 Define and implement new stream_compaction APIs `copy_if` and `drop_null`
 
 
 ## Bug Fixes

cpp/CMakeLists.txt

@@ -392,6 +392,7 @@ add_library(cudf
             src/transform/jit/code/kernel.cpp
             src/transform/legacy/nans_to_nulls.cu
             src/bitmask/legacy/bitmask_ops.cu
+            src/stream_compaction/drop_nulls.cu
             src/stream_compaction/legacy/apply_boolean_mask.cu
             src/stream_compaction/legacy/drop_nulls.cu
             src/stream_compaction/legacy/drop_duplicates.cu

cpp/include/cudf/detail/stream_compaction.hpp

@@ -0,0 +1,71 @@
+/*
+ * Copyright (c) 2019, 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.
+ */
+
+#pragma once
+
+#include <cudf/cudf.h>
+#include <cudf/types.hpp>
+
+namespace cudf {
+namespace experimental {
+namespace detail {
+
+/**
+ * @brief Filters a table to remove null elements.
+ *
+ * Filters the rows of the `input` considering specified columns from
+ * `keys` for validity / null values.
+ *
+ * Given an input table_view, row `i` from the input columns is copied to
+ * the output if the same row `i` of @p keys has at leaast @p keep_threshold
+ * non-null fields.
+ *
+ * This operation is stable: the input order is preserved in the output.
+ *
+ * Any non-nullable column in the input is treated as all non-null.
+ *
+ * @example input   {col1: {1, 2,    3,    null},
+ *                   col2: {4, 5,    null, null},
+ *                   col3: {7, null, null, null}}
+ *          keys = input
+ *          keep_threshold = 2
+ *
+ *          output {col1: {1, 2}
+ *                  col2: {4, 5}
+ *                  col3: {7, null}}
+ *
+ * @note if @p input.num_rows() is zero, or @p keys is empty or has no nulls,
+ * there is no error, and an empty `std::unique_ptr<table>` is returned
+ *
+ * @throws cudf::logic_error if @p keys is non-empty and keys.num_rows() is less
+ * than input.num_rows()
+ *
+ * @param[in] input The input `table_view` to filter.
+ * @param[in] keys The `table_view` to filter `input`.
+ * @param[in] keep_threshold The minimum number of non-null fields in a row
+ *                           required to keep the row.
+ * @param[in] mr Optional, The resource to use for all allocations
+ * @param[in] stream Optional CUDA stream on which to execute kernels
+ * @return unique_ptr<table> Table containing all rows of the `input` with at least @p keep_threshold non-null fields in @p keys.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const& input,
+                 table_view const& keys,
+                 cudf::size_type keep_threshold,
+                 rmm::mr::device_memory_resource *mr = rmm::mr::get_default_resource(),
+                 cudaStream_t stream = 0);
+} // namespace detail
+} // namespace experimental
+} // namespace cudf

cpp/include/cudf/stream_compaction.hpp

@@ -0,0 +1,94 @@
+/*
+ * Copyright (c) 2019, 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.
+ */
+
+#pragma once
+
+#include <cudf/cudf.h>
+#include <cudf/types.hpp>
+
+namespace cudf {
+namespace experimental {
+
+/**
+ * @brief Filters a table to remove null elements.
+ *
+ * Filters the rows of the `input` considering specified columns from
+ * `keys` for validity / null values.
+ *
+ * Given an input table_view, row `i` from the input columns is copied to
+ * the output if the same row `i` of @p keys has at leaast @p keep_threshold
+ * non-null fields.
+ *
+ * This operation is stable: the input order is preserved in the output.
+ * 
+ * Any non-nullable column in the input is treated as all non-null.
+ *
+ * @example input   {col1: {1, 2,    3,    null},
+ *                   col2: {4, 5,    null, null},
+ *                   col3: {7, null, null, null}}
+ *          keys = input
+ *          keep_threshold = 2
+ *
+ *          output {col1: {1, 2}
+ *                  col2: {4, 5}
+ *                  col3: {7, null}}
+ *
+ * @note if @p input.num_rows() is zero, or @p keys is empty or has no nulls,
+ * there is no error, and an empty `std::unique_ptr<table>` is returned
+ * 
+ * @throws cudf::logic_error if @p keys is non-empty and keys.num_rows() is less
+ * than input.num_rows()
+ *
+ * @param[in] input The input `table_view` to filter.
+ * @param[in] keys The `table_view` to filter `input`.
+ * @param[in] keep_threshold The minimum number of non-null fields in a row
+ *                           required to keep the row.
+ * @param[in] mr Optional, The resource to use for all allocations
+ * @return unique_ptr<table> Table containing all rows of the `input` with at least @p keep_threshold non-null fields in @p keys.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const& input,
+                 table_view const& keys,
+                 cudf::size_type keep_threshold,
+                 rmm::mr::device_memory_resource *mr = rmm::mr::get_default_resource());
+
+/**
+ * @brief Filters a table to remove null elements.
+ *
+ * @example input   {col1: {1, 2,    3,    null},
+ *                   col2: {4, 5,    null, null},
+ *                   col3: {7, null, null, null}}
+ *          keys = input
+ *
+ *          output {col1: {1}
+ *                  col2: {4}
+ *                  col3: {7}}
+ *
+ * @overload drop_nulls
+ *
+ * Same as drop_nulls but defaults keep_threshold to the number of columns in
+ * @p keys.
+ *
+ * @param[in] input The input `table_view` to filter.
+ * @param[in] keys The `table_view` to filter `input`.
+ * @param[in] mr Optional, The resource to use for all allocations
+ * @return unique_ptr<table> Table containing all rows of the `input` without nulls in the columns of @p keys.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const &input,
+                 table_view const &keys,
+                 rmm::mr::device_memory_resource *mr = rmm::mr::get_default_resource());
+
+} // namespace experimental
+} // namespace cudf

cpp/include/cudf/table/table_view.hpp

@@ -149,13 +149,14 @@ class table_view : public detail::table_view_base<column_view> {
 class mutable_table_view : public detail::table_view_base<mutable_column_view> {
   using detail::table_view_base<mutable_column_view>::table_view_base;
 
-  mutable_column_view& column(size_type column_index) const noexcept {
-    return const_cast<mutable_column_view&>(table_view_base::column(column_index));
-  }
+  public:
+    mutable_column_view& column(size_type column_index) const noexcept {
+      return const_cast<mutable_column_view&>(table_view_base::column(column_index));
+    }
   /**---------------------------------------------------------------------------*
    * @brief Creates an immutable `table_view` of the columns
    *---------------------------------------------------------------------------**/
-  operator table_view();
+    operator table_view();
 };
 
 inline bool has_nulls(table_view view) {

cpp/include/cudf/wrappers/bool.hpp

@@ -151,7 +151,7 @@ struct bool8 {
   }
 
 private:
-  value_type value{0};
+  value_type value;
 };
 
 // This is necessary for global, constant, non-fundamental types

cpp/include/cudf/wrappers/timestamps.hpp

@@ -35,7 +35,7 @@ namespace detail {
     template <class Duration>
     struct timestamp : time_point<Duration> {
         /// Initialize timestamp to 1/1/1970:00:00:00
-        constexpr timestamp() : time_point<Duration>(Duration(0)) {};
+        constexpr timestamp() : time_point<Duration>(Duration()) {};
         constexpr timestamp(Duration d) : time_point<Duration>(d) {};
         constexpr timestamp(typename Duration::rep t) : time_point<Duration>(Duration(t)) {};
     };

cpp/src/stream_compaction/drop_nulls.cu

@@ -0,0 +1,129 @@
+/*
+ * Copyright (c) 2019, 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/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/detail/stream_compaction.hpp>
+#include <cudf/detail/copy_if.cuh>
+
+namespace {
+
+// Returns true if the mask is true for index i in at least keep_threshold
+// columns
+struct valid_table_filter
+{
+  __device__ inline
+  bool operator()(cudf::size_type i)
+  {
+    auto valid = [i](auto column_device_view) {
+      return column_device_view.is_valid(i);
+    };
+
+    auto count =
+      thrust::count_if(thrust::seq, keys_device_view.begin(), keys_device_view.end(), valid);
+
+    return (count >= keep_threshold);
+  }
+
+  static auto create(cudf::table_device_view const& keys,
+                     cudf::size_type keep_threshold,
+                     cudaStream_t stream = 0)
+  {
+    //auto keys_device_view = cudf::table_device_view::create(keys, stream);
+
+    auto deleter = [stream](valid_table_filter* f) { f->destroy(); };
+    std::unique_ptr<valid_table_filter, decltype(deleter)> p {
+      new valid_table_filter(keys, keys.num_columns(), keep_threshold),
+      deleter
+    };
+
+    return p;
+  }
+
+  __host__ void destroy() {
+    delete this;
+  }
+
+  valid_table_filter() = delete;
+  ~valid_table_filter() = default;
+
+protected:
+
+
+  valid_table_filter(cudf::table_device_view const& keys_device_view,
+                     cudf::size_type num_columns,
+                     cudf::size_type keep_threshold)
+  : keep_threshold(keep_threshold),
+    num_columns(num_columns),
+    keys_device_view(keys_device_view) {}
+
+  cudf::size_type keep_threshold;
+  cudf::size_type num_columns;
+  cudf::table_device_view keys_device_view;
+};
+
+}  // namespace
+
+namespace cudf {
+namespace experimental {
+namespace detail {
+
+/*
+ * Filters a table to remove null elements.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const& input,
+                 table_view const& keys,
+                 cudf::size_type keep_threshold,
+                 rmm::mr::device_memory_resource *mr,
+                 cudaStream_t stream) {
+  if (keys.num_columns() == 0 || keys.num_rows() == 0 ||
+      not cudf::has_nulls(keys)) {
+      return std::make_unique<table>(input, stream, mr);
+  }
+
+  CUDF_EXPECTS(keys.num_rows() <= input.num_rows(),
+               "Column size mismatch");
+
+  auto keys_device_view = cudf::table_device_view::create(keys, stream);
+  auto filter = valid_table_filter::create(*keys_device_view, keep_threshold);
+
+  return cudf::experimental::detail::copy_if(input, *filter.get(), mr, stream);
+}
+
+} //namespace detail
+
+/*
+ * Filters a table to remove null elements.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const& input,
+                 table_view const& keys,
+                 cudf::size_type keep_threshold,
+                 rmm::mr::device_memory_resource *mr) {
+    return cudf::experimental::detail::drop_nulls(input, keys, keep_threshold, mr);
+}
+/*
+ * Filters a table to remove null elements.
+ */
+std::unique_ptr<experimental::table> drop_nulls(table_view const &input,
+                 table_view const &keys,
+                 rmm::mr::device_memory_resource *mr)
+{
+    return cudf::experimental::detail::drop_nulls(input, keys, keys.num_columns(), mr);
+}
+
+} //namespace experimental
+} //namespace cudf

cpp/tests/CMakeLists.txt

@@ -495,7 +495,15 @@ set(LEGACY_MERGE_TEST_SRC
 ConfigureTest(LEGACY_MERGE_TEST "${LEGACY_MERGE_TEST_SRC}")
 
 ###################################################################################################
-# - legacy stream compaction tests ----------------------------------------------------------------
+# - stream compaction tests -----------------------------------------------------------------------
+
+set(STREAM_COMPACTION_TEST_SRC
+    "${CMAKE_CURRENT_SOURCE_DIR}/stream_compaction/drop_nulls_tests.cu")
+
+ConfigureTest(STREAM_COMPACTION_TEST "${STREAM_COMPACTION_TEST_SRC}")
+
+###################################################################################################
+# - legacy stream compaction tests -----------------------------------------------------------------------
 
 set(LEGACY_STREAM_COMPACTION_TEST_SRC
     "${CMAKE_CURRENT_SOURCE_DIR}/stream_compaction/legacy/apply_boolean_mask_tests.cu"