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-2025, 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 <functional>
43 #include <type_traits>
44 
50 namespace CUDF_EXPORT cudf {
51 
61 struct nullate {
62  struct YES : cuda::std::true_type {};
63  struct NO : cuda::std::false_type {};
69  struct DYNAMIC {
70  DYNAMIC() = delete;
79  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
86  CUDF_HOST_DEVICE constexpr operator bool() const noexcept { return value; }
87  bool value;
88  };
89 };
90 
91 namespace detail {
103 class alignas(16) column_device_view_base {
104  public:
105  column_device_view_base() = delete;
106  ~column_device_view_base() = default;
121 
138  template <typename T = void,
139  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
140  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
141  {
142  return static_cast<T const*>(_data);
143  }
144 
160  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
161  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
162  {
163  return head<T>() + _offset;
164  }
165 
171  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
172 
178  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
179 
189  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
190 
200  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
201  {
202  return _null_mask;
203  }
204 
211  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
212 
227  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
228  {
229  return not nullable() or is_valid_nocheck(element_index);
230  }
231 
244  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
245  {
246  return bit_is_set(_null_mask, offset() + element_index);
247  }
248 
262  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
263  {
264  return not is_valid(element_index);
265  }
266 
278  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
279  {
280  return not is_valid_nocheck(element_index);
281  }
282 
292  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
293  {
294  return null_mask()[word_index];
295  }
296 
297  protected:
298  data_type _type{type_id::EMPTY};
299  cudf::size_type _size{};
300  void const* _data{};
301  bitmask_type const* _null_mask{};
303  size_type _offset{};
305 
316  size_type size,
317  void const* data,
318  bitmask_type const* null_mask,
319  size_type offset)
320  : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset}
321  {
322  }
323 
324  template <typename C, typename T, typename = void>
325  struct has_element_accessor_impl : cuda::std::false_type {};
326 
327  template <typename C, typename T>
328  struct has_element_accessor_impl<
329  C,
330  T,
331  void_t<decltype(cuda::std::declval<C>().template element<T>(cuda::std::declval<size_type>()))>>
332  : cuda::std::true_type {};
333 };
334 // @cond
335 // Forward declaration
336 template <typename T>
337 struct value_accessor;
338 template <typename T, typename Nullate>
339 struct optional_accessor;
340 template <typename T, bool has_nulls>
341 struct pair_accessor;
342 template <typename T, bool has_nulls>
343 struct pair_rep_accessor;
344 template <typename T>
345 struct mutable_value_accessor;
346 // @endcond
347 } // namespace detail
348 
356  public:
357  column_device_view() = delete;
358  ~column_device_view() = default;
373 
383  column_device_view(column_view column, void* h_ptr, void* d_ptr);
384 
402  size_type size) const noexcept
403  {
404  return column_device_view{this->type(),
405  size,
406  this->head(),
407  this->null_mask(),
408  this->offset() + offset,
409  d_children,
410  this->num_child_columns()};
411  }
412 
430  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
431  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
432  {
433  return data<T>()[element_index];
434  }
435 
447  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
448  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
449  {
450  size_type index = element_index + offset(); // account for this view's _offset
451  char const* d_strings = static_cast<char const*>(_data);
452  auto const offsets = d_children[strings_column_view::offsets_column_index];
453  auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
454  auto const offset = itr[index];
455  return string_view{d_strings + offset, static_cast<cudf::size_type>(itr[index + 1] - offset)};
456  }
457 
458  private:
464  struct index_element_fn {
465  template <typename IndexType,
466  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_signed_v<IndexType>)>
467  __device__ size_type operator()(column_device_view const& indices, size_type index)
468  {
469  return static_cast<size_type>(indices.element<IndexType>(index));
470  }
471 
472  template <typename IndexType,
473  typename... Args,
474  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_signed_v<IndexType>))>
475  __device__ size_type operator()(Args&&... args)
476  {
477  CUDF_UNREACHABLE("dictionary indices must be a signed integral type");
478  }
479  };
480 
481  public:
506  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
507  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
508  {
509  size_type index = element_index + offset(); // account for this view's _offset
510  auto const indices = d_children[0];
511  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
512  }
513 
524  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
525  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
526  {
527  using namespace numeric;
528  using rep = typename T::rep;
529  auto const scale = scale_type{_type.scale()};
530  return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
531  }
532 
539  template <typename T>
540  CUDF_HOST_DEVICE static constexpr bool has_element_accessor()
541  {
542  return has_element_accessor_impl<column_device_view, T>::value;
543  }
544 
546  using count_it = thrust::counting_iterator<size_type>;
550  template <typename T>
551  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
552 
568  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
569  [[nodiscard]] const_iterator<T> begin() const
570  {
572  }
573 
588  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
589  [[nodiscard]] const_iterator<T> end() const
590  {
591  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
592  }
593 
597  template <typename T, typename Nullate>
599  thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
600 
604  template <typename T, bool has_nulls>
606  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
607 
613  template <typename T, bool has_nulls>
615  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
616 
671  template <typename T,
672  typename Nullate,
673  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
674  auto optional_begin(Nullate has_nulls) const
675  {
678  }
679 
701  template <typename T,
702  bool has_nulls,
703  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
705  {
708  }
709 
733  template <typename T,
734  bool has_nulls,
735  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
737  {
740  }
741 
758  template <typename T,
759  typename Nullate,
760  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
761  auto optional_end(Nullate has_nulls) const
762  {
765  }
766 
778  template <typename T,
779  bool has_nulls,
780  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
782  {
785  }
786 
799  template <typename T,
800  bool has_nulls,
801  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
803  {
806  }
807 
826  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
828 
835  void destroy();
836 
844  static std::size_t extent(column_view const& source_view);
845 
852  [[nodiscard]] __device__ column_device_view child(size_type child_index) const noexcept
853  {
854  return d_children[child_index];
855  }
856 
862  [[nodiscard]] __device__ device_span<column_device_view const> children() const noexcept
863  {
864  return {d_children, static_cast<std::size_t>(_num_children)};
865  }
866 
872  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
873  {
874  return _num_children;
875  }
876 
877  private:
891  size_type size,
892  void const* data,
893  bitmask_type const* null_mask,
894  size_type offset,
895  column_device_view* children,
896  size_type num_children)
897  : column_device_view_base(type, size, data, null_mask, offset),
898  d_children(children),
899  _num_children(num_children)
900  {
901  }
902 
903  protected:
904  column_device_view* d_children{};
908  size_type _num_children{};
909 
921 };
922 
930  public:
931  mutable_column_device_view() = delete;
932  ~mutable_column_device_view() = default;
947 
958 
977  static std::unique_ptr<mutable_column_device_view,
978  std::function<void(mutable_column_device_view*)>>
981 
998  template <typename T = void,
999  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
1000  CUDF_HOST_DEVICE T* head() const noexcept
1001  {
1002  return const_cast<T*>(detail::column_device_view_base::head<T>());
1003  }
1004 
1017  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1018  CUDF_HOST_DEVICE T* data() const noexcept
1019  {
1020  return const_cast<T*>(detail::column_device_view_base::data<T>());
1021  }
1022 
1037  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1038  [[nodiscard]] __device__ T& element(size_type element_index) const noexcept
1039  {
1040  return data<T>()[element_index];
1041  }
1042 
1049  template <typename T>
1050  CUDF_HOST_DEVICE static constexpr bool has_element_accessor()
1051  {
1052  return has_element_accessor_impl<mutable_column_device_view, T>::value;
1053  }
1054 
1063  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept
1064  {
1065  return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
1066  }
1067 
1069  using count_it = thrust::counting_iterator<size_type>;
1073  template <typename T>
1074  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
1075 
1086  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1088  {
1090  }
1091 
1102  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1104  {
1105  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
1106  }
1107 
1114  [[nodiscard]] __device__ mutable_column_device_view child(size_type child_index) const noexcept
1115  {
1116  return d_children[child_index];
1117  }
1118 
1119 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1134  __device__ void set_valid(size_type element_index) const noexcept
1135  {
1136  return set_bit(null_mask(), element_index);
1137  }
1138 
1152  __device__ void set_null(size_type element_index) const noexcept
1153  {
1154  return clear_bit(null_mask(), element_index);
1155  }
1156 
1157 #endif
1158 
1169  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
1170  {
1171  null_mask()[word_index] = new_word;
1172  }
1173 
1182  static std::size_t extent(mutable_column_view source_view);
1183 
1190  void destroy();
1191 
1192  private:
1193  mutable_column_device_view* d_children{};
1197  size_type _num_children{};
1198 
1207  mutable_column_device_view(mutable_column_view source);
1208 };
1209 
1210 namespace detail {
1211 
1212 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1213 
1220 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
1221  size_type destination_word_index,
1222  size_type source_begin_bit,
1223  size_type source_end_bit)
1224 {
1225  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
1226  bitmask_type curr_word = source[source_word_index];
1227  bitmask_type next_word = 0;
1228  if (word_index(source_end_bit - 1) >
1229  word_index(source_begin_bit +
1230  destination_word_index * detail::size_in_bits<bitmask_type>())) {
1231  next_word = source[source_word_index + 1];
1232  }
1233  return __funnelshift_r(curr_word, next_word, source_begin_bit);
1234 }
1235 
1236 #endif
1237 
1252 template <typename T>
1255 
1261  value_accessor(column_device_view const& _col) : col{_col}
1262  {
1263  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1264  }
1265 
1271  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1272 };
1273 
1300 template <typename T, typename Nullate>
1303 
1310  optional_accessor(column_device_view const& _col, Nullate with_nulls)
1311  : col{_col}, has_nulls{with_nulls}
1312  {
1313  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1314  if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1315  }
1316 
1324  __device__ inline cuda::std::optional<T> operator()(cudf::size_type i) const
1325  {
1326  if (has_nulls) {
1327  return (col.is_valid_nocheck(i)) ? cuda::std::optional<T>{col.element<T>(i)}
1328  : cuda::std::optional<T>{cuda::std::nullopt};
1329  }
1330  return cuda::std::optional<T>{col.element<T>(i)};
1331  }
1332 
1333  Nullate has_nulls{};
1334 };
1335 
1355 template <typename T, bool has_nulls = false>
1358 
1364  pair_accessor(column_device_view const& _col) : col{_col}
1365  {
1366  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1367  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1368  }
1369 
1376  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
1377  {
1378  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1379  }
1380 };
1381 
1401 template <typename T, bool has_nulls = false>
1404 
1406 
1412  pair_rep_accessor(column_device_view const& _col) : col{_col}
1413  {
1414  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1415  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1416  }
1417 
1424  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1425  {
1426  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1427  }
1428 
1429  private:
1430  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
1431  [[nodiscard]] __device__ inline auto get_rep(cudf::size_type i) const
1432  {
1433  return col.element<R>(i);
1434  }
1435 
1436  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
1437  [[nodiscard]] __device__ inline auto get_rep(cudf::size_type i) const
1438  {
1439  return col.element<R>(i).value();
1440  }
1441 };
1442 
1454 template <typename T>
1457 
1464  {
1465  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1466  }
1467 
1474  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1475 };
1476 
1502 template <typename ColumnDeviceView, typename ColumnViewIterator>
1503 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1504  ColumnViewIterator child_end,
1505  void* h_ptr,
1506  void* d_ptr)
1507 {
1508  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1509  auto num_children = std::distance(child_begin, child_end);
1510  if (num_children > 0) {
1511  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1512  // struct objects in order for d_children to be used as an array.
1513  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1514  auto d_column = d_children;
1515 
1516  // Any child data is assigned past the end of this array: h_end and d_end.
1517  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1518  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1519  std::for_each(child_begin, child_end, [&](auto const& col) {
1520  // inplace-new each child into host memory
1521  new (h_column) ColumnDeviceView(col, h_end, d_end);
1522  h_column++; // advance to next child
1523  // update the pointers for holding this child column's child data
1524  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1525  h_end += col_child_data_size;
1526  d_end += col_child_data_size;
1527  });
1528  }
1529  return d_children;
1530 }
1531 
1532 } // namespace detail
1533 } // 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: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: host_udf.hpp:37
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:355
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