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 >
	Counting iterator.

template<typename T >
using	const_iterator = thrust::transform_iterator< detail::value_accessor< T >, count_it >
	Iterator for navigating this column.

template<typename T , typename Nullate >
using	const_optional_iterator = thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it >
	Optional iterator for navigating this column.

template<typename T , bool has_nulls>
using	const_pair_iterator = thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it >
	Pair iterator for navigating this column.

template<typename T , bool has_nulls>
using	const_pair_rep_iterator = thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it >
	Pair rep iterator for navigating this column. More...

Public Member Functions
	column_device_view (column_device_view const &)=default
	Copy constructor.

	column_device_view (column_device_view &&)=default
	Move constructor.

column_device_view &	operator= (column_device_view const &)=default
	Copy assignment operator. More...

column_device_view &	operator= (column_device_view &&)=default
	Move assignment operator. More...

	column_device_view (column_view column, void h_ptr, void d_ptr)
	Creates an instance of this class using the specified host memory pointer (h_ptr) to store child objects and the device memory pointer (d_ptr) as a base for any child object pointers. More...

CUDF_HOST_DEVICE column_device_view	slice (size_type offset, size_type size) const noexcept
	Get a new column_device_view which is a slice of this column. More...

template<typename T , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) >
T	element (size_type element_index) const noexcept
	Returns 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. More...

CUDF_HOST_DEVICE size_type	num_child_columns () const noexcept
	Returns the number of child columns. More...

Public Member Functions inherited from cudf::detail::column_device_view_base
	column_device_view_base (column_device_view_base const &)=default
	Copy constructor.

	column_device_view_base (column_device_view_base &&)=default
	Move constructor.

column_device_view_base &	operator= (column_device_view_base const &)=default
	Copy assignment operator. More...

column_device_view_base &	operator= (column_device_view_base &&)=default
	Move assignment operator. More...

template<typename T = void, CUDF_ENABLE_IF(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_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 350 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 609 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.

Parameters

source The 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

T	Type of the elements in the column

Returns: An iterator to the first element of the column

Definition at line 564 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 847 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 857 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_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

Returns: reference to the element at the specified index

Definition at line 426 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 443 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 502 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 520 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 584 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

◆ 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

T	The element type

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

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

Template Parameters

T	The type of elements in the column
Nullate	A cudf::nullate type describing how to check for nulls

Parameters

has_nulls A cudf::nullate type describing how to check for nulls

Returns: An optional iterator to the first element of the column

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

Template Parameters

T	The type of elements in the column
Nullate	A cudf::nullate type describing how to check for nulls

Parameters

has_nulls A 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 756 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.

Returns: A pair iterator to the first element of the column

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

Returns: A pair iterator to the element following the last element of the column

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

Returns: A pair iterator to the first element of the column

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

Returns: A pair iterator to the element following the last element of the column

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

offset	The index of the first element in the slice
size	The number of elements in the slice

Returns: A slice of this column

Definition at line 396 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 899 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()

◆ children()

◆ create()

◆ destroy()

◆ element() [1/4]

◆ element() [2/4]

◆ element() [3/4]

◆ element() [4/4]

◆ end()

◆ extent()

◆ has_element_accessor()

◆ num_child_columns()

◆ operator=() [1/2]

◆ operator=() [2/2]

◆ optional_begin()

◆ optional_end()

◆ pair_begin()

◆ pair_end()

◆ pair_rep_begin()

◆ pair_rep_end()

◆ slice()

Member Data Documentation

◆ d_children