Move table_device_view function definitions from .cuh to .cu

cpp/include/cudf/table/table_device_view.cuh

@@ -131,15 +131,13 @@ class table_device_view_base {
 
  protected:
   /**
-   * @brief Construct a new table device view base object from host table_view
+   * @brief Constructor for subclasses to create a table device view from a host table view
    *
    * @param source_view The host table_view to create table device view from
-   * @param stream The CUDA stream to use for device memory allocation
+   * @param columns Pointer to the array of column device views in device memory corresponding to
+   * the columns of `source_view`
    */
-  table_device_view_base(HostTableView source_view, rmm::cuda_stream_view stream);
-
-  /// Pointer to device memory holding the descendant storage
-  rmm::device_buffer* _descendant_storage{};
+  table_device_view_base(HostTableView source_view, ColumnDeviceView* columns);
 };
 }  // namespace detail
 
@@ -162,19 +160,11 @@ class table_device_view : public detail::table_device_view_base<column_device_vi
    * @return A `unique_ptr` to a `table_device_view` that makes the data from `source_view`
    * available in device memory
    */
-  static auto create(table_view source_view,
-                     rmm::cuda_stream_view stream = cudf::get_default_stream())
-  {
-    auto deleter = [](table_device_view* t) { t->destroy(); };
-    return std::unique_ptr<table_device_view, decltype(deleter)>{
-      new table_device_view(source_view, stream), deleter};
-  }
+  static std::unique_ptr<table_device_view, std::function<void(table_device_view*)>> create(
+    table_view source_view, rmm::cuda_stream_view stream = cudf::get_default_stream());
 
  private:
-  table_device_view(table_view source_view, rmm::cuda_stream_view stream)
-    : detail::table_device_view_base<column_device_view, table_view>(source_view, stream)
-  {
-  }
+  table_device_view(table_view source_view, column_device_view* columns);
 };
 
 /**
@@ -199,20 +189,11 @@ class mutable_table_device_view
    * @return A `unique_ptr` to a `mutable_table_device_view` that makes the data from `source_view`
    * available in device memory
    */
-  static auto create(mutable_table_view source_view,
-                     rmm::cuda_stream_view stream = cudf::get_default_stream())
-  {
-    auto deleter = [](mutable_table_device_view* t) { t->destroy(); };
-    return std::unique_ptr<mutable_table_device_view, decltype(deleter)>{
-      new mutable_table_device_view(source_view, stream), deleter};
-  }
+  static std::unique_ptr<mutable_table_device_view, std::function<void(mutable_table_device_view*)>>
+  create(mutable_table_view source_view, rmm::cuda_stream_view stream = cudf::get_default_stream());
 
  private:
-  mutable_table_device_view(mutable_table_view source_view, rmm::cuda_stream_view stream)
-    : detail::table_device_view_base<mutable_column_device_view, mutable_table_view>(source_view,
-                                                                                     stream)
-  {
-  }
+  mutable_table_device_view(mutable_table_view source_view, mutable_column_device_view* columns);
 };
 
 /**
@@ -225,7 +206,7 @@ class mutable_table_device_view
  * @return tuple of device_buffer and @p ColumnDeviceView device pointer
  */
 template <typename ColumnDeviceView, typename HostTableView>
-std::pair<std::unique_ptr<rmm::device_buffer>, ColumnDeviceView*>
-contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream);
+std::pair<std::unique_ptr<rmm::device_buffer>, ColumnDeviceView*> create_column_device_views(
+  HostTableView source_view, rmm::cuda_stream_view stream);
 
 }  // namespace CUDF_EXPORT cudf

cpp/src/lists/dremel.cu

@@ -210,8 +210,8 @@ dremel_data get_encoding(column_view h_col,
   }
 
   [[maybe_unused]] auto [device_view_owners, d_nesting_levels] =
-    contiguous_copy_column_device_views<column_device_view, host_span<column_view const>>(
-      host_span<column_view const>{nesting_levels}, stream);
+    create_column_device_views<column_device_view>(host_span<column_view const>{nesting_levels},
+                                                   stream);
 
   auto max_def_level = def_at_level.back();
   thrust::exclusive_scan(

cpp/src/strings/copying/concatenate.cu

@@ -67,7 +67,7 @@ auto create_strings_device_views(host_span<column_view const> views, rmm::cuda_s
   CUDF_FUNC_RANGE();
   // Assemble contiguous array of device views
   auto [device_view_owners, device_views_ptr] =
-    contiguous_copy_column_device_views<column_device_view>(views, stream);
+    create_column_device_views<column_device_view>(views, stream);
 
   // Compute the partition offsets and size of offset column
   // Note: Using 64-bit size_t so we can detect overflow of 32-bit size_type

cpp/src/structs/scan/scan_inclusive.cu

@@ -44,9 +44,10 @@ std::unique_ptr<column> scan_inclusive(column_view const& input,
 
   // Gather the children columns of the input column. Must use `get_sliced_child` to properly
   // handle input in case it is a sliced view.
+  auto const structs_view   = structs_column_view{input};
   auto const input_children = [&] {
     auto const it = cudf::detail::make_counting_transform_iterator(
-      0, [structs_view = structs_column_view{input}, &stream](auto const child_idx) {
+      0, [&structs_view, &stream](auto const child_idx) {
         return structs_view.get_sliced_child(child_idx, stream);
       });
     return std::vector<column_view>(it, it + input.num_children());

cpp/src/table/table_device_view.cu

@@ -22,24 +22,14 @@ namespace detail {
 template <typename ColumnDeviceView, typename HostTableView>
 void table_device_view_base<ColumnDeviceView, HostTableView>::destroy()
 {
-  delete _descendant_storage;
   delete this;
 }
 
 template <typename ColumnDeviceView, typename HostTableView>
 table_device_view_base<ColumnDeviceView, HostTableView>::table_device_view_base(
-  HostTableView source_view, rmm::cuda_stream_view stream)
-  : _num_rows{source_view.num_rows()}, _num_columns{source_view.num_columns()}
+  HostTableView source_view, ColumnDeviceView* columns)
+  : _columns{columns}, _num_rows{source_view.num_rows()}, _num_columns{source_view.num_columns()}
 {
-  // The table's columns must be converted to ColumnDeviceView
-  // objects and copied into device memory for the table_device_view's
-  // _columns member.
-  if (source_view.num_columns() > 0) {
-    std::unique_ptr<rmm::device_buffer> descendant_storage_owner;
-    std::tie(descendant_storage_owner, _columns) =
-      contiguous_copy_column_device_views<ColumnDeviceView, HostTableView>(source_view, stream);
-    _descendant_storage = descendant_storage_owner.release();
-  }
 }
 
 // Explicit instantiation for a device table of immutable views
@@ -51,8 +41,8 @@ template class table_device_view_base<mutable_column_device_view, mutable_table_
 }  // namespace detail
 
 template <typename ColumnDeviceView, typename HostTableView>
-std::pair<std::unique_ptr<rmm::device_buffer>, ColumnDeviceView*>
-contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream)
+std::pair<std::unique_ptr<rmm::device_buffer>, ColumnDeviceView*> create_column_device_views(
+  HostTableView source_view, rmm::cuda_stream_view stream)
 {
   // First calculate the size of memory needed to hold the
   // table's ColumnDeviceViews. This is done by calling extent()
@@ -89,7 +79,47 @@ contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_
 }
 
 template std::pair<std::unique_ptr<rmm::device_buffer>, column_device_view*>
-contiguous_copy_column_device_views<column_device_view, host_span<column_view const>>(
+create_column_device_views<column_device_view, host_span<column_view const>>(
   host_span<column_view const> source_view, rmm::cuda_stream_view stream);
 
+table_device_view::table_device_view(table_view source_view, column_device_view* columns)
+  : detail::table_device_view_base<column_device_view, table_view>(source_view, columns)
+{
+}
+
+std::unique_ptr<table_device_view, std::function<void(table_device_view*)>>
+table_device_view::create(table_view source_view, rmm::cuda_stream_view stream)
+{
+  auto [descendant_storage, columns] =
+    create_column_device_views<column_device_view, table_view>(source_view, stream);
+  auto deleter = [ds = descendant_storage.release()](table_device_view* tv) {
+    tv->destroy();
+    delete ds;
+  };
+  std::unique_ptr<table_device_view, decltype(deleter)> result{
+    new table_device_view(source_view, columns), deleter};
+  return result;
+}
+
+mutable_table_device_view::mutable_table_device_view(mutable_table_view source_view,
+                                                     mutable_column_device_view* columns)
+  : detail::table_device_view_base<mutable_column_device_view, mutable_table_view>(source_view,
+                                                                                   columns)
+{
+}
+
+std::unique_ptr<mutable_table_device_view, std::function<void(mutable_table_device_view*)>>
+mutable_table_device_view::create(mutable_table_view source_view, rmm::cuda_stream_view stream)
+{
+  auto [descendant_storage, columns] =
+    create_column_device_views<mutable_column_device_view, mutable_table_view>(source_view, stream);
+  auto deleter = [ds = descendant_storage.release()](mutable_table_device_view* tv) {
+    tv->destroy();
+    delete ds;
+  };
+  std::unique_ptr<mutable_table_device_view, decltype(deleter)> result{
+    new mutable_table_device_view(source_view, columns), deleter};
+  return result;
+}
+
 }  // namespace cudf

cpp/src/transform/row_bit_count.cu

@@ -520,8 +520,8 @@ std::unique_ptr<column> segmented_row_bit_count(table_view const& t,
   }
 
   // create a contiguous block of column_device_views
-  auto d_cols = contiguous_copy_column_device_views<column_device_view>(
-    host_span<column_view const>{cols}, stream);
+  auto d_cols =
+    create_column_device_views<column_device_view>(host_span<column_view const>{cols}, stream);
 
   // move stack info to the gpu
   rmm::device_uvector<column_info> d_info =