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 , 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 >()) >
__device__ 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< T, string_view >::value) >
__device__ 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(std::is_same< T, dictionary32 >::value) >
__device__ 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(std::is_same< T, numeric::decimal32 >::value) >
__device__ T element (size_type element_index) const noexcept
 Returns a numeric::decimal32 element at the specified index for a fixed_point column. More...
 
template<typename T , CUDF_ENABLE_IF(std::is_same< T, numeric::decimal64 >::value) >
__device__ T element (size_type element_index) const noexcept
 Returns a numeric::decimal64 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 , 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 , 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...
 
__device__ column_device_view child (size_type child_index) const noexcept
 Returns the specified child. More...
 
__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
 
 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< T, void >::value or is_rep_layout_compatible< T >()) >
__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 >()) >
__host__ __device__ T const * data () const noexcept
 Returns the underlying data casted to the specified type, plus the offset. More...
 
__host__ __device__ size_type size () const noexcept
 Returns the number of elements in the column.
 
__host__ __device__ data_type type () const noexcept
 Returns the element type.
 
__host__ __device__ bool nullable () const noexcept
 Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask. More...
 
__host__ __device__ bitmask_type const * null_mask () const noexcept
 Returns raw pointer to the underlying bitmask allocation. More...
 
__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>().
 
__device__ bool is_valid (size_type element_index) const noexcept
 Returns whether the specified element holds a valid value (i.e., not null). More...
 
__device__ bool is_valid_nocheck (size_type element_index) const noexcept
 Returns whether the specified element holds a valid value (i.e., not null) More...
 
__device__ bool is_null (size_type element_index) const noexcept
 Returns whether the specified element is null. More...
 
__device__ bool is_null_nocheck (size_type element_index) const noexcept
 Returns whether the specified element is null. More...
 
__device__ 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 276 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 500 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 464 of file column_device_view.cuh.

◆ child()

__device__ 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 645 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/5]

template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) >
__device__ 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 313 of file column_device_view.cuh.

◆ element() [2/5]

template<typename T , CUDF_ENABLE_IF(std::is_same< T, string_view >::value) >
__device__ 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 330 of file column_device_view.cuh.

◆ element() [3/5]

template<typename T , CUDF_ENABLE_IF(std::is_same< T, dictionary32 >::value) >
__device__ 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 391 of file column_device_view.cuh.

◆ element() [4/5]

template<typename T , CUDF_ENABLE_IF(std::is_same< T, numeric::decimal32 >::value) >
__device__ T cudf::column_device_view::element ( size_type  element_index) const
inlinenoexcept

Returns a numeric::decimal32 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::decimal32 representing the element at this index

Definition at line 409 of file column_device_view.cuh.

◆ element() [5/5]

template<typename T , CUDF_ENABLE_IF(std::is_same< T, numeric::decimal64 >::value) >
__device__ T cudf::column_device_view::element ( size_type  element_index) const
inlinenoexcept

Returns a numeric::decimal64 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::decimal64 representing the element at this index

Definition at line 427 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 482 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()

__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 655 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 525 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 575 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 555 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 595 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 658 of file column_device_view.cuh.


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