table_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2022, NVIDIA CORPORATION.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 #pragma once
17 
20 #include <cudf/types.hpp>
21 #include <cudf/utilities/default_stream.hpp>
22 
23 #include <rmm/cuda_stream_view.hpp>
24 #include <rmm/device_buffer.hpp>
25 
26 #include <cassert>
27 #include <memory>
28 #include <numeric>
29 
35 namespace cudf {
36 namespace detail {
37 
54 template <typename ColumnDeviceView, typename HostTableView>
56  public:
57  table_device_view_base() = delete;
58  ~table_device_view_base() = default;
61 
73 
79  __device__ ColumnDeviceView* begin() const noexcept { return _columns; }
80 
89  __device__ ColumnDeviceView* end() const noexcept { return _columns + _num_columns; }
90 
97  __device__ ColumnDeviceView const& column(size_type column_index) const noexcept
98  {
99  assert(column_index >= 0);
100  assert(column_index < _num_columns);
101  return _columns[column_index];
102  }
103 
110  __device__ ColumnDeviceView& column(size_type column_index) noexcept
111  {
112  assert(column_index >= 0);
113  assert(column_index < _num_columns);
114  return _columns[column_index];
115  }
116 
122  [[nodiscard]] __host__ __device__ size_type num_columns() const noexcept { return _num_columns; }
123 
129  [[nodiscard]] __host__ __device__ size_type num_rows() const noexcept { return _num_rows; }
130 
137  void destroy();
138 
139  private:
140  ColumnDeviceView* _columns{};
141  size_type _num_rows{};
142  size_type _num_columns{};
143 
144  protected:
151  table_device_view_base(HostTableView source_view, rmm::cuda_stream_view stream);
152 
155 };
156 } // namespace detail
157 
161 class table_device_view : public detail::table_device_view_base<column_device_view, table_view> {
162  public:
177  static auto create(table_view source_view,
179  {
180  auto deleter = [](table_device_view* t) { t->destroy(); };
181  return std::unique_ptr<table_device_view, decltype(deleter)>{
182  new table_device_view(source_view, stream), deleter};
183  }
184 
185  private:
187  : detail::table_device_view_base<column_device_view, table_view>(source_view, stream)
188  {
189  }
190 };
191 
198  : public detail::table_device_view_base<mutable_column_device_view, mutable_table_view> {
199  public:
214  static auto create(mutable_table_view source_view,
216  {
217  auto deleter = [](mutable_table_device_view* t) { t->destroy(); };
218  return std::unique_ptr<mutable_table_device_view, decltype(deleter)>{
219  new mutable_table_device_view(source_view, stream), deleter};
220  }
221 
222  private:
224  : detail::table_device_view_base<mutable_column_device_view, mutable_table_view>(source_view,
225  stream)
226  {
227  }
228 };
229 
239 template <typename ColumnDeviceView, typename HostTableView>
240 auto contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream)
241 {
242  // First calculate the size of memory needed to hold the
243  // table's ColumnDeviceViews. This is done by calling extent()
244  // for each of the table's ColumnViews columns.
245  std::size_t views_size_bytes = std::accumulate(
246  source_view.begin(), source_view.end(), std::size_t{0}, [](std::size_t init, auto col) {
247  return init + ColumnDeviceView::extent(col);
248  });
249  // pad the allocation for aligning the first pointer
250  auto padded_views_size_bytes = views_size_bytes + std::size_t{alignof(ColumnDeviceView) - 1};
251  // A buffer of CPU memory is allocated to hold the ColumnDeviceView
252  // objects. Once filled, the CPU memory is then copied to device memory
253  // and the pointer is set in the d_columns member.
254  std::vector<int8_t> h_buffer(padded_views_size_bytes);
255  // Each ColumnDeviceView instance may have child objects which may
256  // require setting some internal device pointers before being copied
257  // from CPU to device.
258  // Allocate the device memory to be used in the result.
259  // We need this pointer in order to pass it down when creating the
260  // ColumnDeviceViews so the column can set the pointer(s) for any
261  // of its child objects.
262  // align both h_ptr, d_ptr
263  auto descendant_storage = std::make_unique<rmm::device_buffer>(padded_views_size_bytes, stream);
264  void* h_ptr = detail::align_ptr_for_type<ColumnDeviceView>(h_buffer.data());
265  void* d_ptr = detail::align_ptr_for_type<ColumnDeviceView>(descendant_storage->data());
266  auto d_columns = detail::child_columns_to_device_array<ColumnDeviceView>(
267  source_view.begin(), source_view.end(), h_ptr, d_ptr);
268 
269  CUDF_CUDA_TRY(cudaMemcpyAsync(d_ptr, h_ptr, views_size_bytes, cudaMemcpyDefault, stream.value()));
270  stream.synchronize();
271  return std::make_tuple(std::move(descendant_storage), d_columns);
272 }
273 
274 } // namespace cudf
cudf::detail::table_device_view_base::begin
ColumnDeviceView * begin() const noexcept
Returns an iterator to the first view in the table.
Definition: table_device_view.cuh:79
cudf::mutable_table_view
A set of mutable_column_views of the same size.
Definition: table_view.hpp:255
cudf::detail::table_device_view_base::operator=
table_device_view_base & operator=(table_device_view_base const &)=default
Copy assignment operator.
table_view.hpp
Class definitions for (mutable)_table_view
cudf::detail::table_device_view_base::num_rows
size_type num_rows() const noexcept
Returns the number of rows.
Definition: table_device_view.cuh:129
cudf::mutable_table_device_view::create
static auto create(mutable_table_view source_view, rmm::cuda_stream_view stream=cudf::get_default_stream())
Factory to construct a mutable table device view that is usable in device memory.
Definition: table_device_view.cuh:214
cudf::size_type
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:80
cudf::table_device_view
Table device view that is usable in device memory.
Definition: table_device_view.cuh:161
cudf::column_device_view
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Definition: column_device_view.cuh:349
types.hpp
Type declarations for libcudf.
cudf::table_device_view::create
static auto create(table_view source_view, rmm::cuda_stream_view stream=cudf::get_default_stream())
Factory to construct a table device view that is usable in device memory.
Definition: table_device_view.cuh:177
rmm::cuda_stream_view
cudf::detail::table_device_view_base::column
ColumnDeviceView & column(size_type column_index) noexcept
Returns a reference to the view of the specified column.
Definition: table_device_view.cuh:110
cudf::detail::table_device_view_base::table_device_view_base
table_device_view_base(table_device_view_base const &)=default
Copy constructor.
cudf::detail::table_device_view_base::num_columns
size_type num_columns() const noexcept
Returns the number of columns.
Definition: table_device_view.cuh:122
column_device_view.cuh
Column device view class definitions.
cudf::detail::table_device_view_base::_descendant_storage
rmm::device_buffer * _descendant_storage
Pointer to device memory holding the descendant storage.
Definition: table_device_view.cuh:154
rmm::device_buffer
device_buffer.hpp
cudf::detail::table_device_view_base::destroy
void destroy()
Destroy the table_device_view object.
cudf::table_view
A set of cudf::column_view's of the same size.
Definition: table_view.hpp:187
cudf::contiguous_copy_column_device_views
auto contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream)
Copies the contents of a table_view to a column device view in contiguous device memory.
Definition: table_device_view.cuh:240
cudf::detail::table_device_view_base::end
ColumnDeviceView * end() const noexcept
Returns an iterator one past the last column view in the table.
Definition: table_device_view.cuh:89
cudf::detail::table_device_view_base::operator=
table_device_view_base & operator=(table_device_view_base &&)=default
Move assignment operator.
cudf::mutable_column_device_view
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
Definition: column_device_view.cuh:922
cudf::detail::table_device_view_base::table_device_view_base
table_device_view_base(HostTableView source_view, rmm::cuda_stream_view stream)
Construct a new table device view base object from host table_view.
cudf::detail::table_device_view_base
Base class for a device table of ColumnDeviceViews.
Definition: table_device_view.cuh:55
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::mutable_table_device_view
Mutable table device view that is usable in device memory.
Definition: table_device_view.cuh:198
cudf::detail::table_device_view_base::column
ColumnDeviceView const & column(size_type column_index) const noexcept
Returns a reference to the view of the specified column.
Definition: table_device_view.cuh:97
cudf::get_default_stream
rmm::cuda_stream_view const get_default_stream()
Get the current default stream.
rmm::cuda_stream_view::synchronize
void synchronize() const
rmm::cuda_stream_view::value
constexpr cudaStream_t value() const noexcept
CUDF_CUDA_TRY
#define CUDF_CUDA_TRY(call)
Error checking macro for CUDA runtime API functions.
Definition: error.hpp:232
cudf::detail::table_device_view_base::table_device_view_base
table_device_view_base(table_device_view_base &&)=default
Move constructor.