diff --git a/cpp/include/cudf/table/table_device_view.cuh b/cpp/include/cudf/table/table_device_view.cuh index 163743c3a65..7a3f09ee896 100644 --- a/cpp/include/cudf/table/table_device_view.cuh +++ b/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_basedestroy(); }; - return std::unique_ptr{ - new table_device_view(source_view, stream), deleter}; - } + static std::unique_ptr> 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(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{ - new mutable_table_device_view(source_view, stream), deleter}; - } + static std::unique_ptr> + 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(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 -std::pair, ColumnDeviceView*> -contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream); +std::pair, ColumnDeviceView*> create_column_device_views( + HostTableView source_view, rmm::cuda_stream_view stream); } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/lists/dremel.cu b/cpp/src/lists/dremel.cu index 0db75b984f9..99a30cbb2f3 100644 --- a/cpp/src/lists/dremel.cu +++ b/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>( - host_span{nesting_levels}, stream); + create_column_device_views(host_span{nesting_levels}, + stream); auto max_def_level = def_at_level.back(); thrust::exclusive_scan( diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu index b8f8509b60f..8d90d18c151 100644 --- a/cpp/src/strings/copying/concatenate.cu +++ b/cpp/src/strings/copying/concatenate.cu @@ -67,7 +67,7 @@ auto create_strings_device_views(host_span 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(views, stream); + create_column_device_views(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 diff --git a/cpp/src/structs/scan/scan_inclusive.cu b/cpp/src/structs/scan/scan_inclusive.cu index 2a9753f52b8..3b3671ba3d5 100644 --- a/cpp/src/structs/scan/scan_inclusive.cu +++ b/cpp/src/structs/scan/scan_inclusive.cu @@ -44,9 +44,10 @@ std::unique_ptr 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(it, it + input.num_children()); diff --git a/cpp/src/table/table_device_view.cu b/cpp/src/table/table_device_view.cu index 3e145c2a999..03da0d065de 100644 --- a/cpp/src/table/table_device_view.cu +++ b/cpp/src/table/table_device_view.cu @@ -22,24 +22,14 @@ namespace detail { template void table_device_view_base::destroy() { - delete _descendant_storage; delete this; } template table_device_view_base::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 descendant_storage_owner; - std::tie(descendant_storage_owner, _columns) = - contiguous_copy_column_device_views(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 -std::pair, ColumnDeviceView*> -contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream) +std::pair, 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, column_device_view*> -contiguous_copy_column_device_views>( +create_column_device_views>( host_span 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(source_view, columns) +{ +} + +std::unique_ptr> +table_device_view::create(table_view source_view, rmm::cuda_stream_view stream) +{ + auto [descendant_storage, columns] = + create_column_device_views(source_view, stream); + auto deleter = [ds = descendant_storage.release()](table_device_view* tv) { + tv->destroy(); + delete ds; + }; + std::unique_ptr 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(source_view, + columns) +{ +} + +std::unique_ptr> +mutable_table_device_view::create(mutable_table_view source_view, rmm::cuda_stream_view stream) +{ + auto [descendant_storage, columns] = + create_column_device_views(source_view, stream); + auto deleter = [ds = descendant_storage.release()](mutable_table_device_view* tv) { + tv->destroy(); + delete ds; + }; + std::unique_ptr result{ + new mutable_table_device_view(source_view, columns), deleter}; + return result; +} + } // namespace cudf diff --git a/cpp/src/transform/row_bit_count.cu b/cpp/src/transform/row_bit_count.cu index 2b944f380d0..2e204d812ba 100644 --- a/cpp/src/transform/row_bit_count.cu +++ b/cpp/src/transform/row_bit_count.cu @@ -520,8 +520,8 @@ std::unique_ptr segmented_row_bit_count(table_view const& t, } // create a contiguous block of column_device_views - auto d_cols = contiguous_copy_column_device_views( - host_span{cols}, stream); + auto d_cols = + create_column_device_views(host_span{cols}, stream); // move stack info to the gpu rmm::device_uvector d_info =