List element Equality comparator

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

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 ---------------------------------------------------------------------------

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
+                  });

cpp/include/cudf/column/column_device_view.cuh

@@ -110,7 +110,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);
   }
@@ -131,20 +131,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
@@ -155,7 +155,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.
@@ -164,7 +164,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;
   }
@@ -173,7 +173,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
@@ -268,11 +268,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}
   {
   }
@@ -328,6 +328,28 @@ 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{
+      _type, size, _data, _null_mask, _offset + offset, d_children, _num_children};
+  }
+
   /**
    * @brief Returns reference to element at the specified index.
    *
@@ -345,7 +367,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];
   }
@@ -364,9 +386,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};
@@ -762,11 +783,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.
@@ -851,7 +898,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>());
   }
@@ -869,7 +916,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>());
   }
@@ -911,7 +958,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());
   }

cpp/include/cudf/detail/iterator.cuh

@@ -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.
@@ -65,7 +65,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);
 }
@@ -115,26 +116,39 @@ 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.
- *
- * @throws cudf::logic_error if the column is not nullable.
+ * A unary functor that returns validity at index `i`.
  */
+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);
+    }
+  }
 };
 
 /**
@@ -287,16 +301,18 @@ 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
  *
  * @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});
 }
 
 /**