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

Public Types

using count_it = thrust::counting_iterator< size_type >
 Iterator for navigating this column.
 
template<typename T >
using const_iterator = thrust::transform_iterator< detail::value_accessor< T >, count_it >
 
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
 
 column_device_view (column_device_view &&)=default
 
column_device_viewoperator= (column_device_view const &)=default
 
column_device_viewoperator= (column_device_view &&)=default
 
 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...
 
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_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.
 
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
 
 column_device_view_base (column_device_view_base &&)=default
 
column_device_view_baseoperator= (column_device_view_base const &)=default
 
column_device_view_baseoperator= (column_device_view_base &&)=default
 
template<typename T = void, CUDF_ENABLE_IF(std::is_same_v< T, void > or is_rep_layout_compatible< T >()) >
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 >()) >
T const * data () const noexcept
 Returns the underlying data casted to the specified type, plus the offset. More...
 
size_type size () const noexcept
 Returns the number of elements in the column.
 
data_type type () const noexcept
 Returns the element type.
 
bool nullable () const noexcept
 Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask. More...
 
bitmask_type const * null_mask () const noexcept
 Returns raw pointer to the underlying bitmask allocation. More...
 
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>().
 
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 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.
 
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. 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
 column_device_view_base (data_type type, size_type size, void const *data, bitmask_type const *null_mask, size_type offset)
 

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 311 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 523 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.

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.

Definition at line 480 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 747 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 = rmm::cuda_stream_default 
)
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

Definition at line 348 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 365 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};

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 424 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 442 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.

Definition at line 498 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

◆ num_child_columns()

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

◆ 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 thrust::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()});
}

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

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.

Definition at line 578 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 thrust::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.

Definition at line 657 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.

Definition at line 606 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.

Definition at line 677 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.

Definition at line 636 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.

Definition at line 697 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 771 of file column_device_view.cuh.


The documentation for this class was generated from the following file:
cudf::nullate::DYNAMIC
Definition: column_device_view.cuh:58
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:300
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::nullate::NO
Definition: column_device_view.cuh:56
cudf::detail::column_view_base::has_nulls
bool has_nulls() const
Indicates if the column contains null elements, i.e., null_count() > 0
Definition: column_view.hpp:186
cudf::nullate::YES
Definition: column_device_view.cuh:54