Support create lists column from a list_scalar

conda/recipes/libcudf/meta.yaml

@@ -136,6 +136,7 @@ test:
     - test -f $PREFIX/include/cudf/join.hpp
     - test -f $PREFIX/include/cudf/lists/detail/concatenate.hpp
     - test -f $PREFIX/include/cudf/lists/detail/copying.hpp
+    - test -f $PREFIX/include/cudf/lists/lists_column_factories.hpp
     - test -f $PREFIX/include/cudf/lists/detail/drop_list_duplicates.hpp
     - test -f $PREFIX/include/cudf/lists/detail/interleave_columns.hpp
     - test -f $PREFIX/include/cudf/lists/detail/sorting.hpp

cpp/include/cudf/column/column_factories.hpp

@@ -541,7 +541,8 @@ std::unique_ptr<cudf::column> make_structs_column(
  *
  * The output column will have the same type as `s.type()`
  * The output column will contain all null rows if `s.invalid()==false`
- * The output column will be empty if `size==0`.
+ * The output column will be empty if `size==0`. For LIST scalars, the column hierarchy
+ * from @p s is preserved.
  *
  * @param[in] s The scalar to use for values in the column.
  * @param[in] size The number of rows for the output column.

cpp/include/cudf/lists/lists_column_factories.hpp

@@ -0,0 +1,42 @@
+/*
+ * 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.hpp>
+#include <cudf/scalar/scalar.hpp>
+#include <cudf/types.hpp>
+
+namespace cudf {
+namespace lists {
+namespace detail {
+
+/**
+ * @brief Internal API to construct a lists column from a `list_scalar`, for public
+ * use, use `cudf::make_column_from_scalar`.
+ *
+ * @param[in] value The `list_scalar` to construct from
+ * @param[in] size The number of rows for the output column.
+ * @param[in] stream CUDA stream used for device memory operations and kernel launches.
+ * @param[in] mr Device memory resource used to allocate the returned column's device memory.
+ */
+std::unique_ptr<cudf::column> make_lists_column_from_scalar(
+  list_scalar const& value,
+  size_type size,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+}  // namespace detail
+}  // namespace lists
+}  // namespace cudf

cpp/src/column/column_factories.cpp

@@ -20,6 +20,7 @@
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/dictionary/dictionary_factories.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
+#include <cudf/lists/lists_column_factories.hpp>
 #include <cudf/scalar/scalar.hpp>
 #include <cudf/scalar/scalar_factories.hpp>
 #include <cudf/strings/detail/fill.hpp>
@@ -166,6 +167,7 @@ struct column_from_scalar_dispatch {
                                            rmm::cuda_stream_view stream,
                                            rmm::mr::device_memory_resource* mr) const
   {
+    if (size == 0) {return make_empty_column(value.type());}
     if (!value.is_valid())
       return make_fixed_width_column(value.type(), size, mask_state::ALL_NULL, stream, mr);
     auto output_column =
@@ -183,6 +185,7 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::stri
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr) const
 {
+  if (size == 0) {return make_empty_column(value.type());}
   auto null_mask = detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr);
 
   if (!value.is_valid())
@@ -221,7 +224,8 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::list
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr) const
 {
-  CUDF_FAIL("TODO");
+  auto lv = static_cast<list_scalar const*>(&value);
+  return lists::detail::make_lists_column_from_scalar(*lv, size, stream, mr);
 }
 
 template <>
@@ -239,7 +243,6 @@ std::unique_ptr<column> make_column_from_scalar(scalar const& s,
                                                 rmm::cuda_stream_view stream,
                                                 rmm::mr::device_memory_resource* mr)
 {
-  if (size == 0) return make_empty_column(s.type());
   return type_dispatcher(s.type(), column_from_scalar_dispatch{}, s, size, stream, mr);
 }
 

cpp/src/lists/lists_column_factories.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.
@@ -16,10 +16,68 @@
 
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/gather.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/sequence.h>
 
 namespace cudf {
+namespace lists {
+namespace detail {
+
+std::unique_ptr<cudf::column> make_lists_column_from_scalar(list_scalar const& value,
+                                                            size_type size,
+                                                            rmm::cuda_stream_view stream,
+                                                            rmm::mr::device_memory_resource* mr)
+{
+  if (size == 0) {
+    return make_lists_column(0,
+                             make_empty_column(data_type{type_to_id<offset_type>()}),
+                             empty_like(value.view()),
+                             0,
+                             cudf::create_null_mask(0, mask_state::UNALLOCATED),
+                             stream,
+                             mr);
+  }
+  auto mr_final = size == 1 ? mr : rmm::mr::get_current_device_resource();
+
+  // Handcraft a 1-row column
+  auto offsets = make_numeric_column(
+    data_type{type_to_id<offset_type>()}, 2, mask_state::UNALLOCATED, stream, mr_final);
+  auto m_offsets = offsets->mutable_view();
+  thrust::sequence(rmm::exec_policy(stream),
+                   m_offsets.begin<size_type>(),
+                   m_offsets.end<size_type>(),
+                   0,
+                   value.view().size());
+  auto child           = std::make_unique<column>(value.view(), stream, mr_final);
+  size_type null_count = value.is_valid(stream) ? 0 : 1;
+  auto null_mask_state = null_count ? mask_state::ALL_NULL : mask_state::UNALLOCATED;
+  auto null_mask       = cudf::detail::create_null_mask(1, null_mask_state, stream, mr_final);
+
+  if (size == 1) {
+    return make_lists_column(
+      1, std::move(offsets), std::move(child), null_count, std::move(null_mask), stream, mr);
+  }
+
+  auto one_row_col = make_lists_column(
+    1, std::move(offsets), std::move(child), null_count, std::move(null_mask), stream);
+
+  auto begin = thrust::make_constant_iterator(0);
+  auto res   = cudf::detail::gather(table_view({one_row_col->view()}),
+                                  begin,
+                                  begin + size,
+                                  out_of_bounds_policy::DONT_CHECK,
+                                  stream,
+                                  mr);
+  return std::move(res->release()[0]);
+}
+
+}  // namespace detail
+}  // namespace lists
 
 /**
  * @copydoc cudf::make_lists_column