column_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2022, NVIDIA CORPORATION.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 #pragma once
17 
19 #include <cudf/detail/utilities/alignment.hpp>
21 #include <cudf/lists/list_view.hpp>
22 #include <cudf/strings/string_view.cuh>
25 #include <cudf/types.hpp>
26 #include <cudf/utilities/bit.hpp>
27 #include <cudf/utilities/default_stream.hpp>
28 #include <cudf/utilities/span.hpp>
31 
32 #include <rmm/cuda_stream_view.hpp>
33 
34 #include <thrust/iterator/counting_iterator.h>
35 #include <thrust/iterator/transform_iterator.h>
36 #include <thrust/optional.h>
37 #include <thrust/pair.h>
38 
39 #include <algorithm>
40 
46 namespace cudf {
47 
55 struct nullate {
56  struct YES : std::bool_constant<true> {
57  };
58  struct NO : std::bool_constant<false> {
59  };
65  struct DYNAMIC {
66  DYNAMIC() = delete;
75  constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
82  constexpr operator bool() const noexcept { return value; }
83  bool value;
84  };
85 };
86 
87 namespace detail {
99 class alignas(16) column_device_view_base {
100  public:
101  column_device_view_base() = delete;
102  ~column_device_view_base() = default;
105 
117 
134  template <typename T = void,
135  CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
136  [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept
137  {
138  return static_cast<T const*>(_data);
139  }
140 
156  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
157  [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept
158  {
159  return head<T>() + _offset;
160  }
161 
167  [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; }
168 
174  [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; }
175 
185  [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; }
186 
196  [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept
197  {
198  return _null_mask;
199  }
200 
207  [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; }
208 
223  [[nodiscard]] __device__ bool is_valid(size_type element_index) const noexcept
224  {
225  return not nullable() or is_valid_nocheck(element_index);
226  }
227 
240  [[nodiscard]] __device__ bool is_valid_nocheck(size_type element_index) const noexcept
241  {
242  return bit_is_set(_null_mask, offset() + element_index);
243  }
244 
258  [[nodiscard]] __device__ bool is_null(size_type element_index) const noexcept
259  {
260  return not is_valid(element_index);
261  }
262 
274  [[nodiscard]] __device__ bool is_null_nocheck(size_type element_index) const noexcept
275  {
276  return not is_valid_nocheck(element_index);
277  }
278 
288  [[nodiscard]] __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
289  {
290  return null_mask()[word_index];
291  }
292 
293  protected:
294  data_type _type{type_id::EMPTY};
296  void const* _data{};
298  size_type _offset{};
300 
312  size_type size,
313  void const* data,
314  bitmask_type const* null_mask,
317  {
318  }
319 
320  template <typename C, typename T, typename = void>
321  struct has_element_accessor_impl : std::false_type {
322  };
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 };
332 // @cond
333 // Forward declaration
334 template <typename T>
335 struct value_accessor;
336 template <typename T, typename Nullate>
337 struct optional_accessor;
338 template <typename T, bool has_nulls>
339 struct pair_accessor;
340 template <typename T, bool has_nulls>
341 struct pair_rep_accessor;
342 template <typename T>
343 struct mutable_value_accessor;
344 // @endcond
345 } // namespace detail
346 
354  public:
355  column_device_view() = delete;
356  ~column_device_view() = default;
359 
371 
381  column_device_view(column_view column, void* h_ptr, void* d_ptr);
382 
399  [[nodiscard]] CUDF_HOST_DEVICE column_device_view slice(size_type offset,
400  size_type size) const noexcept
401  {
402  return column_device_view{this->type(),
403  size,
404  this->head(),
405  this->null_mask(),
406  this->offset() + offset,
407  d_children,
408  this->num_child_columns()};
409  }
410 
428  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
429  [[nodiscard]] __device__ T element(size_type element_index) const noexcept
430  {
431  return data<T>()[element_index];
432  }
433 
445  template <typename T, CUDF_ENABLE_IF(std::is_same_v<T, string_view>)>
446  __device__ T element(size_type element_index) const noexcept
447  {
448  size_type index = element_index + offset(); // account for this view's _offset
449  const auto* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
450  const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
451  size_type offset = d_offsets[index];
452  return string_view{d_strings + offset, d_offsets[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__ 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__ 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  {
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(
824  column_view source_view, rmm::cuda_stream_view stream = cudf::default_stream_value);
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  {
862  }
863 
869  [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
870  {
871  return _num_children;
872  }
873 
874  private:
887  CUDF_HOST_DEVICE column_device_view(data_type type,
888  size_type size,
889  void const* data,
890  bitmask_type const* null_mask,
893  size_type num_children)
894  : column_device_view_base(type, size, data, null_mask, offset),
896  _num_children(num_children)
897  {
898  }
899 
900  protected:
906 
918 };
919 
927  public:
928  mutable_column_device_view() = delete;
929  ~mutable_column_device_view() = default;
932 
944 
955 
974  static std::unique_ptr<mutable_column_device_view,
975  std::function<void(mutable_column_device_view*)>>
977  rmm::cuda_stream_view stream = cudf::default_stream_value);
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__ 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  {
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  {
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__
1117 
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:
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
cudf::bit_is_set
CUDF_HOST_DEVICE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
Definition: bit.hpp:126
cudf::detail::column_device_view_base::_data
void const * _data
Pointer to device memory containing elements.
Definition: column_device_view.cuh:296
cudf::column_device_view::operator=
column_device_view & operator=(column_device_view const &)=default
Copy assignment operator.
cudf::column_device_view::optional_end
auto optional_end(Nullate has_nulls) const
Return an optional iterator to the element following the last element of the column.
Definition: column_device_view.cuh:758
cudf::detail::mutable_value_accessor::mutable_value_accessor
mutable_value_accessor(mutable_column_device_view &_col)
Constructor.
Definition: column_device_view.cuh:1460
bit.hpp
Utilities for bit and bitmask operations.
fixed_point.hpp
Class definition for fixed point data type.
cudf::detail::column_device_view_base::head
CUDF_HOST_DEVICE T const * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
Definition: column_device_view.cuh:136
cudf::column
A container of nullable device data as a column of elements.
Definition: column.hpp:48
cudf::mutable_column_device_view::mutable_column_device_view
mutable_column_device_view(mutable_column_device_view &&)=default
Move constructor.
cudf::detail::pair_accessor::operator()
thrust::pair< T, bool > operator()(cudf::size_type i) const
Pair accessor.
Definition: column_device_view.cuh:1373
cudf::data_type::id
constexpr type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:280
cudf::detail::pair_rep_accessor::pair_rep_accessor
pair_rep_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1409
strings_column_view.hpp
Class definition for cudf::strings_column_view.
cudf::detail::column_device_view_base::_offset
size_type _offset
Definition: column_device_view.cuh:299
cudf::detail::column_device_view_base::operator=
column_device_view_base & operator=(column_device_view_base &&)=default
Move assignment operator.
cudf::mutable_column_device_view::operator=
mutable_column_device_view & operator=(mutable_column_device_view &&)=default
Move assignment operator.
cudf::nullate::DYNAMIC
nullate::DYNAMIC defers the determination of nullability to run time rather than compile time....
Definition: column_device_view.cuh:65
cudf::size_type
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:80
cudf::column_view
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
Definition: column_view.hpp:322
cudf::column_device_view::num_child_columns
CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
Returns the number of child columns.
Definition: column_device_view.cuh:869
cudf::nullate
Indicates the presence of nulls at compile-time or runtime.
Definition: column_device_view.cuh:55
cudf::mutable_column_device_view::destroy
void destroy()
Destroy the mutable_column_device_view object.
cudf::column_device_view
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Definition: column_device_view.cuh:353
types.hpp
Type declarations for libcudf.
cudf::column_device_view::const_optional_iterator
thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > const_optional_iterator
Optional iterator for navigating this column.
Definition: column_device_view.cuh:596
cudf::detail::column_device_view_base::is_valid
bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
Definition: column_device_view.cuh:223
cudf::column_device_view::operator=
column_device_view & operator=(column_device_view &&)=default
Move assignment operator.
cudf::string_view
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
Definition: string_view.hpp:42
cudf::mutable_column_view
A non-owning, mutable view of device data as a column of elements, some of which may be null as indic...
Definition: column_view.hpp:485
rmm::cuda_stream_view
cudf::column_device_view::pair_end
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
Definition: column_device_view.cuh:778
cudf::data_type::scale
constexpr int32_t scale() const noexcept
Returns the scale (for fixed_point types)
Definition: types.hpp:287
cudf::column_device_view::create
static std::unique_ptr< column_device_view, std::function< void(column_device_view *)> > create(column_view source_view, rmm::cuda_stream_view stream=cudf::default_stream_value)
Factory to construct a column view that is usable in device memory.
cudf::detail::column_device_view_base::null_mask
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
Definition: column_device_view.cuh:196
cudf::detail::column_device_view_base::is_null_nocheck
bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:274
cudf::column_device_view::optional_begin
auto optional_begin(Nullate has_nulls) const
Return an optional iterator to the first element of the column.
Definition: column_device_view.cuh:671
cudf::column_device_view::column_device_view
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
cudf::bitmask_type
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
Definition: types.hpp:81
cudf::column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Counting iterator.
Definition: column_device_view.cuh:543
cudf::device_storage_type_t
std::conditional_t< std::is_same_v< numeric::decimal32, T >, int32_t, std::conditional_t< std::is_same_v< numeric::decimal64, T >, int64_t, std::conditional_t< std::is_same_v< numeric::decimal128, T >, __int128_t, T > >> device_storage_type_t
"Returns" the corresponding type that is stored on the device when using cudf::column
Definition: type_dispatcher.hpp:114
cudf::detail::pair_rep_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1400
CUDF_ENABLE_IF
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:50
cudf::detail::value_accessor::operator()
T operator()(cudf::size_type i) const
Returns the value of element at index i
Definition: column_device_view.cuh:1268
cudf::column_device_view::slice
CUDF_HOST_DEVICE column_device_view slice(size_type offset, size_type size) const noexcept
Get a new column_device_view which is a slice of this column.
Definition: column_device_view.cuh:399
cudf::detail::mutable_value_accessor
Mutable value accessor of column without null bitmask.
Definition: column_device_view.cuh:1452
cudf::column_device_view::column_device_view
column_device_view(column_device_view const &)=default
Copy constructor.
cudf::detail::pair_accessor::pair_accessor
pair_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1361
cudf::detail::column_device_view_base::offset
CUDF_HOST_DEVICE size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
Definition: column_device_view.cuh:207
cudf::dictionary_wrapper
A strongly typed wrapper for indices in a DICTIONARY type column.
Definition: dictionary.hpp:48
cudf::detail::child_columns_to_device_array
ColumnDeviceView * child_columns_to_device_array(ColumnViewIterator child_begin, ColumnViewIterator child_end, void *h_ptr, void *d_ptr)
Helper function for use by column_device_view and mutable_column_device_view constructors to build de...
Definition: column_device_view.cuh:1500
cudf::column_device_view::const_pair_rep_iterator
thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > const_pair_rep_iterator
Pair rep iterator for navigating this column.
Definition: column_device_view.cuh:612
cudf::mutable_column_device_view::mutable_column_device_view
mutable_column_device_view(mutable_column_view column, void *h_ptr, void *d_ptr)
Creates an instance of this class using the specified host memory pointer (h_ptr) to store child obje...
cudf::type_dispatcher
CUDF_HOST_DEVICE constexpr decltype(auto) __forceinline__ type_dispatcher(cudf::data_type dtype, Functor f, Ts &&... args)
Invokes an operator() template with the type instantiation based on the specified cudf::data_type's i...
Definition: type_dispatcher.hpp:440
cudf::mutable_column_device_view::end
iterator< T > end()
Return one past the last element after underlying data is casted to the specified type.
Definition: column_device_view.cuh:1100
cudf::mutable_column_device_view::element
T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:1035
cudf::mutable_column_device_view::iterator
thrust::transform_iterator< detail::mutable_value_accessor< T >, count_it > iterator
Iterator for navigating this column.
Definition: column_device_view.cuh:1071
cudf::mutable_column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Counting iterator.
Definition: column_device_view.cuh:1066
cudf::strings_column_view::offsets_column_index
static constexpr size_type offsets_column_index
Child index of the offsets column.
Definition: strings_column_view.hpp:60
numeric::scaled_integer
Helper struct for constructing fixed_point when value is already shifted.
Definition: fixed_point.hpp:191
cudf::nullate::NO
Definition: column_device_view.cuh:58
cudf::mutable_column_device_view::create
static std::unique_ptr< mutable_column_device_view, std::function< void(mutable_column_device_view *)> > create(mutable_column_view source_view, rmm::cuda_stream_view stream=cudf::default_stream_value)
Factory to construct a column view that is usable in device memory.
cudf::detail::column_device_view_base::get_mask_word
bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the the specified bitmask word from the null_mask().
Definition: column_device_view.cuh:288
cudf::strings_column_view::chars_column_index
static constexpr size_type chars_column_index
Child index of the characters column.
Definition: strings_column_view.hpp:61
cudf::detail::pair_accessor
pair accessor of column with/without null bitmask
Definition: column_device_view.cuh:1353
cudf::column_device_view::pair_rep_end
const_pair_rep_iterator< T, has_nulls > pair_rep_end() const
Return a pair iterator to the element following the last element of the column.
Definition: column_device_view.cuh:799
cudf::detail::pair_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1354
cudf::detail::column_device_view_base::is_null
bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:258
cudf::data_type
Indicator for the logical data type of an element in a column.
Definition: types.hpp:236
cudf::mutable_column_device_view::extent
static std::size_t extent(mutable_column_view source_view)
Return the size in bytes of the amount of memory needed to hold a device view of the specified column...
cudf::column_device_view::column_device_view
column_device_view(column_device_view &&)=default
Move constructor.
cudf::mutable_column_device_view
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
Definition: column_device_view.cuh:926
cudf::detail::value_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1251
list_view.hpp
Class definition for cudf::list_view.
cudf::column_device_view::pair_rep_begin
const_pair_rep_iterator< T, has_nulls > pair_rep_begin() const
Return a pair iterator to the first element of the column.
Definition: column_device_view.cuh:733
cudf::nullate::YES
Definition: column_device_view.cuh:56
cudf::mutable_column_device_view::operator=
mutable_column_device_view & operator=(mutable_column_device_view const &)=default
Copy assignment operator.
cudf::detail::column_device_view_base::nullable
CUDF_HOST_DEVICE bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
Definition: column_device_view.cuh:185
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::column_device_view::end
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
Definition: column_device_view.cuh:586
cudf::detail::column_device_view_base::_type
data_type _type
Element type.
Definition: column_device_view.cuh:294
cudf::column_device_view::child
column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:849
cudf::detail::column_device_view_base::column_device_view_base
column_device_view_base(column_device_view_base const &)=default
Copy constructor.
cudf::column_device_view::children
device_span< column_device_view const > children() const noexcept
Returns a span containing the children of this column.
Definition: column_device_view.cuh:859
cudf::detail::value_accessor
value accessor of column without null bitmask
Definition: column_device_view.cuh:1250
cudf::column_device_view::begin
const_iterator< T > begin() const
Return an iterator to the first element of the column.
Definition: column_device_view.cuh:566
cudf::detail::optional_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1299
cudf::detail::column_device_view_base::_null_mask
bitmask_type const * _null_mask
Definition: column_device_view.cuh:297
cudf::detail::column_device_view_base
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Definition: column_device_view.cuh:99
cudf::has_nulls
bool has_nulls(table_view const &view)
Returns True if the table has nulls in any of its columns.
Definition: table_view.hpp:318
cudf::mutable_column_device_view::mutable_column_device_view
mutable_column_device_view(mutable_column_device_view const &)=default
Copy constructor.
cudf::mutable_column_device_view::child
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:1111
cudf::detail::column_device_view_base::type
CUDF_HOST_DEVICE data_type type() const noexcept
Returns the element type.
Definition: column_device_view.cuh:174
cudf::detail::pair_rep_accessor
pair accessor of column with/without null bitmask
Definition: column_device_view.cuh:1399
cudf::column_device_view::d_children
column_device_view * d_children
Definition: column_device_view.cuh:901
cudf::detail::column_device_view_base::operator=
column_device_view_base & operator=(column_device_view_base const &)=default
Copy assignment operator.
cudf::mutable_column_device_view::has_element_accessor
static constexpr bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
Definition: column_device_view.cuh:1047
cudf::column_device_view::column_device_view
column_device_view(column_view column, void *h_ptr, void *d_ptr)
Creates an instance of this class using the specified host memory pointer (h_ptr) to store child obje...
CUDF_EXPECTS
#define CUDF_EXPECTS(cond, reason)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
Definition: error.hpp:113
cudf::column_device_view::pair_begin
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
Definition: column_device_view.cuh:701
cudf::detail::mutable_value_accessor::operator()
T & operator()(cudf::size_type i)
Accessor.
Definition: column_device_view.cuh:1471
cudf::column_device_view::destroy
void destroy()
Destroy the column_device_view object.
cudf::detail::value_accessor::value_accessor
value_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1258
cudf::mutable_column_device_view::set_mask_word
void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
Updates the specified bitmask word in the null_mask() with a new word.
Definition: column_device_view.cuh:1166
cudf::column_device_view::has_element_accessor
static constexpr bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
Definition: column_device_view.cuh:537
cudf::detail::pair_rep_accessor::rep_type
device_storage_type_t< T > rep_type
representation type
Definition: column_device_view.cuh:1402
cudf::detail::pair_rep_accessor::operator()
thrust::pair< rep_type, bool > operator()(cudf::size_type i) const
Pair accessor.
Definition: column_device_view.cuh:1421
cudf::detail::column_device_view_base::_size
cudf::size_type _size
Number of elements.
Definition: column_device_view.cuh:295
cudf::detail::column_device_view_base::size
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
Definition: column_device_view.cuh:167
cudf::mutable_column_device_view::begin
iterator< T > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
Definition: column_device_view.cuh:1084
cudf::distance
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
Definition: types.hpp:95
cudf::device_span
Device version of C++20 std::span with reduced feature set.
Definition: span.hpp:285
cudf::column_device_view::extent
static std::size_t extent(column_view const &source_view)
Return the size in bytes of the amount of memory needed to hold a device view of the specified column...
type_dispatcher.hpp
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
traits.hpp
cudf::void_t
void void_t
Utility metafunction that maps a sequence of any types to the type void.
Definition: traits.hpp:37
cudf::column_device_view::const_iterator
thrust::transform_iterator< detail::value_accessor< T >, count_it > const_iterator
Iterator for navigating this column.
Definition: column_device_view.cuh:548
cudf::detail::optional_accessor
optional accessor of a column
Definition: column_device_view.cuh:1298
cudf::detail::optional_accessor::optional_accessor
optional_accessor(column_device_view const &_col, Nullate with_nulls)
Constructor.
Definition: column_device_view.cuh:1307
cudf::detail::column_device_view_base::column_device_view_base
column_device_view_base(column_device_view_base &&)=default
Move constructor.
column_view.hpp
column view class definitions
struct_view.hpp
Class definition for cudf::struct_view.
cudf::detail::optional_accessor::operator()
thrust::optional< T > operator()(cudf::size_type i) const
Returns a thrust::optional of column[i].
Definition: column_device_view.cuh:1321
cudf::column_device_view::_num_children
size_type _num_children
The number of child columns.
Definition: column_device_view.cuh:905
cudf::detail::mutable_value_accessor::col
mutable_column_device_view col
mutable column view of column in device
Definition: column_device_view.cuh:1453
cudf::nullate::DYNAMIC::DYNAMIC
constexpr DYNAMIC(bool b) noexcept
Create a runtime nullate object.
Definition: column_device_view.cuh:75
cudf::detail::column_device_view_base::data
CUDF_HOST_DEVICE T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Definition: column_device_view.cuh:157
cudf::detail::column_device_view_base::is_valid_nocheck
bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
Definition: column_device_view.cuh:240
cudf::column_device_view::element
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:429
numeric::scale_type
scale_type
The scale type for fixed_point.
Definition: fixed_point.hpp:38
numeric
fixed_point and supporting types
Definition: fixed_point.hpp:35
cudf::word_index
constexpr CUDF_HOST_DEVICE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
Definition: bit.hpp:72
cudf::nullate::DYNAMIC::value
bool value
True if nulls are expected.
Definition: column_device_view.cuh:83
cudf::column_device_view::const_pair_iterator
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.
Definition: column_device_view.cuh:603