column_device_view.cuh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2019-2020, 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 
18 #include <algorithm>
20 #include <cudf/detail/utilities/alignment.hpp>
22 #include <cudf/lists/list_view.cuh>
23 #include <cudf/strings/string_view.cuh>
26 #include <cudf/types.hpp>
27 #include <cudf/utilities/bit.hpp>
30 
31 #include <rmm/cuda_stream_view.hpp>
32 
33 #include <thrust/iterator/counting_iterator.h>
34 #include <thrust/iterator/transform_iterator.h>
35 
41 namespace cudf {
42 namespace detail {
54 class alignas(16) column_device_view_base {
55  public:
56  column_device_view_base() = delete;
57  ~column_device_view_base() = default;
60  column_device_view_base& operator=(column_device_view_base const&) = default;
61  column_device_view_base& operator=(column_device_view_base&&) = default;
62 
76  template <typename T = void>
77  __host__ __device__ T const* head() const noexcept
78  {
79  return static_cast<T const*>(_data);
80  }
81 
94  template <typename T>
95  __host__ __device__ T const* data() const noexcept
96  {
97  return head<T>() + _offset;
98  }
99 
103  __host__ __device__ size_type size() const noexcept { return _size; }
104 
108  __host__ __device__ data_type type() const noexcept { return _type; }
109 
119  __host__ __device__ bool nullable() const noexcept { return nullptr != _null_mask; }
120 
128  __host__ __device__ bitmask_type const* null_mask() const noexcept { return _null_mask; }
129 
134  __host__ __device__ size_type offset() const noexcept { return _offset; }
135 
150  __device__ bool is_valid(size_type element_index) const noexcept
151  {
152  return not nullable() or is_valid_nocheck(element_index);
153  }
154 
167  __device__ bool is_valid_nocheck(size_type element_index) const noexcept
168  {
169  return bit_is_set(_null_mask, offset() + element_index);
170  }
171 
185  __device__ bool is_null(size_type element_index) const noexcept
186  {
187  return not is_valid(element_index);
188  }
189 
201  __device__ bool is_null_nocheck(size_type element_index) const noexcept
202  {
203  return not is_valid_nocheck(element_index);
204  }
205 
215  __device__ bitmask_type get_mask_word(size_type word_index) const noexcept
216  {
217  return null_mask()[word_index];
218  }
219 
220  protected:
221  data_type _type{type_id::EMPTY};
222  cudf::size_type _size{};
223  void const* _data{};
224  bitmask_type const* _null_mask{};
225  size_type _offset{};
227 
229  column_device_view_base(data_type type,
230  size_type size,
231  void const* data,
232  bitmask_type const* null_mask,
233  size_type offset)
235  {
236  }
237 };
238 
239 // Forward declaration
240 template <typename T>
241 struct value_accessor;
242 template <typename T, bool has_nulls>
243 struct pair_accessor;
244 template <typename T>
245 struct mutable_value_accessor;
246 } // namespace detail
247 
255  public:
256  column_device_view() = delete;
257  ~column_device_view() = default;
258  column_device_view(column_device_view const&) = default;
260  column_device_view& operator=(column_device_view const&) = default;
261  column_device_view& operator=(column_device_view&&) = default;
262 
272  column_device_view(column_view column, void* h_ptr, void* d_ptr);
273 
286  template <typename T>
287  __device__ T const element(size_type element_index) const noexcept
288  {
289  return data<T>()[element_index];
290  }
291 
295  using count_it = thrust::counting_iterator<size_type>;
296  template <typename T>
297  using const_iterator = thrust::transform_iterator<detail::value_accessor<T>, count_it>;
298 
308  template <typename T>
309  const_iterator<T> begin() const
310  {
311  return const_iterator<T>{count_it{0}, detail::value_accessor<T>{*this}};
312  }
313 
323  template <typename T>
324  const_iterator<T> end() const
325  {
326  return const_iterator<T>{count_it{size()}, detail::value_accessor<T>{*this}};
327  }
328 
332  template <typename T, bool has_nulls>
334  thrust::transform_iterator<detail::pair_accessor<T, has_nulls>, count_it>;
335 
352  template <typename T, bool has_nulls>
354  {
357  }
358 
367  template <typename T, bool has_nulls>
369  {
372  }
373 
392  static std::unique_ptr<column_device_view, std::function<void(column_device_view*)>> create(
393  column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
394 
401  void destroy();
402 
410  static std::size_t extent(column_view const& source_view);
411 
418  __device__ column_device_view child(size_type child_index) const noexcept
419  {
420  return d_children[child_index];
421  }
422 
423  protected:
425  size_type _num_children{};
429 
439 };
440 
448  public:
449  mutable_column_device_view() = delete;
450  ~mutable_column_device_view() = default;
453  mutable_column_device_view& operator=(mutable_column_device_view const&) = default;
455 
466 
485  static std::unique_ptr<mutable_column_device_view,
486  std::function<void(mutable_column_device_view*)>>
487  create(mutable_column_view source_view, rmm::cuda_stream_view stream = rmm::cuda_stream_default);
488 
502  template <typename T = void>
503  __host__ __device__ T* head() const noexcept
504  {
505  return const_cast<T*>(detail::column_device_view_base::head<T>());
506  }
507 
519  template <typename T>
520  __host__ __device__ T* data() const noexcept
521  {
522  return const_cast<T*>(detail::column_device_view_base::data<T>());
523  }
524 
533  template <typename T>
534  __device__ T& element(size_type element_index) noexcept
535  {
536  return data<T>()[element_index];
537  }
538 
546  __host__ __device__ bitmask_type* null_mask() const noexcept
547  {
548  return const_cast<bitmask_type*>(detail::column_device_view_base::null_mask());
549  }
550 
554  using count_it = thrust::counting_iterator<size_type>;
555  template <typename T>
556  using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>, count_it>;
557 
565  template <typename T>
566  std::enable_if_t<is_fixed_width<T>(), iterator<T>> begin()
567  {
568  return iterator<T>{count_it{0}, detail::mutable_value_accessor<T>{*this}};
569  }
570 
578  template <typename T>
579  std::enable_if_t<is_fixed_width<T>(), iterator<T>> end()
580  {
581  return iterator<T>{count_it{size()}, detail::mutable_value_accessor<T>{*this}};
582  }
583 
590  __device__ mutable_column_device_view child(size_type child_index) const noexcept
591  {
592  return d_children[child_index];
593  }
594 
609  __device__ void set_valid(size_type element_index) const noexcept
610  {
611  return set_bit(null_mask(), element_index);
612  }
613 
627  __device__ void set_null(size_type element_index) const noexcept
628  {
629  return clear_bit(null_mask(), element_index);
630  }
631 
642  __device__ void set_mask_word(size_type word_index, bitmask_type new_word) const noexcept
643  {
644  null_mask()[word_index] = new_word;
645  }
646 
653  static std::size_t extent(mutable_column_view source_view);
654 
661  void destroy();
662 
663  private:
664  mutable_column_device_view* d_children{};
665  size_type _num_children{};
669 
678  mutable_column_device_view(mutable_column_view source);
679 };
680 
692 template <>
693 __device__ inline string_view const column_device_view::element<string_view>(
694  size_type element_index) const noexcept
695 {
696  size_type index = element_index + offset(); // account for this view's _offset
697  const int32_t* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
698  const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
699  size_type offset = d_offsets[index];
700  return string_view{d_strings + offset, d_offsets[index + 1] - offset};
701 }
702 
709  template <
710  typename IndexType,
711  std::enable_if_t<is_index_type<IndexType>() and std::is_unsigned<IndexType>::value>* = nullptr>
712  __device__ size_type operator()(column_device_view const& input, size_type index)
713  {
714  return static_cast<size_type>(input.element<IndexType>(index));
715  }
716  template <typename IndexType,
717  typename... Args,
718  std::enable_if_t<not(is_index_type<IndexType>() and
719  std::is_unsigned<IndexType>::value)>* = nullptr>
720  __device__ size_type operator()(Args&&... args)
721  {
722  release_assert(false and "dictionary indices must be an unsigned integral type");
723  return 0;
724  }
725 };
726 
751 template <>
752 __device__ inline dictionary32 const column_device_view::element<dictionary32>(
753  size_type element_index) const noexcept
754 {
755  size_type index = element_index + offset(); // account for this view's _offset
756  auto const indices = d_children[0];
757  return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)};
758 }
759 
769 template <>
770 __device__ inline numeric::decimal32 const column_device_view::element<numeric::decimal32>(
771  size_type element_index) const noexcept
772 {
773  using namespace numeric;
774  auto const scale = scale_type{_type.scale()};
775  return decimal32{scaled_integer<int32_t>{data<int32_t>()[element_index], scale}};
776 }
777 
787 template <>
788 __device__ inline numeric::decimal64 const column_device_view::element<numeric::decimal64>(
789  size_type element_index) const noexcept
790 {
791  using namespace numeric;
792  auto const scale = scale_type{_type.scale()};
793  return decimal64{scaled_integer<int64_t>{data<int64_t>()[element_index], scale}};
794 }
795 
796 namespace detail {
797 
804 __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restrict__ source,
805  size_type destination_word_index,
806  size_type source_begin_bit,
807  size_type source_end_bit)
808 {
809  size_type source_word_index = destination_word_index + word_index(source_begin_bit);
810  bitmask_type curr_word = source[source_word_index];
811  bitmask_type next_word = 0;
812  if (word_index(source_end_bit) >
813  word_index(source_begin_bit +
814  destination_word_index * detail::size_in_bits<bitmask_type>())) {
815  next_word = source[source_word_index + 1];
816  }
817  return __funnelshift_r(curr_word, next_word, source_begin_bit);
818 }
819 
834 template <typename T>
837 
842  value_accessor(column_device_view const& _col) : col{_col}
843  {
844  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
845  }
846 
847  __device__ T operator()(cudf::size_type i) const { return col.element<T>(i); }
848 };
849 
868 template <typename T, bool has_nulls = false>
871 
876  pair_accessor(column_device_view const& _col) : col{_col}
877  {
878  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
879  if (has_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
880  }
881 
882  CUDA_DEVICE_CALLABLE
883  thrust::pair<T, bool> operator()(cudf::size_type i) const
884  {
885  return {col.element<T>(i), (has_nulls ? col.is_valid_nocheck(i) : true)};
886  }
887 };
888 
889 template <typename T>
892 
898  {
899  CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
900  }
901 
902  __device__ T& operator()(cudf::size_type i) { return col.element<T>(i); }
903 };
904 
930 template <typename ColumnDeviceView, typename ColumnViewIterator>
931 ColumnDeviceView* child_columns_to_device_array(ColumnViewIterator child_begin,
932  ColumnViewIterator child_end,
933  void* h_ptr,
934  void* d_ptr)
935 {
936  ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
937  auto num_children = std::distance(child_begin, child_end);
938  if (num_children > 0) {
939  // The beginning of the memory must be the fixed-sized ColumnDeviceView
940  // struct objects in order for d_children to be used as an array.
941  auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
942  auto d_column = d_children;
943 
944  // Any child data is assigned past the end of this array: h_end and d_end.
945  auto h_end = reinterpret_cast<int8_t*>(h_column + num_children);
946  auto d_end = reinterpret_cast<int8_t*>(d_column + num_children);
947  std::for_each(child_begin, child_end, [&](auto const& col) {
948  // inplace-new each child into host memory
949  new (h_column) ColumnDeviceView(col, h_end, d_end);
950  h_column++; // advance to next child
951  // update the pointers for holding this child column's child data
952  auto col_child_data_size = ColumnDeviceView::extent(col) - sizeof(ColumnDeviceView);
953  h_end += col_child_data_size;
954  d_end += col_child_data_size;
955  });
956  }
957  return d_children;
958 }
959 
960 } // namespace detail
961 } // namespace cudf
cudf::detail::column_device_view_base::_data
void const * _data
Pointer to device memory containing elements.
Definition: column_device_view.cuh:223
cudf::detail::mutable_value_accessor::mutable_value_accessor
mutable_value_accessor(mutable_column_device_view &_col)
constructor
Definition: column_device_view.cuh:897
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:401
bit.hpp
Utilities for bit and bitmask operations.
fixed_point.hpp
Class definition for fixed point data type.
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:503
cudf::column
A container of nullable device data as a column of elements.
Definition: column.hpp:45
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:418
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:226
numeric::fixed_point
A type for representing a number with a fixed amount of precision.
Definition: fixed_point.hpp:212
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:281
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:128
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:254
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:388
rmm::cuda_stream_view
cudf::mutable_column_device_view::element
__device__ T & element(size_type element_index) noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:534
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:642
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:215
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::column_device_view::element
__device__ T const element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:287
cudf::mutable_column_device_view::set_valid
__device__ void set_valid(size_type element_index) const noexcept
Updates the null mask to indicate that the specified element is valid.
Definition: column_device_view.cuh:609
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:119
cudf::column_device_view::count_it
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
Definition: column_device_view.cuh:295
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:201
cudf::detail::mutable_value_accessor
Definition: column_device_view.cuh:890
cudf::detail::pair_accessor::pair_accessor
pair_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:876
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:931
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:77
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:368
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:150
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:554
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:520
numeric::scaled_integer
Helper struct for constructing fixed_point when value is already shifted.
Definition: fixed_point.hpp:196
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:167
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:869
cudf::detail::pair_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:870
cudf::data_type
Indicator for the logical data type of an element in a column.
Definition: types.hpp:235
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::end
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
Definition: column_device_view.cuh:324
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:447
cudf::detail::get_mask_offset_word
__device__ bitmask_type get_mask_offset_word(bitmask_type const *__restrict__ source, size_type destination_word_index, size_type source_begin_bit, size_type source_end_bit)
Convenience function to get offset word from a bitmask.
Definition: column_device_view.cuh:804
cudf::detail::value_accessor::col
column_device_view const col
column view of column in device
Definition: column_device_view.cuh:836
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:103
cudf
cuDF interfaces
Definition: aggregation.hpp:34
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:353
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:185
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:309
cudf::detail::column_device_view_base::_type
data_type _type
Element type.
Definition: column_device_view.cuh:221
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:95
cudf::detail::value_accessor
value accessor of column without null bitmask A unary functor returns scalar value at id....
Definition: column_device_view.cuh:835
cudf::detail::column_device_view_base::_null_mask
bitmask_type const * _null_mask
Definition: column_device_view.cuh:224
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:54
cudf::mutable_column_device_view::set_null
__device__ void set_null(size_type element_index) const noexcept
Updates the null mask to indicate that the specified element is null.
Definition: column_device_view.cuh:627
list_view.cuh
Class definition for cudf::list_view.
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:546
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:424
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:134
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::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:590
cudf::detail::column_device_view_base::type
__host__ __device__ data_type type() const noexcept
Returns the element type.
Definition: column_device_view.cuh:108
cudf::detail::value_accessor::value_accessor
value_accessor(column_device_view const &_col)
constructor
Definition: column_device_view.cuh:842
cudf::mutable_column_device_view::end
std::enable_if_t< is_fixed_width< T >), iterator< T > > end()
Return one past the last element after underlying data is casted to the specified type.
Definition: column_device_view.cuh:579
cudf::detail::column_device_view_base::_size
cudf::size_type _size
Number of elements.
Definition: column_device_view.cuh:222
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:104
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
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:428
cudf::detail::mutable_value_accessor::col
mutable_column_device_view col
mutable column view of column in device
Definition: column_device_view.cuh:891
cudf::mutable_column_device_view::begin
std::enable_if_t< is_fixed_width< T >), iterator< T > > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
Definition: column_device_view.cuh:566
cudf::data_type::id
CUDA_HOST_DEVICE_CALLABLE type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:265
numeric
fixed_point and supporting types
Definition: fixed_point.hpp:32
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:334
cudf::index_element_fn
Dispatch functor for resolving the index value for a dictionary element.
Definition: column_device_view.cuh:708