column_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2023, 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/offsets_iterator.cuh>
20 #include <cudf/detail/utilities/alignment.hpp>
22 #include <cudf/lists/list_view.hpp>
23 #include <cudf/strings/string_view.cuh>
26 #include <cudf/types.hpp>
27 #include <cudf/utilities/bit.hpp>
28 #include <cudf/utilities/default_stream.hpp>
29 #include <cudf/utilities/span.hpp>
32 
33 #include <rmm/cuda_stream_view.hpp>
34 
35 #include <thrust/iterator/counting_iterator.h>
36 #include <thrust/iterator/transform_iterator.h>
37 #include <thrust/optional.h>
38 #include <thrust/pair.h>
39 
40 #include <algorithm>
41 
47 namespace cudf {
48 
56 struct nullate {
57  struct YES : std::bool_constant<true> {};
58  struct NO : std::bool_constant<false> {};
64  struct DYNAMIC {
65  DYNAMIC() = delete;
74  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
81  constexpr operator bool() const noexcept { return value; }
82  bool value;
83  };
84 };
85 
86 namespace detail {
98 class alignas(16) column_device_view_base {
99  public:
100  column_device_view_base() = delete;
101  ~column_device_view_base() = default;
116 
133  template <typename T = void,
134  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
135  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
136  {
137  return static_cast<T const*>(_data);
138  }
139 
155  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
156  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
157  {
158  return head<T>() + _offset;
159  }
160 
166  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
167 
173  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
174 
184  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
185 
195  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
196  {
197  return _null_mask;
198  }
199 
206  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
207 
222  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
223  {
224  return not nullable() or is_valid_nocheck(element_index);
225  }
226 
239  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
240  {
241  return bit_is_set(_null_mask, offset() + element_index);
242  }
243 
257  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
258  {
259  return not is_valid(element_index);
260  }
261 
273  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
274  {
275  return not is_valid_nocheck(element_index);
276  }
277 
287  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
288  {
289  return null_mask()[word_index];
290  }
291 
292  protected:
295  void const* _data{};
300 
311  size_type size,
312  void const* data,
313  bitmask_type const* null_mask,
316  {
317  }
318 
319  template <typename C, typename T, typename = void>
320  struct has_element_accessor_impl : std::false_type {};
321 
322  template <typename C, typename T>
323  struct has_element_accessor_impl<
324  C,
325  T,
326  void_t<decltype(std::declval<C>().template element<T>(std::declval<size_type>()))>>
327  : std::true_type {};
328 };
329 // @cond
330 // Forward declaration
331 template <typename T>
332 struct value_accessor;
333 template <typename T, typename Nullate>
334 struct optional_accessor;
335 template <typename T, bool has_nulls>
336 struct pair_accessor;
337 template <typename T, bool has_nulls>
338 struct pair_rep_accessor;
339 template <typename T>
340 struct mutable_value_accessor;
341 // @endcond
342 } // namespace detail
343 
351  public:
352  column_device_view() = delete;
353  ~column_device_view() = default;
368 
378  column_device_view(column_view column, void* h_ptr, void* d_ptr);
379 
396  [[nodiscard]] CUDF_HOST_DEVICE column_device_view slice(size_type offset,
397  size_type size) const noexcept
398  {
399  return column_device_view{this->type(),
400  size,
401  this->head(),
402  this->null_mask(),
403  this->offset() + offset,
404  d_children,
405  this->num_child_columns()};
406  }
407 
425  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
426  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
427  {
428  return data<T>()[element_index];
429  }
430 
442  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
443  __device__ T element(size_type element_index) const noexcept
444  {
445  size_type index = element_index + offset(); // account for this view's _offset
446  char const* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
448  auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
449  auto const offset = itr[index];
450  return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
451  }
452 
453  private:
459  struct index_element_fn {
460  template <typename IndexType,
461  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>)>
462  __device__ size_type operator()(column_device_view const& indices, size_type index)
463  {
464  return static_cast<size_type>(indices.element<IndexType>(index));
465  }
466 
467  template <typename IndexType,
468  typename... Args,
469  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>))>
470  __device__ size_type operator()(Args&&... args)
471  {
472  CUDF_UNREACHABLE("dictionary indices must be an unsigned integral type");
473  }
474  };
475 
476  public:
501  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
502  __device__ T element(size_type element_index) const noexcept
503  {
504  size_type index = element_index + offset(); // account for this view's _offset
505  auto const indices = d_children[0];
506  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
507  }
508 
519  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
520  __device__ T element(size_type element_index) const noexcept
521  {
522  using namespace numeric;
523  using rep = typename T::rep;
524  auto const scale = scale_type{_type.scale()};
525  return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
526  }
527 
534  template <typename T>
535  static constexpr bool has_element_accessor()
536  {
537  return has_element_accessor_impl<column_device_view, T>::value;
538  }
539 
541  using count_it = thrust::counting_iterator<size_type>;
545  template <typename T>
546  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
547 
563  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
564  [[nodiscard]] const_iterator<T> begin() const
565  {
567  }
568 
583  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
584  [[nodiscard]] const_iterator<T> end() const
585  {
587  }
588 
592  template <typename T, typename Nullate>
594  thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
595 
599  template <typename T, bool has_nulls>
601  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
602 
608  template <typename T, bool has_nulls>
610  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
611 
666  template <typename T,
667  typename Nullate,
668  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
669  auto optional_begin(Nullate has_nulls) const
670  {
673  }
674 
696  template <typename T,
697  bool has_nulls,
698  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
700  {
703  }
704 
728  template <typename T,
729  bool has_nulls,
730  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
732  {
735  }
736 
753  template <typename T,
754  typename Nullate,
755  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
756  auto optional_end(Nullate has_nulls) const
757  {
760  }
761 
773  template <typename T,
774  bool has_nulls,
775  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
777  {
780  }
781 
794  template <typename T,
795  bool has_nulls,
796  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
798  {
801  }
802 
821  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
823 
830  void destroy();
831 
839  static std::size_t extent(column_view const& source_view);
840 
847  [[nodiscard]] __device__ column_device_view child(size_type child_index) const noexcept
848  {
849  return d_children[child_index];
850  }
851 
857  [[nodiscard]] __device__ device_span<column_device_view const> children() const noexcept
858  {
860  }
861 
867  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
868  {
869  return _num_children;
870  }
871 
872  private:
885  CUDF_HOST_DEVICE column_device_view(data_type type,
886  size_type size,
887  void const* data,
888  bitmask_type const* null_mask,
891  size_type num_children)
892  : column_device_view_base(type, size, data, null_mask, offset),
894  _num_children(num_children)
895  {
896  }
897 
898  protected:
904 
916 };
917 
925  public:
926  mutable_column_device_view() = delete;
927  ~mutable_column_device_view() = default;
942 
953 
972  static std::unique_ptr<mutable_column_device_view,
973  std::function<void(mutable_column_device_view*)>>
976 
993  template <typename T = void,
994  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
995  CUDF_HOST_DEVICE T* head() const noexcept
996  {
997  return const_cast<T*>(detail::column_device_view_base::head<T>());
998  }
999 
1012  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1013  CUDF_HOST_DEVICE T* data() const noexcept
1014  {
1015  return const_cast<T*>(detail::column_device_view_base::data<T>());
1016  }
1017 
1032  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1033  __device__ T& element(size_type element_index) const noexcept
1034  {
1035  return data<T>()[element_index];
1036  }
1037 
1044  template <typename T>
1045  static constexpr bool has_element_accessor()
1046  {
1047  return has_element_accessor_impl<mutable_column_device_view, T>::value;
1048  }
1049 
1058  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
1059  {
1061  }
1062 
1064  using count_it = thrust::counting_iterator<size_type>;
1068  template <typename T>
1069  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
1070 
1081  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1083  {
1085  }
1086 
1097  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1099  {
1101  }
1102 
1109  [[nodiscard]] __device__ mutable_column_device_view child(size_type child_index) const noexcept
1110  {
1111  return d_children[child_index];
1112  }
1113 
1114 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1129  __device__ void set_valid(size_type element_index) const noexcept
1130  {
1131  return set_bit(null_mask(), element_index);
1132  }
1133 
1147  __device__ void set_null(size_type element_index) const noexcept
1148  {
1149  return clear_bit(null_mask(), element_index);
1150  }
1151 
1152 #endif
1153 
1164  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
1165  {
1166  null_mask()[word_index] = new_word;
1167  }
1168 
1177  static std::size_t extent(mutable_column_view source_view);
1178 
1185  void destroy();
1186 
1187  private:
1193 
1202  mutable_column_device_view(mutable_column_view source);
1203 };
1204 
1205 namespace detail {
1206 
1207 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1208 
1215 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
1216  size_type destination_word_index,
1217  size_type source_begin_bit,
1218  size_type source_end_bit)
1219 {
1220  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
1221  bitmask_type curr_word = source[source_word_index];
1222  bitmask_type next_word = 0;
1223  if (word_index(source_end_bit - 1) >
1224  word_index(source_begin_bit +
1225  destination_word_index * detail::size_in_bits<bitmask_type>())) {
1226  next_word = source[source_word_index + 1];
1227  }
1228  return __funnelshift_r(curr_word, next_word, source_begin_bit);
1229 }
1230 
1231 #endif
1232 
1247 template <typename T>
1250 
1256  value_accessor(column_device_view const& _col) : col{_col}
1257  {
1258  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1259  }
1260 
1266  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1267 };
1268 
1295 template <typename T, typename Nullate>
1298 
1305  optional_accessor(column_device_view const& _col, Nullate with_nulls)
1306  : col{_col}, has_nulls{with_nulls}
1307  {
1308  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1309  if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1310  }
1311 
1319  __device__ inline thrust::optional<T> operator()(cudf::size_type i) const
1320  {
1321  if (has_nulls) {
1322  return (col.is_valid_nocheck(i)) ? thrust::optional<T>{col.element<T>(i)}
1323  : thrust::optional<T>{thrust::nullopt};
1324  }
1325  return thrust::optional<T>{col.element<T>(i)};
1326  }
1327 
1328  Nullate has_nulls{};
1329 };
1330 
1350 template <typename T, bool has_nulls = false>
1353 
1359  pair_accessor(column_device_view const& _col) : col{_col}
1360  {
1361  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1362  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1363  }
1364 
1371  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
1372  {
1373  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1374  }
1375 };
1376 
1396 template <typename T, bool has_nulls = false>
1399 
1401 
1407  pair_rep_accessor(column_device_view const& _col) : col{_col}
1408  {
1409  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1410  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1411  }
1412 
1419  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1420  {
1421  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1422  }
1423 
1424  private:
1425  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
1426  __device__ inline auto get_rep(cudf::size_type i) const
1427  {
1428  return col.element<R>(i);
1429  }
1430 
1431  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
1432  __device__ inline auto get_rep(cudf::size_type i) const
1433  {
1434  return col.element<R>(i).value();
1435  }
1436 };
1437 
1449 template <typename T>
1452 
1459  {
1460  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1461  }
1462 
1469  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1470 };
1471 
1497 template <typename ColumnDeviceView, typename ColumnViewIterator>
1498 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1499  ColumnViewIterator child_end,
1500  void* h_ptr,
1501  void* d_ptr)
1502 {
1503  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1504  auto num_children = std::distance(child_begin, child_end);
1505  if (num_children > 0) {
1506  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1507  // struct objects in order for d_children to be used as an array.
1508  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1509  auto d_column = d_children;
1510 
1511  // Any child data is assigned past the end of this array: h_end and d_end.
1512  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1513  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1514  std::for_each(child_begin, child_end, [&](auto const& col) {
1515  // inplace-new each child into host memory
1516  new (h_column) ColumnDeviceView(col, h_end, d_end);
1517  h_column++; // advance to next child
1518  // update the pointers for holding this child column's child data
1519  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1520  h_end += col_child_data_size;
1521  d_end += col_child_data_size;
1522  });
1523  }
1524  return d_children;
1525 }
1526 
1527 } // namespace detail
1528 } // namespace cudf
Utilities for bit and bitmask operations.
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
static std::unique_ptr< column_device_view, std::function< void(column_device_view *)> > create(column_view source_view, rmm::cuda_stream_view stream=cudf::get_default_stream())
Factory to construct a column view that is usable in device memory.
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
column_device_view child(size_type child_index) const noexcept
Returns the specified child.
void destroy()
Destroy the column_device_view object.
column_device_view & operator=(column_device_view &&)=default
Move assignment operator.
thrust::counting_iterator< size_type > count_it
Counting iterator.
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
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.
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
static constexpr bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
column_device_view & operator=(column_device_view const &)=default
Copy assignment operator.
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.
thrust::transform_iterator< detail::value_accessor< T >, count_it > const_iterator
Iterator for navigating this column.
auto optional_end(Nullate has_nulls) const
Return an optional iterator to the element following the last element of the column.
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.
size_type _num_children
The number of child columns.
const_pair_rep_iterator< T, has_nulls > pair_rep_begin() const
Return a pair iterator to the first element of the column.
thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > const_optional_iterator
Optional iterator for navigating this column.
column_device_view * d_children
CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
Returns the number of child columns.
auto optional_begin(Nullate has_nulls) const
Return an optional iterator to the first element of the column.
column_device_view(column_device_view const &)=default
Copy constructor.
column_device_view(column_device_view &&)=default
Move constructor.
const_iterator< T > begin() const
Return an iterator to the first element of the column.
thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > const_pair_rep_iterator
Pair rep iterator for navigating this column.
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...
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...
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
device_span< column_device_view const > children() const noexcept
Returns a span containing the children of this column.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
A container of nullable device data as a column of elements.
Definition: column.hpp:48
Indicator for the logical data type of an element in a column.
Definition: types.hpp:227
constexpr int32_t scale() const noexcept
Returns the scale (for fixed_point types)
Definition: types.hpp:278
constexpr type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:271
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
CUDF_HOST_DEVICE T const * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
CUDF_HOST_DEVICE data_type type() const noexcept
Returns the element type.
column_device_view_base & operator=(column_device_view_base &&)=default
Move assignment operator.
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the specified bitmask word from the null_mask().
column_device_view_base(column_device_view_base &&)=default
Move constructor.
column_device_view_base(column_device_view_base const &)=default
Copy constructor.
void const * _data
Pointer to device memory containing elements.
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
cudf::size_type _size
Number of elements.
bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
CUDF_HOST_DEVICE T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
column_device_view_base & operator=(column_device_view_base const &)=default
Copy assignment operator.
bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
CUDF_HOST_DEVICE size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
CUDF_HOST_DEVICE bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
void destroy()
Destroy the mutable_column_device_view object.
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.
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...
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...
thrust::counting_iterator< size_type > count_it
Counting iterator.
static constexpr bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
mutable_column_device_view(mutable_column_device_view const &)=default
Copy constructor.
T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
iterator< T > end()
Return one past the last element after underlying data is casted to the specified type.
mutable_column_device_view & operator=(mutable_column_device_view &&)=default
Move assignment operator.
mutable_column_device_view(mutable_column_device_view &&)=default
Move constructor.
mutable_column_device_view & operator=(mutable_column_device_view const &)=default
Copy assignment operator.
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=cudf::get_default_stream())
Factory to construct a column view that is usable in device memory.
thrust::transform_iterator< detail::mutable_value_accessor< T >, count_it > iterator
Iterator for navigating this column.
iterator< T > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
A non-owning, mutable view of device data as a column of elements, some of which may be null as indic...
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
Definition: string_view.hpp:44
static constexpr size_type chars_column_index
Child index of the characters column.
static constexpr size_type offsets_column_index
Child index of the offsets column.
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...
column view class definitions
Class definition for fixed point data type.
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:72
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:126
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...
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
#define CUDF_EXPECTS(...)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
Definition: error.hpp:170
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:80
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
Definition: types.hpp:81
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
Definition: types.hpp:94
void void_t
Utility metafunction that maps a sequence of any types to the type void.
Definition: traits.hpp:37
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:50
@ EMPTY
Always null with no underlying data.
Class definition for cudf::list_view.
cuDF interfaces
Definition: aggregation.hpp:34
rmm::cuda_stream_view const get_default_stream()
Get the current default stream.
bool has_nulls(table_view const &view)
Returns True if the table has nulls in any of its columns.
Definition: table_view.hpp:318
fixed_point and supporting types
Definition: fixed_point.hpp:32
scale_type
The scale type for fixed_point.
Definition: fixed_point.hpp:35
Class definition for cudf::strings_column_view.
Class definition for cudf::struct_view.
Mutable value accessor of column without null bitmask.
T & operator()(cudf::size_type i)
Accessor.
mutable_value_accessor(mutable_column_device_view &_col)
Constructor.
mutable_column_device_view col
mutable column view of column in device
optional accessor of a column
thrust::optional< T > operator()(cudf::size_type i) const
Returns a thrust::optional of column[i].
column_device_view const col
column view of column in device
optional_accessor(column_device_view const &_col, Nullate with_nulls)
Constructor.
pair accessor of column with/without null bitmask
column_device_view const col
column view of column in device
thrust::pair< T, bool > operator()(cudf::size_type i) const
Pair accessor.
pair_accessor(column_device_view const &_col)
constructor
pair accessor of column with/without null bitmask
thrust::pair< rep_type, bool > operator()(cudf::size_type i) const
Pair accessor.
column_device_view const col
column view of column in device
device_storage_type_t< T > rep_type
representation type
pair_rep_accessor(column_device_view const &_col)
constructor
value accessor of column without null bitmask
column_device_view const col
column view of column in device
value_accessor(column_device_view const &_col)
constructor
T operator()(cudf::size_type i) const
Returns the value of element at index i
Device version of C++20 std::span with reduced feature set.
Definition: span.hpp:277
A strongly typed wrapper for indices in a DICTIONARY type column.
Definition: dictionary.hpp:48
nullate::DYNAMIC defers the determination of nullability to run time rather than compile time....
bool value
True if nulls are expected.
constexpr DYNAMIC(bool b) noexcept
Create a runtime nullate object.
Indicates the presence of nulls at compile-time or runtime.
Helper struct for constructing fixed_point when value is already shifted.
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
Type declarations for libcudf.