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, CUDF_ENABLE_IF(std::is_same< T, void >::value or is_rep_layout_compatible< T >()) >
__host__ __device__ T * 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 * data () const noexcept
 Returns the underlying data casted to the specified type, plus the offset. 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...
 
__host__ __device__ bitmask_type * null_mask () const noexcept
 Returns raw pointer to the underlying bitmask allocation. More...
 
template<typename T , CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor< T >()) >
iterator< T > begin ()
 Return first element (accounting for offset) after underlying data is casted to the specified type. More...
 
template<typename T , CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor< 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_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, 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

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...
 
template<typename T >
static constexpr bool has_element_accessor ()
 For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
 
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 681 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 , CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor< 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.

This function does not participate in overload resolution if mutable_column_device_view::has_element_accessor<T>() is false.

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

Definition at line 824 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 851 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 , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) >
__host__ __device__ T* cudf::mutable_column_device_view::data ( ) const
inlinenoexcept

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

This function does not participate in overload resolution if is_rep_layout_compatible<T> is false.

Note
If offset() == 0, then head<T>() == data<T>()
Template Parameters
TThe type to cast to
Returns
T* Typed pointer to underlying data, including the offset

Definition at line 759 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 , CUDF_ENABLE_IF(is_rep_layout_compatible< T >()) >
__device__ T& cudf::mutable_column_device_view::element ( size_type  element_index) const
inlinenoexcept

Returns reference to element at the specified index.

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 778 of file column_device_view.cuh.

◆ end()

template<typename T , CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor< 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.

This function does not participate in overload resolution if mutable_column_device_view::has_element_accessor<T>() is false.

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

Definition at line 840 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, CUDF_ENABLE_IF(std::is_same< T, void >::value or is_rep_layout_compatible< T >()) >
__host__ __device__ T* cudf::mutable_column_device_view::head ( ) const
inlinenoexcept

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

This function will only participate in overload resolution if is_rep_layout_compatible<T>() or std::is_same<T,void>::value are true.

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 741 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 801 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 906 of file column_device_view.cuh.


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