Classes | Public Types | Public Member Functions | Static Public Member Functions | Protected Member Functions | Protected Attributes | List of all members
cudf::column_device_view Class Reference

An immutable, non-owning view of device data as a column of elements that is trivially copyable and usable in CUDA device code. More...

Inheritance diagram for cudf::column_device_view:
cudf::detail::column_device_view_base cudf::detail::lists_column_device_view cudf::detail::structs_column_device_view

Public Types

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_viewoperator= (column_device_view const &)=default
 Copy assignment operator. More...
 
column_device_viewoperator= (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 >()) >
element (size_type element_index) const noexcept
 Returns reference to element at the specified index. More...
 
template<typename T , CUDF_ENABLE_IF(std::is_same_v< T, string_view >) >
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(std::is_same_v< T, dictionary32 >) >
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(cudf::is_fixed_point< 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(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_nullspair_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_nullspair_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_nullspair_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_nullspair_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::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_baseoperator= (column_device_view_base const &)=default
 Copy assignment operator. More...
 
column_device_view_baseoperator= (column_device_view_base &&)=default
 Move assignment operator. More...
 
template<typename T = void, CUDF_ENABLE_IF(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 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 it's children. More...
 

Protected Member Functions

 column_device_view (column_view source)
 Construct's a column_device_view from a column_view populating all but the children. 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

column_device_viewd_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 {}
 

Detailed Description

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 352 of file column_device_view.cuh.

Member Typedef Documentation

◆ const_pair_rep_iterator

template<typename T , bool has_nulls>
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 611 of file column_device_view.cuh.

Constructor & Destructor Documentation

◆ column_device_view() [1/2]

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.

Parameters
columnColumn view from which to create this instance.
h_ptrHost memory pointer on which to place any child data.
d_ptrDevice memory pointer on which to base any child pointers.

◆ column_device_view() [2/2]

cudf::column_device_view::column_device_view ( column_view  source)
protected

Construct's a column_device_view from a column_view populating all but the children.

Note
This constructor is for internal use only. To create a column_device_view from a column_view, the column_device_view::create() function should be used.
Parameters
sourceThe column_view to use for this construction

Member Function Documentation

◆ begin()

template<typename T , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_iterator<T> cudf::column_device_view::begin ( ) const
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.

Template Parameters
TType of the elements in the column
Returns
An iterator to the first element of the column

Definition at line 566 of file column_device_view.cuh.

◆ child()

column_device_view cudf::column_device_view::child ( size_type  child_index) const
inlinenoexcept

Returns the specified child.

Parameters
child_indexThe index of the desired child
Returns
column_view The requested child column_view

Definition at line 849 of file column_device_view.cuh.

◆ children()

device_span<column_device_view const> cudf::column_device_view::children ( ) const
inlinenoexcept

Returns a span containing the children of this column.

Returns
A span containing the children of this column

Definition at line 859 of file column_device_view.cuh.

◆ create()

static std::unique_ptr<column_device_view, std::function<void(column_device_view*)> > cudf::column_device_view::create ( column_view  source_view,
rmm::cuda_stream_view  stream = cudf::get_default_stream() 
)
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.

Parameters
source_viewThe column_view to make usable in device code
streamCUDA stream used for device memory operations for children columns.
Returns
A unique_ptr to a column_device_view that makes the data from source_view available in device memory.

◆ destroy()

void cudf::column_device_view::destroy ( )

Destroy the column_device_view object.

Note
Does not free the column data, simply frees the device memory allocated to hold the child views.

◆ element() [1/4]

template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) >
T cudf::column_device_view::element ( size_type  element_index) const
inlinenoexcept

Returns reference to 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.

Template Parameters
TThe element type
Parameters
element_indexPosition of the desired element
Returns
reference to the element at the specified index

Definition at line 428 of file column_device_view.cuh.

◆ element() [2/4]

template<typename T , CUDF_ENABLE_IF(std::is_same_v< T, string_view >) >
T cudf::column_device_view::element ( size_type  element_index) const
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.

Parameters
element_indexPosition of the desired string element
Returns
string_view instance representing this element at this index

Definition at line 445 of file column_device_view.cuh.

◆ element() [3/4]

template<typename T , CUDF_ENABLE_IF(std::is_same_v< T, dictionary32 >) >
T cudf::column_device_view::element ( size_type  element_index) const
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:

{c++}
keys: {"foo", "bar", "baz"}
indices: {2, 0, 2, 1, 0}
d.element<dictionary32>(0) == dictionary32{2};
d.element<dictionary32>(1) == dictionary32{0};
dictionary_wrapper< int32_t > dictionary32
32-bit integer indexed dictionary wrapper
Definition: dictionary.hpp:217

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.

Parameters
element_indexPosition of the desired element
Returns
dictionary32 instance representing this element at this index

Definition at line 504 of file column_device_view.cuh.

◆ element() [4/4]

template<typename T , CUDF_ENABLE_IF(cudf::is_fixed_point< T >()) >
T cudf::column_device_view::element ( size_type  element_index) const
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.

Parameters
element_indexPosition of the desired element
Returns
numeric::fixed_point representing the element at this index

Definition at line 522 of file column_device_view.cuh.

◆ end()

template<typename T , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_iterator<T> cudf::column_device_view::end ( ) const
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.

Returns
An iterator to the element following the last element of the column

Definition at line 586 of file column_device_view.cuh.

◆ extent()

static std::size_t cudf::column_device_view::extent ( column_view const &  source_view)
static

Return the size in bytes of the amount of memory needed to hold a device view of the specified column and it's children.

Parameters
source_viewThe column_view to use for this calculation.
Returns
number of bytes to store device view in GPU memory

◆ has_element_accessor()

template<typename T >
static constexpr bool cudf::column_device_view::has_element_accessor ( )
inlinestaticconstexpr

For a given T, indicates if column_device_view::element<T>() has a valid overload.

Template Parameters
TThe element type
Returns
true if column_device_view::element<T>() has a valid overload, false otherwise

Definition at line 537 of file column_device_view.cuh.

◆ num_child_columns()

CUDF_HOST_DEVICE size_type cudf::column_device_view::num_child_columns ( ) const
inlinenoexcept

Returns the number of child columns.

Returns
The number of child columns

Definition at line 869 of file column_device_view.cuh.

◆ operator=() [1/2]

column_device_view& cudf::column_device_view::operator= ( column_device_view &&  )
default

Move assignment operator.

Returns
Reference to this object (after transferring ownership)

◆ operator=() [2/2]

column_device_view& cudf::column_device_view::operator= ( column_device_view const &  )
default

Copy assignment operator.

Returns
Reference to this object

◆ optional_begin()

template<typename T , typename Nullate , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
auto cudf::column_device_view::optional_begin ( Nullate  has_nulls) const
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.

template<typename T>
void some_function(cudf::column_view<T> const& col_view){
auto d_col = cudf::column_device_view::create(col_view);
// Create a `DYNAMIC` optional iterator
auto optional_iterator =
d_col->optional_begin<T>(cudf::nullate::DYNAMIC{col_view.has_nulls()});
}
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.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
bool has_nulls() const
Indicates if the column contains null elements, i.e., null_count() > 0
nullate::DYNAMIC defers the determination of nullability to run time rather than 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.

template<typename T, bool has_nulls>
void some_function(cudf::column_view<T> const& col_view){
auto d_col = cudf::column_device_view::create(col_view);
if constexpr(has_nulls) {
auto optional_iterator = d_col->optional_begin<T>(cudf::nullate::YES{});
//use optional_iterator
} else {
auto optional_iterator = d_col->optional_begin<T>(cudf::nullate::NO{});
//use optional_iterator
}
}
bool has_nulls(table_view const &view)
Returns True if the table has nulls in any of its columns.

This function does not participate in overload resolution if column_device_view::has_element_accessor<T>() is false.

Exceptions
cudf::logic_errorif the column is not nullable and has_nulls evaluates to true.
cudf::logic_errorif column datatype and Element type mismatch.
Template Parameters
TThe type of elements in the column
NullateA cudf::nullate type describing how to check for nulls
Parameters
has_nullsA cudf::nullate type describing how to check for nulls
Returns
An optional iterator to the first element of the column

Definition at line 671 of file column_device_view.cuh.

◆ optional_end()

template<typename T , typename Nullate , CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
auto cudf::column_device_view::optional_end ( Nullate  has_nulls) const
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.

Exceptions
cudf::logic_errorif the column is not nullable and has_nulls is true
cudf::logic_errorif column datatype and Element type mismatch.
Template Parameters
TThe type of elements in the column
NullateA cudf::nullate type describing how to check for nulls
Parameters
has_nullsA cudf::nullate type describing how to check for nulls
Returns
An optional iterator to the element following the last element of the column

Definition at line 758 of file column_device_view.cuh.

◆ pair_begin()

template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_pair_iterator<T, has_nulls> cudf::column_device_view::pair_begin ( ) const
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.

Exceptions
cudf::logic_errorif tparam has_nulls == true and nullable() == false
cudf::logic_errorif column datatype and Element type mismatch.
Returns
A pair iterator to the first element of the column

Definition at line 701 of file column_device_view.cuh.

◆ pair_end()

template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_pair_iterator<T, has_nulls> cudf::column_device_view::pair_end ( ) const
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.

Exceptions
cudf::logic_errorif tparam has_nulls == true and nullable() == false
cudf::logic_errorif column datatype and Element type mismatch.
Returns
A pair iterator to the element following the last element of the column

Definition at line 778 of file column_device_view.cuh.

◆ pair_rep_begin()

template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_pair_rep_iterator<T, has_nulls> cudf::column_device_view::pair_rep_begin ( ) const
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.

Exceptions
cudf::logic_errorif tparam has_nulls == true and nullable() == false
cudf::logic_errorif column datatype and Element type mismatch.
Returns
A pair iterator to the first element of the column

Definition at line 733 of file column_device_view.cuh.

◆ pair_rep_end()

template<typename T , bool has_nulls, CUDF_ENABLE_IF(column_device_view::has_element_accessor< T >()) >
const_pair_rep_iterator<T, has_nulls> cudf::column_device_view::pair_rep_end ( ) const
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.

Exceptions
cudf::logic_errorif tparam has_nulls == true and nullable() == false
cudf::logic_errorif column datatype and Element type mismatch.
Returns
A pair iterator to the element following the last element of the column

Definition at line 799 of file column_device_view.cuh.

◆ slice()

CUDF_HOST_DEVICE column_device_view cudf::column_device_view::slice ( size_type  offset,
size_type  size 
) const
inlinenoexcept

Get a new column_device_view which is a slice of this column.

Example:

// column = column_device_view([1, 2, 3, 4, 5, 6, 7])
auto c = column.slice(1, 3);
// c = column_device_view([2, 3, 4])
auto c1 = column.slice(2, 3);
// c1 = column_device_view([3, 4, 5])
Parameters
offsetThe index of the first element in the slice
sizeThe number of elements in the slice
Returns
A slice of this column

Definition at line 398 of file column_device_view.cuh.

Member Data Documentation

◆ d_children

column_device_view* cudf::column_device_view::d_children {}
protected

Array of column_device_view objects in device memory. Based on element type, children may contain additional data

Definition at line 901 of file column_device_view.cuh.


The documentation for this class was generated from the following file: