column_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2021, 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.cuh>
22 #include <cudf/strings/string_view.cuh>
25 #include <cudf/types.hpp>
26 #include <cudf/utilities/bit.hpp>
29 
30 #include <rmm/cuda_stream_view.hpp>
31 
32 #include <thrust/iterator/counting_iterator.h>
33 #include <thrust/iterator/transform_iterator.h>
34 
35 #include <algorithm>
36 
42 namespace cudf {
43 namespace detail {
55 class alignas(16) column_device_view_base {
56  public:
57  column_device_view_base() = delete;
58  ~column_device_view_base() = default;
61  column_device_view_base& operator=(column_device_view_base const&) = default;
62  column_device_view_base& operator=(column_device_view_base&&) = default;
63 
80  template <typename T = void,
81  CUDF_ENABLE_IF(std::is_same<T, void>::value or is_rep_layout_compatible<T>())>
82  __host__ __device__ T const* head() const noexcept
83  {
84  return static_cast<T const*>(_data);
85  }
86 
102  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
103  __host__ __device__ T const* data() const noexcept
104  {
105  return head<T>() + _offset;
106  }
107 
111  __host__ __device__ size_type size() const noexcept { return _size; }
112 
116  __host__ __device__ data_type type() const noexcept { return _type; }
117 
127  __host__ __device__ bool nullable() const noexcept { return nullptr != _null_mask; }
128 
136  __host__ __device__ bitmask_type const* null_mask() const noexcept { return _null_mask; }
137 
142  __host__ __device__ size_type offset() const noexcept { return _offset; }
143 
158  __device__ bool is_valid(size_type element_index) const noexcept
159  {
160  return not nullable() or is_valid_nocheck(element_index);
161  }
162 
175  __device__ bool is_valid_nocheck(size_type element_index) const noexcept
176  {
177  return bit_is_set(_null_mask, offset() + element_index);
178  }
179 
193  __device__ bool is_null(size_type element_index) const noexcept
194  {
195  return not is_valid(element_index);
196  }
197 
209  __device__ bool is_null_nocheck(size_type element_index) const noexcept
210  {
211  return not is_valid_nocheck(element_index);
212  }
213 
223  __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
224  {
225  return null_mask()[word_index];
226  }
227 
228  protected:
229  data_type _type{type_id::EMPTY};
230  cudf::size_type _size{};
231  void const* _data{};
232  bitmask_type const* _null_mask{};
233  size_type _offset{};
235 
237  column_device_view_base(data_type type,
238  size_type size,
239  void const* data,
240  bitmask_type const* null_mask,
241  size_type offset)
243  {
244  }
245 
246  template <typename C, typename T, typename = void>
247  struct has_element_accessor_impl : std::false_type {
248  };
249 
250  template <typename C, typename T>
252  C,
253  T,
254  void_t<decltype(std::declval<C>().template element<T>(std::declval<size_type>()))>>
255  : std::true_type {
256  };
257 };
258 
259 // Forward declaration
260 template <typename T>
261 struct value_accessor;
262 template <typename T, bool has_nulls>
263 struct pair_accessor;
264 template <typename T, bool has_nulls>
265 struct pair_rep_accessor;
266 template <typename T>
268 } // namespace detail
269 
277  public:
278  column_device_view() = delete;
279  ~column_device_view() = default;
280  column_device_view(column_device_view const&) = default;
282  column_device_view& operator=(column_device_view const&) = default;
283  column_device_view& operator=(column_device_view&&) = default;
284 
294  column_device_view(column_view column, void* h_ptr, void* d_ptr);
295 
312  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
313  __device__ T element(size_type element_index) const noexcept
314  {
315  return data<T>()[element_index];
316  }
317 
329  template <typename T, CUDF_ENABLE_IF(std::is_same<T, string_view>::value)>
330  __device__ T element(size_type element_index) const noexcept
331  {
332  size_type index = element_index + offset(); // account for this view's _offset
333  const int32_t* d_offsets =
334  d_children[strings_column_view::offsets_column_index].data<int32_t>();
335  const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
336  size_type offset = d_offsets[index];
337  return string_view{d_strings + offset, d_offsets[index + 1] - offset};
338  }
339 
340  private:
346  struct index_element_fn {
347  template <typename IndexType,
348  CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_unsigned<IndexType>::value)>
349  __device__ size_type operator()(column_device_view const& indices, size_type index)
350  {
351  return static_cast<size_type>(indices.element<IndexType>(index));
352  }
353 
354  template <typename IndexType,
355  typename... Args,
356  CUDF_ENABLE_IF(not(is_index_type<IndexType>() and
357  std::is_unsigned<IndexType>::value))>
358  __device__ size_type operator()(Args&&... args)
359  {
360  cudf_assert(false and "dictionary indices must be an unsigned integral type");
361  return 0;
362  }
363  };
364 
365  public:
390  template <typename T, CUDF_ENABLE_IF(std::is_same<T, dictionary32>::value)>
391  __device__ T element(size_type element_index) const noexcept
392  {
393  size_type index = element_index + offset(); // account for this view's _offset
394  auto const indices = d_children[0];
395  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
396  }
397 
408  template <typename T, CUDF_ENABLE_IF(std::is_same<T, numeric::decimal32>::value)>
409  __device__ T element(size_type element_index) const noexcept
410  {
411  using namespace numeric;
412  auto const scale = scale_type{_type.scale()};
413  return decimal32{scaled_integer<int32_t>{data<int32_t>()[element_index], scale}};
414  }
415 
426  template <typename T, CUDF_ENABLE_IF(std::is_same<T, numeric::decimal64>::value)>
427  __device__ T element(size_type element_index) const noexcept
428  {
429  using namespace numeric;
430  auto const scale = scale_type{_type.scale()};
431  return decimal64{scaled_integer<int64_t>{data<int64_t>()[element_index], scale}};
432  }
433 
438  template <typename T>
439  static constexpr bool has_element_accessor()
440  {
442  }
443 
447  using count_it = thrust::counting_iterator<size_type>;
448  template <typename T>
449  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
450 
463  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
464  const_iterator<T> begin() const
465  {
466  return const_iterator<T>{count_it{0}, detail::value_accessor<T>{*this}};
467  }
468 
481  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
482  const_iterator<T> end() const
483  {
484  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
485  }
486 
490  template <typename T, bool has_nulls>
492  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
493 
499  template <typename T, bool has_nulls>
501  thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it>;
502 
522  template <typename T,
523  bool has_nulls,
524  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
526  {
529  }
530 
552  template <typename T,
553  bool has_nulls,
554  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
556  {
559  }
560 
572  template <typename T,
573  bool has_nulls,
574  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
576  {
579  }
580 
592  template <typename T,
593  bool has_nulls,
594  CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
596  {
599  }
600 
619  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
620  column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
621 
628  void destroy();
629 
637  static std::size_t extent(column_view const& source_view);
638 
645  __device__ column_device_view child(size_type child_index) const noexcept
646  {
647  return d_children[child_index];
648  }
649 
655  __host__ __device__ size_type num_child_columns() const noexcept { return _num_children; }
656 
657  protected:
659  size_type _num_children{};
663 
673 };
674 
682  public:
683  mutable_column_device_view() = delete;
684  ~mutable_column_device_view() = default;
687  mutable_column_device_view& operator=(mutable_column_device_view const&) = default;
689 
700 
719  static std::unique_ptr<mutable_column_device_view,
720  std::function<void(mutable_column_device_view*)>>
721  create(mutable_column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
722 
739  template <typename T = void,
740  CUDF_ENABLE_IF(std::is_same<T, void>::value or is_rep_layout_compatible<T>())>
741  __host__ __device__ T* head() const noexcept
742  {
743  return const_cast<T*>(detail::column_device_view_base::head<T>());
744  }
745 
758  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
759  __host__ __device__ T* data() const noexcept
760  {
761  return const_cast<T*>(detail::column_device_view_base::data<T>());
762  }
763 
777  template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
778  __device__ T& element(size_type element_index) const noexcept
779  {
780  return data<T>()[element_index];
781  }
782 
788  template <typename T>
789  static constexpr bool has_element_accessor()
790  {
792  }
793 
801  __host__ __device__ bitmask_type* null_mask() const noexcept
802  {
803  return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
804  }
805 
809  using count_it = thrust::counting_iterator<size_type>;
810  template <typename T>
811  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
812 
823  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
824  iterator<T> begin()
825  {
826  return iterator<T>{count_it{0}, detail::mutable_value_accessor<T>{*this}};
827  }
828 
839  template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
840  iterator<T> end()
841  {
842  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
843  }
844 
851  __device__ mutable_column_device_view child(size_type child_index) const noexcept
852  {
853  return d_children[child_index];
854  }
855 
856 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
857 
871  __device__ void set_valid(size_type element_index) const noexcept
872  {
873  return set_bit(null_mask(), element_index);
874  }
875 
889  __device__ void set_null(size_type element_index) const noexcept
890  {
891  return clear_bit(null_mask(), element_index);
892  }
893 
894 #endif
895 
906  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
907  {
908  null_mask()[word_index] = new_word;
909  }
910 
917  static std::size_t extent(mutable_column_view source_view);
918 
925  void destroy();
926 
927  private:
929  size_type _num_children{};
933 
942  mutable_column_device_view(mutable_column_view source);
943 };
944 
945 namespace detail {
946 
947 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
948 
955 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
956  size_type destination_word_index,
957  size_type source_begin_bit,
958  size_type source_end_bit)
959 {
960  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
961  bitmask_type curr_word = source[source_word_index];
962  bitmask_type next_word = 0;
963  if (word_index(source_end_bit) >
964  word_index(source_begin_bit +
965  destination_word_index * detail::size_in_bits<bitmask_type>())) {
966  next_word = source[source_word_index + 1];
967  }
968  return __funnelshift_r(curr_word, next_word, source_begin_bit);
969 }
970 
971 #endif
972 
986 template <typename T>
989 
994  value_accessor(column_device_view const& _col) : col{_col}
995  {
996  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
997  }
998 
999  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
1000 };
1001 
1020 template <typename T, bool has_nulls = false>
1023 
1028  pair_accessor(column_device_view const& _col) : col{_col}
1029  {
1030  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1031  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1032  }
1033 
1034  CUDA_DEVICE_CALLABLE
1035  thrust::pair<T, bool> operator()(cudf::size_type i) const
1036  {
1037  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1038  }
1039 };
1040 
1059 template <typename T, bool has_nulls = false>
1062 
1063  using rep_type = device_storage_type_t<T>;
1064 
1069  pair_rep_accessor(column_device_view const& _col) : col{_col}
1070  {
1071  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1072  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
1073  }
1074 
1075  CUDA_DEVICE_CALLABLE
1076  thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
1077  {
1078  return {get_rep<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
1079  }
1080 
1081  private:
1082  template <typename R, std::enable_if_t<std::is_same<R, rep_type>::value, void>* = nullptr>
1083  CUDA_DEVICE_CALLABLE auto get_rep(cudf::size_type i) const
1084  {
1085  return col.element<R>(i);
1086  }
1087 
1088  template <typename R, std::enable_if_t<not std::is_same<R, rep_type>::value, void>* = nullptr>
1089  CUDA_DEVICE_CALLABLE auto get_rep(cudf::size_type i) const
1090  {
1091  return col.element<R>(i).value();
1092  }
1093 };
1094 
1095 template <typename T>
1098 
1104  {
1105  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
1106  }
1107 
1108  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
1109 };
1110 
1136 template <typename ColumnDeviceView, typename ColumnViewIterator>
1137 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
1138  ColumnViewIterator child_end,
1139  void* h_ptr,
1140  void* d_ptr)
1141 {
1142  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1143  auto num_children = std::distance(child_begin, child_end);
1144  if (num_children > 0) {
1145  // The beginning of the memory must be the fixed-sized ColumnDeviceView
1146  // struct objects in order for d_children to be used as an array.
1147  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1148  auto d_column = d_children;
1149 
1150  // Any child data is assigned past the end of this array: h_end and d_end.
1151  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
1152  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
1153  std::for_each(child_begin, child_end, [&](auto const& col) {
1154  // inplace-new each child into host memory
1155  new (h_column) ColumnDeviceView(col, h_end, d_end);
1156  h_column++; // advance to next child
1157  // update the pointers for holding this child column's child data
1158  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
1159  h_end += col_child_data_size;
1160  d_end += col_child_data_size;
1161  });
1162  }
1163  return d_children;
1164 }
1165 
1166 } // namespace detail
1167 } // namespace cudf
cudf::detail::column_device_view_base::_data
void const * _data
Pointer to device memory containing elements.
Definition: column_device_view.cuh:231
cudf::detail::mutable_value_accessor::mutable_value_accessor
mutable_value_accessor(mutable_column_device_view &_col)
constructor
Definition: column_device_view.cuh:1103
cudf::type_dispatcher
constexpr decltype(auto) CUDA_HOST_DEVICE_CALLABLE 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:414
bit.hpp
Utilities for bit and bitmask operations.
fixed_point.hpp
Class definition for fixed point data type.
cudf::column
A container of nullable device data as a column of elements.
Definition: column.hpp:45
cudf::data_type::id
constexpr type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:273
cudf::detail::pair_rep_accessor::pair_rep_accessor
pair_rep_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1069
cudf::column_device_view::child
__device__ column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:645
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:234
numeric::fixed_point
A type for representing a number with a fixed amount of precision.
Definition: fixed_point.hpp:213
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:294
cudf::mutable_column_device_view::destroy
void destroy()
Destroy the mutable_column_device_view object.
cudf::detail::column_device_view_base::null_mask
__host__ __device__ bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
Definition: column_device_view.cuh:136
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:276
types.hpp
Type declarations for libcudf.
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:63
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:401
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:575
cudf::mutable_column_device_view::set_mask_word
__device__ 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:906
cudf::detail::column_device_view_base::get_mask_word
__device__ 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:223
cudf::data_type::scale
constexpr int32_t scale() const noexcept
Returns the scale (for fixed_point types)
Definition: types.hpp:278
cudf::detail::column_device_view_base::has_element_accessor_impl
Definition: column_device_view.cuh:247
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::mutable_column_device_view::element
__device__ T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:778
cudf::detail::column_device_view_base::nullable
__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:127
cudf::column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
Definition: column_device_view.cuh:447
cudf::detail::column_device_view_base::is_null_nocheck
__device__ bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:209
cudf::detail::pair_rep_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1061
CUDF_ENABLE_IF
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Definition: traits.hpp:48
cudf::detail::mutable_value_accessor
Definition: column_device_view.cuh:1096
cudf::detail::pair_accessor::pair_accessor
pair_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:1028
cudf::device_storage_type_t
std::conditional_t< std::is_same< numeric::decimal32, T >::value, int32_t, std::conditional_t< std::is_same< numeric::decimal64, T >::value, int64_t, T > > device_storage_type_t
"Returns" the corresponding type that is stored on the device when using cudf::column
Definition: type_dispatcher.hpp:102
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:1137
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:501
cudf::detail::column_device_view_base::head
__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:82
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::detail::column_device_view_base::is_valid
__device__ 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:158
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:840
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=rmm::cuda_stream_default)
Factory to construct a column view that is usable in device memory.
cudf::mutable_column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
Definition: column_device_view.cuh:809
numeric::scaled_integer
Helper struct for constructing fixed_point when value is already shifted.
Definition: fixed_point.hpp:197
cudf::detail::column_device_view_base::is_valid_nocheck
__device__ 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:175
cudf::detail::pair_accessor
pair accessor of column with/without null bitmask A unary functor returns pair with scalar value at i...
Definition: column_device_view.cuh:1021
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:595
cudf::detail::pair_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:1022
cudf::data_type
Indicator for the logical data type of an element in a column.
Definition: types.hpp:243
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::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:681
cudf::detail::value_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:988
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:555
cudf::detail::column_device_view_base::size
__host__ __device__ size_type size() const noexcept
Returns the number of elements in the column.
Definition: column_device_view.cuh:111
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::detail::column_device_view_base::is_null
__device__ bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:193
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:482
cudf::detail::column_device_view_base::_type
data_type _type
Element type.
Definition: column_device_view.cuh:229
cudf::detail::value_accessor
value accessor of column without null bitmask A unary functor returns scalar value at id....
Definition: column_device_view.cuh:987
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:464
cudf::column_device_view::element
__device__ T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:313
cudf::detail::column_device_view_base::_null_mask
bitmask_type const * _null_mask
Definition: column_device_view.cuh:232
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:55
list_view.cuh
Class definition for cudf::list_view.
cudf::column_device_view::num_child_columns
__host__ __device__ size_type num_child_columns() const noexcept
Returns the number of child columns.
Definition: column_device_view.cuh:655
cudf::mutable_column_device_view::null_mask
__host__ __device__ bitmask_type * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
Definition: column_device_view.cuh:801
cudf::detail::pair_rep_accessor
pair accessor of column with/without null bitmask A unary functor returns pair with representative sc...
Definition: column_device_view.cuh:1060
cudf::bit_is_set
CUDA_HOST_DEVICE_CALLABLE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
Definition: bit.hpp:111
cudf::column_device_view::d_children
column_device_view * d_children
Definition: column_device_view.cuh:658
cudf::detail::column_device_view_base::offset
__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:142
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:789
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:62
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:525
cudf::column_device_view::destroy
void destroy()
Destroy the column_device_view object.
cudf::mutable_column_device_view::child
__device__ mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:851
cudf::detail::column_device_view_base::type
__host__ __device__ data_type type() const noexcept
Returns the element type.
Definition: column_device_view.cuh:116
cudf::detail::value_accessor::value_accessor
value_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:994
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:439
cudf::detail::column_device_view_base::_size
cudf::size_type _size
Number of elements.
Definition: column_device_view.cuh:230
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:824
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:103
cudf::detail::column_device_view_base::data
__host__ __device__ T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Definition: column_device_view.cuh:103
cudf::word_index
constexpr CUDA_HOST_DEVICE_CALLABLE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
Definition: bit.hpp:61
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=rmm::cuda_stream_default)
Factory to construct a column view that is usable in device memory.
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::mutable_column_device_view::data
__host__ __device__ T * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Definition: column_device_view.cuh:759
column_view.hpp
column view class definitons
struct_view.hpp
Class definition for cudf::struct_view.
cudf::column_device_view::_num_children
size_type _num_children
The number of child columns.
Definition: column_device_view.cuh:662
cudf::detail::mutable_value_accessor::col
mutable_column_device_view col
mutable column view of column in device
Definition: column_device_view.cuh:1097
numeric
fixed_point and supporting types
Definition: fixed_point.hpp:33
cudf::mutable_column_device_view::head
__host__ __device__ T * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
Definition: column_device_view.cuh:741
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:492