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 , 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 >()) >
__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_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< 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_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 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

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 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_index The 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_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/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

T	The element type

Parameters

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

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

◆ element() [2/5]

◆ element() [3/5]

◆ element() [4/5]

◆ element() [5/5]

◆ end()

◆ extent()

◆ num_child_columns()

◆ pair_begin()

◆ pair_end()

◆ pair_rep_begin()

◆ pair_rep_end()

Member Data Documentation

◆ d_children