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 <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_EXPORT cudf {
48 
58 struct nullate {
59  struct YES : std::bool_constant<true> {};
60  struct NO : std::bool_constant<false> {};
66  struct DYNAMIC {
67  DYNAMIC() = delete;
76  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
83  constexpr operator bool() const noexcept { return value; }
84  bool value;
85  };
86 };
87 
88 namespace detail {
100 class alignas(16) column_device_view_base {
101  public:
102  column_device_view_base() = delete;
103  ~column_device_view_base() = default;
118 
135  template <typename T = void,
136  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
137  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
138  {
139  return static_cast<T const*>(_data);
140  }
141 
157  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
158  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
159  {
160  return head<T>() + _offset;
161  }
162 
168  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
169 
175  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
176 
186  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
187 
197  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
198  {
199  return _null_mask;
200  }
201 
208  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
209 
224  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
225  {
226  return not nullable() or is_valid_nocheck(element_index);
227  }
228 
241  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
242  {
243  return bit_is_set(_null_mask, offset() + element_index);
244  }
245 
259  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
260  {
261  return not is_valid(element_index);
262  }
263 
275  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
276  {
277  return not is_valid_nocheck(element_index);
278  }
279 
289  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
290  {
291  return null_mask()[word_index];
292  }
293 
294  protected:
295  data_type _type{type_id::EMPTY};
296  cudf::size_type _size{};
297  void const* _data{};
298  bitmask_type const* _null_mask{};
300  size_type _offset{};
302 
313  size_type size,
314  void const* data,
315  bitmask_type const* null_mask,
316  size_type offset)
317  : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset}
318  {
319  }
320 
321  template <typename C, typename T, typename = void>
322  struct has_element_accessor_impl : std::false_type {};
323 
324  template <typename C, typename T>
325  struct has_element_accessor_impl<
326  C,
327  T,
328  void_t<decltype(std::declval<C>().template element<T>(std::declval<size_type>()))>>
329  : std::true_type {};
330 };
331 // @cond
332 // Forward declaration
333 template <typename T>
334 struct value_accessor;
335 template <typename T, typename Nullate>
336 struct optional_accessor;
337 template <typename T, bool has_nulls>
338 struct pair_accessor;
339 template <typename T, bool has_nulls>
340 struct pair_rep_accessor;
341 template <typename T>
342 struct mutable_value_accessor;
343 // @endcond
344 } // namespace detail
345 
353  public:
354  column_device_view() = delete;
355  ~column_device_view() = default;
370 
380  column_device_view(column_view column, void* h_ptr, void* d_ptr);
381 
399  size_type size) const noexcept
400  {
401  return column_device_view{this->type(),
402  size,
403  this->head(),
404  this->null_mask(),
405  this->offset() + offset,
406  d_children,
407  this->num_child_columns()};
408  }
409 
427  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
428  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
429  {
430  return data<T>()[element_index];
431  }
432 
444  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
445  __device__ [[nodiscard]] T element(size_type element_index) const noexcept
446  {
447  size_type index = element_index + offset(); // account for this view's _offset
448  char const* d_strings = static_cast<char const*>(_data);
449  auto const offsets = d_children[strings_column_view::offsets_column_index];
450  auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
451  auto const offset = itr[index];
452  return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
453  }
454 
455  private:
461  struct index_element_fn {
462  template <typename IndexType,
463  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>)>
464  __device__ size_type operator()(column_device_view const& indices, size_type index)
465  {
466  return static_cast<size_type>(indices.element<IndexType>(index));
467  }
468 
469  template <typename IndexType,
470  typename... Args,
471  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>))>
472  __device__ size_type operator()(Args&&... args)
473  {
474  CUDF_UNREACHABLE("dictionary indices must be an unsigned integral type");
475  }
476  };
477 
478  public:
503  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
504  __device__ [[nodiscard]] T element(size_type element_index) const noexcept
505  {
506  size_type index = element_index + offset(); // account for this view's _offset
507  auto const indices = d_children[0];
508  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
509  }
510 
521  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
522  __device__ [[nodiscard]] T element(size_type element_index) const noexcept
523  {
524  using namespace numeric;
525  using rep = typename T::rep;
526  auto const scale = scale_type{_type.scale()};
527  return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
528  }
529 
536  template <typename T>
537  static constexpr bool has_element_accessor()
538  {
539  return has_element_accessor_impl<column_device_view, T>::value;
540  }
541 
543  using count_it = thrust::counting_iterator<size_type>;
547  template <typename T>
548  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
549 
565  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
566  [[nodiscard]] const_iterator<T> begin() const
567  {
569  }
570 
585  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
586  [[nodiscard]] const_iterator<T> end() const
587  {
588  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
589  }
590 
594  template <typename T, typename Nullate>
596  thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
597 
601  template <typename T, bool has_nulls>
603  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
604 
610  template <typename T, bool has_nulls>
612  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
613 
668  template <typename T,
669  typename Nullate,
670  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
671  auto optional_begin(Nullate has_nulls) const
672  {
675  }
676 
698  template <typename T,
699  bool has_nulls,
700  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
702  {
705  }
706 
730  template <typename T,
731  bool has_nulls,
732  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
734  {
737  }
738 
755  template <typename T,
756  typename Nullate,
757  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
758  auto optional_end(Nullate has_nulls) const
759  {
762  }
763 
775  template <typename T,
776  bool has_nulls,
777  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
779  {
782  }
783 
796  template <typename T,
797  bool has_nulls,
798  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
800  {
803  }
804 
823  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
825 
832  void destroy();
833 
841  static std::size_t extent(column_view const& source_view);
842 
849  [[nodiscard]] __device__ column_device_view child(size_type child_index) const noexcept
850  {
851  return d_children[child_index];
852  }
853 
859  [[nodiscard]] __device__ device_span<column_device_view const> children() const noexcept
860  {
861  return {d_children, static_cast<std::size_t>(_num_children)};
862  }
863 
869  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
870  {
871  return _num_children;
872  }
873 
874  private:
888  size_type size,
889  void const* data,
890  bitmask_type const* null_mask,
891  size_type offset,
892  column_device_view* children,
893  size_type num_children)
894  : column_device_view_base(type, size, data, null_mask, offset),
895  d_children(children),
896  _num_children(num_children)
897  {
898  }
899 
900  protected:
901  column_device_view* d_children{};
905  size_type _num_children{};
906 
918 };
919 
927  public:
928  mutable_column_device_view() = delete;
929  ~mutable_column_device_view() = default;
944 
955 
974  static std::unique_ptr<mutable_column_device_view,
975  std::function<void(mutable_column_device_view*)>>
978 
995  template <typename T = void,
996  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
997  CUDF_HOST_DEVICE T* head() const noexcept
998  {
999  return const_cast<T*>(detail::column_device_view_base::head<T>());
1000  }
1001 
1014  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1015  CUDF_HOST_DEVICE T* data() const noexcept
1016  {
1017  return const_cast<T*>(detail::column_device_view_base::data<T>());
1018  }
1019 
1034  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1035  __device__ [[nodiscard]] T& element(size_type element_index) const noexcept
1036  {
1037  return data<T>()[element_index];
1038  }
1039 
1046  template <typename T>
1047  static constexpr bool has_element_accessor()
1048  {
1049  return has_element_accessor_impl<mutable_column_device_view, T>::value;
1050  }
1051 
1060  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
1061  {
1062  return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
1063  }
1064 
1066  using count_it = thrust::counting_iterator<size_type>;
1070  template <typename T>
1071  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
1072 
1083  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1085  {
1087  }
1088 
1099  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1101  {
1102  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
1103  }
1104 
1111  [[nodiscard]] __device__ mutable_column_device_view child(size_type child_index) const noexcept
1112  {
1113  return d_children[child_index];
1114  }
1115 
1116 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1131  __device__ void set_valid(size_type element_index) const noexcept
1132  {
1133  return set_bit(null_mask(), element_index);
1134  }
1135 
1149  __device__ void set_null(size_type element_index) const noexcept
1150  {
1151  return clear_bit(null_mask(), element_index);
1152  }
1153 
1154 #endif
1155 
1166  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
1167  {
1168  null_mask()[word_index] = new_word;
1169  }
1170 
1179  static std::size_t extent(mutable_column_view source_view);
1180 
1187  void destroy();
1188 
1189  private:
1190  mutable_column_device_view* d_children{};
1194  size_type _num_children{};
1195 
1204  mutable_column_device_view(mutable_column_view source);
1205 };
1206 
1207 namespace detail {
1208 
1209 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1210 
1217 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
1218  size_type destination_word_index,
1219  size_type source_begin_bit,
1220  size_type source_end_bit)
1221 {
1222  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
1223  bitmask_type curr_word = source[source_word_index];
1224  bitmask_type next_word = 0;
1225  if (word_index(source_end_bit - 1) >
1226  word_index(source_begin_bit +
1227  destination_word_index * detail::size_in_bits<bitmask_type>())) {
1228  next_word = source[source_word_index + 1];
1229  }
1230  return __funnelshift_r(curr_word, next_word, source_begin_bit);
1231 }
1232 
1233 #endif
1234 
1249 template <typename T>
1252 
1258  value_accessor(column_device_view const& _col) : col{_col}
1259  {
1260  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1261  }
1262 
1268  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1269 };
1270 
1297 template <typename T, typename Nullate>
1300 
1307  optional_accessor(column_device_view const& _col, Nullate with_nulls)
1308  : col{_col}, has_nulls{with_nulls}
1309  {
1310  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1311  if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1312  }
1313 
1321  __device__ inline thrust::optional<T> operator()(cudf::size_type i) const
1322  {
1323  if (has_nulls) {
1324  return (col.is_valid_nocheck(i)) ? thrust::optional<T>{col.element<T>(i)}
1325  : thrust::optional<T>{thrust::nullopt};
1326  }
1327  return thrust::optional<T>{col.element<T>(i)};
1328  }
1329 
1330  Nullate has_nulls{};
1331 };
1332 
1352 template <typename T, bool has_nulls = false>
1355 
1361  pair_accessor(column_device_view const& _col) : col{_col}
1362  {
1363  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1364  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1365  }
1366 
1373  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
1374  {
1375  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1376  }
1377 };
1378 
1398 template <typename T, bool has_nulls = false>
1401 
1403 
1409  pair_rep_accessor(column_device_view const& _col) : col{_col}
1410  {
1411  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1412  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1413  }
1414 
1421  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1422  {
1423  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1424  }
1425 
1426  private:
1427  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
1428  __device__ inline auto get_rep(cudf::size_type i) const
1429  {
1430  return col.element<R>(i);
1431  }
1432 
1433  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
1434  __device__ inline auto get_rep(cudf::size_type i) const
1435  {
1436  return col.element<R>(i).value();
1437  }
1438 };
1439 
1451 template <typename T>
1454 
1461  {
1462  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1463  }
1464 
1471  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1472 };
1473 
1499 template <typename ColumnDeviceView, typename ColumnViewIterator>
1500 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1501  ColumnViewIterator child_end,
1502  void* h_ptr,
1503  void* d_ptr)
1504 {
1505  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1506  auto num_children = std::distance(child_begin, child_end);
1507  if (num_children > 0) {
1508  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1509  // struct objects in order for d_children to be used as an array.
1510  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1511  auto d_column = d_children;
1512 
1513  // Any child data is assigned past the end of this array: h_end and d_end.
1514  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1515  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1516  std::for_each(child_begin, child_end, [&](auto const& col) {
1517  // inplace-new each child into host memory
1518  new (h_column) ColumnDeviceView(col, h_end, d_end);
1519  h_column++; // advance to next child
1520  // update the pointers for holding this child column's child data
1521  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1522  h_end += col_child_data_size;
1523  d_end += col_child_data_size;
1524  });
1525  }
1526  return d_children;
1527 }
1528 
1529 } // namespace detail
1530 } // 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.
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.
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:48
Indicator for the logical data type of an element in a column.
Definition: types.hpp:243
constexpr 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...
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.
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=rmm::mr::get_current_device_resource())
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:37
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:50
Class definition for cudf::list_view.
cuDF interfaces
Definition: aggregation.hpp:35
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
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:328
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