[REVIEW] Initial support for struct columns

cpp/CMakeLists.txt

@@ -562,6 +562,8 @@ add_library(cudf
             src/lists/lists_column_view.cu
             src/lists/copying/concatenate.cu
             src/lists/copying/gather.cu
+            src/structs/structs_column_view.cu
+            src/structs/structs_column_factories.cu
             src/text/detokenize.cu
             src/text/generate_ngrams.cu
             src/text/normalize.cu

cpp/include/cudf/column/column_device_view.cuh

@@ -19,6 +19,7 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <cudf/column/column_view.hpp>
 #include <cudf/lists/list_view.cuh>
+#include <cudf/structs/struct_view.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/strings/strings_column_view.hpp>
 #include <cudf/types.hpp>

cpp/include/cudf/column/column_factories.hpp

@@ -502,6 +502,14 @@ std::unique_ptr<cudf::column> make_lists_column(
   cudaStream_t stream                 = 0,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
 
+std::unique_ptr<cudf::column> make_structs_column(
+  size_type num_rows,
+  std::vector<std::unique_ptr<column>>&& child_column,
+  size_type null_count,
+  rmm::device_buffer&& null_mask,
+  cudaStream_t stream                 = 0,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+
 /**
  * @brief Return a column with size elements that are all equal to the
  * given scalar.

cpp/include/cudf/column/column_view.hpp

@@ -335,6 +335,9 @@ class column_view : public detail::column_view_base {
    **/
   size_type num_children() const noexcept { return _children.size(); }
 
+  auto child_begin() const noexcept { return _children.begin(); }
+  auto child_end() const noexcept { return _children.end(); }
+
  private:
   std::vector<column_view> _children{};  ///< Based on element type, children
                                          ///< may contain additional data

cpp/include/cudf/detail/gather.cuh

@@ -357,6 +357,19 @@ struct column_gatherer_impl<list_view, MapItRoot> {
   }
 };
 
+template <typename MapItRoot>
+struct column_gatherer_impl<struct_view, MapItRoot> {
+  std::unique_ptr<column> operator()(column_view const& column,
+                                    MapItRoot gather_map_begin,
+                                    MapItRoot gather_map_end,
+                                    bool nullify_out_of_bounds,
+                                    cudaStream_t stream,
+                                    rmm::mr::device_memory_resource* mr)
+  {
+    CUDF_FAIL("Gather not yet supported on struct_view.");
+  }
+};
+
 /**
  * @brief Function object for gathering a type-erased
  * column. To be used with the cudf::type_dispatcher.

cpp/include/cudf/detail/null_mask.hpp

@@ -39,6 +39,21 @@ std::vector<size_type> segmented_count_unset_bits(bitmask_type const* bitmask,
                                                   std::vector<size_type> const& indices,
                                                   cudaStream_t stream = 0);
 
+/**
+ * @brief Returns a bitwise AND of the specified bitmasks
+ * 
+ * @param masks The list of data pointers of the bitmasks to be ANDed
+ * @param begin_bits The bit offsets from which each mask is to be ANDed
+ * @param mask_size The number of bits to be ANDed in each mask
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @param mr Device memory resource used to allocate the returned device_buffer
+ * @return rmm::device_buffer Output bitmask 
+ */
+rmm::device_buffer bitmask_and(std::vector<bitmask_type const *> const &masks,
+                               std::vector<size_type> const &begin_bits,
+                               size_type mask_size,
+                               cudaStream_t stream,
+                               rmm::mr::device_memory_resource *mr);
 }  // namespace detail
 
 }  // namespace cudf

cpp/include/cudf/structs/struct_view.hpp

@@ -0,0 +1,35 @@
+/*
+ * Copyright (c) 2020, 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
+
+/**
+ * @file struct_view.cuh
+ * @brief Class definition for cudf::struct_view.
+ */
+
+namespace cudf {
+
+/**
+ * @brief A non-owning, immutable view of device data that represents
+ * a struct with fields of arbitrary types (including primitives, lists, 
+ * and other structs)
+ *
+ */
+class struct_view {
+
+};
+
+}  // namespace cudf

cpp/include/cudf/structs/structs_column_view.hpp

@@ -0,0 +1,48 @@
+/*
+ * Copyright (c) 2020, 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.hpp>
+#include <cudf/column/column_view.hpp>
+
+namespace cudf {
+
+
+class structs_column_view : private column_view 
+{
+
+    public:
+
+        // Foundation members:
+        structs_column_view(structs_column_view const&) = default;
+        structs_column_view(structs_column_view &&) = default;
+        ~structs_column_view() = default;
+        structs_column_view& operator=(structs_column_view const&) = default;
+        structs_column_view& operator=(structs_column_view &&) = default;
+
+        explicit structs_column_view(column_view const& rhs);
+
+        using column_view::has_nulls;
+        using column_view::null_count;
+        using column_view::null_mask;
+        using column_view::offset;
+        using column_view::size;
+        using column_view::child_begin;
+        using column_view::child_end;
+
+}; // class structs_column_view;
+
+} // namespace cudf;

cpp/include/cudf/types.hpp

@@ -59,6 +59,7 @@ class column_view;
 class mutable_column_view;
 class string_view;
 class list_view;
+class struct_view;
 
 class scalar;
 template <typename T>
@@ -79,6 +80,8 @@ class duration_scalar_device_view;
 
 class list_scalar;
 
+class struct_scalar;
+
 class table;
 class table_view;
 class mutable_table_view;
@@ -207,6 +210,7 @@ enum class type_id : int32_t {
   DICTIONARY32,            ///< Dictionary type using int32 indices
   STRING,                  ///< String elements
   LIST,                    ///< List elements
+  STRUCT,                  ///< Struct elements
   // `NUM_TYPE_IDS` must be last!
   NUM_TYPE_IDS  ///< Total number of type ids
 };

cpp/include/cudf/utilities/traits.hpp

@@ -24,6 +24,7 @@
 #include <cudf/wrappers/timestamps.hpp>
 
 #include <type_traits>
+#include "cudf/structs/struct_view.hpp"
 
 namespace cudf {
 
@@ -455,7 +456,7 @@ template <typename T>
 constexpr inline bool is_compound()
 {
   return std::is_same<T, cudf::string_view>::value or std::is_same<T, cudf::dictionary32>::value or
-         std::is_same<T, cudf::list_view>::value;
+         std::is_same<T, cudf::list_view>::value or std::is_same<T, cudf::struct_view>::value;
 }
 
 struct is_compound_impl {
@@ -497,7 +498,8 @@ constexpr inline bool is_compound(data_type type)
 template <typename T>
 constexpr inline bool is_nested()
 {
-  return std::is_same<T, cudf::list_view>::value;
+  return std::is_same<T, cudf::list_view>::value 
+      || std::is_same<T, cudf::struct_view>::value;
 }
 
 struct is_nested_impl {

cpp/include/cudf/utilities/type_dispatcher.hpp

@@ -132,6 +132,7 @@ CUDF_TYPE_MAPPING(cudf::duration_us, type_id::DURATION_MICROSECONDS);
 CUDF_TYPE_MAPPING(cudf::duration_ns, type_id::DURATION_NANOSECONDS);
 CUDF_TYPE_MAPPING(dictionary32, type_id::DICTIONARY32);
 CUDF_TYPE_MAPPING(cudf::list_view, type_id::LIST);
+CUDF_TYPE_MAPPING(cudf::struct_view, type_id::STRUCT);
 
 template <typename T>
 struct type_to_scalar_type_impl {
@@ -184,6 +185,12 @@ struct type_to_scalar_type_impl<cudf::list_view> {
   // using ScalarDeviceType = cudf::list_scalar_device_view;
 };
 
+template <>  // TODO: Ditto, likewise.
+struct type_to_scalar_type_impl<cudf::struct_view> {
+  using ScalarType = cudf::struct_scalar;
+  // using ScalarDeviceType = cudf::struct_scalar_device_view; // CALEB: TODO!
+};
+
 #ifndef MAP_TIMESTAMP_SCALAR
 #define MAP_TIMESTAMP_SCALAR(Type)                                     \
   template <>                                                          \
@@ -400,6 +407,9 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty
     case type_id::LIST:
       return f.template operator()<typename IdTypeMap<type_id::LIST>::type>(
         std::forward<Ts>(args)...);
+    case type_id::STRUCT:
+      return f.template operator()<typename IdTypeMap<type_id::STRUCT>::type>(
+        std::forward<Ts>(args)...);
     default: {
 #ifndef __CUDA_ARCH__
       CUDF_FAIL("Unsupported type_id.");

cpp/src/bitmask/null_mask.cu

@@ -341,6 +341,36 @@ __global__ void offset_bitmask_and(bitmask_type *__restrict__ destination,
   }
 }
 
+// convert [first_bit_index,last_bit_index) to
+// [first_word_index,last_word_index)
+struct to_word_index : public thrust::unary_function<size_type, size_type> {
+  const bool _inclusive                 = false;
+  size_type const *const _d_bit_indices = nullptr;
+
+  /**
+   * @brief Constructor of a functor that converts bit indices to bitmask word
+   * indices.
+   *
+   * @param[in] inclusive Flag that indicates whether bit indices are inclusive
+   * or exclusive.
+   * @param[in] d_bit_indices Pointer to an array of bit indices
+   */
+  __host__ to_word_index(bool inclusive, size_type const *d_bit_indices)
+    : _inclusive(inclusive), _d_bit_indices(d_bit_indices)
+  {
+  }
+
+  __device__ size_type operator()(const size_type &i) const
+  {
+    auto bit_index = _d_bit_indices[i];
+    return word_index(bit_index) + ((_inclusive || intra_word_index(bit_index) == 0) ? 0 : 1);
+  }
+};
+
+}  // namespace
+
+namespace detail {
+
 // Bitwise AND of the masks
 rmm::device_buffer bitmask_and(std::vector<bitmask_type const *> const &masks,
                                std::vector<size_type> const &begin_bits,
@@ -378,35 +408,6 @@ rmm::device_buffer bitmask_and(std::vector<bitmask_type const *> const &masks,
   return dest_mask;
 }
 
-// convert [first_bit_index,last_bit_index) to
-// [first_word_index,last_word_index)
-struct to_word_index : public thrust::unary_function<size_type, size_type> {
-  const bool _inclusive                 = false;
-  size_type const *const _d_bit_indices = nullptr;
-
-  /**
-   * @brief Constructor of a functor that converts bit indices to bitmask word
-   * indices.
-   *
-   * @param[in] inclusive Flag that indicates whether bit indices are inclusive
-   * or exclusive.
-   * @param[in] d_bit_indices Pointer to an array of bit indices
-   */
-  __host__ to_word_index(bool inclusive, size_type const *d_bit_indices)
-    : _inclusive(inclusive), _d_bit_indices(d_bit_indices)
-  {
-  }
-
-  __device__ size_type operator()(const size_type &i) const
-  {
-    auto bit_index = _d_bit_indices[i];
-    return word_index(bit_index) + ((_inclusive || intra_word_index(bit_index) == 0) ? 0 : 1);
-  }
-};
-
-}  // namespace
-
-namespace detail {
 cudf::size_type count_set_bits(bitmask_type const *bitmask,
                                size_type start,
                                size_type stop,
@@ -662,7 +663,7 @@ rmm::device_buffer bitmask_and(table_view const &view,
     }
   }
 
-  if (masks.size() > 0) { return bitmask_and(masks, offsets, view.num_rows(), stream, mr); }
+  if (masks.size() > 0) { return cudf::detail::bitmask_and(masks, offsets, view.num_rows(), stream, mr); }
 
   return null_mask;
 }

cpp/src/column/column.cu

@@ -235,6 +235,13 @@ struct create_column_from_view {
   {
     CUDF_FAIL("list_view not supported yet");
   }
+
+  template <typename ColumnType,
+            std::enable_if_t<std::is_same<ColumnType, cudf::struct_view>::value> * = nullptr>
+  std::unique_ptr<column> operator()()
+  {
+    CUDF_FAIL("struct_view not supported yet");
+  }
 };
 }  // anonymous namespace
 

cpp/src/column/column_factories.cpp

@@ -190,6 +190,16 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::list
   CUDF_FAIL("TODO");
 }
 
+template <>
+std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::struct_view>(
+  scalar const& value,
+  size_type size,
+  rmm::mr::device_memory_resource* mr,
+  cudaStream_t stream) const
+{
+  CUDF_FAIL("TODO. struct_view currently not supported.");
+}
+
 std::unique_ptr<column> make_column_from_scalar(scalar const& s,
                                                 size_type size,
                                                 rmm::mr::device_memory_resource* mr,

cpp/src/copying/copy.cu

@@ -115,6 +115,21 @@ struct copy_if_else_functor_impl<list_view, Left, Right, Filter> {
   }
 };
 
+template <typename Left, typename Right, typename Filter>
+struct copy_if_else_functor_impl<struct_view, Left, Right, Filter> {
+  std::unique_ptr<column> operator()(Left const& lhs,
+                                     Right const& rhs,
+                                     size_type size,
+                                     bool left_nullable,
+                                     bool right_nullable,
+                                     Filter filter,
+                                     rmm::mr::device_memory_resource* mr,
+                                     cudaStream_t stream)
+  {
+    CUDF_FAIL("copy_if_else not supported for struct_view yet");
+  }
+};
+
 /**
  * @brief Functor called by the `type_dispatcher` to invoke copy_if_else on combinations
  *        of column_view and scalar