list_device_view.cuh
1 /*
2  * Copyright (c) 2020-2024, NVIDIA CORPORATION.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 #pragma once
17 
18 #include <cudf/detail/iterator.cuh>
19 #include <cudf/lists/lists_column_device_view.cuh>
20 #include <cudf/types.hpp>
22 
23 #include <cuda_runtime.h>
24 #include <thrust/iterator/counting_iterator.h>
25 #include <thrust/iterator/transform_iterator.h>
26 #include <thrust/pair.h>
27 
28 namespace CUDF_EXPORT cudf {
29 
36 
37  public:
38  list_device_view() = default;
39 
46  __device__ inline list_device_view(lists_column_device_view const& lists_column,
47  size_type const& row_index)
48  : lists_column(lists_column), _row_index(row_index)
49  {
50  column_device_view const& offsets = lists_column.offsets();
51  cudf_assert(row_index >= 0 && row_index < lists_column.size() && row_index < offsets.size() &&
52  "row_index out of bounds");
53 
54  begin_offset = offsets.element<size_type>(row_index + lists_column.offset());
55  cudf_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() &&
56  "begin_offset out of bounds.");
57  _size = offsets.element<size_type>(row_index + 1 + lists_column.offset()) - begin_offset;
58  }
59 
60  ~list_device_view() = default;
61 
86  [[nodiscard]] __device__ inline size_type element_offset(size_type idx) const
87  {
88  cudf_assert(idx >= 0 && idx < size() && "idx out of bounds");
89  return begin_offset + idx;
90  }
91 
99  template <typename T>
100  __device__ inline T element(size_type idx) const
101  {
102  return lists_column.child().element<T>(element_offset(idx));
103  }
104 
111  [[nodiscard]] __device__ inline bool is_null(size_type idx) const
112  {
113  cudf_assert(idx >= 0 && idx < size() && "Index out of bounds.");
114  auto element_offset = begin_offset + idx;
115  return lists_column.child().is_null(element_offset);
116  }
117 
123  [[nodiscard]] __device__ inline bool is_null() const { return lists_column.is_null(_row_index); }
124 
130  [[nodiscard]] __device__ inline size_type size() const { return _size; }
131 
137  [[nodiscard]] __device__ inline size_type row_index() const { return _row_index; }
138 
144  [[nodiscard]] __device__ inline lists_column_device_view const& get_column() const
145  {
146  return lists_column;
147  }
148 
149  template <typename T>
150  struct pair_accessor;
151 
152  template <typename T>
153  struct pair_rep_accessor;
154 
156  template <typename T>
158  thrust::transform_iterator<pair_accessor<T>, thrust::counting_iterator<cudf::size_type>>;
159 
161  template <typename T>
163  thrust::transform_iterator<pair_rep_accessor<T>, thrust::counting_iterator<cudf::size_type>>;
164 
181  template <typename T>
182  [[nodiscard]] __device__ inline const_pair_iterator<T> pair_begin() const
183  {
184  return const_pair_iterator<T>{thrust::counting_iterator<size_type>(0), pair_accessor<T>{*this}};
185  }
186 
194  template <typename T>
195  [[nodiscard]] __device__ inline const_pair_iterator<T> pair_end() const
196  {
197  return const_pair_iterator<T>{thrust::counting_iterator<size_type>(size()),
198  pair_accessor<T>{*this}};
199  }
200 
219  template <typename T>
220  [[nodiscard]] __device__ inline const_pair_rep_iterator<T> pair_rep_begin() const
221  {
222  return const_pair_rep_iterator<T>{thrust::counting_iterator<size_type>(0),
223  pair_rep_accessor<T>{*this}};
224  }
225 
233  template <typename T>
234  [[nodiscard]] __device__ inline const_pair_rep_iterator<T> pair_rep_end() const
235  {
236  return const_pair_rep_iterator<T>{thrust::counting_iterator<size_type>(size()),
237  pair_rep_accessor<T>{*this}};
238  }
239 
240  private:
241  lists_column_device_view const& lists_column;
242  size_type _row_index{}; // Row index in the Lists column vector.
243  size_type _size{}; // Number of elements in *this* list row.
244 
245  size_type begin_offset; // Offset in list_column_device_view where this list begins.
246 
256  template <typename T>
257  struct pair_accessor {
259 
265  explicit CUDF_HOST_DEVICE inline pair_accessor(list_device_view const& _list) : list{_list} {}
266 
273  __device__ inline thrust::pair<T, bool> operator()(cudf::size_type i) const
274  {
275  return {list.element<T>(i), !list.is_null(i)};
276  }
277  };
278 
291  template <typename T>
294 
296 
302  explicit CUDF_HOST_DEVICE inline pair_rep_accessor(list_device_view const& _list) : list{_list}
303  {
304  }
305 
312  __device__ inline thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
313  {
314  return {get_rep<T>(i), !list.is_null(i)};
315  }
316 
317  private:
318  template <typename R, std::enable_if_t<std::is_same_v<R, rep_type>, void>* = nullptr>
319  __device__ inline rep_type get_rep(cudf::size_type i) const
320  {
321  return list.element<R>(i);
322  }
323 
324  template <typename R, std::enable_if_t<not std::is_same_v<R, rep_type>, void>* = nullptr>
325  __device__ inline rep_type get_rep(cudf::size_type i) const
326  {
327  return list.element<R>(i).value();
328  }
329  };
330 };
331 
344  : d_column(d_col)
345  {
346  }
353  __device__ inline size_type operator()(size_type idx)
354  {
355  if (d_column.is_null(idx)) return size_type{0};
356  return d_column.offset_at(idx + 1) - d_column.offset_at(idx);
357  }
358 };
359 
376 {
377  return detail::make_counting_transform_iterator(0, list_size_functor{c});
378 }
379 
380 } // namespace CUDF_EXPORT cudf
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
CUDF_HOST_DEVICE size_type size() const noexcept
Returns the number of elements in the column.
Given a column_device_view, an instance of this class provides a wrapper on this compound column for ...
column_device_view offsets() const
Fetches the offsets column of the underlying list column.
size_type offset_at(size_type idx) const
Fetches the list offset value at a given row index while taking column offset into account.
column_device_view child() const
Fetches the child column of the underlying list column.
A non-owning, immutable view of device data that represents a list of elements of arbitrary type (inc...
const_pair_rep_iterator< T > pair_rep_end() const
Fetcher for a pair iterator to one position past the last element in the list_device_view.
T element(size_type idx) const
Fetches the element at the specified index within the list row.
size_type element_offset(size_type idx) const
Fetches the offset in the list column's child that corresponds to the element at the specified list i...
bool is_null() const
Checks whether this list row is null.
const_pair_iterator< T > pair_end() const
Fetcher for a pair iterator to one position past the last element in the list_device_view.
size_type row_index() const
Returns the row index of this list in the original lists column.
bool is_null(size_type idx) const
Checks whether the element is null at the specified index in the list.
const_pair_rep_iterator< T > pair_rep_begin() const
Fetcher for a pair iterator to the first element in the list_device_view.
thrust::transform_iterator< pair_accessor< T >, thrust::counting_iterator< cudf::size_type > > const_pair_iterator
const pair iterator for the list
thrust::transform_iterator< pair_rep_accessor< T >, thrust::counting_iterator< cudf::size_type > > const_pair_rep_iterator
const pair iterator type for the list
list_device_view(lists_column_device_view const &lists_column, size_type const &row_index)
Constructs a list_device_view from a list column and index.
const_pair_iterator< T > pair_begin() const
Fetcher for a pair iterator to the first element in the list_device_view.
lists_column_device_view const & get_column() const
Fetches the lists_column_device_view that contains this list.
size_type size() const
Fetches the number of elements in this list row.
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
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:95
cuDF interfaces
Definition: host_udf.hpp:39
CUDF_HOST_DEVICE auto make_list_size_iterator(detail::lists_column_device_view const &c)
Makes an iterator that returns size of the list by row index.
pair accessor for elements in a list_device_view
CUDF_HOST_DEVICE pair_accessor(list_device_view const &_list)
constructor
list_device_view const & list
The list_device_view to access.
thrust::pair< T, bool > operator()(cudf::size_type i) const
Accessor for the {data, validity} pair at the specified index.
pair rep accessor for elements in a list_device_view
list_device_view const & list
The list_device_view whose rows are being accessed.
device_storage_type_t< T > rep_type
The type used to store the value on the device.
CUDF_HOST_DEVICE pair_rep_accessor(list_device_view const &_list)
constructor
thrust::pair< rep_type, bool > operator()(cudf::size_type i) const
Accessor for the {rep_data, validity} pair at the specified index.
Returns the size of the list by row index.
detail::lists_column_device_view const d_column
The list column to access.
CUDF_HOST_DEVICE list_size_functor(detail::lists_column_device_view const &d_col)
Constructor.
size_type operator()(size_type idx)
Returns size of the list by row index.
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.
Definition: types.hpp:32