Add conversions between column_view and device_span<T const>. (#10302)

This PR adds an implicit conversion operator from `column_view` to `device_span<T const>`. The immediate purpose of this PR is to make it possible to use the API `segmented_reduce(column_view data, device_span<size_type> offsets, ...)` in PR #9621.

This PR also resolves #9656 by adding a `column_view` constructor from `device_span<T const>`.

More broadly, this PR should make it easier to refactor instances where `column.data()` is used with counting iterators to build transform iterators, or other patterns that require a length (e.g. vector factories to copy to host).

Authors:
  - Bradley Dice (https://github.com/bdice)

Approvers:
  - Mark Harris (https://github.com/harrism)
  - Jake Hemstad (https://github.com/jrhemstad)
  - David Wendt (https://github.com/davidwendt)

URL: #10302

cpp/include/cudf/column/column_view.hpp

@@ -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.
@@ -16,8 +16,13 @@
 #pragma once
 
 #include <cudf/types.hpp>
+#include <cudf/utilities/error.hpp>
+#include <cudf/utilities/span.hpp>
 #include <cudf/utilities/traits.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
 
+#include <limits>
+#include <type_traits>
 #include <vector>
 
 /**
@@ -375,6 +380,43 @@ class column_view : public detail::column_view_base {
    */
   auto child_end() const noexcept { return _children.cend(); }
 
+  /**
+   * @brief Construct a column view from a device_span<T>.
+   *
+   * Only numeric and chrono types are supported.
+   *
+   * @tparam T The device span type. Must be const and match the column view's type.
+   * @param data A typed device span containing the column view's data.
+   */
+  template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() or cudf::is_chrono<T>())>
+  column_view(device_span<T const> data)
+    : column_view(
+        cudf::data_type{cudf::type_to_id<T>()}, data.size(), data.data(), nullptr, 0, 0, {})
+  {
+    CUDF_EXPECTS(data.size() < std::numeric_limits<cudf::size_type>::max(),
+                 "Data exceeds the maximum size of a column view.");
+  }
+
+  /**
+   * @brief Converts a column view into a device span.
+   *
+   * Only numeric and chrono data types are supported. The column view must not
+   * be nullable.
+   *
+   * @tparam T The device span type. Must be const and match the column view's type.
+   * @throws cudf::logic_error if the column view type does not match the span type.
+   * @throws cudf::logic_error if the column view is nullable.
+   * @return A typed device span of the column view's data.
+   */
+  template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() or cudf::is_chrono<T>())>
+  [[nodiscard]] operator device_span<T const>() const
+  {
+    CUDF_EXPECTS(type() == cudf::data_type{cudf::type_to_id<T>()},
+                 "Device span type must match column view type.");
+    CUDF_EXPECTS(!nullable(), "A nullable column view cannot be converted to a device span.");
+    return device_span<T const>(data<T>(), size());
+  }
+
  private:
   friend column_view bit_cast(column_view const& input, data_type type);
 

cpp/include/cudf/utilities/span.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-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.

cpp/tests/CMakeLists.txt

@@ -48,8 +48,13 @@ endfunction()
 # ##################################################################################################
 # * column tests ----------------------------------------------------------------------------------
 ConfigureTest(
-  COLUMN_TEST column/bit_cast_test.cpp column/column_view_shallow_test.cpp column/column_test.cu
-  column/column_device_view_test.cu column/compound_test.cu
+  COLUMN_TEST
+  column/bit_cast_test.cpp
+  column/column_device_view_test.cu
+  column/column_test.cu
+  column/column_view_device_span_test.cpp
+  column/column_view_shallow_test.cpp
+  column/compound_test.cu
 )
 
 # ##################################################################################################

cpp/tests/column/column_view_device_span_test.cpp

@@ -0,0 +1,75 @@
+/*
+ * 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 <cudf/column/column_view.hpp>
+#include <cudf/null_mask.hpp>
+#include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
+#include <cudf/utilities/traits.hpp>
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_utilities.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/cudf_gtest.hpp>
+#include <cudf_test/type_lists.hpp>
+
+#include <thrust/iterator/counting_iterator.h>
+
+#include <memory>
+
+template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() or cudf::is_chrono<T>())>
+std::unique_ptr<cudf::column> example_column()
+{
+  auto begin = thrust::make_counting_iterator(1);
+  auto end   = thrust::make_counting_iterator(16);
+  return cudf::test::fixed_width_column_wrapper<T>(begin, end).release();
+}
+
+template <typename T>
+struct ColumnViewDeviceSpanTests : public cudf::test::BaseFixture {
+};
+
+using DeviceSpanTypes = cudf::test::FixedWidthTypesWithoutFixedPoint;
+TYPED_TEST_SUITE(ColumnViewDeviceSpanTests, DeviceSpanTypes);
+
+TYPED_TEST(ColumnViewDeviceSpanTests, conversion_round_trip)
+{
+  auto col      = example_column<TypeParam>();
+  auto col_view = cudf::column_view{*col};
+
+  // Test implicit conversion, round trip
+  cudf::device_span<const TypeParam> device_span_from_col_view = col_view;
+  cudf::column_view col_view_from_device_span                  = device_span_from_col_view;
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(col_view, col_view_from_device_span);
+}
+
+struct ColumnViewDeviceSpanErrorTests : public cudf::test::BaseFixture {
+};
+
+TEST_F(ColumnViewDeviceSpanErrorTests, type_mismatch)
+{
+  auto col      = example_column<int32_t>();
+  auto col_view = cudf::column_view{*col};
+  EXPECT_THROW((void)cudf::device_span<const float>{col_view}, cudf::logic_error);
+}
+
+TEST_F(ColumnViewDeviceSpanErrorTests, nullable_column)
+{
+  auto col = example_column<int32_t>();
+  col->set_null_mask(cudf::create_null_mask(col->size(), cudf::mask_state::ALL_NULL), col->size());
+  auto col_view = cudf::column_view{*col};
+  EXPECT_THROW((void)cudf::device_span<const int32_t>{col_view}, cudf::logic_error);
+}