An immutable, non-owning view of device data as a column of elements that is trivially copyable and usable in CUDA device code. More...
Public Types | |
| using | base = column_device_view_core |
| Base type. | |
| using | count_it = thrust::counting_iterator< size_type > |
| Counting iterator. | |
| template<typename T > | |
| using | const_iterator = thrust::transform_iterator< detail::value_accessor< T >, count_it > |
| Iterator for navigating this column. | |
| template<typename T , typename Nullate > | |
| using | const_optional_iterator = thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > |
| Optional iterator for navigating this column. | |
| template<typename T , bool has_nulls> | |
| using | const_pair_iterator = thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > |
| Pair iterator for navigating this column. | |
| template<typename T , bool has_nulls> | |
| using | const_pair_rep_iterator = thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > |
| Pair rep iterator for navigating this column. More... | |
Public Member Functions | |
| column_device_view (column_device_view const &)=default | |
| Copy constructor. | |
| column_device_view (column_device_view &&)=default | |
| Move constructor. | |
| column_device_view & | operator= (column_device_view const &)=default |
| Copy assignment operator. More... | |
| column_device_view & | operator= (column_device_view &&)=default |
| Move assignment operator. More... | |
| 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 objects and the device memory pointer (d_ptr) as a base for any child object pointers. More... | |
| 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. More... | |
| template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) > | |
| T | element (size_type element_index) const noexcept |
| Returns a copy of the element at the specified index. More... | |
| template<typename T , CUDF_ENABLE_IF(cuda::std::is_same_v< T, string_view >) > | |
| T | element (size_type element_index) const noexcept |
Returns string_view to the string element at the specified index. More... | |
| template<typename T , CUDF_ENABLE_IF(cudf::is_fixed_point< T >()) > | |
| T | element (size_type element_index) const noexcept |
Returns a numeric::fixed_point element at the specified index for a fixed_point column. More... | |
| template<typename T , CUDF_ENABLE_IF(cuda::std::is_same_v< T, dictionary32 >) > | |
| T | element (size_type element_index) const noexcept |
Returns dictionary32 element at the specified index for a dictionary column. More... | |
| template<typename T , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| const_iterator< T > | begin () const |
| Return an iterator to the first element of the column. More... | |
| template<typename T , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| const_iterator< T > | end () const |
| Returns an iterator to the element following the last element of the column. More... | |
| template<typename T , typename Nullate , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| auto | optional_begin (Nullate has_nulls) const |
| Return an optional iterator to the first element of the column. More... | |
| template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| const_pair_iterator< T, has_nulls > | pair_begin () const |
| Return a pair iterator to the first element of the column. More... | |
| template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| const_pair_rep_iterator< T, has_nulls > | pair_rep_begin () const |
| Return a pair iterator to the first element of the column. More... | |
| template<typename T , typename Nullate , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| auto | optional_end (Nullate has_nulls) const |
| Return an optional iterator to the element following the last element of the column. More... | |
| template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| const_pair_iterator< T, has_nulls > | pair_end () const |
| Return a pair iterator to the element following the last element of the column. More... | |
| template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) > | |
| 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. More... | |
| void | destroy () |
Destroy the column_device_view object. More... | |
| column_device_view | child (size_type child_index) const noexcept |
| Returns the specified child. More... | |
| device_span< column_device_view const > | children () const noexcept |
| Returns a span containing the children of this column. More... | |
| CUDF_HOST_DEVICE size_type | num_child_columns () const noexcept |
| Returns the number of child columns. More... | |
Public Member Functions inherited from cudf::column_device_view_core | |
| column_device_view_core (column_device_view_core const &)=default | |
| Copy constructor. | |
| column_device_view_core (column_device_view_core &&)=default | |
| Move constructor. | |
| column_device_view_core & | operator= (column_device_view_core const &)=default |
| Copy assignment operator. More... | |
| column_device_view_core & | operator= (column_device_view_core &&)=default |
| Move assignment operator. More... | |
| column_device_view_core (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 objects and the device memory pointer (d_ptr) as a base for any child object pointers. More... | |
| CUDF_HOST_DEVICE column_device_view_core | slice (size_type offset, size_type size) const noexcept |
| Get a new raw_column_device_view which is a slice of this column. More... | |
| template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) > | |
| T | element (size_type element_index) const noexcept |
| Returns a copy of the element at the specified index. More... | |
| template<typename T , CUDF_ENABLE_IF(cuda::std::is_same_v< T, string_view >) > | |
| T | element (size_type element_index) const noexcept |
Returns string_view to the string element at the specified index. More... | |
| template<typename T , CUDF_ENABLE_IF(cudf::is_fixed_point< T >()) > | |
| T | element (size_type element_index) const noexcept |
Returns a numeric::fixed_point element at the specified index for a fixed_point column. More... | |
| column_device_view_core | child (size_type child_index) const noexcept |
| Returns the specified child. More... | |
| CUDF_HOST_DEVICE size_type | num_child_columns () const noexcept |
| Returns the number of child columns. More... | |
Public Member Functions inherited from cudf::detail::column_device_view_base | |
| column_device_view_base (column_device_view_base const &)=default | |
| Copy constructor. | |
| column_device_view_base (column_device_view_base &&)=default | |
| Move constructor. | |
| column_device_view_base & | operator= (column_device_view_base const &)=default |
| Copy assignment operator. More... | |
| column_device_view_base & | operator= (column_device_view_base &&)=default |
| Move assignment operator. More... | |
| template<typename T = void, CUDF_ENABLE_IF(cuda::std::is_same_v< T, void > or is_rep_layout_compatible< T >()) > | |
| CUDF_HOST_DEVICE T const * | head () const noexcept |
| Returns pointer to the base device memory allocation casted to the specified type. More... | |
| template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) > | |
| CUDF_HOST_DEVICE T const * | data () const noexcept |
| Returns the underlying data casted to the specified type, plus the offset. More... | |
| CUDF_HOST_DEVICE size_type | size () const noexcept |
| Returns the number of elements in the column. More... | |
| CUDF_HOST_DEVICE data_type | type () const noexcept |
| Returns the element type. More... | |
| CUDF_HOST_DEVICE bool | nullable () const noexcept |
| Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask. More... | |
| CUDF_HOST_DEVICE bitmask_type const * | null_mask () const noexcept |
| Returns raw pointer to the underlying bitmask allocation. More... | |
| CUDF_HOST_DEVICE size_type | offset () const noexcept |
Returns the index of the first element relative to the base memory allocation, i.e., what is returned from head<T>(). More... | |
| bool | is_valid (size_type element_index) const noexcept |
| Returns whether the specified element holds a valid value (i.e., not null). More... | |
| bool | is_valid_nocheck (size_type element_index) const noexcept |
| Returns whether the specified element holds a valid value (i.e., not null) More... | |
| bool | is_null (size_type element_index) const noexcept |
| Returns whether the specified element is null. More... | |
| bool | is_null_nocheck (size_type element_index) const noexcept |
| Returns whether the specified element is null. More... | |
| bitmask_type | get_mask_word (size_type word_index) const noexcept |
Returns the specified bitmask word from the null_mask(). More... | |
Static Public Member Functions | |
| template<typename T > | |
| static constexpr CUDF_HOST_DEVICE bool | has_element_accessor () |
For a given T, indicates if column_device_view::element<T>() has a valid overload. More... | |
| static std::unique_ptr< column_device_view, std::function< void(column_device_view *)> > | create (column_view source_view, rmm::cuda_stream_view stream=cudf::get_default_stream()) |
| Factory to construct a column view that is usable in device memory. More... | |
| 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 and its children. More... | |
Additional Inherited Members | |
Static Public Attributes inherited from cudf::detail::column_device_view_base | |
| static constexpr size_type | offsets_column_index {0} |
| Child index of the offsets column. | |
Protected Member Functions inherited from cudf::column_device_view_core | |
| CUDF_HOST_DEVICE | column_device_view_core (data_type type, size_type size, void const *data, bitmask_type const *null_mask, size_type offset, column_device_view_core *children, size_type num_children) |
| Creates an instance of this class using pre-existing device memory pointers to data, nullmask, and offset. More... | |
Protected Member Functions inherited from cudf::detail::column_device_view_base | |
| CUDF_HOST_DEVICE | column_device_view_base (data_type type, size_type size, void const *data, bitmask_type const *null_mask, size_type offset) |
| Constructs a column with the specified type, size, data, nullmask and offset. More... | |
Protected Attributes inherited from cudf::column_device_view_core | |
| column_device_view_core * | d_children {} |
| size_type | _num_children {} |
| The number of child columns. | |
Protected Attributes inherited from cudf::detail::column_device_view_base | |
| data_type | _type {type_id::EMPTY} |
| Element type. | |
| cudf::size_type | _size {} |
| Number of elements. | |
| void const * | _data {} |
| Pointer to device memory containing elements. | |
| bitmask_type const * | _null_mask {} |
| size_type | _offset {} |
An immutable, non-owning view of device data as a column of elements that is trivially copyable and usable in CUDA device code.
Definition at line 38 of file column_device_view.cuh.
| using cudf::column_device_view::const_pair_rep_iterator = thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>, count_it> |
Pair rep iterator for navigating this column.
Each row value is accessed in its representative form.
Definition at line 291 of file column_device_view.cuh.
| cudf::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 objects and the device memory pointer (d_ptr) as a base for any child object pointers.
| column | Column view from which to create this instance. |
| h_ptr | Host memory pointer on which to place any child data. |
| d_ptr | Device memory pointer on which to base any child pointers. |
|
inline |
Return an iterator to the first element of the column.
This iterator only supports columns where has_nulls() == false. Using it with columns where has_nulls() == true will result in undefined behavior when accessing null elements.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
For columns with null elements, use make_null_replacement_iterator.
| T | Type of the elements in the column |
Definition at line 246 of file column_device_view.cuh.
|
inlinenoexcept |
Returns the specified child.
| child_index | The index of the desired child |
column_view Definition at line 529 of file column_device_view.cuh.
|
inlinenoexcept |
Returns a span containing the children of this column.
Definition at line 539 of file column_device_view.cuh.
|
static |
Factory to construct a column view that is usable in device memory.
Allocates and copies views of source_view's children to device memory to make them accessible in device code.
If source_view.num_children() == 0, then no device memory is allocated.
Returns a std::unique_ptr<column_device_view> with a custom deleter to free the device memory allocated for the children.
A column_device_view should be passed by value into GPU kernels.
| source_view | The column_view to make usable in device code |
| stream | CUDA stream used for device memory operations for children columns. |
unique_ptr to a column_device_view that makes the data from source_view available in device memory. | void cudf::column_device_view::destroy | ( | ) |
Destroy the column_device_view object.
|
inlinenoexcept |
Returns a copy of the element at the specified index.
If the element at the specified index is NULL, i.e., is_null(element_index) == true, then any attempt to use the result will lead to undefined behavior.
This function accounts for the offset.
This function does not participate in overload resolution if is_rep_layout_compatible<T> is false. Specializations of this function may exist for types T where is_rep_layout_compatible<T> is false.
| T | The element type |
| element_index | Position of the desired element |
Definition at line 116 of file column_device_view.cuh.
|
inlinenoexcept |
Returns string_view to the string element at the specified index.
If the element at the specified index is NULL, i.e., is_null(element_index) == true, then any attempt to use the result will lead to undefined behavior.
This function accounts for the offset.
| element_index | Position of the desired string element |
Definition at line 133 of file column_device_view.cuh.
|
inlinenoexcept |
Returns a numeric::fixed_point element at the specified index for a fixed_point column.
If the element at the specified index is NULL, i.e., is_null(element_index) == true, then any attempt to use the result will lead to undefined behavior.
| element_index | Position of the desired element |
Definition at line 149 of file column_device_view.cuh.
|
inlinenoexcept |
Returns dictionary32 element at the specified index for a dictionary column.
dictionary32 is a strongly typed wrapper around an int32_t value that holds the offset into the dictionary keys for the specified element.
For example, given a dictionary column d with:
If the element at the specified index is NULL, i.e., is_null(element_index) == true, then any attempt to use the result will lead to undefined behavior.
This function accounts for the offset.
| element_index | Position of the desired element |
Definition at line 203 of file column_device_view.cuh.
|
inline |
Returns an iterator to the element following the last element of the column.
This iterator only supports columns where has_nulls() == false. Using it with columns where has_nulls() == true will result in undefined behavior when accessing null elements.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
For columns with null elements, use make_null_replacement_iterator.
Definition at line 266 of file column_device_view.cuh.
|
static |
Return the size in bytes of the amount of memory needed to hold a device view of the specified column and its children.
| source_view | The column_view to use for this calculation. |
|
inlinestaticconstexpr |
For a given T, indicates if column_device_view::element<T>() has a valid overload.
| T | The element type |
true if column_device_view::element<T>() has a valid overload, false otherwise Definition at line 217 of file column_device_view.cuh.
|
inlinenoexcept |
Returns the number of child columns.
Definition at line 549 of file column_device_view.cuh.
|
default |
Move assignment operator.
|
default |
Copy assignment operator.
|
inline |
Return an optional iterator to the first element of the column.
Dereferencing the returned iterator returns a cuda::std::optional<T>.
The element of this iterator contextually converts to bool. The conversion returns true if the object contains a value and false if it does not contain a value.
Calling this method with nullate::DYNAMIC defers the assumption of nullability to runtime with the caller indicating if the column has nulls. The nullate::DYNAMIC is useful when an algorithm is going to execute on multiple iterators and all the combinations of iterator types are not required at compile time.
Calling this method with nullate::YES means that the column supports nulls and the optional returned might not contain a value.
Calling this method with nullate::NO means that the column has no null values and the optional returned will always contain a value.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if the column is not nullable and has_nulls evaluates to true. |
| cudf::logic_error | if column datatype and Element type mismatch. |
| T | The type of elements in the column |
| Nullate | A cudf::nullate type describing how to check for nulls |
| has_nulls | A cudf::nullate type describing how to check for nulls |
Definition at line 351 of file column_device_view.cuh.
|
inline |
Return an optional iterator to the element following the last element of the column.
The returned iterator represents a cuda::std::optional<T> element.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if the column is not nullable and has_nulls is true |
| cudf::logic_error | if column datatype and Element type mismatch. |
| T | The type of elements in the column |
| Nullate | A cudf::nullate type describing how to check for nulls |
| has_nulls | A cudf::nullate type describing how to check for nulls |
Definition at line 438 of file column_device_view.cuh.
|
inline |
Return a pair iterator to the first element of the column.
Dereferencing the returned iterator returns a thrust::pair<T, bool>.
If an element at position i is valid (or has_nulls == false), then for p = *(iter + i), p.first contains the value of the element at i and p.second == true.
Else, if the element at i is null, then the value of p.first is undefined and p.second == false.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if tparam has_nulls == true and nullable() == false |
| cudf::logic_error | if column datatype and Element type mismatch. |
Definition at line 381 of file column_device_view.cuh.
|
inline |
Return a pair iterator to the element following the last element of the column.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if tparam has_nulls == true and nullable() == false |
| cudf::logic_error | if column datatype and Element type mismatch. |
Definition at line 458 of file column_device_view.cuh.
|
inline |
Return a pair iterator to the first element of the column.
Dereferencing the returned iterator returns a thrust::pair<rep_type, bool>, where rep_type is device_storage_type<T>, the type used to store the value on the device.
If an element at position i is valid (or has_nulls == false), then for p = *(iter + i), p.first contains the value of the element at i and p.second == true.
Else, if the element at i is null, then the value of p.first is undefined and p.second == false.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if tparam has_nulls == true and nullable() == false |
| cudf::logic_error | if column datatype and Element type mismatch. |
Definition at line 413 of file column_device_view.cuh.
|
inline |
Return a pair iterator to the element following the last element of the column.
This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.
| cudf::logic_error | if tparam has_nulls == true and nullable() == false |
| cudf::logic_error | if column datatype and Element type mismatch. |
Definition at line 479 of file column_device_view.cuh.
|
inlinenoexcept |
Get a new column_device_view which is a slice of this column.
Example:
| offset | The index of the first element in the slice |
| size | The number of elements in the slice |
Definition at line 86 of file column_device_view.cuh.