19 #include <cudf/detail/offsets_iterator.cuh>
20 #include <cudf/detail/utilities/alignment.hpp>
23 #include <cudf/strings/string_view.cuh>
35 #include <cuda/std/optional>
36 #include <cuda/std/type_traits>
37 #include <thrust/iterator/counting_iterator.h>
38 #include <thrust/iterator/transform_iterator.h>
39 #include <thrust/pair.h>
42 #include <type_traits>
49 namespace CUDF_EXPORT
cudf {
61 struct YES : cuda::std::bool_constant<true> {};
62 struct NO : cuda::std::bool_constant<false> {};
78 constexpr
explicit DYNAMIC(
bool b) noexcept : value{b} {}
137 template <
typename T = void,
138 CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
141 return static_cast<T const*
>(_data);
159 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
162 return head<T>() + _offset;
228 return not
nullable() or is_valid_nocheck(element_index);
245 return bit_is_set(_null_mask, offset() + element_index);
279 return not is_valid_nocheck(element_index);
319 : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset}
323 template <
typename C,
typename T,
typename =
void>
324 struct has_element_accessor_impl : cuda::std::false_type {};
326 template <
typename C,
typename T>
327 struct has_element_accessor_impl<
330 void_t<decltype(cuda::std::declval<C>().template element<T>(cuda::std::declval<size_type>()))>>
331 : cuda::std::true_type {};
335 template <
typename T>
336 struct value_accessor;
337 template <
typename T,
typename Nullate>
338 struct optional_accessor;
339 template <
typename T,
bool has_nulls>
340 struct pair_accessor;
341 template <
typename T,
bool has_nulls>
342 struct pair_rep_accessor;
343 template <
typename T>
344 struct mutable_value_accessor;
407 this->offset() + offset,
409 this->num_child_columns()};
429 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
432 return data<T>()[element_index];
446 template <
typename T, CUDF_ENABLE_IF(std::is_same_v<T,
string_view>)>
449 size_type index = element_index + offset();
450 char const* d_strings =
static_cast<char const*
>(_data);
451 auto const offsets = d_children[strings_column_view::offsets_column_index];
452 auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
453 auto const offset = itr[index];
463 struct index_element_fn {
464 template <
typename IndexType,
465 CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_signed_v<IndexType>)>
471 template <
typename IndexType,
473 CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_signed_v<IndexType>))>
474 __device__
size_type operator()(Args&&... args)
476 CUDF_UNREACHABLE(
"dictionary indices must be a signed integral type");
505 template <
typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
508 size_type index = element_index + offset();
509 auto const indices = d_children[0];
523 template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
527 using rep =
typename T::rep;
538 template <
typename T>
541 return has_element_accessor_impl<column_device_view, T>::value;
545 using count_it = thrust::counting_iterator<size_type>;
549 template <
typename T>
567 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
587 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
596 template <
typename T,
typename Nullate>
598 thrust::transform_iterator<detail::optional_accessor<T, Nullate>,
count_it>;
603 template <
typename T,
bool has_nulls>
605 thrust::transform_iterator<detail::pair_accessor<T, has_nulls>,
count_it>;
612 template <
typename T,
bool has_nulls>
614 thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>,
count_it>;
670 template <
typename T,
700 template <
typename T,
732 template <
typename T,
757 template <
typename T,
777 template <
typename T,
798 template <
typename T,
853 return d_children[child_index];
863 return {d_children,
static_cast<std::size_t
>(_num_children)};
873 return _num_children;
896 : column_device_view_base(type, size, data, null_mask, offset),
897 d_children(children),
898 _num_children(num_children)
997 template <
typename T = void,
998 CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
1001 return const_cast<T*
>(detail::column_device_view_base::head<T>());
1016 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1019 return const_cast<T*
>(detail::column_device_view_base::data<T>());
1036 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
1039 return data<T>()[element_index];
1048 template <
typename T>
1051 return has_element_accessor_impl<mutable_column_device_view, T>::value;
1064 return const_cast<bitmask_type*
>(detail::column_device_view_base::null_mask());
1072 template <
typename T>
1085 template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1101 template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1115 return d_children[child_index];
1133 __device__ void set_valid(size_type element_index) const noexcept
1135 return set_bit(null_mask(), element_index);
1151 __device__
void set_null(
size_type element_index)
const noexcept
1153 return clear_bit(null_mask(), element_index);
1206 mutable_column_device_view(mutable_column_view source);
1229 destination_word_index * detail::size_in_bits<bitmask_type>())) {
1230 next_word = source[source_word_index + 1];
1232 return __funnelshift_r(curr_word, next_word, source_begin_bit);
1251 template <
typename T>
1262 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1299 template <
typename T,
typename Nullate>
1312 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1327 : cuda::std::optional<T>{cuda::std::nullopt};
1329 return cuda::std::optional<T>{col.
element<T>(i)};
1354 template <
typename T,
bool has_nulls = false>
1365 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1400 template <
typename T,
bool has_nulls = false>
1413 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1429 template <
typename R, std::enable_if_t<std::is_same_v<R, rep_type>,
void>* =
nullptr>
1435 template <
typename R, std::enable_if_t<not std::is_same_v<R, rep_type>,
void>* =
nullptr>
1438 return col.
element<R>(i).value();
1453 template <
typename T>
1464 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1501 template <
typename ColumnDeviceView,
typename ColumnViewIterator>
1503 ColumnViewIterator child_end,
1507 ColumnDeviceView* d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1509 if (num_children > 0) {
1512 auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1513 auto d_column = d_children;
1516 auto h_end =
reinterpret_cast<int8_t*
>(h_column + num_children);
1517 auto d_end =
reinterpret_cast<int8_t*
>(d_column + num_children);
1518 std::for_each(child_begin, child_end, [&](
auto const& col) {
1520 new (h_column) ColumnDeviceView(col, h_end, d_end);
1523 auto col_child_data_size = ColumnDeviceView::extent(col) -
sizeof(ColumnDeviceView);
1524 h_end += col_child_data_size;
1525 d_end += col_child_data_size;
Utilities for bit and bitmask operations.
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
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.
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
column_device_view child(size_type child_index) const noexcept
Returns the specified child.
void destroy()
Destroy the column_device_view object.
column_device_view & operator=(column_device_view &&)=default
Move assignment operator.
thrust::counting_iterator< size_type > count_it
Counting iterator.
static constexpr CUDF_HOST_DEVICE bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
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.
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
column_device_view & operator=(column_device_view const &)=default
Copy assignment operator.
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.
thrust::transform_iterator< detail::value_accessor< T >, count_it > const_iterator
Iterator for navigating this column.
auto optional_end(Nullate has_nulls) const
Return an optional iterator to the element following the last element of the column.
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.
const_pair_rep_iterator< T, has_nulls > pair_rep_begin() const
Return a pair iterator to the first element of the column.
thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > const_optional_iterator
Optional iterator for navigating this column.
CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
Returns the number of child columns.
auto optional_begin(Nullate has_nulls) const
Return an optional iterator to the first element of the column.
column_device_view(column_device_view const &)=default
Copy constructor.
column_device_view(column_device_view &&)=default
Move constructor.
const_iterator< T > begin() const
Return an iterator to the first element of the column.
thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > const_pair_rep_iterator
Pair rep iterator for navigating this column.
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...
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...
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
device_span< column_device_view const > children() const noexcept
Returns a span containing the children of this column.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
A container of nullable device data as a column of elements.
Indicator for the logical data type of an element in a column.
constexpr CUDF_HOST_DEVICE type_id id() const noexcept
Returns the type identifier.
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
CUDF_HOST_DEVICE T const * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
CUDF_HOST_DEVICE data_type type() const noexcept
Returns the element type.
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.
column_device_view_base & operator=(column_device_view_base &&)=default
Move assignment operator.
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the specified bitmask word from the null_mask().
column_device_view_base(column_device_view_base &&)=default
Move constructor.
column_device_view_base(column_device_view_base const &)=default
Copy constructor.
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
CUDF_HOST_DEVICE T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
column_device_view_base & operator=(column_device_view_base const &)=default
Copy assignment operator.
bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
CUDF_HOST_DEVICE size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
CUDF_HOST_DEVICE bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
static constexpr CUDF_HOST_DEVICE bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
void destroy()
Destroy the mutable_column_device_view object.
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.
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...
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...
thrust::counting_iterator< size_type > count_it
Counting iterator.
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
CUDF_HOST_DEVICE bitmask_type * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
CUDF_HOST_DEVICE T * head() const noexcept
Returns pointer to the base device memory allocation casted to the specified type.
CUDF_HOST_DEVICE T * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
mutable_column_device_view(mutable_column_device_view const &)=default
Copy constructor.
T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
iterator< T > end()
Return one past the last element after underlying data is casted to the specified type.
mutable_column_device_view & operator=(mutable_column_device_view &&)=default
Move assignment operator.
mutable_column_device_view(mutable_column_device_view &&)=default
Move constructor.
mutable_column_device_view & operator=(mutable_column_device_view const &)=default
Copy assignment operator.
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=cudf::get_default_stream())
Factory to construct a column view that is usable in device memory.
thrust::transform_iterator< detail::mutable_value_accessor< T >, count_it > iterator
Iterator for navigating this column.
iterator< T > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
A non-owning, mutable view of device data as a column of elements, some of which may be null as indic...
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
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...
column view class definitions
Class definition for fixed point data type.
rmm::cuda_stream_view const get_default_stream()
Get the current default stream.
scale_type
The scale type for fixed_point.
constexpr CUDF_HOST_DEVICE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
CUDF_HOST_DEVICE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
CUDF_HOST_DEVICE constexpr decltype(auto) __forceinline__ 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...
std::conditional_t< std::is_same_v< numeric::decimal32, T >, int32_t, std::conditional_t< std::is_same_v< numeric::decimal64, T >, int64_t, std::conditional_t< std::is_same_v< numeric::decimal128, T >, __int128_t, T > >> device_storage_type_t
"Returns" the corresponding type that is stored on the device when using cudf::column
#define CUDF_EXPECTS(...)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
int32_t size_type
Row index type for columns and tables.
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
void void_t
Utility metafunction that maps a sequence of any types to the type void.
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
Class definition for cudf::list_view.
bool nullable(table_view const &view)
Returns True if any of the columns in the table is nullable. (not entire hierarchy)
bool has_nulls(table_view const &view)
Returns True if the table has nulls in any of its columns.
fixed_point and supporting types
Class definition for cudf::strings_column_view.
Class definition for cudf::struct_view.
Mutable value accessor of column without null bitmask.
T & operator()(cudf::size_type i)
Accessor.
mutable_value_accessor(mutable_column_device_view &_col)
Constructor.
mutable_column_device_view col
mutable column view of column in device
optional accessor of a column
cuda::std::optional< T > operator()(cudf::size_type i) const
Returns a cuda::std::optional of column[i].
column_device_view const col
column view of column in device
optional_accessor(column_device_view const &_col, Nullate with_nulls)
Constructor.
pair accessor of column with/without null bitmask
column_device_view const col
column view of column in device
thrust::pair< T, bool > operator()(cudf::size_type i) const
Pair accessor.
pair_accessor(column_device_view const &_col)
constructor
pair accessor of column with/without null bitmask
thrust::pair< rep_type, bool > operator()(cudf::size_type i) const
Pair accessor.
column_device_view const col
column view of column in device
device_storage_type_t< T > rep_type
representation type
pair_rep_accessor(column_device_view const &_col)
constructor
value accessor of column without null bitmask
column_device_view const col
column view of column in device
value_accessor(column_device_view const &_col)
constructor
T operator()(cudf::size_type i) const
Returns the value of element at index i
Device version of C++20 std::span with reduced feature set.
A strongly typed wrapper for indices in a DICTIONARY type column.
nullate::DYNAMIC defers the determination of nullability to run time rather than compile time....
bool value
True if nulls are expected.
constexpr DYNAMIC(bool b) noexcept
Create a runtime nullate object.
Indicates the presence of nulls at compile-time or runtime.
Helper struct for constructing fixed_point when value is already shifted.
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
Type declarations for libcudf.
#define CUDF_HOST_DEVICE
Indicates that the function or method is usable on host and device.