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:

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_view &	operator= (column_device_view const &)=default

column_device_view &	operator= (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 >()) >
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 >) >
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_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(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(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_base &	operator= (column_device_view_base const &)=default

column_device_view_base &	operator= (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_view *	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 {}

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

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.

◆ 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_index The 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_view	The `column_view` to make usable in device code
stream	CUDA 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

T	The element type

Parameters

element_index Position 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_index Position 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_index Position 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_index Position 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_view The 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_error	if the column is not nullable and `has_nulls` evaluates to true.
cudf::logic_error	if 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_error	if the column is not nullable and `has_nulls` is true
cudf::logic_error	if 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_error	if tparam `has_nulls == true` and `nullable() == false`
cudf::logic_error	if 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_error	if tparam `has_nulls == true` and `nullable() == false`
cudf::logic_error	if 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_error	if tparam `has_nulls == true` and `nullable() == false`
cudf::logic_error	if 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_error	if tparam `has_nulls == true` and `nullable() == false`
cudf::logic_error	if 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:

column_device_view.cuh

Public Types

Public Member Functions

Static Public Member Functions

Protected Member Functions

Protected Attributes

Detailed Description

Member Typedef Documentation

◆ const_pair_rep_iterator

Constructor & Destructor Documentation

◆ column_device_view() [1/2]

◆ column_device_view() [2/2]

Member Function Documentation

◆ begin()

◆ child()

◆ create()

◆ destroy()

◆ element() [1/4]

◆ element() [2/4]

◆ element() [3/4]

◆ element() [4/4]

◆ end()

◆ extent()

◆ num_child_columns()

◆ optional_begin()

◆ optional_end()

◆ pair_begin()

◆ pair_end()

◆ pair_rep_begin()

◆ pair_rep_end()

Member Data Documentation

◆ d_children