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 
22 #include <rmm/cuda_stream_view.hpp>
23 #include <rmm/device_buffer.hpp>
24 
25 #include <cassert>
26 #include <memory>
27 #include <numeric>
28 
34 namespace cudf {
35 namespace detail {
36 template <typename ColumnDeviceView, typename HostTableView>
38  public:
39  table_device_view_base() = delete;
40  ~table_device_view_base() = default;
43  table_device_view_base& operator=(table_device_view_base const&) = default;
44  table_device_view_base& operator=(table_device_view_base&&) = default;
45 
46  __device__ ColumnDeviceView* begin() const noexcept { return _columns; }
47 
48  __device__ ColumnDeviceView* end() const noexcept { return _columns + _num_columns; }
49 
50  __device__ ColumnDeviceView const& column(size_type column_index) const noexcept
51  {
52  assert(column_index >= 0);
53  assert(column_index < _num_columns);
54  return _columns[column_index];
55  }
56 
57  __device__ ColumnDeviceView& column(size_type column_index) noexcept
58  {
59  assert(column_index >= 0);
60  assert(column_index < _num_columns);
61  return _columns[column_index];
62  }
63 
64  [[nodiscard]] __host__ __device__ size_type num_columns() const noexcept { return _num_columns; }
65 
66  [[nodiscard]] __host__ __device__ size_type num_rows() const noexcept { return _num_rows; }
67 
68  void destroy();
69 
70  private:
71  ColumnDeviceView* _columns{};
72  size_type _num_rows{};
73  size_type _num_columns{};
74 
75  protected:
76  table_device_view_base(HostTableView source_view, rmm::cuda_stream_view stream);
77 
78  rmm::device_buffer* _descendant_storage{};
79 };
80 } // namespace detail
81 
82 class table_device_view : public detail::table_device_view_base<column_device_view, table_view> {
83  public:
84  static auto create(table_view source_view,
85  rmm::cuda_stream_view stream = rmm::cuda_stream_default)
86  {
87  auto deleter = [](table_device_view* t) { t->destroy(); };
88  return std::unique_ptr<table_device_view, decltype(deleter)>{
89  new table_device_view(source_view, stream), deleter};
90  }
91 
92  private:
95  {
96  }
97 };
98 
100  : public detail::table_device_view_base<mutable_column_device_view, mutable_table_view> {
101  public:
102  static auto create(mutable_table_view source_view,
103  rmm::cuda_stream_view stream = rmm::cuda_stream_default)
104  {
105  auto deleter = [](mutable_table_device_view* t) { t->destroy(); };
106  return std::unique_ptr<mutable_table_device_view, decltype(deleter)>{
107  new mutable_table_device_view(source_view, stream), deleter};
108  }
109 
110  private:
113  stream)
114  {
115  }
116 };
117 
118 template <typename ColumnDeviceView, typename HostTableView>
119 auto contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_stream_view stream)
120 {
121  // First calculate the size of memory needed to hold the
122  // table's ColumnDeviceViews. This is done by calling extent()
123  // for each of the table's ColumnViews columns.
124  std::size_t views_size_bytes = std::accumulate(
125  source_view.begin(), source_view.end(), std::size_t{0}, [](std::size_t init, auto col) {
126  return init + ColumnDeviceView::extent(col);
127  });
128  // pad the allocation for aligning the first pointer
129  auto padded_views_size_bytes = views_size_bytes + std::size_t{alignof(ColumnDeviceView) - 1};
130  // A buffer of CPU memory is allocated to hold the ColumnDeviceView
131  // objects. Once filled, the CPU memory is then copied to device memory
132  // and the pointer is set in the d_columns member.
133  std::vector<int8_t> h_buffer(padded_views_size_bytes);
134  // Each ColumnDeviceView instance may have child objects which may
135  // require setting some internal device pointers before being copied
136  // from CPU to device.
137  // Allocate the device memory to be used in the result.
138  // We need this pointer in order to pass it down when creating the
139  // ColumnDeviceViews so the column can set the pointer(s) for any
140  // of its child objects.
141  // align both h_ptr, d_ptr
142  auto descendant_storage = std::make_unique<rmm::device_buffer>(padded_views_size_bytes, stream);
143  void* h_ptr = detail::align_ptr_for_type<ColumnDeviceView>(h_buffer.data());
144  void* d_ptr = detail::align_ptr_for_type<ColumnDeviceView>(descendant_storage->data());
145  auto d_columns = detail::child_columns_to_device_array<ColumnDeviceView>(
146  source_view.begin(), source_view.end(), h_ptr, d_ptr);
147 
148  CUDF_CUDA_TRY(cudaMemcpyAsync(d_ptr, h_ptr, views_size_bytes, cudaMemcpyDefault, stream.value()));
149  stream.synchronize();
150  return std::make_tuple(std::move(descendant_storage), d_columns);
151 }
152 
153 } // namespace cudf
cudf::mutable_table_view
A set of mutable_column_views of the same size.
Definition: table_view.hpp:222
table_view.hpp
Class definitions for (mutable)_table_view
cudf::column
A container of nullable device data as a column of elements.
Definition: column.hpp:45
cudf::size_type
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:84
cudf::table_device_view
Definition: table_device_view.cuh:82
types.hpp
Type declarations for libcudf.
rmm::cuda_stream_view
column_device_view.cuh
Column device view class definitions.
rmm::device_buffer
device_buffer.hpp
cudf::table_view
A set of cudf::column_view's of the same size.
Definition: table_view.hpp:154
cudf::detail::table_device_view_base
Definition: table_device_view.cuh:37
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::mutable_table_device_view
Definition: table_device_view.cuh:100
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:142