list_device_view.cuh
1 /*
2  * Copyright (c) 2020-2021, 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 <cuda_runtime.h>
19 #include <cudf/lists/lists_column_device_view.cuh>
20 #include <cudf/types.hpp>
22 
23 namespace cudf {
24 
31 
32  public:
33  list_device_view() = default;
34 
35  CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column,
36  size_type const& row_index)
37  : lists_column(lists_column), _row_index(row_index)
38  {
39  column_device_view const& offsets = lists_column.offsets();
40  cudf_assert(row_index >= 0 && row_index < lists_column.size() && row_index < offsets.size() &&
41  "row_index out of bounds");
42 
43  begin_offset = offsets.element<size_type>(row_index + lists_column.offset());
44  cudf_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() &&
45  "begin_offset out of bounds.");
46  _size = offsets.element<size_type>(row_index + 1 + lists_column.offset()) - begin_offset;
47  }
48 
49  ~list_device_view() = default;
50 
72  CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const
73  {
74  cudf_assert(idx >= 0 && idx < size() && "idx out of bounds");
75  return begin_offset + idx;
76  }
77 
85  template <typename T>
86  CUDA_DEVICE_CALLABLE T element(size_type idx) const
87  {
88  return lists_column.child().element<T>(element_offset(idx));
89  }
90 
94  CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const
95  {
96  cudf_assert(idx >= 0 && idx < size() && "Index out of bounds.");
97  auto element_offset = begin_offset + idx;
98  return lists_column.child().is_null(element_offset);
99  }
100 
104  CUDA_DEVICE_CALLABLE bool is_null() const { return lists_column.is_null(_row_index); }
105 
109  CUDA_DEVICE_CALLABLE size_type size() const { return _size; }
110 
114  CUDA_DEVICE_CALLABLE lists_column_device_view const& get_column() const { return lists_column; }
115 
116  template <typename T>
117  struct pair_accessor;
118 
119  template <typename T>
120  struct pair_rep_accessor;
121 
122  template <typename T>
123  using const_pair_iterator =
124  thrust::transform_iterator<pair_accessor<T>, thrust::counting_iterator<cudf::size_type>>;
125 
126  template <typename T>
127  using const_pair_rep_iterator =
128  thrust::transform_iterator<pair_rep_accessor<T>, thrust::counting_iterator<cudf::size_type>>;
129 
143  template <typename T>
144  CUDA_DEVICE_CALLABLE const_pair_iterator<T> pair_begin() const
145  {
146  return const_pair_iterator<T>{thrust::counting_iterator<size_type>(0), pair_accessor<T>{*this}};
147  }
148 
153  template <typename T>
154  CUDA_DEVICE_CALLABLE const_pair_iterator<T> pair_end() const
155  {
156  return const_pair_iterator<T>{thrust::counting_iterator<size_type>(size()),
157  pair_accessor<T>{*this}};
158  }
159 
175  template <typename T>
176  CUDA_DEVICE_CALLABLE const_pair_rep_iterator<T> pair_rep_begin() const
177  {
178  return const_pair_rep_iterator<T>{thrust::counting_iterator<size_type>(0),
179  pair_rep_accessor<T>{*this}};
180  }
181 
186  template <typename T>
187  CUDA_DEVICE_CALLABLE const_pair_rep_iterator<T> pair_rep_end() const
188  {
189  return const_pair_rep_iterator<T>{thrust::counting_iterator<size_type>(size()),
190  pair_rep_accessor<T>{*this}};
191  }
192 
193  private:
194  lists_column_device_view const& lists_column;
195  size_type _row_index{}; // Row index in the Lists column vector.
196  size_type _size{}; // Number of elements in *this* list row.
197 
198  size_type begin_offset; // Offset in list_column_device_view where this list begins.
199 
209  template <typename T>
210  struct pair_accessor {
211  list_device_view const& list;
212 
218  explicit CUDA_HOST_DEVICE_CALLABLE pair_accessor(list_device_view const& _list) : list{_list} {}
219 
226  CUDA_DEVICE_CALLABLE
227  thrust::pair<T, bool> operator()(cudf::size_type i) const
228  {
229  return {list.element<T>(i), !list.is_null(i)};
230  }
231  };
232 
245  template <typename T>
247  list_device_view const& list;
248 
249  using rep_type = device_storage_type_t<T>;
250 
256  explicit CUDA_HOST_DEVICE_CALLABLE pair_rep_accessor(list_device_view const& _list)
257  : list{_list}
258  {
259  }
260 
267  CUDA_DEVICE_CALLABLE
268  thrust::pair<rep_type, bool> operator()(cudf::size_type i) const
269  {
270  return {get_rep<T>(i), !list.is_null(i)};
271  }
272 
273  private:
274  template <typename R, std::enable_if_t<std::is_same<R, rep_type>::value, void>* = nullptr>
275  CUDA_DEVICE_CALLABLE rep_type get_rep(cudf::size_type i) const
276  {
277  return list.element<R>(i);
278  }
279 
280  template <typename R, std::enable_if_t<not std::is_same<R, rep_type>::value, void>* = nullptr>
281  CUDA_DEVICE_CALLABLE rep_type get_rep(cudf::size_type i) const
282  {
283  return list.element<R>(i).value();
284  }
285  };
286 };
287 
293  column_device_view const d_column;
294  CUDA_HOST_DEVICE_CALLABLE list_size_functor(column_device_view const& d_col) : d_column(d_col)
295  {
296 #if defined(__CUDA_ARCH__)
297  cudf_assert(d_col.type().id() == type_id::LIST && "Only list type column is supported");
298 #else
299  CUDF_EXPECTS(d_col.type().id() == type_id::LIST, "Only list type column is supported");
300 #endif
301  }
302  CUDA_DEVICE_CALLABLE size_type operator()(size_type idx)
303  {
304  if (d_column.is_null(idx)) return size_type{0};
305  auto d_offsets =
306  d_column.child(lists_column_view::offsets_column_index).data<size_type>() + d_column.offset();
307  return d_offsets[idx + 1] - d_offsets[idx];
308  }
309 };
310 
311 } // namespace cudf
cudf::data_type::id
constexpr type_id id() const noexcept
Returns the type identifier.
Definition: types.hpp:272
cudf::column_device_view::child
__device__ column_device_view child(size_type child_index) const noexcept
Returns the specified child.
Definition: column_device_view.cuh:851
cudf::detail::lists_column_device_view::size
CUDA_HOST_DEVICE_CALLABLE cudf::size_type size() const
Fetches number of rows in the lists column.
Definition: lists_column_device_view.cuh:49
cudf::list_device_view::pair_accessor
pair accessor for elements in a list_device_view
Definition: list_device_view.cuh:210
cudf::column_device_view
An immutable, non-owning view of device data as a column of elements that is trivially copyable and u...
Definition: column_device_view.cuh:300
types.hpp
Type declarations for libcudf.
cudf::list_device_view::element_offset
CUDA_DEVICE_CALLABLE 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...
Definition: list_device_view.cuh:72
cudf::list_size_functor
returns size of the list by row index
Definition: list_device_view.cuh:292
cudf::list_device_view::element
CUDA_DEVICE_CALLABLE T element(size_type idx) const
Fetches the element at the specified index within the list row.
Definition: list_device_view.cuh:86
cudf::list_device_view::pair_begin
CUDA_DEVICE_CALLABLE const_pair_iterator< T > pair_begin() const
Fetcher for a pair iterator to the first element in the list_device_view.
Definition: list_device_view.cuh:144
cudf::detail::lists_column_device_view
Given a column-device-view, an instance of this class provides a wrapper on this compound column for ...
Definition: lists_column_device_view.cuh:32
cudf::list_device_view::is_null
CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const
Checks whether element is null at specified index in the list row.
Definition: list_device_view.cuh:94
cudf::device_storage_type_t
std::conditional_t< std::is_same< numeric::decimal32, T >::value, int32_t, std::conditional_t< std::is_same< numeric::decimal64, T >::value, int64_t, T > > device_storage_type_t
"Returns" the corresponding type that is stored on the device when using cudf::column
Definition: type_dispatcher.hpp:102
cudf::list_device_view::pair_rep_begin
CUDA_DEVICE_CALLABLE const_pair_rep_iterator< T > pair_rep_begin() const
Fetcher for a pair iterator to the first element in the list_device_view.
Definition: list_device_view.cuh:176
cudf::detail::lists_column_device_view::offsets
CUDA_DEVICE_CALLABLE column_device_view offsets() const
Fetches the offsets column of the underlying list column.
Definition: lists_column_device_view.cuh:54
cudf::list_device_view::pair_rep_end
CUDA_DEVICE_CALLABLE 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.
Definition: list_device_view.cuh:187
cudf::detail::lists_column_device_view::offset
CUDA_DEVICE_CALLABLE size_type offset() const
Fetches the offset of the underlying column_device_view, in case it is a sliced/offset column.
Definition: lists_column_device_view.cuh:82
cudf::detail::column_device_view_base::size
__host__ __device__ size_type size() const noexcept
Returns the number of elements in the column.
Definition: column_device_view.cuh:134
cudf
cuDF interfaces
Definition: aggregation.hpp:34
cudf::detail::column_device_view_base::is_null
__device__ bool is_null(size_type element_index) const noexcept
Returns whether the specified element is null.
Definition: column_device_view.cuh:216
cudf::list_device_view::size
CUDA_DEVICE_CALLABLE size_type size() const
Fetches the number of elements in this list row.
Definition: list_device_view.cuh:109
cudf::list_device_view::pair_rep_accessor::pair_rep_accessor
CUDA_HOST_DEVICE_CALLABLE pair_rep_accessor(list_device_view const &_list)
constructor
Definition: list_device_view.cuh:256
cudf::column_device_view::element
__device__ T element(size_type element_index) const noexcept
Returns reference to element at the specified index.
Definition: column_device_view.cuh:337
cudf::list_device_view::pair_rep_accessor
pair rep accessor for elements in a list_device_view
Definition: list_device_view.cuh:246
cudf::detail::lists_column_device_view::child
CUDA_DEVICE_CALLABLE column_device_view child() const
Fetches the child column of the underlying list column.
Definition: lists_column_device_view.cuh:62
cudf::detail::column_device_view_base::offset
__host__ __device__ size_type offset() const noexcept
Returns the index of the first element relative to the base memory allocation, i.e....
Definition: column_device_view.cuh:165
CUDF_EXPECTS
#define CUDF_EXPECTS(cond, reason)
Macro for checking (pre-)conditions that throws an exception when a condition is violated.
Definition: error.hpp:62
cudf::detail::column_device_view_base::type
__host__ __device__ data_type type() const noexcept
Returns the element type.
Definition: column_device_view.cuh:139
cudf::list_device_view::pair_rep_accessor::operator()
CUDA_DEVICE_CALLABLE thrust::pair< rep_type, bool > operator()(cudf::size_type i) const
Accessor for the {rep_data, validity} pair at the specified index.
Definition: list_device_view.cuh:268
cudf::list_device_view::get_column
CUDA_DEVICE_CALLABLE lists_column_device_view const & get_column() const
Fetches the lists_column_device_view that contains this list.
Definition: list_device_view.cuh:114
cudf::list_device_view::pair_accessor::operator()
CUDA_DEVICE_CALLABLE thrust::pair< T, bool > operator()(cudf::size_type i) const
Accessor for the {data, validity} pair at the specified index.
Definition: list_device_view.cuh:227
cudf::detail::lists_column_device_view::is_null
CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const
Indicates whether the row (i.e. list) at the specified index is null.
Definition: lists_column_device_view.cuh:76
cudf::list_device_view::is_null
CUDA_DEVICE_CALLABLE bool is_null() const
Checks whether this list row is null.
Definition: list_device_view.cuh:104
cudf::detail::column_device_view_base::data
__host__ __device__ T const * data() const noexcept
Returns the underlying data casted to the specified type, plus the offset.
Definition: column_device_view.cuh:126
type_dispatcher.hpp
Defines the mapping between cudf::type_id runtime type information and concrete C++ types.
cudf::list_device_view::pair_end
CUDA_DEVICE_CALLABLE const_pair_iterator< T > pair_end() const
Fetcher for a pair iterator to one position past the last element in the list_device_view.
Definition: list_device_view.cuh:154
cudf::list_device_view::pair_accessor::pair_accessor
CUDA_HOST_DEVICE_CALLABLE pair_accessor(list_device_view const &_list)
constructor
Definition: list_device_view.cuh:218
cudf::list_device_view
A non-owning, immutable view of device data that represents a list of elements of arbitrary type (inc...
Definition: list_device_view.cuh:29