Support get_element from LIST column

conda/recipes/libcudf/meta.yaml

@@ -78,6 +78,7 @@ test:
     - test -f $PREFIX/include/cudf/detail/groupby/sort_helper.hpp
     - test -f $PREFIX/include/cudf/detail/hashing.hpp
     - test -f $PREFIX/include/cudf/detail/interop.hpp
+    - test -f $PREFIX/include/cudf/detail/is_element_valid.hpp
     - test -f $PREFIX/include/cudf/detail/null_mask.hpp
     - test -f $PREFIX/include/cudf/detail/nvtx/nvtx3.hpp
     - test -f $PREFIX/include/cudf/detail/nvtx/ranges.hpp

cpp/CMakeLists.txt

@@ -155,6 +155,7 @@ add_library(cudf
     src/binaryop/compiled/binary_ops.cu
     src/labeling/label_bins.cu
     src/bitmask/null_mask.cu
+    src/bitmask/is_element_valid.cpp
     src/column/column.cu
     src/column/column_device_view.cu
     src/column/column_factories.cpp

cpp/docs/DEVELOPER_GUIDE.md

@@ -255,6 +255,11 @@ currently supported by cudf. Each type of value is represented by a separate typ
 which are all derived from `cudf::scalar`. e.g. A `numeric_scalar` holds a single numerical value, 
 a `string_scalar` holds a single string. The data for the stored value resides in device memory.
 
+A `list_scalar` holds the underlying data of a single list. This means the underlying data can be any type
+that cudf supports. For example, a `list_scalar` representing a list of integers stores a `cudf::column`
+of type `INT32`. A `list_scalar` representing a list of lists of integers stores a `cudf::column` of
+type `LIST`, which in turn stores a column of type `INT32`.
+
 |Value type|Scalar class|Notes|
 |-|-|-|
 |fixed-width|`fixed_width_scalar<T>`| `T` can be any fixed-width type|
@@ -263,6 +268,7 @@ a `string_scalar` holds a single string. The data for the stored value resides i
 |timestamp|`timestamp_scalar<T>` | `T` can be `timestamp_D`, `timestamp_s`, etc.|
 |duration|`duration_scalar<T>` | `T` can be `duration_D`, `duration_s`, etc.|
 |string|`string_scalar`| This class object is immutable|
+|list|`list_scalar`| Underlying data can be any type supported by cudf |
 
 ### Construction
 `scalar`s can be created using either their respective constructors or using factory functions like 
@@ -285,11 +291,16 @@ auto s1 = static_cast<ScalarType *>(s.get());
 ```
 
 ### Passing to device
-Each scalar type has a corresponding non-owning device view class which allows access to the value 
-and its validity from the device. This can be obtained using the function 
+Each scalar type, except `list_scalar`, has a corresponding non-owning device view class which allows
+access to the value and its validity from the device. This can be obtained using the function 
 `get_scalar_device_view(ScalarType s)`. Note that a device view is not provided for a base scalar 
 object, only for the derived typed scalar class objects.
 
+The underlying data for `list_scalar` can be accessed via `view()` method. For non-nested data,
+the device view can be obtained via function `column_device_view::create(column_view)`. For nested
+data, a specialized device view for list columns can be constructed via
+`lists_column_device_view(column_device_view)`.
+
 # libcudf++ API and Implementation
 
 ## Streams

cpp/include/cudf/detail/is_element_valid.hpp

@@ -0,0 +1,46 @@
+/*
+ * Copyright (c) 2021, 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/column/column_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+namespace cudf {
+namespace detail {
+
+/**
+ * @brief Return validity of a row
+ *
+ * Retrieves the validity (NULL or non-NULL) of the specified row from device memory.
+ *
+ * @note Synchronizes `stream`.
+ *
+ * @throw cudf::logic_error if `element_index < 0 or >= col_view.size()`
+ *
+ * @param col_view The column to retrieve the validity from.
+ * @param element_index The index of the row to retrieve.
+ * @param stream The stream to use for copying the validity to the host.
+ * @return Host boolean that indicates the validity of the row.
+ */
+
+bool is_element_valid_sync(column_view const& col_view,
+                           size_type element_index,
+                           rmm::cuda_stream_view stream = rmm::cuda_stream_default);
+
+}  // namespace detail
+}  // namespace cudf

cpp/include/cudf/scalar/scalar_factories.hpp

@@ -156,5 +156,17 @@ std::unique_ptr<scalar> make_fixed_point_scalar(
   return std::make_unique<scalar_type_t<T>>(value, scale, true, stream, mr);
 }
 
+/**
+ * @brief Construct scalar using the given value of fixed_point type
+ *
+ * @param elements Elements of the list
+ * @param stream CUDA stream used for device memory operations.
+ * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool.
+ */
+std::unique_ptr<scalar> make_list_scalar(
+  column_view elements,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /** @} */  // end of group
 }  // namespace cudf

cpp/src/bitmask/is_element_valid.cpp

@@ -0,0 +1,47 @@
+
+/*
+ * Copyright (c) 2021, 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/column/column_view.hpp>
+#include <cudf/utilities/bit.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+namespace cudf {
+namespace detail {
+
+bool is_element_valid_sync(column_view const& col_view,
+                           size_type element_index,
+                           rmm::cuda_stream_view stream)
+{
+  CUDF_EXPECTS(element_index >= 0 and element_index < col_view.size(), "invalid index.");
+  if (!col_view.nullable()) { return true; }
+
+  bitmask_type word;
+  // null_mask() returns device ptr to bitmask without offset
+  size_type index = element_index + col_view.offset();
+  CUDA_TRY(cudaMemcpyAsync(&word,
+                           col_view.null_mask() + word_index(index),
+                           sizeof(bitmask_type),
+                           cudaMemcpyDeviceToHost,
+                           stream.value()));
+  stream.synchronize();
+  return static_cast<bool>(word & (bitmask_type{1} << intra_word_index(index)));
+}
+
+}  // namespace detail
+}  // namespace cudf

cpp/src/copying/get_element.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,8 +17,11 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/copying.hpp>
 #include <cudf/detail/indexalator.cuh>
+#include <cudf/detail/is_element_valid.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/lists/detail/copying.hpp>
+#include <cudf/lists/lists_column_view.hpp>
 #include <cudf/scalar/scalar_device_view.cuh>
 #include <cudf/scalar/scalar_factories.hpp>
 
@@ -122,7 +125,19 @@ struct get_element_functor {
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource *mr = rmm::mr::get_current_device_resource())
   {
-    CUDF_FAIL("get_element_functor not supported for list_view");
+    bool valid = is_element_valid_sync(input, index, stream);
+
+    if (valid) {
+      lists_column_view lcv(input);
+      // Make a copy of the row
+      auto row_slice_contents =
+        lists::detail::copy_slice(lcv, index, index + 1, stream, mr)->release();
+      // Construct scalar with row data
+      return std::make_unique<list_scalar>(
+        std::move(*row_slice_contents.children[1]), valid, stream, mr);
+    } else {
+      return make_default_constructed_scalar(data_type(type_id::LIST));
+    }
   }
 
   template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()> *p = nullptr>

cpp/src/scalar/scalar.cpp

@@ -440,7 +440,7 @@ list_scalar::list_scalar(cudf::column&& data,
                          bool is_valid,
                          rmm::cuda_stream_view stream,
                          rmm::mr::device_memory_resource* mr)
-  : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(std::move(data))
+  : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(data)
 {
 }
 

cpp/src/scalar/scalar_factories.cpp

@@ -97,6 +97,13 @@ std::unique_ptr<scalar> make_fixed_width_scalar(data_type type,
   return type_dispatcher(type, scalar_construction_helper{}, stream, mr);
 }
 
+std::unique_ptr<scalar> make_list_scalar(column_view elements,
+                                         rmm::cuda_stream_view stream,
+                                         rmm::mr::device_memory_resource* mr)
+{
+  return std::make_unique<list_scalar>(elements, true, stream, mr);
+}
+
 namespace {
 struct default_scalar_functor {
   template <typename T>
@@ -125,7 +132,7 @@ template <>
 std::unique_ptr<cudf::scalar> default_scalar_functor::operator()<list_view>(
   rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr)
 {
-  CUDF_FAIL("list_view type not supported");
+  return std::make_unique<list_scalar>(column(), false, stream, mr);
 }
 
 template <>

cpp/tests/CMakeLists.txt

@@ -368,7 +368,8 @@ ConfigureTest(TEXT_TEST
 ConfigureTest(BITMASK_TEST
     bitmask/valid_if_tests.cu
     bitmask/set_nullmask_tests.cu
-    bitmask/bitmask_tests.cu)
+    bitmask/bitmask_tests.cu
+    bitmask/is_element_valid_tests.cpp)
 
 
 ###################################################################################################

cpp/tests/bitmask/is_element_valid_tests.cpp

@@ -0,0 +1,92 @@
+/*
+ * Copyright (c) 2021, 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_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
+
+#include <cudf/copying.hpp>
+#include <cudf/detail/is_element_valid.hpp>
+#include <cudf/detail/iterator.cuh>
+
+#include <thrust/iterator/counting_iterator.h>
+
+namespace cudf {
+namespace test {
+
+struct IsElementValidTest : public BaseFixture {
+};
+
+TEST_F(IsElementValidTest, IsElementValidBasic)
+{
+  fixed_width_column_wrapper<int32_t> col({1, 1, 1, 1, 1}, {1, 0, 0, 0, 1});
+  EXPECT_TRUE(cudf::detail::is_element_valid_sync(col, 0));
+  EXPECT_FALSE(cudf::detail::is_element_valid_sync(col, 1));
+  EXPECT_FALSE(cudf::detail::is_element_valid_sync(col, 2));
+  EXPECT_FALSE(cudf::detail::is_element_valid_sync(col, 3));
+  EXPECT_TRUE(cudf::detail::is_element_valid_sync(col, 4));
+}
+
+TEST_F(IsElementValidTest, IsElementValidLarge)
+{
+  auto filter        = [](auto i) { return static_cast<bool>(i % 3); };
+  auto val           = thrust::make_counting_iterator(0);
+  auto valid         = cudf::detail::make_counting_transform_iterator(0, filter);
+  size_type num_rows = 1000;
+
+  fixed_width_column_wrapper<int32_t> col(val, val + num_rows, valid);
+
+  for (int i = 0; i < num_rows; i++) {
+    EXPECT_EQ(cudf::detail::is_element_valid_sync(col, i), filter(i));
+  }
+}
+
+TEST_F(IsElementValidTest, IsElementValidOffset)
+{
+  fixed_width_column_wrapper<int32_t> col({1, 1, 1, 1, 1}, {1, 0, 0, 0, 1});
+  {
+    auto offset_col = slice(col, {1, 5}).front();
+    EXPECT_FALSE(cudf::detail::is_element_valid_sync(offset_col, 0));
+    EXPECT_FALSE(cudf::detail::is_element_valid_sync(offset_col, 1));
+    EXPECT_FALSE(cudf::detail::is_element_valid_sync(offset_col, 2));
+    EXPECT_TRUE(cudf::detail::is_element_valid_sync(offset_col, 3));
+  }
+  {
+    auto offset_col = slice(col, {2, 5}).front();
+    EXPECT_FALSE(cudf::detail::is_element_valid_sync(offset_col, 0));
+    EXPECT_FALSE(cudf::detail::is_element_valid_sync(offset_col, 1));
+    EXPECT_TRUE(cudf::detail::is_element_valid_sync(offset_col, 2));
+  }
+}
+
+TEST_F(IsElementValidTest, IsElementValidOffsetLarge)
+{
+  auto filter        = [](auto i) { return static_cast<bool>(i % 3); };
+  size_type offset   = 37;
+  auto val           = thrust::make_counting_iterator(0);
+  auto valid         = cudf::detail::make_counting_transform_iterator(0, filter);
+  size_type num_rows = 1000;
+
+  fixed_width_column_wrapper<int32_t> col(val, val + num_rows, valid);
+  auto offset_col = slice(col, {offset, num_rows}).front();
+
+  for (int i = 0; i < offset_col.size(); i++) {
+    EXPECT_EQ(cudf::detail::is_element_valid_sync(offset_col, i), filter(i + offset));
+  }
+}
+
+}  // namespace test
+
+}  // namespace cudf