Go to the documentation of this file.
20 #include <cudf/detail/utilities/alignment.hpp>
23 #include <cudf/strings/string_view.cuh>
31 #include <rmm/cuda_stream_view.hpp>
33 #include <thrust/iterator/counting_iterator.h>
34 #include <thrust/iterator/transform_iterator.h>
76 template <
typename T =
void>
77 __host__ __device__ T
const*
head() const noexcept
79 return static_cast<T const*
>(
_data);
95 __host__ __device__ T
const*
data() const noexcept
103 __host__ __device__ size_type
size() const noexcept {
return _size; }
150 __device__
bool is_valid(size_type element_index)
const noexcept
185 __device__
bool is_null(size_type element_index)
const noexcept
240 template <
typename T>
241 struct value_accessor;
242 template <
typename T,
bool has_nulls>
243 struct pair_accessor;
244 template <
typename T>
245 struct mutable_value_accessor;
286 template <
typename T>
287 __device__ T
const element(size_type element_index)
const noexcept
289 return data<T>()[element_index];
295 using count_it = thrust::counting_iterator<size_type>;
296 template <
typename T>
297 using const_iterator = thrust::transform_iterator<detail::value_accessor<T>,
count_it>;
308 template <
typename T>
323 template <
typename T>
324 const_iterator<T>
end()
const
332 template <
typename T,
bool has_nulls>
334 thrust::transform_iterator<detail::pair_accessor<T, has_nulls>,
count_it>;
352 template <
typename T,
bool has_nulls>
367 template <
typename T,
bool has_nulls>
502 template <
typename T =
void>
503 __host__ __device__ T*
head() const noexcept
505 return const_cast<T*
>(detail::column_device_view_base::head<T>());
519 template <
typename T>
520 __host__ __device__ T*
data() const noexcept
522 return const_cast<T*
>(detail::column_device_view_base::data<T>());
533 template <
typename T>
534 __device__ T&
element(size_type element_index) noexcept
536 return data<T>()[element_index];
546 __host__ __device__ bitmask_type*
null_mask() const noexcept
554 using count_it = thrust::counting_iterator<size_type>;
555 template <
typename T>
556 using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>,
count_it>;
565 template <
typename T>
566 std::enable_if_t<is_fixed_width<T>(), iterator<T>>
begin()
578 template <
typename T>
579 std::enable_if_t<is_fixed_width<T>(), iterator<T>>
end()
592 return d_children[child_index];
609 __device__
void set_valid(size_type element_index)
const noexcept
611 return set_bit(
null_mask(), element_index);
627 __device__
void set_null(size_type element_index)
const noexcept
629 return clear_bit(
null_mask(), element_index);
665 size_type _num_children{};
678 mutable_column_device_view(mutable_column_view source);
693 __device__
inline string_view const column_device_view::element<string_view>(
694 size_type element_index)
const noexcept
696 size_type index = element_index + offset();
697 const int32_t* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
698 const char* d_strings = d_children[strings_column_view::chars_column_index].data<
char>();
699 size_type offset = d_offsets[index];
700 return string_view{d_strings + offset, d_offsets[index + 1] - offset};
711 std::enable_if_t<is_index_type<IndexType>() and std::is_unsigned<IndexType>::value>* =
nullptr>
714 return static_cast<size_type
>(input.
element<IndexType>(index));
716 template <
typename IndexType,
718 std::enable_if_t<not(is_index_type<IndexType>() and
719 std::is_unsigned<IndexType>::value)>* =
nullptr>
720 __device__ size_type
operator()(Args&&... args)
722 release_assert(
false and
"dictionary indices must be an unsigned integral type");
752 __device__
inline dictionary32 const column_device_view::element<dictionary32>(
753 size_type element_index)
const noexcept
755 size_type index = element_index + offset();
756 auto const indices = d_children[0];
771 size_type element_index)
const noexcept
774 auto const scale = scale_type{_type.scale()};
789 size_type element_index)
const noexcept
792 auto const scale = scale_type{_type.scale()};
805 size_type destination_word_index,
806 size_type source_begin_bit,
807 size_type source_end_bit)
809 size_type source_word_index = destination_word_index +
word_index(source_begin_bit);
810 bitmask_type curr_word = source[source_word_index];
811 bitmask_type next_word = 0;
814 destination_word_index * detail::size_in_bits<bitmask_type>())) {
815 next_word = source[source_word_index + 1];
817 return __funnelshift_r(curr_word, next_word, source_begin_bit);
834 template <
typename T>
844 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
847 __device__ T operator()(cudf::size_type i)
const {
return col.
element<T>(i); }
868 template <
typename T,
bool has_nulls = false>
878 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
879 if (has_nulls) {
CUDF_EXPECTS(_col.nullable(),
"Unexpected non-nullable column."); }
883 thrust::pair<T, bool> operator()(cudf::size_type i)
const
889 template <
typename T>
899 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
902 __device__ T& operator()(cudf::size_type i) {
return col.
element<T>(i); }
930 template <
typename ColumnDeviceView,
typename ColumnViewIterator>
932 ColumnViewIterator child_end,
936 ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
938 if (num_children > 0) {
941 auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
942 auto d_column = d_children;
945 auto h_end =
reinterpret_cast<int8_t*
>(h_column + num_children);
946 auto d_end =
reinterpret_cast<int8_t*
>(d_column + num_children);
947 std::for_each(child_begin, child_end, [&](
auto const& col) {
949 new (h_column) ColumnDeviceView(col, h_end, d_end);
952 auto col_child_data_size = ColumnDeviceView::extent(col) -
sizeof(ColumnDeviceView);
953 h_end += col_child_data_size;
954 d_end += col_child_data_size;
void const * _data
Pointer to device memory containing elements.
mutable_value_accessor(mutable_column_device_view &_col)
constructor
constexpr decltype(auto) CUDA_HOST_DEVICE_CALLABLE type_dispatcher(cudf::data_type dtype, Functor f, Ts &&... args)
Invokes an operator() template with the type instantiation based on the specified cudf::data_type's i...
Utilities for bit and bitmask operations.
Class definition for fixed point data type.
__host__ __device__ T * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
A container of nullable device data as a column of elements.
__device__ column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Class definition for cudf::strings_column_view.
A type for representing a number with a fixed amount of precision.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
void destroy()
Destroy the mutable_column_device_view object.
__host__ __device__ bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Type declarations for libcudf.
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
A non-owning, mutable view of device data as a column of elements, some of which may be null as indic...
__device__ T & element(size_type element_index) noexcept
Returns reference to element at the specified index.
__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.
__device__ bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the the specified bitmask word from the null_mask().
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
__device__ T const element(size_type element_index) const noexcept
Returns reference to element at the specified index.
__device__ void set_valid(size_type element_index) const noexcept
Updates the null mask to indicate that the specified element is valid.
__host__ __device__ bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
__device__ bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
pair_accessor(column_device_view const &_col)
constructor
A strongly typed wrapper for indices in a DICTIONARY type column.
ColumnDeviceView * child_columns_to_device_array(ColumnViewIterator child_begin, ColumnViewIterator child_end, void *h_ptr, void *d_ptr)
Helper function for use by column_device_view and mutable_column_device_view constructors to build de...
__host__ __device__ T const * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
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 obje...
__device__ bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
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.
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
__host__ __device__ T * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Helper struct for constructing fixed_point when value is already shifted.
__device__ bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
pair accessor of column with/without null bitmask A unary functor returns pair with scalar value at i...
column_device_view const col
column view of column in device
Indicator for the logical data type of an element in a column.
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...
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
__device__ bitmask_type get_mask_offset_word(bitmask_type const *__restrict__ source, size_type destination_word_index, size_type source_begin_bit, size_type source_end_bit)
Convenience function to get offset word from a bitmask.
column_device_view const col
column view of column in device
__host__ __device__ size_type size() const noexcept
Returns the number of elements in the column.
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
__device__ bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
const_iterator< T > begin() const
Return an iterator to the first element of the column.
data_type _type
Element type.
__host__ __device__ T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
value accessor of column without null bitmask A unary functor returns scalar value at id....
bitmask_type const * _null_mask
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
__device__ void set_null(size_type element_index) const noexcept
Updates the null mask to indicate that the specified element is null.
Class definition for cudf::list_view.
__host__ __device__ bitmask_type * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
CUDA_HOST_DEVICE_CALLABLE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
column_device_view * d_children
__host__ __device__ size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
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 obje...
#define CUDF_EXPECTS(cond, reason)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
void destroy()
Destroy the column_device_view object.
__device__ mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
__host__ __device__ data_type type() const noexcept
Returns the element type.
value_accessor(column_device_view const &_col)
constructor
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.
cudf::size_type _size
Number of elements.
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
constexpr CUDA_HOST_DEVICE_CALLABLE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
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.
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...
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
column view class definitons
Class definition for cudf::struct_view.
size_type _num_children
The number of child columns.
mutable_column_device_view col
mutable column view of column in device
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.
CUDA_HOST_DEVICE_CALLABLE type_id id() const noexcept
Returns the type identifier.
fixed_point and supporting types
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.
Dispatch functor for resolving the index value for a dictionary element.