expressions.hpp
Go to the documentation of this file.
1 /*
2  * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
3  * SPDX-License-Identifier: Apache-2.0
4  */
5 #pragma once
6 
8 #include <cudf/column/scalar_column_view.hpp>
10 #include <cudf/scalar/scalar.hpp>
13 #include <cudf/types.hpp>
14 #include <cudf/utilities/error.hpp>
15 
16 #include <cuda/stream>
17 
18 #include <cstdint>
19 #include <functional>
20 #include <initializer_list>
21 #include <memory>
22 #include <optional>
23 #include <vector>
24 
30 namespace CUDF_EXPORT cudf {
31 
32 namespace detail {
33 namespace row_ir {
34 
44 struct node;
45 
49 struct ast_converter;
50 
51 } // namespace row_ir
52 } // namespace detail
53 
54 namespace ast {
60 // Forward declaration.
61 namespace detail {
62 class expression_parser;
63 class expression_transformer;
64 } // namespace detail
65 
72 struct [[nodiscard]] expression {
79  virtual cudf::size_type accept(detail::expression_parser& visitor) const = 0;
80 
87  virtual std::reference_wrapper<expression const> accept(
88  detail::expression_transformer& visitor) const = 0;
89 
96  [[nodiscard]] virtual std::unique_ptr<cudf::detail::row_ir::node> accept(
97  cudf::detail::row_ir::ast_converter& visitor) const = 0;
98 
106  [[nodiscard]] bool may_evaluate_null(table_view const& left, cuda::stream_ref stream) const
107  {
108  return may_evaluate_null(left, left, stream);
109  }
110 
119  [[nodiscard]] virtual bool may_evaluate_null(table_view const& left,
120  table_view const& right,
121  cuda::stream_ref stream) const = 0;
122 
123  virtual ~expression() {}
124 };
125 
131 enum class table_reference {
132  LEFT,
133  RIGHT,
134  OUTPUT
135 };
136 
141  public:
148  template <typename T>
149  __device__ T const value() const noexcept
150  {
151  if constexpr (std::is_same_v<T, cudf::string_view>) {
152  return string_view(static_cast<char const*>(_data), _size);
153  }
154  if constexpr (cudf::is_fixed_point<T>()) {
155  using rep_type = typename T::rep;
156  auto const rep = *static_cast<rep_type const*>(_data);
157  auto const scale = numeric::scale_type{type().scale()};
158  return T{numeric::scaled_integer<rep_type>{rep, scale}};
159  }
160  return *static_cast<T const*>(_data);
161  }
162 
167  template <typename T>
169  : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
170  {
171  }
172 
177  template <typename T>
179  : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
180  {
181  }
182 
187  template <typename T>
189  : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
190  {
191  }
192 
198  : generic_scalar_device_view(s.type(), s.data(), s.validity_data(), s.size())
199  {
200  }
201 
206  template <typename T>
208  : generic_scalar_device_view{s.type(), s.data(), s.validity_data()}
209  {
210  }
211 
212  protected:
213  void const* _data{};
214  size_type const _size{};
215 
224  generic_scalar_device_view(data_type type, void const* data, bool* is_valid)
225  : cudf::detail::scalar_device_view_base(type, is_valid), _data(data)
226  {
227  }
228 
237  generic_scalar_device_view(data_type type, void const* data, bool* is_valid, size_type size)
238  : cudf::detail::scalar_device_view_base(type, is_valid), _data(data), _size(size)
239  {
240  }
241 };
242 
246 class literal : public expression {
247  public:
254  template <typename T>
256  : scalar{ast_scalar{std::ref(value), generic_scalar_device_view(value)}}
257  {
258  }
259 
266  template <typename T>
268  : scalar{ast_scalar{std::ref(value), generic_scalar_device_view(value)}}
269  {
270  }
271 
278  template <typename T>
280  : scalar{ast_scalar{std::ref(value), generic_scalar_device_view(value)}}
281  {
282  }
283 
290  : scalar{ast_scalar{std::ref(value), generic_scalar_device_view(value)}}
291  {
292  }
293 
299  template <typename T>
301  : scalar{ast_scalar{std::ref(value), generic_scalar_device_view(value)}}
302  {
303  }
304 
310  literal(scalar_column_view value) : scalar{std::move(value)} {}
311 
317  [[nodiscard]] cudf::data_type get_data_type() const
318  {
319  return std::visit(
320  [](auto const& value) {
321  if constexpr (std::is_same_v<std::decay_t<decltype(value)>, ast_scalar>) {
322  return value.value.type();
323  } else {
324  return value.type();
325  }
326  },
327  scalar);
328  }
329 
335  [[nodiscard]] bool is_scalar_column_view() const noexcept
336  {
337  return std::holds_alternative<scalar_column_view>(scalar);
338  }
339 
345  [[nodiscard]] generic_scalar_device_view get_value() const
346  {
347  return std::get<ast_scalar>(scalar).value;
348  }
349 
355  [[nodiscard]] cudf::scalar const& get_scalar() const
356  {
357  return std::get<ast_scalar>(scalar).scalar.get();
358  }
359 
365  [[nodiscard]] scalar_column_view const& get_scalar_column_view() const
366  {
367  return std::get<scalar_column_view>(scalar);
368  }
369 
373  cudf::size_type accept(detail::expression_parser& visitor) const override;
374 
378  std::reference_wrapper<expression const> accept(
379  detail::expression_transformer& visitor) const override;
380 
384  [[nodiscard]] std::unique_ptr<cudf::detail::row_ir::node> accept(
385  cudf::detail::row_ir::ast_converter& visitor) const override;
386 
387  [[nodiscard]] bool may_evaluate_null(table_view const& left,
388  table_view const& right,
389  cuda::stream_ref stream) const override
390  {
391  return !is_valid(stream);
392  }
393 
400  [[nodiscard]] bool is_valid(cuda::stream_ref stream) const
401  {
402  if (auto* s = std::get_if<ast_scalar>(&scalar)) {
403  return s->scalar.get().is_valid(stream);
404  } else {
405  auto& c = std::get<scalar_column_view>(scalar);
406  return c.null_count() == 0;
407  }
408  }
409 
410  private:
411  struct ast_scalar {
412  std::reference_wrapper<cudf::scalar const> scalar;
414  };
415 
416  std::variant<ast_scalar, scalar_column_view> scalar;
417 };
418 
422 class column_reference : public expression {
423  public:
432  table_reference table_source = table_reference::LEFT)
433  : column_index(column_index), table_source(table_source)
434  {
435  }
436 
442  [[nodiscard]] cudf::size_type get_column_index() const { return column_index; }
443 
449  [[nodiscard]] table_reference get_table_source() const { return table_source; }
450 
457  [[nodiscard]] cudf::data_type get_data_type(table_view const& table) const;
458 
466  [[nodiscard]] cudf::data_type get_data_type(table_view const& left_table,
467  table_view const& right_table) const;
468 
472  cudf::size_type accept(detail::expression_parser& visitor) const override;
473 
477  std::reference_wrapper<expression const> accept(
478  detail::expression_transformer& visitor) const override;
479 
480  [[nodiscard]] bool may_evaluate_null(table_view const& left,
481  table_view const& right,
482  cuda::stream_ref stream) const override
483  {
484  return (table_source == table_reference::LEFT ? left : right).column(column_index).has_nulls();
485  }
486 
490  [[nodiscard]] std::unique_ptr<cudf::detail::row_ir::node> accept(
491  cudf::detail::row_ir::ast_converter& visitor) const override;
492 
493  private:
494  cudf::size_type column_index;
495  table_reference table_source;
496 };
497 
501 class operation : public expression {
502  public:
509  operation(ast_operator op, expression const& input);
510 
518  operation(ast_operator op, expression const& left, expression const& right);
519 
520  // operation only stores references to expressions, so it does not accept r-value
521  // references: the calling code must own the expressions.
522  operation(ast_operator op, expression&& input) = delete;
523  operation(ast_operator op, expression&& left, expression&& right) = delete;
524  operation(ast_operator op, expression&& left, expression const& right) = delete;
525  operation(ast_operator op, expression const& left, expression&& right) = delete;
526 
532  [[nodiscard]] ast_operator get_operator() const { return op; }
533 
539  [[nodiscard]] std::vector<std::reference_wrapper<expression const>> const& get_operands() const
540  {
541  return operands;
542  }
543 
547  cudf::size_type accept(detail::expression_parser& visitor) const override;
548 
552  std::reference_wrapper<expression const> accept(
553  detail::expression_transformer& visitor) const override;
554 
555  [[nodiscard]] bool may_evaluate_null(table_view const& left,
556  table_view const& right,
557  cuda::stream_ref stream) const override;
558 
562  [[nodiscard]] std::unique_ptr<cudf::detail::row_ir::node> accept(
563  cudf::detail::row_ir::ast_converter& visitor) const override;
564 
565  private:
566  ast_operator op;
567  std::vector<std::reference_wrapper<expression const>> operands;
568 };
569 
570 namespace detail {
571 
576 class predicate : public expression {
577  public:
582  predicate(expression const& source) : source_{source} {}
583 
587  cudf::size_type accept(detail::expression_parser& visitor) const override;
588 
592  std::reference_wrapper<expression const> accept(
593  detail::expression_transformer& visitor) const override;
594 
595  [[nodiscard]] bool may_evaluate_null(table_view const& left,
596  table_view const& right,
597  cuda::stream_ref stream) const override;
598 
602  [[nodiscard]] std::unique_ptr<cudf::detail::row_ir::node> accept(
603  cudf::detail::row_ir::ast_converter& visitor) const override;
604 
609  [[nodiscard]] expression const& get_operand() const { return source_; }
610 
611  private:
612  std::reference_wrapper<expression const> source_;
613 };
614 
615 } // namespace detail
616 
621  public:
628  column_name_reference(std::string column_name) : column_name(std::move(column_name)) {}
629 
635  [[nodiscard]] std::string get_column_name() const { return column_name; }
636 
640  cudf::size_type accept(detail::expression_parser& visitor) const override;
641 
645  std::reference_wrapper<expression const> accept(
646  detail::expression_transformer& visitor) const override;
647 
648  [[nodiscard]] bool may_evaluate_null(table_view const& left,
649  table_view const& right,
650  cuda::stream_ref stream) const override
651  {
652  return true;
653  }
654 
658  [[nodiscard]] std::unique_ptr<cudf::detail::row_ir::node> accept(
659  cudf::detail::row_ir::ast_converter& visitor) const override;
660 
661  private:
662  std::string column_name;
663 };
664 
669 class tree {
670  public:
674  tree() = default;
675 
679  tree(tree&&) = default;
680 
685  tree& operator=(tree&&) = default;
686 
687  ~tree() = default;
688 
689  // the tree is not copyable
690  tree(tree const&) = delete;
691  tree& operator=(tree const&) = delete;
692 
698  template <typename Expr, typename... Args>
699  std::enable_if_t<std::is_base_of_v<expression, Expr>, Expr const&> emplace(Args&&... args)
700  {
701  auto expr = std::make_unique<Expr>(std::forward<Args>(args)...);
702  Expr const& expr_ref = *expr;
703  expressions.emplace_back(std::move(expr));
704  return expr_ref;
705  }
706 
712  template <typename Expr>
713  decltype(auto) push(Expr expr)
714  {
715  return emplace<Expr>(std::move(expr));
716  }
717 
722  [[nodiscard]] expression const& front() const { return *expressions.front(); }
723 
728  [[nodiscard]] expression const& back() const { return *expressions.back(); }
729 
734  [[nodiscard]] size_t size() const { return expressions.size(); }
735 
741  expression const& at(size_t index) { return *expressions.at(index); }
742 
748  expression const& operator[](size_t index) const { return *expressions[index]; }
749 
750  private:
751  // TODO: use better ownership semantics, the unique_ptr here is redundant. Consider using a bump
752  // allocator with type-erased deleters.
753  std::vector<std::unique_ptr<expression>> expressions;
754 };
755 
756 namespace jit {
757 
761 enum class op : uint8_t {
762  // Identity operators
763  IDENTITY,
764 
765  // Null handling operators
766  IS_NULL,
767 
768  COALESCE,
769  PREDICATE,
770 
772  ADD,
773  SUB,
774  MUL,
775  DIV,
776  NEG,
777  ABS,
778  MOD,
779  PYMOD,
780  TRUE_DIV,
781  FLOOR_DIV,
782 
784  ADD_OVERFLOW,
785  SUB_OVERFLOW,
786  MUL_OVERFLOW,
787  DIV_OVERFLOW,
788  NEG_OVERFLOW,
789  ABS_OVERFLOW,
790  MOD_OVERFLOW,
791  CHECK_PRECISION,
792 
794  BITWISE_AND,
795  BITWISE_INVERT,
796  BITWISE_OR,
797  BITWISE_XOR,
798  BITWISE_SHIFT_LEFT,
799  BITWISE_SHIFT_RIGHT,
800 
803  CAST_TO_INT8,
804  CAST_TO_INT16,
805  CAST_TO_INT32,
807  CAST_TO_UINT8,
808  CAST_TO_UINT16,
809  CAST_TO_UINT32,
811  CAST_TO_FLOAT32,
813  CAST_TO_DECIMAL32,
814  CAST_TO_DECIMAL64,
815  CAST_TO_DECIMAL128,
816  RESCALE,
817 
819  EQUAL,
820  NOT_EQUAL,
821  GREATER,
823  LESS,
824  LESS_EQUAL,
825  NULL_EQUAL,
828  LOGICAL_AND,
829  LOGICAL_OR,
830  LOGICAL_NOT,
831  IF_ELSE,
832 
834  CBRT,
835  CEIL,
836  FLOOR,
837  RINT,
838  SQRT,
839  POW,
840  EXP,
841  LOG,
842 
844  ARCCOS,
845  ARCCOSH,
846  ARCSIN,
847  ARCSINH,
848  ARCTAN,
849  ARCTANH,
850  COS,
851  COSH,
852  SIN,
853  SINH,
854  TAN,
855  TANH,
856 };
857 
869  op operator_id,
870  std::vector<std::reference_wrapper<expression const>> const& args,
872  std::optional<int32_t> target_scale = std::nullopt);
873 
885  op operator_id,
886  std::initializer_list<std::reference_wrapper<expression const>> args,
888  std::optional<int32_t> target_scale = std::nullopt);
889 
890 } // namespace jit
891  // end of group
893 } // namespace ast
894 } // namespace CUDF_EXPORT cudf
Enum defining the supported AST operators.
A expression referring to data from a column in a table.
bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const override
Returns true if the expression may evaluate to null.
cudf::size_type accept(detail::expression_parser &visitor) const override
Accepts a visitor class.
std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const override
Accepts a visitor class.
std::string get_column_name() const
Get the column name.
std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const override
Accepts a visitor class.
column_name_reference(std::string column_name)
Construct a new column name reference object.
A expression referring to data from a column in a table.
cudf::data_type get_data_type(table_view const &left_table, table_view const &right_table) const
Get the data type.
std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const override
Accepts a visitor class.
cudf::size_type accept(detail::expression_parser &visitor) const override
Accepts a visitor class.
cudf::data_type get_data_type(table_view const &table) const
Get the data type.
bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const override
Returns true if the expression may evaluate to null.
std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const override
Accepts a visitor class.
table_reference get_table_source() const
Get the table source.
column_reference(cudf::size_type column_index, table_reference table_source=table_reference::LEFT)
Construct a new column reference object.
cudf::size_type get_column_index() const
Get the column index.
An expression that represents a predicate.
std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const override
Accepts a visitor class.
bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const override
Returns true if the expression may evaluate to null.
predicate(expression const &source)
Construct a new filter predicate object.
expression const & get_operand() const
Get the operand expression.
std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const override
Accepts a visitor class.
cudf::size_type accept(detail::expression_parser &visitor) const override
Accepts a visitor class.
A type-erased scalar_device_view where the value is a fixed width type or a string.
generic_scalar_device_view(duration_scalar< T > &s)
Construct a new generic scalar device view object from a duration scalar.
generic_scalar_device_view(cudf::fixed_point_scalar< T > &s)
Construct a new generic scalar device view object from a fixed-point scalar.
generic_scalar_device_view(data_type type, void const *data, bool *is_valid)
Construct a new fixed width scalar device view object.
generic_scalar_device_view(string_scalar &s)
Construct a new generic scalar device view object from a string scalar.
generic_scalar_device_view(timestamp_scalar< T > &s)
Construct a new generic scalar device view object from a timestamp scalar.
generic_scalar_device_view(data_type type, void const *data, bool *is_valid, size_type size)
Construct a new string scalar device view object.
T const value() const noexcept
Returns the stored value.
generic_scalar_device_view(numeric_scalar< T > &s)
Construct a new generic scalar device view object from a numeric scalar.
A literal value used in an abstract syntax tree.
literal(cudf::numeric_scalar< T > &value)
Construct a new literal object.
std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const override
Accepts a visitor class.
cudf::size_type accept(detail::expression_parser &visitor) const override
Accepts a visitor class.
bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const override
Returns true if the expression may evaluate to null.
generic_scalar_device_view get_value() const
Get the value object.
bool is_scalar_column_view() const noexcept
Check whether the literal is backed by a scalar column view.
cudf::data_type get_data_type() const
Get the data type.
literal(cudf::string_scalar &value)
Construct a new literal object.
cudf::scalar const & get_scalar() const
Get the scalar.
literal(cudf::duration_scalar< T > &value)
Construct a new literal object.
scalar_column_view const & get_scalar_column_view() const
Get scalar column view.
bool is_valid(cuda::stream_ref stream) const
Check if the underlying scalar is valid.
std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const override
Accepts a visitor class.
literal(scalar_column_view value)
Construct a new literal object.
literal(cudf::fixed_point_scalar< T > &value)
Construct a new literal object.
literal(cudf::timestamp_scalar< T > &value)
Construct a new literal object.
An operation expression holds an operator and zero or more operands.
std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const override
Accepts a visitor class.
std::vector< std::reference_wrapper< expression const > > const & get_operands() const
Get the operands.
ast_operator get_operator() const
Get the operator.
operation(ast_operator op, expression const &left, expression const &right)
Construct a new binary operation object.
bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const override
Returns true if the expression may evaluate to null.
std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const override
Accepts a visitor class.
operation(ast_operator op, expression const &input)
Construct a new unary operation object.
cudf::size_type accept(detail::expression_parser &visitor) const override
Accepts a visitor class.
An AST expression tree. It owns and contains multiple dependent expressions. All the expressions are ...
tree(tree &&)=default
Moves the ast tree.
tree & operator=(tree &&)=default
move-assigns the AST tree
size_t size() const
get the number of expressions added to the tree
std::enable_if_t< std::is_base_of_v< expression, Expr >, Expr const & > emplace(Args &&... args)
Add an expression to the AST tree.
tree()=default
construct an empty ast tree
expression const & front() const
get the first expression in the tree
expression const & at(size_t index)
get the expression at an index in the tree. Index is checked.
expression const & back() const
get the last expression in the tree
expression const & operator[](size_t index) const
get the expression at an index in the tree. Index is unchecked.
Indicator for the logical data type of an element in a column.
Definition: types.hpp:279
A non-owning view of scalar from device that is trivially copyable and usable in CUDA device code.
An owning class to represent a duration value in device memory.
Definition: scalar.hpp:752
An owning class to represent a fixed_point number in device memory.
Definition: scalar.hpp:309
An owning class to represent a numerical value in device memory.
Definition: scalar.hpp:230
An owning class to represent a singular value.
Definition: scalar.hpp:42
An owning class to represent a string in device memory.
Definition: scalar.hpp:450
A non-owning, immutable view of device data that is a variable length char array representing a UTF-8...
Definition: string_view.hpp:35
A set of cudf::column_view's of the same size.
Definition: table_view.hpp:206
A set of cudf::column's of the same size.
Definition: table.hpp:31
An owning class to represent a timestamp value in device memory.
Definition: scalar.hpp:677
Exception types and error-checking macros used throughout libcudf.
op
JIT operation kinds for cudf::ast::jit::operation.
@ ADD_OVERFLOW
Overflow-checking Arithmetic functions. raise errors on overflow, division by zero,...
@ CAST_TO_BOOL8
Type conversion/scaling operators.
Class definition for fixed point data type.
table_reference
Enum of table references.
ast_operator
Enum of supported operators.
@ RIGHT
Column index in the right table.
@ OUTPUT
Column index in the output table.
@ LEFT
Column index in the left table.
@ TANH
Hyperbolic tangent.
@ DIV
operator / using common type of lhs and rhs
@ CBRT
Cube-root (x^(1.0/3))
@ ARCSINH
Hyperbolic sine inverse.
@ SQRT
Square-root (x^0.5)
@ PYMOD
operator % using Python's sign rules for negatives
@ LOG
Natural Logarithm (base e)
@ FLOOR
largest integer value not greater than arg
@ ARCTAN
Trigonometric tangent inverse.
@ SIN
Trigonometric sine.
@ CEIL
Smallest integer value not less than arg.
@ ARCSIN
Trigonometric sine inverse.
@ RINT
Rounds the floating-point argument arg to an integer value.
@ TAN
Trigonometric tangent.
@ ARCCOS
Trigonometric cosine inverse.
@ ABS
Absolute value.
@ ARCTANH
Hyperbolic tangent inverse.
@ GREATER_EQUAL
operator >=
@ EXP
Exponential (base e, Euler number)
@ CAST_TO_FLOAT64
Cast value to double.
@ TRUE_DIV
operator / after promoting type to floating point
@ SINH
Hyperbolic sine.
@ CAST_TO_UINT64
Cast value to uint64_t.
@ COSH
Hyperbolic cosine.
@ ARCCOSH
Hyperbolic cosine inverse.
@ IDENTITY
Identity function.
@ IS_NULL
Check if operand is null.
@ CAST_TO_INT64
Cast value to int64_t.
@ COS
Trigonometric cosine.
scale_type
The scale type for fixed_point.
Definition: fixed_point.hpp:38
std::unique_ptr< cudf::column > is_valid(cudf::column_view const &input, cuda::stream_ref stream=cudf::get_default_stream(), rmm::device_async_resource_ref mr=cudf::get_current_device_resource_ref())
Creates a column of type_id::BOOL8 elements where for every element in input true indicates the value...
int32_t size_type
Row index type for columns and tables.
Definition: types.hpp:76
error_policy
Indicates whether a function nullifies its output on error.
Definition: types.hpp:260
@ PROPAGATE
The function propagates errors.
cuDF interfaces
Definition: host_udf.hpp:27
Class definitions for cudf::scalar.
Scalar device view class definitions.
A generic expression that can be evaluated to return a value.
Definition: expressions.hpp:72
virtual std::unique_ptr< cudf::detail::row_ir::node > accept(cudf::detail::row_ir::ast_converter &visitor) const =0
Accepts an row_ir::ast_converter class.
virtual cudf::size_type accept(detail::expression_parser &visitor) const =0
Accepts a visitor class.
bool may_evaluate_null(table_view const &left, cuda::stream_ref stream) const
Returns true if the expression may evaluate to null.
virtual std::reference_wrapper< expression const > accept(detail::expression_transformer &visitor) const =0
Accepts a visitor class.
virtual bool may_evaluate_null(table_view const &left, table_view const &right, cuda::stream_ref stream) const =0
Returns true if the expression may evaluate to null.
A non-owning, immutable view of device data as a column of elements, some of which may be null as ind...
Helper struct for constructing fixed_point when value is already shifted.
Class definitions for (mutable)_table_view
Type declarations for libcudf.