column_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 
19 #include <cudf/detail/utilities/alignment.hpp>
21 #include <cudf/lists/list_view.hpp>
22 #include <cudf/strings/string_view.cuh>
25 #include <cudf/types.hpp>
26 #include <cudf/utilities/bit.hpp>
27 #include <cudf/utilities/span.hpp>
30 
31 #include <rmm/cuda_stream_view.hpp>
32 
33 #include <thrust/iterator/counting_iterator.h>
34 #include <thrust/iterator/transform_iterator.h>
35 #include <thrust/optional.h>
36 #include <thrust/pair.h>
37 
38 #include <algorithm>
39 
45 namespace cudf {
46 
54 struct nullate {
55  struct YES : std::bool_constant<true> {
56  };
57  struct NO : std::bool_constant<false> {
58  };
59  struct DYNAMIC {
60  DYNAMIC() = delete;
69  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
70  constexpr operator bool() const noexcept { return value; }
71  bool value;
72  };
73 };
74 
75 namespace detail {
87 class alignas(16) column_device_view_base {
88  public:
89  column_device_view_base() = delete;
90  ~column_device_view_base() = default;
93  column_device_view_base& operator=(column_device_view_base const&) = default;
94  column_device_view_base& operator=(column_device_view_base&&) = default;
95 
112  template <typename T = void,
113  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
114  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
115  {
116  return static_cast<T const*>(_data);
117  }
118 
134  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
135  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
136  {
137  return head<T>() + _offset;
138  }
139 
143  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
144 
148  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
149 
159  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
160 
168  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
169  {
170  return _null_mask;
171  }
172 
177  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
178 
193  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
194  {
195  return not nullable() or is_valid_nocheck(element_index);
196  }
197 
210  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
211  {
212  return bit_is_set(_null_mask, offset() + element_index);
213  }
214 
228  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
229  {
230  return not is_valid(element_index);
231  }
232 
244  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
245  {
246  return not is_valid_nocheck(element_index);
247  }
248 
258  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
259  {
260  return null_mask()[word_index];
261  }
262 
263  protected:
264  data_type _type{type_id::EMPTY};
266  void const* _data{};
268  size_type _offset{};
270 
272  CUDF_HOST_DEVICE column_device_view_base(data_type type,
273  size_type size,
274  void const* data,
275  bitmask_type const* null_mask,
278  {
279  }
280 
281  template <typename C, typename T, typename = void>
282  struct has_element_accessor_impl : std::false_type {
283  };
284 
285  template <typename C, typename T>
286  struct has_element_accessor_impl<
287  C,
288  T,
289  void_t<decltype(std::declval<C>().template element<T>(std::declval<size_type>()))>>
290  : std::true_type {
291  };
292 };
293 // Forward declaration
294 template <typename T>
295 struct value_accessor;
296 template <typename T, typename Nullate>
297 struct optional_accessor;
298 template <typename T, bool has_nulls>
299 struct pair_accessor;
300 template <typename T, bool has_nulls>
301 struct pair_rep_accessor;
302 template <typename T>
303 struct mutable_value_accessor;
304 } // namespace detail
305 
313  public:
314  column_device_view() = delete;
315  ~column_device_view() = default;
316  column_device_view(column_device_view const&) = default;
318  column_device_view& operator=(column_device_view const&) = default;
319  column_device_view& operator=(column_device_view&&) = default;
320 
330  column_device_view(column_view column, void* h_ptr, void* d_ptr);
331 
347  [[nodiscard]] CUDF_HOST_DEVICE column_device_view slice(size_type offset,
348  size_type size) const noexcept
349  {
350  return column_device_view{this->type(),
351  size,
352  this->head(),
353  this->null_mask(),
354  this->offset() + offset,
355  d_children,
356  this->num_child_columns()};
357  }
358 
375  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
376  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
377  {
378  return data<T>()[element_index];
379  }
380 
392  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
393  __device__ T element(size_type element_index) const noexcept
394  {
395  size_type index = element_index + offset(); // account for this view's _offset
396  const auto* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
397  const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
398  size_type offset = d_offsets[index];
399  return string_view{d_strings + offset, d_offsets[index + 1] - offset};
400  }
401 
402  private:
408  struct index_element_fn {
409  template <typename IndexType,
410  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>)>
411  __device__ size_type operator()(column_device_view const& indices, size_type index)
412  {
413  return static_cast<size_type>(indices.element<IndexType>(index));
414  }
415 
416  template <typename IndexType,
417  typename... Args,
418  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>))>
419  __device__ size_type operator()(Args&&... args)
420  {
421  CUDF_UNREACHABLE("dictionary indices must be an unsigned integral type");
422  }
423  };
424 
425  public:
450  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
451  __device__ T element(size_type element_index) const noexcept
452  {
453  size_type index = element_index + offset(); // account for this view's _offset
454  auto const indices = d_children[0];
455  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
456  }
457 
468  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
469  __device__ T element(size_type element_index) const noexcept
470  {
471  using namespace numeric;
472  using rep = typename T::rep;
473  auto const scale = scale_type{_type.scale()};
474  return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
475  }
476 
481  template <typename T>
482  static constexpr bool has_element_accessor()
483  {
484  return has_element_accessor_impl<column_device_view, T>::value;
485  }
486 
490  using count_it = thrust::counting_iterator<size_type>;
491  template <typename T>
492  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
493 
506  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
507  [[nodiscard]] const_iterator<T> begin() const
508  {
509  return const_iterator<T>{count_it{0}, detail::value_accessor<T>{*this}};
510  }
511 
524  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
525  [[nodiscard]] const_iterator<T> end() const
526  {
527  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
528  }
529 
533  template <typename T, typename Nullate>
535  thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
536 
540  template <typename T, bool has_nulls>
542  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
543 
549  template <typename T, bool has_nulls>
551  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
552 
602  template <typename T,
603  typename Nullate,
604  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
605  auto optional_begin(Nullate has_nulls) const
606  {
608  count_it{0}, detail::optional_accessor<T, Nullate>{*this, has_nulls}};
609  }
610 
630  template <typename T,
631  bool has_nulls,
632  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
634  {
637  }
638 
660  template <typename T,
661  bool has_nulls,
662  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
664  {
667  }
668 
681  template <typename T,
682  typename Nullate,
683  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
684  auto optional_end(Nullate has_nulls) const
685  {
687  count_it{size()}, detail::optional_accessor<T, Nullate>{*this, has_nulls}};
688  }
689 
701  template <typename T,
702  bool has_nulls,
703  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
705  {
708  }
709 
721  template <typename T,
722  bool has_nulls,
723  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
725  {
728  }
729 
748  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
749  column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
750 
757  void destroy();
758 
766  static std::size_t extent(column_view const& source_view);
767 
774  [[nodiscard]] __device__ column_device_view child(size_type child_index) const noexcept
775  {
776  return d_children[child_index];
777  }
778 
782  [[nodiscard]] __device__ device_span<column_device_view const> children() const noexcept
783  {
785  }
786 
792  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
793  {
794  return _num_children;
795  }
796 
797  private:
810  CUDF_HOST_DEVICE column_device_view(data_type type,
811  size_type size,
812  void const* data,
813  bitmask_type const* null_mask,
816  size_type num_children)
817  : column_device_view_base(type, size, data, null_mask, offset),
819  _num_children(num_children)
820  {
821  }
822 
823  protected:
829 
839 };
840 
848  public:
849  mutable_column_device_view() = delete;
850  ~mutable_column_device_view() = default;
853  mutable_column_device_view& operator=(mutable_column_device_view const&) = default;
855 
866 
885  static std::unique_ptr<mutable_column_device_view,
886  std::function<void(mutable_column_device_view*)>>
887  create(mutable_column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
888 
905  template <typename T = void,
906  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
907  CUDF_HOST_DEVICE T* head() const noexcept
908  {
909  return const_cast<T*>(detail::column_device_view_base::head<T>());
910  }
911 
924  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
925  CUDF_HOST_DEVICE T* data() const noexcept
926  {
927  return const_cast<T*>(detail::column_device_view_base::data<T>());
928  }
929 
943  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
944  __device__ T& element(size_type element_index) const noexcept
945  {
946  return data<T>()[element_index];
947  }
948 
954  template <typename T>
955  static constexpr bool has_element_accessor()
956  {
957  return has_element_accessor_impl<mutable_column_device_view, T>::value;
958  }
959 
967  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
968  {
970  }
971 
975  using count_it = thrust::counting_iterator<size_type>;
976  template <typename T>
977  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
978 
989  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
990  iterator<T> begin()
991  {
992  return iterator<T>{count_it{0}, detail::mutable_value_accessor<T>{*this}};
993  }
994 
1005  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1006  iterator<T> end()
1007  {
1008  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
1009  }
1010 
1017  [[nodiscard]] __device__ mutable_column_device_view child(size_type child_index) const noexcept
1018  {
1019  return d_children[child_index];
1020  }
1021 
1022 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1023 
1037  __device__ void set_valid(size_type element_index) const noexcept
1038  {
1039  return set_bit(null_mask(), element_index);
1040  }
1041 
1055  __device__ void set_null(size_type element_index) const noexcept
1056  {
1057  return clear_bit(null_mask(), element_index);
1058  }
1059 
1060 #endif
1061 
1072  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
1073  {
1074  null_mask()[word_index] = new_word;
1075  }
1076 
1083  static std::size_t extent(mutable_column_view source_view);
1084 
1091  void destroy();
1092 
1093  private:
1099 
1108  mutable_column_device_view(mutable_column_view source);
1109 };
1110 
1111 namespace detail {
1112 
1113 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1114 
1121 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
1122  size_type destination_word_index,
1123  size_type source_begin_bit,
1124  size_type source_end_bit)
1125 {
1126  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
1127  bitmask_type curr_word = source[source_word_index];
1128  bitmask_type next_word = 0;
1129  if (word_index(source_end_bit - 1) >
1130  word_index(source_begin_bit +
1131  destination_word_index * detail::size_in_bits<bitmask_type>())) {
1132  next_word = source[source_word_index + 1];
1133  }
1134  return __funnelshift_r(curr_word, next_word, source_begin_bit);
1135 }
1136 
1137 #endif
1138 
1152 template <typename T>
1155 
1160  value_accessor(column_device_view const& _col) : col{_col}
1161  {
1162  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1163  }
1164 
1165  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1166 };
1167 
1194 template <typename T, typename Nullate>
1197 
1204  optional_accessor(column_device_view const& _col, Nullate with_nulls)
1205  : col{_col}, has_nulls{with_nulls}
1206  {
1207  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1208  if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1209  }
1210 
1211  __device__ inline thrust::optional<T> operator()(cudf::size_type i) const
1212  {
1213  if (has_nulls) {
1214  return (col.is_valid_nocheck(i)) ? thrust::optional<T>{col.element<T>(i)}
1215  : thrust::optional<T>{thrust::nullopt};
1216  }
1217  return thrust::optional<T>{col.element<T>(i)};
1218  }
1219 
1220  Nullate has_nulls{};
1221 };
1222 
1241 template <typename T, bool has_nulls = false>
1244 
1249  pair_accessor(column_device_view const& _col) : col{_col}
1250  {
1251  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1252  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1253  }
1254 
1255  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
1256  {
1257  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1258  }
1259 };
1260 
1279 template <typename T, bool has_nulls = false>
1282 
1283  using rep_type = device_storage_type_t<T>;
1284 
1289  pair_rep_accessor(column_device_view const& _col) : col{_col}
1290  {
1291  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1292  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1293  }
1294 
1295  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1296  {
1297  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1298  }
1299 
1300  private:
1301  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
1302  __device__ inline auto get_rep(cudf::size_type i) const
1303  {
1304  return col.element<R>(i);
1305  }
1306 
1307  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
1308  __device__ inline auto get_rep(cudf::size_type i) const
1309  {
1310  return col.element<R>(i).value();
1311  }
1312 };
1313 
1314 template <typename T>
1317 
1323  {
1324  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1325  }
1326 
1327  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1328 };
1329 
1355 template <typename ColumnDeviceView, typename ColumnViewIterator>
1356 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1357  ColumnViewIterator child_end,
1358  void* h_ptr,
1359  void* d_ptr)
1360 {
1361  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1362  auto num_children = std::distance(child_begin, child_end);
1363  if (num_children > 0) {
1364  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1365  // struct objects in order for d_children to be used as an array.
1366  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1367  auto d_column = d_children;
1368 
1369  // Any child data is assigned past the end of this array: h_end and d_end.
1370  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1371  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1372  std::for_each(child_begin, child_end, [&](auto const& col) {
1373  // inplace-new each child into host memory
1374  new (h_column) ColumnDeviceView(col, h_end, d_end);
1375  h_column++; // advance to next child
1376  // update the pointers for holding this child column's child data
1377  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1378  h_end += col_child_data_size;
1379  d_end += col_child_data_size;
1380  });
1381  }
1382  return d_children;
1383 }
1384 
1385 } // namespace detail
1386 } // namespace cudf
cudf::bit_is_set
CUDF_HOST_DEVICE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
Definition: bit.hpp:112
cudf::detail::column_device_view_base::_data
void const * _data
Pointer to device memory containing elements.
Definition: column_device_view.cuh:266
cudf::column_device_view::optional_end
auto optional_end(Nullate has_nulls) const
Return an optional iterator to the element following the last element of the column.
Definition: column_device_view.cuh:684
cudf::detail::mutable_value_accessor::mutable_value_accessor
mutable_value_accessor(mutable_column_device_view &_col)
constructor
Definition: column_device_view.cuh:1322
bit.hpp
Utilities for bit and bitmask operations.
fixed_point.hpp
Class definition for fixed point data type.
cudf::detail::column_device_view_base::head
CUDF_HOST_DEVICE T const * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
Definition: column_device_view.cuh:114
cudf::column
A container of nullable device data as a column of elements.
Definition: column.hpp:45
cudf::data_type::id
constexpr type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:284
cudf::detail::pair_rep_accessor::pair_rep_accessor
pair_rep_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1289
strings_column_view.hpp
Class definition for cudf::strings_column_view.
cudf::detail::column_device_view_base::_offset
size_type _offset
Definition: column_device_view.cuh:269
cudf::nullate::DYNAMIC
Definition: column_device_view.cuh:59
cudf::size_type
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:84
cudf::column_view
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
Definition: column_view.hpp:300
cudf::column_device_view::num_child_columns
CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
Returns the number of child columns.
Definition: column_device_view.cuh:792
cudf::nullate
Indicates the presence of nulls at compile-time or runtime.
Definition: column_device_view.cuh:54
cudf::mutable_column_device_view::destroy
void destroy()
Destroy the mutable_column_device_view object.
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:312
types.hpp
Type declarations for libcudf.
cudf::column_device_view::const_optional_iterator
thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > const_optional_iterator
Optional iterator for navigating this column.
Definition: column_device_view.cuh:535
cudf::detail::column_device_view_base::is_valid
bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
Definition: column_device_view.cuh:193
cudf::string_view
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
Definition: string_view.hpp:49
cudf::mutable_column_view
A non-owning, mutable view of device data as a column of elements, some of which may be null as indic...
Definition: column_view.hpp:448
rmm::cuda_stream_view
cudf::column_device_view::pair_end
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
Definition: column_device_view.cuh:704
cudf::data_type::scale
constexpr int32_t scale() const noexcept
Returns the scale (for fixed_point types)
Definition: types.hpp:291
cudf::detail::column_device_view_base::null_mask
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
Definition: column_device_view.cuh:168
cudf::detail::column_device_view_base::is_null_nocheck
bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:244
cudf::column_device_view::optional_begin
auto optional_begin(Nullate has_nulls) const
Return an optional iterator to the first element of the column.
Definition: column_device_view.cuh:605
cudf::column_device_view::column_device_view
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
cudf::bitmask_type
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
Definition: types.hpp:85
cudf::column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
Definition: column_device_view.cuh:490
cudf::device_storage_type_t
std::conditional_t< std::is_same_v< numeric::decimal32, T >, int32_t, std::conditional_t< std::is_same_v< numeric::decimal64, T >, int64_t, std::conditional_t< std::is_same_v< numeric::decimal128, T >, __int128_t, T > >> device_storage_type_t
"Returns" the corresponding type that is stored on the device when using cudf::column
Definition: type_dispatcher.hpp:104
cudf::detail::pair_rep_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1281
CUDF_ENABLE_IF
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:48
cudf::column_device_view::slice
CUDF_HOST_DEVICE column_device_view slice(size_type offset, size_type size) const noexcept
Get a new column_device_view which is a slice of this column.
Definition: column_device_view.cuh:347
cudf::detail::mutable_value_accessor
Definition: column_device_view.cuh:1315
cudf::detail::pair_accessor::pair_accessor
pair_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1249
cudf::detail::column_device_view_base::offset
CUDF_HOST_DEVICE size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
Definition: column_device_view.cuh:177
cudf::dictionary_wrapper
A strongly typed wrapper for indices in a DICTIONARY type column.
Definition: dictionary.hpp:48
cudf::detail::child_columns_to_device_array
ColumnDeviceView * child_columns_to_device_array(ColumnViewIterator child_begin, ColumnViewIterator child_end, void *h_ptr, void *d_ptr)
Helper function for use by column_device_view and mutable_column_device_view constructors to build de...
Definition: column_device_view.cuh:1356
cudf::column_device_view::const_pair_rep_iterator
thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > const_pair_rep_iterator
Pair rep iterator for navigating this column.
Definition: column_device_view.cuh:551
cudf::mutable_column_device_view::mutable_column_device_view
mutable_column_device_view(mutable_column_view column, void *h_ptr, void *d_ptr)
Creates an instance of this class using the specified host memory pointer (h_ptr) to store child obje...
cudf::type_dispatcher
CUDF_HOST_DEVICE constexpr decltype(auto) __forceinline__ type_dispatcher(cudf::data_type dtype, Functor f, Ts &&... args)
Invokes an operator() template with the type instantiation based on the specified cudf::data_type's i...
Definition: type_dispatcher.hpp:438
cudf::mutable_column_device_view::end
iterator< T > end()
Return one past the last element after underlying data is casted to the specified type.
Definition: column_device_view.cuh:1006
cudf::mutable_column_device_view::element
T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:944
cudf::column_device_view::create
static std::unique_ptr< column_device_view, std::function< void(column_device_view *)> > create(column_view source_view, rmm::cuda_stream_view stream=rmm::cuda_stream_default)
Factory to construct a column view that is usable in device memory.
cudf::mutable_column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
Definition: column_device_view.cuh:975
numeric::scaled_integer
Helper struct for constructing fixed_point when value is already shifted.
Definition: fixed_point.hpp:177
cudf::nullate::NO
Definition: column_device_view.cuh:57
cudf::detail::column_device_view_base::get_mask_word
bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the the specified bitmask word from the null_mask().
Definition: column_device_view.cuh:258
cudf::detail::pair_accessor
pair accessor of column with/without null bitmask A unary functor returns pair with scalar value at i...
Definition: column_device_view.cuh:1242
cudf::column_device_view::pair_rep_end
const_pair_rep_iterator< T, has_nulls > pair_rep_end() const
Return a pair iterator to the element following the last element of the column.
Definition: column_device_view.cuh:724
cudf::detail::pair_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1243
cudf::detail::column_device_view_base::is_null
bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:228
cudf::data_type
Indicator for the logical data type of an element in a column.
Definition: types.hpp:240
cudf::mutable_column_device_view::extent
static std::size_t extent(mutable_column_view source_view)
Return the size in bytes of the amount of memory needed to hold a device view of the specified column...
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:847
cudf::detail::value_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1154
list_view.hpp
Class definition for cudf::list_view.
cudf::column_device_view::pair_rep_begin
const_pair_rep_iterator< T, has_nulls > pair_rep_begin() const
Return a pair iterator to the first element of the column.
Definition: column_device_view.cuh:663
cudf::nullate::YES
Definition: column_device_view.cuh:55
cudf::detail::column_device_view_base::nullable
CUDF_HOST_DEVICE bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
Definition: column_device_view.cuh:159
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::column_device_view::end
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
Definition: column_device_view.cuh:525
cudf::detail::column_device_view_base::_type
data_type _type
Element type.
Definition: column_device_view.cuh:264
cudf::column_device_view::child
column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:774
cudf::column_device_view::children
device_span< column_device_view const > children() const noexcept
Returns a span containing the children of this column.
Definition: column_device_view.cuh:782
cudf::detail::value_accessor
value accessor of column without null bitmask A unary functor returns scalar value at id....
Definition: column_device_view.cuh:1153
cudf::column_device_view::begin
const_iterator< T > begin() const
Return an iterator to the first element of the column.
Definition: column_device_view.cuh:507
cudf::detail::optional_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1196
cudf::detail::column_device_view_base::_null_mask
bitmask_type const * _null_mask
Definition: column_device_view.cuh:267
cudf::detail::column_device_view_base
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Definition: column_device_view.cuh:87
cudf::mutable_column_device_view::child
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:1017
cudf::detail::column_device_view_base::type
CUDF_HOST_DEVICE data_type type() const noexcept
Returns the element type.
Definition: column_device_view.cuh:148
cudf::detail::pair_rep_accessor
pair accessor of column with/without null bitmask A unary functor returns pair with representative sc...
Definition: column_device_view.cuh:1280
cudf::column_device_view::d_children
column_device_view * d_children
Definition: column_device_view.cuh:824
cudf::mutable_column_device_view::has_element_accessor
static constexpr bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
Definition: column_device_view.cuh:955
cudf::column_device_view::column_device_view
column_device_view(column_view column, void *h_ptr, void *d_ptr)
Creates an instance of this class using the specified host memory pointer (h_ptr) to store child obje...
CUDF_EXPECTS
#define CUDF_EXPECTS(cond, reason)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
Definition: error.hpp:91
cudf::column_device_view::pair_begin
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
Definition: column_device_view.cuh:633
cudf::column_device_view::destroy
void destroy()
Destroy the column_device_view object.
cudf::detail::value_accessor::value_accessor
value_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1160
cudf::mutable_column_device_view::set_mask_word
void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
Updates the specified bitmask word in the null_mask() with a new word.
Definition: column_device_view.cuh:1072
cudf::column_device_view::has_element_accessor
static constexpr bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
Definition: column_device_view.cuh:482
cudf::detail::column_device_view_base::_size
cudf::size_type _size
Number of elements.
Definition: column_device_view.cuh:265
cudf::detail::column_device_view_base::size
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
Definition: column_device_view.cuh:143
cudf::mutable_column_device_view::begin
iterator< T > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
Definition: column_device_view.cuh:990
cudf::distance
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
Definition: types.hpp:99
cudf::mutable_column_device_view::create
static std::unique_ptr< mutable_column_device_view, std::function< void(mutable_column_device_view *)> > create(mutable_column_view source_view, rmm::cuda_stream_view stream=rmm::cuda_stream_default)
Factory to construct a column view that is usable in device memory.
cudf::device_span
Definition: span.hpp:194
cudf::column_device_view::extent
static std::size_t extent(column_view const &source_view)
Return the size in bytes of the amount of memory needed to hold a device view of the specified column...
type_dispatcher.hpp
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
traits.hpp
cudf::detail::optional_accessor
optional accessor of a column
Definition: column_device_view.cuh:1195
cudf::detail::optional_accessor::optional_accessor
optional_accessor(column_device_view const &_col, Nullate with_nulls)
Constructor.
Definition: column_device_view.cuh:1204
column_view.hpp
column view class definitions
struct_view.hpp
Class definition for cudf::struct_view.
cudf::column_device_view::_num_children
size_type _num_children
The number of child columns.
Definition: column_device_view.cuh:828
cudf::detail::mutable_value_accessor::col
mutable_column_device_view col
mutable column view of column in device
Definition: column_device_view.cuh:1316
cudf::nullate::DYNAMIC::DYNAMIC
constexpr DYNAMIC(bool b) noexcept
Create a runtime nullate object.
Definition: column_device_view.cuh:69
cudf::detail::column_device_view_base::data
CUDF_HOST_DEVICE T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Definition: column_device_view.cuh:135
cudf::detail::column_device_view_base::is_valid_nocheck
bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
Definition: column_device_view.cuh:210
cudf::column_device_view::element
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:376
numeric
fixed_point and supporting types
Definition: fixed_point.hpp:34
cudf::word_index
constexpr CUDF_HOST_DEVICE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
Definition: bit.hpp:61
cudf::nullate::DYNAMIC::value
bool value
True if nulls are expected.
Definition: column_device_view.cuh:71
cudf::column_device_view::const_pair_iterator
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.
Definition: column_device_view.cuh:542