Public Types | Public Member Functions | Static Public Member Functions | List of all members
cudf::mutable_column_device_view Class Reference

A mutable, 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::mutable_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 iterator = thrust::transform_iterator< detail::mutable_value_accessor< T >, count_it >
 

Public Member Functions

 mutable_column_device_view (mutable_column_device_view const &)=default
 
 mutable_column_device_view (mutable_column_device_view &&)=default
 
mutable_column_device_viewoperator= (mutable_column_device_view const &)=default
 
mutable_column_device_viewoperator= (mutable_column_device_view &&)=default
 
 mutable_column_device_view (mutable_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 = void>
__host__ __device__ T * head () const noexcept
 Returns pointer to the base device memory allocation casted to the specified type. More...
 
template<typename T >
__host__ __device__ T * data () const noexcept
 Returns the underlying data casted to the specified type, plus the offset. More...
 
template<typename T >
__device__ T & element (size_type element_index) noexcept
 Returns reference to element at the specified index. More...
 
__host__ __device__ bitmask_type * null_mask () const noexcept
 Returns raw pointer to the underlying bitmask allocation. More...
 
template<typename T >
std::enable_if_t< is_fixed_width< T >), iterator< T > > begin ()
 Return first element (accounting for offset) after underlying data is casted to the specified type. More...
 
template<typename T >
std::enable_if_t< is_fixed_width< T >), iterator< T > > end ()
 Return one past the last element after underlying data is casted to the specified type. More...
 
__device__ mutable_column_device_view child (size_type child_index) const noexcept
 Returns the specified child. More...
 
__device__ void set_valid (size_type element_index) const noexcept
 Updates the null mask to indicate that the specified element is valid. More...
 
__device__ void set_null (size_type element_index) const noexcept
 Updates the null mask to indicate that the specified element is null. More...
 
__device__ void set_mask_word (size_type word_index, bitmask_type new_word) const noexcept
 Updates the specified bitmask word in the null_mask() with a new word. More...
 
void destroy ()
 Destroy the mutable_column_device_view object. 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>
__host__ __device__ T const * head () const noexcept
 Returns pointer to the base device memory allocation casted to the specified type. More...
 
template<typename 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

static std::unique_ptr< mutable_column_device_view, std::function< void(mutable_column_device_view *)> > create (mutable_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 (mutable_column_view 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...
 

Additional Inherited Members

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

A mutable, non-owning view of device data as a column of elements that is trivially copyable and usable in CUDA device code.

Definition at line 447 of file column_device_view.cuh.

Constructor & Destructor Documentation

◆ mutable_column_device_view()

cudf::mutable_column_device_view::mutable_column_device_view ( mutable_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.

Member Function Documentation

◆ begin()

template<typename T >
std::enable_if_t<is_fixed_width<T>), iterator<T> > cudf::mutable_column_device_view::begin ( )
inline

Return first element (accounting for offset) after underlying data is casted to the specified type.

Template Parameters
TThe desired type
Returns
T* Pointer to the first element after casting

Definition at line 566 of file column_device_view.cuh.

◆ child()

__device__ mutable_column_device_view cudf::mutable_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 590 of file column_device_view.cuh.

◆ create()

static std::unique_ptr<mutable_column_device_view, std::function<void(mutable_column_device_view*)> > cudf::mutable_column_device_view::create ( mutable_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<mutable_column_device_view> with a custom deleter to free the device memory allocated for the children.

A mutable_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 mutable_column_device_view that makes the data from source_view available in device memory.

◆ data()

template<typename T >
__host__ __device__ T* cudf::mutable_column_device_view::data ( ) const
inlinenoexcept

Returns the underlying data casted to the specified type, plus the offset.

Note
If offset() == 0, then head<T>() == data<T>()

This pointer is undefined for columns with children.

Template Parameters
TThe type to cast to
Returns
T* Typed pointer to underlying data, including the offset

Definition at line 520 of file column_device_view.cuh.

◆ destroy()

void cudf::mutable_column_device_view::destroy ( )

Destroy the mutable_column_device_view object.

Note
Does not free the column data, simply frees the device memory allocated to hold the child views.

◆ element()

template<typename T >
__device__ T& cudf::mutable_column_device_view::element ( size_type  element_index)
inlinenoexcept

Returns reference to element at the specified index.

This function accounts for the offset.

Template Parameters
TThe element type
Parameters
element_indexPosition of the desired element

Definition at line 534 of file column_device_view.cuh.

◆ end()

template<typename T >
std::enable_if_t<is_fixed_width<T>), iterator<T> > cudf::mutable_column_device_view::end ( )
inline

Return one past the last element after underlying data is casted to the specified type.

Template Parameters
TThe desired type
Returns
T const* Pointer to one past the last element after casting

Definition at line 579 of file column_device_view.cuh.

◆ extent()

static std::size_t cudf::mutable_column_device_view::extent ( mutable_column_view  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.

◆ head()

template<typename T = void>
__host__ __device__ T* cudf::mutable_column_device_view::head ( ) const
inlinenoexcept

Returns pointer to the base device memory allocation casted to the specified type.

Note
If offset() == 0, then head<T>() == data<T>()
It should be rare to need to access the head<T>() allocation of a column, and instead, accessing the elements should be done via data<T>().
Template Parameters
Thetype to cast to
Returns
T* Typed pointer to underlying data

Definition at line 503 of file column_device_view.cuh.

◆ null_mask()

__host__ __device__ bitmask_type* cudf::mutable_column_device_view::null_mask ( ) const
inlinenoexcept

Returns raw pointer to the underlying bitmask allocation.

Note
This function does not account for the offset().
If null_count() == 0, this may return nullptr.

Definition at line 546 of file column_device_view.cuh.

◆ set_mask_word()

__device__ void cudf::mutable_column_device_view::set_mask_word ( size_type  word_index,
bitmask_type  new_word 
) const
inlinenoexcept

Updates the specified bitmask word in the null_mask() with a new word.

Note
It is undefined behavior to call this function if nullable() == false.
Parameters
element_indexThe index of the element to update
new_elementThe new bitmask element

Definition at line 642 of file column_device_view.cuh.

◆ set_null()

__device__ void cudf::mutable_column_device_view::set_null ( size_type  element_index) const
inlinenoexcept

Updates the null mask to indicate that the specified element is null.

Note
This operation requires a global atomic operation. Therefore, it is not recommended to use this function in performance critical regions. When possible, it is more efficient to compute and update an entire word at once using set_mask_word.
It is undefined behavior to call this function if nullable() == false.
Parameters
element_indexThe index of the element to update

Definition at line 627 of file column_device_view.cuh.

◆ set_valid()

__device__ void cudf::mutable_column_device_view::set_valid ( size_type  element_index) const
inlinenoexcept

Updates the null mask to indicate that the specified element is valid.

Note
This operation requires a global atomic operation. Therefore, it is not recommended to use this function in performance critical regions. When possible, it is more efficient to compute and update an entire word at once using set_mask_word.
It is undefined behavior to call this function if nullable() == false.
Parameters
element_indexThe index of the element to update

Definition at line 609 of file column_device_view.cuh.


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