All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Modules Pages
column_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2024, 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>
29 #include <cudf/utilities/span.hpp>
32 
33 #include <rmm/cuda_stream_view.hpp>
34 
35 #include <cuda/std/optional>
36 #include <cuda/std/type_traits>
37 #include <thrust/iterator/counting_iterator.h>
38 #include <thrust/iterator/transform_iterator.h>
39 #include <thrust/pair.h>
40 
41 #include <algorithm>
42 #include <type_traits>
43 
49 namespace CUDF_EXPORT cudf {
50 
60 struct nullate {
61  struct YES : cuda::std::bool_constant<true> {};
62  struct NO : cuda::std::bool_constant<false> {};
68  struct DYNAMIC {
69  DYNAMIC() = delete;
78  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
85  CUDF_HOST_DEVICE constexpr operator bool() const noexcept { return value; }
86  bool value;
87  };
88 };
89 
90 namespace detail {
102 class alignas(16) column_device_view_base {
103  public:
104  column_device_view_base() = delete;
105  ~column_device_view_base() = default;
120 
137  template <typename T = void,
138  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
139  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
140  {
141  return static_cast<T const*>(_data);
142  }
143 
159  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
160  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
161  {
162  return head<T>() + _offset;
163  }
164 
170  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
171 
177  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
178 
188  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
189 
199  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
200  {
201  return _null_mask;
202  }
203 
210  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
211 
226  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
227  {
228  return not nullable() or is_valid_nocheck(element_index);
229  }
230 
243  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
244  {
245  return bit_is_set(_null_mask, offset() + element_index);
246  }
247 
261  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
262  {
263  return not is_valid(element_index);
264  }
265 
277  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
278  {
279  return not is_valid_nocheck(element_index);
280  }
281 
291  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
292  {
293  return null_mask()[word_index];
294  }
295 
296  protected:
297  data_type _type{type_id::EMPTY};
298  cudf::size_type _size{};
299  void const* _data{};
300  bitmask_type const* _null_mask{};
302  size_type _offset{};
304 
315  size_type size,
316  void const* data,
317  bitmask_type const* null_mask,
318  size_type offset)
319  : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset}
320  {
321  }
322 
323  template <typename C, typename T, typename = void>
324  struct has_element_accessor_impl : cuda::std::false_type {};
325 
326  template <typename C, typename T>
327  struct has_element_accessor_impl<
328  C,
329  T,
330  void_t<decltype(cuda::std::declval<C>().template element<T>(cuda::std::declval<size_type>()))>>
331  : cuda::std::true_type {};
332 };
333 // @cond
334 // Forward declaration
335 template <typename T>
336 struct value_accessor;
337 template <typename T, typename Nullate>
338 struct optional_accessor;
339 template <typename T, bool has_nulls>
340 struct pair_accessor;
341 template <typename T, bool has_nulls>
342 struct pair_rep_accessor;
343 template <typename T>
344 struct mutable_value_accessor;
345 // @endcond
346 } // namespace detail
347 
355  public:
356  column_device_view() = delete;
357  ~column_device_view() = default;
372 
382  column_device_view(column_view column, void* h_ptr, void* d_ptr);
383 
401  size_type size) const noexcept
402  {
403  return column_device_view{this->type(),
404  size,
405  this->head(),
406  this->null_mask(),
407  this->offset() + offset,
408  d_children,
409  this->num_child_columns()};
410  }
411 
429  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
430  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
431  {
432  return data<T>()[element_index];
433  }
434 
446  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
447  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
448  {
449  size_type index = element_index + offset(); // account for this view's _offset
450  char const* d_strings = static_cast<char const*>(_data);
451  auto const offsets = d_children[strings_column_view::offsets_column_index];
452  auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
453  auto const offset = itr[index];
454  return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
455  }
456 
457  private:
463  struct index_element_fn {
464  template <typename IndexType,
465  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_signed_v<IndexType>)>
466  __device__ size_type operator()(column_device_view const& indices, size_type index)
467  {
468  return static_cast<size_type>(indices.element<IndexType>(index));
469  }
470 
471  template <typename IndexType,
472  typename... Args,
473  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_signed_v<IndexType>))>
474  __device__ size_type operator()(Args&&... args)
475  {
476  CUDF_UNREACHABLE("dictionary indices must be a signed integral type");
477  }
478  };
479 
480  public:
505  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
506  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
507  {
508  size_type index = element_index + offset(); // account for this view's _offset
509  auto const indices = d_children[0];
510  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
511  }
512 
523  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
524  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
525  {
526  using namespace numeric;
527  using rep = typename T::rep;
528  auto const scale = scale_type{_type.scale()};
529  return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
530  }
531 
538  template <typename T>
539  CUDF_HOST_DEVICE static constexpr bool has_element_accessor()
540  {
541  return has_element_accessor_impl<column_device_view, T>::value;
542  }
543 
545  using count_it = thrust::counting_iterator<size_type>;
549  template <typename T>
550  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
551 
567  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
568  [[nodiscard]] const_iterator<T> begin() const
569  {
571  }
572 
587  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
588  [[nodiscard]] const_iterator<T> end() const
589  {
590  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
591  }
592 
596  template <typename T, typename Nullate>
598  thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
599 
603  template <typename T, bool has_nulls>
605  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
606 
612  template <typename T, bool has_nulls>
614  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
615 
670  template <typename T,
671  typename Nullate,
672  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
673  auto optional_begin(Nullate has_nulls) const
674  {
677  }
678 
700  template <typename T,
701  bool has_nulls,
702  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
704  {
707  }
708 
732  template <typename T,
733  bool has_nulls,
734  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
736  {
739  }
740 
757  template <typename T,
758  typename Nullate,
759  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
760  auto optional_end(Nullate has_nulls) const
761  {
764  }
765 
777  template <typename T,
778  bool has_nulls,
779  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
781  {
784  }
785 
798  template <typename T,
799  bool has_nulls,
800  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
802  {
805  }
806 
825  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
827 
834  void destroy();
835 
843  static std::size_t extent(column_view const& source_view);
844 
851  [[nodiscard]] __device__ column_device_view child(size_type child_index) const noexcept
852  {
853  return d_children[child_index];
854  }
855 
861  [[nodiscard]] __device__ device_span<column_device_view const> children() const noexcept
862  {
863  return {d_children, static_cast<std::size_t>(_num_children)};
864  }
865 
871  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
872  {
873  return _num_children;
874  }
875 
876  private:
890  size_type size,
891  void const* data,
892  bitmask_type const* null_mask,
893  size_type offset,
894  column_device_view* children,
895  size_type num_children)
896  : column_device_view_base(type, size, data, null_mask, offset),
897  d_children(children),
898  _num_children(num_children)
899  {
900  }
901 
902  protected:
903  column_device_view* d_children{};
907  size_type _num_children{};
908 
920 };
921 
929  public:
930  mutable_column_device_view() = delete;
931  ~mutable_column_device_view() = default;
946 
957 
976  static std::unique_ptr<mutable_column_device_view,
977  std::function<void(mutable_column_device_view*)>>
980 
997  template <typename T = void,
998  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
999  CUDF_HOST_DEVICE T* head() const noexcept
1000  {
1001  return const_cast<T*>(detail::column_device_view_base::head<T>());
1002  }
1003 
1016  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1017  CUDF_HOST_DEVICE T* data() const noexcept
1018  {
1019  return const_cast<T*>(detail::column_device_view_base::data<T>());
1020  }
1021 
1036  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1037  [[nodiscard]] __device__ T& element(size_type element_index) const noexcept
1038  {
1039  return data<T>()[element_index];
1040  }
1041 
1048  template <typename T>
1049  CUDF_HOST_DEVICE static constexpr bool has_element_accessor()
1050  {
1051  return has_element_accessor_impl<mutable_column_device_view, T>::value;
1052  }
1053 
1062  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
1063  {
1064  return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
1065  }
1066 
1068  using count_it = thrust::counting_iterator<size_type>;
1072  template <typename T>
1073  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
1074 
1085  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1087  {
1089  }
1090 
1101  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1103  {
1104  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
1105  }
1106 
1113  [[nodiscard]] __device__ mutable_column_device_view child(size_type child_index) const noexcept
1114  {
1115  return d_children[child_index];
1116  }
1117 
1118 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1133  __device__ void set_valid(size_type element_index) const noexcept
1134  {
1135  return set_bit(null_mask(), element_index);
1136  }
1137 
1151  __device__ void set_null(size_type element_index) const noexcept
1152  {
1153  return clear_bit(null_mask(), element_index);
1154  }
1155 
1156 #endif
1157 
1168  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
1169  {
1170  null_mask()[word_index] = new_word;
1171  }
1172 
1181  static std::size_t extent(mutable_column_view source_view);
1182 
1189  void destroy();
1190 
1191  private:
1192  mutable_column_device_view* d_children{};
1196  size_type _num_children{};
1197 
1206  mutable_column_device_view(mutable_column_view source);
1207 };
1208 
1209 namespace detail {
1210 
1211 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1212 
1219 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
1220  size_type destination_word_index,
1221  size_type source_begin_bit,
1222  size_type source_end_bit)
1223 {
1224  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
1225  bitmask_type curr_word = source[source_word_index];
1226  bitmask_type next_word = 0;
1227  if (word_index(source_end_bit - 1) >
1228  word_index(source_begin_bit +
1229  destination_word_index * detail::size_in_bits<bitmask_type>())) {
1230  next_word = source[source_word_index + 1];
1231  }
1232  return __funnelshift_r(curr_word, next_word, source_begin_bit);
1233 }
1234 
1235 #endif
1236 
1251 template <typename T>
1254 
1260  value_accessor(column_device_view const& _col) : col{_col}
1261  {
1262  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1263  }
1264 
1270  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1271 };
1272 
1299 template <typename T, typename Nullate>
1302 
1309  optional_accessor(column_device_view const& _col, Nullate with_nulls)
1310  : col{_col}, has_nulls{with_nulls}
1311  {
1312  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1313  if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1314  }
1315 
1323  __device__ inline cuda::std::optional<T> operator()(cudf::size_type i) const
1324  {
1325  if (has_nulls) {
1326  return (col.is_valid_nocheck(i)) ? cuda::std::optional<T>{col.element<T>(i)}
1327  : cuda::std::optional<T>{cuda::std::nullopt};
1328  }
1329  return cuda::std::optional<T>{col.element<T>(i)};
1330  }
1331 
1332  Nullate has_nulls{};
1333 };
1334 
1354 template <typename T, bool has_nulls = false>
1357 
1363  pair_accessor(column_device_view const& _col) : col{_col}
1364  {
1365  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1366  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1367  }
1368 
1375  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
1376  {
1377  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1378  }
1379 };
1380 
1400 template <typename T, bool has_nulls = false>
1403 
1405 
1411  pair_rep_accessor(column_device_view const& _col) : col{_col}
1412  {
1413  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1414  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1415  }
1416 
1423  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1424  {
1425  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1426  }
1427 
1428  private:
1429  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
1430  [[nodiscard]] __device__ inline auto get_rep(cudf::size_type i) const
1431  {
1432  return col.element<R>(i);
1433  }
1434 
1435  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
1436  [[nodiscard]] __device__ inline auto get_rep(cudf::size_type i) const
1437  {
1438  return col.element<R>(i).value();
1439  }
1440 };
1441 
1453 template <typename T>
1456 
1463  {
1464  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1465  }
1466 
1473  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1474 };
1475 
1501 template <typename ColumnDeviceView, typename ColumnViewIterator>
1502 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1503  ColumnViewIterator child_end,
1504  void* h_ptr,
1505  void* d_ptr)
1506 {
1507  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1508  auto num_children = std::distance(child_begin, child_end);
1509  if (num_children > 0) {
1510  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1511  // struct objects in order for d_children to be used as an array.
1512  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1513  auto d_column = d_children;
1514 
1515  // Any child data is assigned past the end of this array: h_end and d_end.
1516  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1517  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1518  std::for_each(child_begin, child_end, [&](auto const& col) {
1519  // inplace-new each child into host memory
1520  new (h_column) ColumnDeviceView(col, h_end, d_end);
1521  h_column++; // advance to next child
1522  // update the pointers for holding this child column's child data
1523  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1524  h_end += col_child_data_size;
1525  d_end += col_child_data_size;
1526  });
1527  }
1528  return d_children;
1529 }
1530 
1531 } // namespace detail
1532 } // namespace CUDF_EXPORT 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.
static constexpr CUDF_HOST_DEVICE bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
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.
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.
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.
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:47
Indicator for the logical data type of an element in a column.
Definition: types.hpp:243
constexpr CUDF_HOST_DEVICE type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:287
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.
CUDF_HOST_DEVICE column_device_view_base(data_type type, size_type size, void const *data, bitmask_type const *null_mask, size_type offset)
Constructs a column with the specified type, size, data, nullmask and offset.
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.
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
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...
static constexpr CUDF_HOST_DEVICE bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
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.
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
CUDF_HOST_DEVICE bitmask_type * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
CUDF_HOST_DEVICE T * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
CUDF_HOST_DEVICE T * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
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
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.
rmm::cuda_stream_view const get_default_stream()
Get the current default stream.
scale_type
The scale type for fixed_point.
Definition: fixed_point.hpp:43
std::unique_ptr< cudf::column > is_valid(cudf::column_view const &input, rmm::cuda_stream_view stream=cudf::get_default_stream(), rmm::device_async_resource_ref mr=cudf::get_current_device_resource_ref())
Creates a column of type_id::BOOL8 elements where for every element in input true indicates the value...
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:74
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:128
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:178
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:95
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
Definition: types.hpp:96
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
Definition: types.hpp:110
void void_t
Utility metafunction that maps a sequence of any types to the type void.
Definition: traits.hpp:35
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:48
Class definition for cudf::list_view.
cuDF interfaces
Definition: host_udf.hpp:39
bool nullable(table_view const &view)
Returns True if any of the columns in the table is nullable. (not entire hierarchy)
bool has_nulls(table_view const &view)
Returns True if the table has nulls in any of its columns.
fixed_point and supporting types
Definition: fixed_point.hpp:33
APIs for spans.
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
cuda::std::optional< T > operator()(cudf::size_type i) const
Returns a cuda::std::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:346
A strongly typed wrapper for indices in a DICTIONARY type column.
Definition: dictionary.hpp:49
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.
#define CUDF_HOST_DEVICE
Indicates that the function or method is usable on host and device.
Definition: types.hpp:32