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