List element Equality comparator (#10289)

This PR implements equality comparator for LIST columns. This only supports "self" comparison for now, meaning the two rows to be compared should belong to the same table. A comparator that works on rows of two different tables will be implemented in another PR. This works only on "sanitized" list columns. See #10291 for details. This will partially support #10186. Authors: - Devavret Makkar (https://github.com/devavret) Approvers: - Robert Maynard (https://github.com/robertmaynard) - Vyas Ramasubramani (https://github.com/vyasr) - Mike Wilson (https://github.com/hyperbolic2346) - Jake Hemstad (https://github.com/jrhemstad) - Jordan Jacobelli (https://github.com/Ethyling) URL: #10289
rapidsai · Apr 13, 2022 · 0ea6f8e · 0ea6f8e
1 parent c9e16c7
commit 0ea6f8e
Show file tree

Hide file tree

Showing 15 changed files with 1,106 additions and 188 deletions.
diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml
@@ -110,6 +110,7 @@ outputs:
         - test -f $PREFIX/include/cudf/detail/transpose.hpp
         - test -f $PREFIX/include/cudf/detail/unary.hpp
         - test -f $PREFIX/include/cudf/detail/utilities/alignment.hpp
+        - test -f $PREFIX/include/cudf/detail/utilities/column.hpp
         - test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp
         - test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h
         - test -f $PREFIX/include/cudf/detail/utilities/vector_factories.hpp

diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt
@@ -181,7 +181,7 @@ ConfigureBench(
   REDUCTION_BENCH reduction/anyall.cpp reduction/dictionary.cpp reduction/minmax.cpp
   reduction/reduce.cpp reduction/scan.cpp
 )
-ConfigureNVBench(REDUCTION_NVBENCH reduction/segment_reduce.cu)
+ConfigureNVBench(REDUCTION_NVBENCH reduction/segment_reduce.cu reduction/rank.cpp)
 
 # ##################################################################################################
 # * reduction benchmark ---------------------------------------------------------------------------

diff --git a/cpp/benchmarks/reduction/rank.cpp b/cpp/benchmarks/reduction/rank.cpp
@@ -0,0 +1,64 @@
+/*
+ * 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 <benchmarks/common/generate_input.hpp>
+#include <benchmarks/fixture/rmm_pool_raii.hpp>
+
+#include <cudf/detail/scan.hpp>
+#include <cudf/filling.hpp>
+#include <cudf/lists/list_view.cuh>
+
+#include <nvbench/nvbench.cuh>
+
+template <typename type>
+static void nvbench_reduction_scan(nvbench::state& state, nvbench::type_list<type>)
+{
+  cudf::rmm_pool_raii pool_raii;
+
+  auto const dtype = cudf::type_to_id<type>();
+
+  double const null_frequency = state.get_float64("null_frequency");
+  size_t const size           = state.get_int64("data_size");
+
+  data_profile table_data_profile;
+  table_data_profile.set_distribution_params(dtype, distribution_id::UNIFORM, 0, 5);
+  table_data_profile.set_null_frequency(null_frequency);
+
+  auto const table = create_random_table({dtype}, table_size_bytes{size / 2}, table_data_profile);
+
+  auto const new_tbl = cudf::repeat(table->view(), 2);
+  cudf::column_view input(new_tbl->view().column(0));
+
+  state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
+    rmm::cuda_stream_view stream_view{launch.get_stream()};
+    auto result = cudf::detail::inclusive_dense_rank_scan(
+      input, stream_view, rmm::mr::get_current_device_resource());
+  });
+}
+
+using data_type = nvbench::type_list<int32_t, cudf::list_view>;
+
+NVBENCH_BENCH_TYPES(nvbench_reduction_scan, NVBENCH_TYPE_AXES(data_type))
+  .set_name("rank_scan")
+  .add_float64_axis("null_frequency", {0, 0.1, 0.5, 0.9})
+  .add_int64_axis("data_size",
+                  {
+                    10000,      // 10k
+                    100000,     // 100k
+                    1000000,    // 1M
+                    10000000,   // 10M
+                    100000000,  // 100M
+                  });
diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh
@@ -111,7 +111,7 @@ class alignas(16) column_device_view_base {
    */
   template <typename T = void,
             CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
-  __host__ __device__ T const* head() const noexcept
+  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
   {
     return static_cast<T const*>(_data);
   }
@@ -132,20 +132,20 @@ class alignas(16) column_device_view_base {
    * @return T const* Typed pointer to underlying data, including the offset
    */
   template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
-  __host__ __device__ T const* data() const noexcept
+  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
   {
     return head<T>() + _offset;
   }
 
   /**
    * @brief Returns the number of elements in the column.
    */
-  [[nodiscard]] __host__ __device__ size_type size() const noexcept { return _size; }
+  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
 
   /**
    * @brief Returns the element type
    */
-  [[nodiscard]] __host__ __device__ data_type type() const noexcept { return _type; }
+  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
 
   /**
    * @brief Indicates whether the column can contain null elements, i.e., if it
@@ -156,7 +156,7 @@ class alignas(16) column_device_view_base {
    * @return true The bitmask is allocated
    * @return false The bitmask is not allocated
    */
-  [[nodiscard]] __host__ __device__ bool nullable() const noexcept { return nullptr != _null_mask; }
+  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
 
   /**
    * @brief Returns raw pointer to the underlying bitmask allocation.
@@ -165,7 +165,7 @@ class alignas(16) column_device_view_base {
    *
    * @note If `null_count() == 0`, this may return `nullptr`.
    */
-  [[nodiscard]] __host__ __device__ bitmask_type const* null_mask() const noexcept
+  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
   {
     return _null_mask;
   }
@@ -174,7 +174,7 @@ class alignas(16) column_device_view_base {
    * @brief Returns the index of the first element relative to the base memory
    * allocation, i.e., what is returned from `head<T>()`.
    */
-  [[nodiscard]] __host__ __device__ size_type offset() const noexcept { return _offset; }
+  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
 
   /**
    * @brief Returns whether the specified element holds a valid value (i.e., not
@@ -269,11 +269,11 @@ class alignas(16) column_device_view_base {
   size_type _offset{};               ///< Index position of the first element.
                                      ///< Enables zero-copy slicing
 
-  column_device_view_base(data_type type,
-                          size_type size,
-                          void const* data,
-                          bitmask_type const* null_mask,
-                          size_type offset)
+  CUDF_HOST_DEVICE column_device_view_base(data_type type,
+                                           size_type size,
+                                           void const* data,
+                                           bitmask_type const* null_mask,
+                                           size_type offset)
     : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset}
   {
   }
@@ -329,6 +329,33 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
    */
   column_device_view(column_view column, void* h_ptr, void* d_ptr);
 
+  /**
+   * @brief Get a new column_device_view which is a slice of this column.
+   *
+   * Example:
+   * @code{.cpp}
+   * // column = column_device_view([1, 2, 3, 4, 5, 6, 7])
+   * auto c = column.slice(1, 3);
+   * // c = column_device_view([2, 3, 4])
+   * auto c1 = column.slice(2, 3);
+   * // c1 = column_device_view([3, 4, 5])
+   * @endcode
+   *
+   * @param offset The index of the first element in the slice
+   * @param size The number of elements in the slice
+   */
+  [[nodiscard]] CUDF_HOST_DEVICE column_device_view slice(size_type offset,
+                                                          size_type size) const noexcept
+  {
+    return column_device_view{this->type(),
+                              size,
+                              this->head(),
+                              this->null_mask(),
+                              this->offset() + offset,
+                              d_children,
+                              this->num_child_columns()};
+  }
+
   /**
    * @brief Returns reference to element at the specified index.
    *
@@ -346,7 +373,7 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
    * @param element_index Position of the desired element
    */
   template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
-  __device__ T element(size_type element_index) const noexcept
+  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
   {
     return data<T>()[element_index];
   }
@@ -365,9 +392,8 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
   template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
   __device__ T element(size_type element_index) const noexcept
   {
-    size_type index = element_index + offset();  // account for this view's _offset
-    const int32_t* d_offsets =
-      d_children[strings_column_view::offsets_column_index].data<int32_t>();
+    size_type index       = element_index + offset();  // account for this view's _offset
+    const auto* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
     const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
     size_type offset      = d_offsets[index];
     return string_view{d_strings + offset, d_offsets[index + 1] - offset};
@@ -763,11 +789,37 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
    *
    * @return The number of child columns
    */
-  [[nodiscard]] __host__ __device__ size_type num_child_columns() const noexcept
+  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
   {
     return _num_children;
   }
 
+ private:
+  /**
+   * @brief Creates an instance of this class using pre-existing device memory pointers to data,
+   * nullmask, and offset.
+   *
+   * @param type The type of the column
+   * @param size The number of elements in the column
+   * @param data Pointer to the device memory containing the data
+   * @param null_mask Pointer to the device memory containing the null bitmask
+   * @param offset The index of the first element in the column
+   * @param children Pointer to the device memory containing child data
+   * @param num_children The number of child columns
+   */
+  CUDF_HOST_DEVICE column_device_view(data_type type,
+                                      size_type size,
+                                      void const* data,
+                                      bitmask_type const* null_mask,
+                                      size_type offset,
+                                      column_device_view* children,
+                                      size_type num_children)
+    : column_device_view_base(type, size, data, null_mask, offset),
+      d_children(children),
+      _num_children(num_children)
+  {
+  }
+
  protected:
   column_device_view* d_children{};  ///< Array of `column_device_view`
                                      ///< objects in device memory.
@@ -852,7 +904,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view
    */
   template <typename T = void,
             CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
-  __host__ __device__ T* head() const noexcept
+  CUDF_HOST_DEVICE T* head() const noexcept
   {
     return const_cast<T*>(detail::column_device_view_base::head<T>());
   }
@@ -870,7 +922,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view
    * @return T* Typed pointer to underlying data, including the offset
    */
   template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
-  __host__ __device__ T* data() const noexcept
+  CUDF_HOST_DEVICE T* data() const noexcept
   {
     return const_cast<T*>(detail::column_device_view_base::data<T>());
   }
@@ -912,7 +964,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view
    *
    * @note If `null_count() == 0`, this may return `nullptr`.
    */
-  [[nodiscard]] __host__ __device__ bitmask_type* null_mask() const noexcept
+  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
   {
     return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
   }

diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
@@ -67,7 +67,8 @@ namespace detail {
  * @return A transform iterator that applies `f` to a counting iterator
  */
 template <typename UnaryFunction>
-inline auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f)
+CUDF_HOST_DEVICE inline auto make_counting_transform_iterator(cudf::size_type start,
+                                                              UnaryFunction f)
 {
   return thrust::make_transform_iterator(thrust::make_counting_iterator(start), f);
 }
@@ -117,26 +118,42 @@ struct null_replaced_value_accessor {
 
 /**
  * @brief validity accessor of column with null bitmask
- * A unary functor returns validity at `id`.
- * `operator() (cudf::size_type id)` computes validity flag at `id`
- * This functor is only allowed for nullable columns.
+ * A unary functor that returns validity at index `i`.
  *
- * @throws cudf::logic_error if the column is not nullable.
+ * @tparam safe If false, the accessor with throw logic_error if the column is not nullable. If
+ * true, the accessor checks for nullability and if col is not nullable, returns true.
  */
+template <bool safe = false>
 struct validity_accessor {
   column_device_view const col;
 
   /**
    * @brief constructor
+   *
+   * @throws cudf::logic_error if not safe and `col` does not have a validity bitmask
+   *
    * @param[in] _col column device view of cudf column
    */
-  validity_accessor(column_device_view const& _col) : col{_col}
+  CUDF_HOST_DEVICE validity_accessor(column_device_view const& _col) : col{_col}
   {
-    // verify valid is non-null, otherwise, is_valid() will crash
-    CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column.");
+    if constexpr (not safe) {
+      // verify col is nullable, otherwise, is_valid_nocheck() will crash
+#if defined(__CUDA_ARCH__)
+      cudf_assert(_col.nullable() && "Unexpected non-nullable column.");
+#else
+      CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column.");
+#endif
+    }
   }
 
-  __device__ inline bool operator()(cudf::size_type i) const { return col.is_valid_nocheck(i); }
+  __device__ inline bool operator()(cudf::size_type i) const
+  {
+    if constexpr (safe) {
+      return col.is_valid(i);
+    } else {
+      return col.is_valid_nocheck(i);
+    }
+  }
 };
 
 /**
@@ -289,16 +306,20 @@ auto make_pair_rep_iterator(column_device_view const& column)
  *
  * Dereferencing the returned iterator for element `i` will return the validity
  * of `column[i]`
- * This iterator is only allowed for nullable columns.
+ * This iterator is only allowed for nullable columns if `safe` = false
+ * When safe = true, if the column is not nullable then the validity is always true.
  *
- * @throws cudf::logic_error if the column is not nullable.
+ * @throws cudf::logic_error if the column is not nullable when safe = false
  *
+ * @tparam safe If false, the accessor with throw logic_error if the column is not nullable. If
+ * true, the accessor checks for nullability and if col is not nullable, returns true.
  * @param column The column to iterate
  * @return auto Iterator that returns validities of column elements.
  */
-auto inline make_validity_iterator(column_device_view const& column)
+template <bool safe = false>
+CUDF_HOST_DEVICE auto inline make_validity_iterator(column_device_view const& column)
 {
-  return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor{column});
+  return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor<safe>{column});
 }
 
 /**