Go to the documentation of this file.
19 #include <cudf/detail/utilities/alignment.hpp>
22 #include <cudf/strings/string_view.cuh>
27 #include <cudf/utilities/span.hpp>
31 #include <rmm/cuda_stream_view.hpp>
33 #include <thrust/iterator/counting_iterator.h>
34 #include <thrust/iterator/transform_iterator.h>
35 #include <thrust/optional.h>
36 #include <thrust/pair.h>
55 struct YES : std::bool_constant<true> {
57 struct NO : std::bool_constant<false> {
70 constexpr
operator bool() const noexcept {
return value; }
112 template <
typename T = void,
113 CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
114 [[nodiscard]] CUDF_HOST_DEVICE T
const*
head() const noexcept
116 return static_cast<T const*
>(
_data);
134 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
135 [[nodiscard]] CUDF_HOST_DEVICE T
const*
data()
const noexcept
281 template <
typename C,
typename T,
typename =
void>
282 struct has_element_accessor_impl : std::false_type {
285 template <
typename C,
typename T>
286 struct has_element_accessor_impl<
289 void_t<decltype(std::declval<C>().template element<T>(std::declval<size_type>()))>>
294 template <
typename T>
295 struct value_accessor;
296 template <
typename T,
typename Nullate>
297 struct optional_accessor;
298 template <
typename T,
bool has_nulls>
299 struct pair_accessor;
300 template <
typename T,
bool has_nulls>
301 struct pair_rep_accessor;
302 template <
typename T>
303 struct mutable_value_accessor;
375 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
378 return data<T>()[element_index];
392 template <
typename T, CUDF_ENABLE_IF(std::is_same_v<T,
string_view>)>
396 const auto* d_offsets =
d_children[strings_column_view::offsets_column_index].
data<int32_t>();
397 const char* d_strings =
d_children[strings_column_view::chars_column_index].
data<
char>();
408 struct index_element_fn {
409 template <
typename IndexType,
410 CUDF_ENABLE_IF(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>)>
416 template <
typename IndexType,
418 CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>))>
419 __device__
size_type operator()(Args&&... args)
421 CUDF_UNREACHABLE(
"dictionary indices must be an unsigned integral type");
450 template <
typename T, CUDF_ENABLE_IF(std::is_same_v<T, dictionary32>)>
468 template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
472 using rep =
typename T::rep;
481 template <
typename T>
484 return has_element_accessor_impl<column_device_view, T>::value;
490 using count_it = thrust::counting_iterator<size_type>;
491 template <
typename T>
492 using const_iterator = thrust::transform_iterator<detail::value_accessor<T>,
count_it>;
506 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
507 [[nodiscard]] const_iterator<T>
begin()
const
524 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
525 [[nodiscard]] const_iterator<T>
end()
const
533 template <
typename T,
typename Nullate>
535 thrust::transform_iterator<detail::optional_accessor<T, Nullate>,
count_it>;
540 template <
typename T,
bool has_nulls>
542 thrust::transform_iterator<detail::pair_accessor<T, has_nulls>,
count_it>;
549 template <
typename T,
bool has_nulls>
551 thrust::transform_iterator<detail::pair_rep_accessor<T, has_nulls>,
count_it>;
602 template <
typename T,
630 template <
typename T,
660 template <
typename T,
681 template <
typename T,
701 template <
typename T,
721 template <
typename T,
905 template <
typename T = void,
906 CUDF_ENABLE_IF(std::is_same_v<T, void> or is_rep_layout_compatible<T>())>
907 CUDF_HOST_DEVICE T*
head() const noexcept
909 return const_cast<T*
>(detail::column_device_view_base::head<T>());
924 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
925 CUDF_HOST_DEVICE T*
data()
const noexcept
927 return const_cast<T*
>(detail::column_device_view_base::data<T>());
943 template <typename T, CUDF_ENABLE_IF(is_rep_layout_compatible<T>())>
946 return data<T>()[element_index];
954 template <
typename T>
957 return has_element_accessor_impl<mutable_column_device_view, T>::value;
975 using count_it = thrust::counting_iterator<size_type>;
976 template <
typename T>
977 using iterator = thrust::transform_iterator<detail::mutable_value_accessor<T>,
count_it>;
989 template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1005 template <typename T, CUDF_ENABLE_IF(mutable_column_device_view::has_element_accessor<T>())>
1022 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1037 __device__
void set_valid(
size_type element_index)
const noexcept
1039 return set_bit(
null_mask(), element_index);
1055 __device__
void set_null(
size_type element_index)
const noexcept
1057 return clear_bit(
null_mask(), element_index);
1108 mutable_column_device_view(mutable_column_view source);
1113 #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__
1131 destination_word_index * detail::size_in_bits<bitmask_type>())) {
1132 next_word = source[source_word_index + 1];
1134 return __funnelshift_r(curr_word, next_word, source_begin_bit);
1152 template <
typename T>
1162 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1194 template <
typename T,
typename Nullate>
1205 : col{_col}, has_nulls{with_nulls}
1207 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1208 if (with_nulls) {
CUDF_EXPECTS(_col.nullable(),
"Unexpected non-nullable column."); }
1211 __device__
inline thrust::optional<T> operator()(
cudf::size_type i)
const
1215 : thrust::optional<T>{thrust::nullopt};
1217 return thrust::optional<T>{col.
element<T>(i)};
1220 Nullate has_nulls{};
1241 template <
typename T,
bool has_nulls = false>
1251 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1252 if (has_nulls) {
CUDF_EXPECTS(_col.nullable(),
"Unexpected non-nullable column."); }
1255 __device__
inline thrust::pair<T, bool> operator()(
cudf::size_type i)
const
1279 template <
typename T,
bool has_nulls = false>
1291 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1292 if (has_nulls) {
CUDF_EXPECTS(_col.nullable(),
"Unexpected non-nullable column."); }
1295 __device__
inline thrust::pair<rep_type, bool> operator()(
cudf::size_type i)
const
1301 template <
typename R, std::enable_if_t<std::is_same_v<R, rep_type>,
void>* =
nullptr>
1307 template <
typename R, std::enable_if_t<not std::is_same_v<R, rep_type>,
void>* =
nullptr>
1310 return col.
element<R>(i).value();
1314 template <
typename T>
1324 CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.
type().
id()),
"the data type mismatch");
1355 template <
typename ColumnDeviceView,
typename ColumnViewIterator>
1357 ColumnViewIterator child_end,
1361 ColumnDeviceView*
d_children = detail::align_ptr_for_type<ColumnDeviceView>(d_ptr);
1363 if (num_children > 0) {
1366 auto h_column = detail::align_ptr_for_type<ColumnDeviceView>(h_ptr);
1370 auto h_end =
reinterpret_cast<int8_t*
>(h_column + num_children);
1371 auto d_end =
reinterpret_cast<int8_t*
>(d_column + num_children);
1372 std::for_each(child_begin, child_end, [&](
auto const& col) {
1374 new (h_column) ColumnDeviceView(col, h_end, d_end);
1377 auto col_child_data_size = ColumnDeviceView::extent(col) -
sizeof(ColumnDeviceView);
1378 h_end += col_child_data_size;
1379 d_end += col_child_data_size;
CUDF_HOST_DEVICE bool bit_is_set(bitmask_type const *bitmask, size_type bit_index)
Indicates whether the specified bit is set to 1
void const * _data
Pointer to device memory containing elements.
auto optional_end(Nullate has_nulls) const
Return an optional iterator to the element following the last element of the column.
mutable_value_accessor(mutable_column_device_view &_col)
constructor
Utilities for bit and bitmask operations.
Class definition for fixed point data type.
CUDF_HOST_DEVICE T const * 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.
constexpr type_id id() const noexcept
Returns the type identifier.
pair_rep_accessor(column_device_view const &_col)
constructor
Class definition for cudf::strings_column_view.
int32_t size_type
Row index type for columns and tables.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
CUDF_HOST_DEVICE size_type num_child_columns() const noexcept
Returns the number of child columns.
Indicates the presence of nulls at compile-time or runtime.
void destroy()
Destroy the mutable_column_device_view object.
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Type declarations for libcudf.
thrust::transform_iterator< detail::optional_accessor< T, Nullate >, count_it > const_optional_iterator
Optional iterator for navigating this column.
bool is_valid(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null).
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...
const_pair_iterator< T, has_nulls > pair_end() const
Return a pair iterator to the element following the last element of the column.
constexpr int32_t scale() const noexcept
Returns the scale (for fixed_point types)
CUDF_HOST_DEVICE bitmask_type const * null_mask() const noexcept
Returns raw pointer to the underlying bitmask allocation.
bool is_null_nocheck(size_type element_index) const noexcept
Returns whether the specified element is null.
auto optional_begin(Nullate has_nulls) const
Return an optional iterator to the first element of the column.
column_device_view(column_view source)
Construct's a column_device_view from a column_view populating all but the children.
uint32_t bitmask_type
Bitmask type stored as 32-bit unsigned integer.
thrust::counting_iterator< size_type > count_it
Iterator for navigating this column.
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
column_device_view const col
column view of column in device
#define CUDF_ENABLE_IF(...)
Convenience macro for SFINAE as an unnamed template parameter.
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.
pair_accessor(column_device_view const &_col)
constructor
CUDF_HOST_DEVICE size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
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...
thrust::transform_iterator< detail::pair_rep_accessor< T, has_nulls >, count_it > const_pair_rep_iterator
Pair rep iterator for navigating this 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...
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...
iterator< T > end()
Return one past the last element after underlying data is casted to the specified type.
T & element(size_type element_index) const noexcept
Returns reference to element at the specified index.
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.
Helper struct for constructing fixed_point when value is already shifted.
bitmask_type get_mask_word(size_type word_index) const noexcept
Returns the the specified bitmask word from the null_mask().
pair accessor of column with/without null bitmask A unary functor returns pair with scalar value at i...
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.
column_device_view const col
column view of column in device
bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
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...
A mutable, non-owning view of device data as a column of elements that is trivially copyable and usab...
column_device_view const col
column view of column in device
Class definition for cudf::list_view.
const_pair_rep_iterator< T, has_nulls > pair_rep_begin() const
Return a pair iterator to the first element of the column.
CUDF_HOST_DEVICE bool nullable() const noexcept
Indicates whether the column can contain null elements, i.e., if it has an allocated bitmask.
const_iterator< T > end() const
Returns an iterator to the element following the last element of the column.
data_type _type
Element type.
column_device_view child(size_type child_index) const noexcept
Returns the specified child.
device_span< column_device_view const > children() const noexcept
Returns a span containing the children of this column.
value accessor of column without null bitmask A unary functor returns scalar value at id....
const_iterator< T > begin() const
Return an iterator to the first element of the column.
column_device_view const col
column view of column in device
bitmask_type const * _null_mask
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
mutable_column_device_view child(size_type child_index) const noexcept
Returns the specified child.
CUDF_HOST_DEVICE data_type type() const noexcept
Returns the element type.
pair accessor of column with/without null bitmask A unary functor returns pair with representative sc...
column_device_view * d_children
static constexpr bool has_element_accessor()
For a given T, indicates if mutable_column_device_view::element<T>() has a valid overload.
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.
const_pair_iterator< T, has_nulls > pair_begin() const
Return a pair iterator to the first element of the column.
void destroy()
Destroy the column_device_view object.
value_accessor(column_device_view const &_col)
constructor
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.
static constexpr bool has_element_accessor()
For a given T, indicates if column_device_view::element<T>() has a valid overload.
cudf::size_type _size
Number of elements.
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
iterator< T > begin()
Return first element (accounting for offset) after underlying data is casted to the specified type.
size_type distance(T f, T l)
Similar to std::distance but returns cudf::size_type and performs static_cast
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.
optional accessor of a column
optional_accessor(column_device_view const &_col, Nullate with_nulls)
Constructor.
column view class definitions
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
constexpr DYNAMIC(bool b) noexcept
Create a runtime nullate object.
CUDF_HOST_DEVICE T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
bool is_valid_nocheck(size_type element_index) const noexcept
Returns whether the specified element holds a valid value (i.e., not null)
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
fixed_point and supporting types
constexpr CUDF_HOST_DEVICE size_type word_index(size_type bit_index)
Returns the index of the word containing the specified bit.
bool value
True if nulls are expected.
thrust::transform_iterator< detail::pair_accessor< T, has_nulls >, count_it > const_pair_iterator
Pair iterator for navigating this column.