diff --git a/cpp/include/cudf/ast/detail/expression_evaluator.cuh b/cpp/include/cudf/ast/detail/expression_evaluator.cuh index 428b84eb03bc..1a95bf16382f 100644 --- a/cpp/include/cudf/ast/detail/expression_evaluator.cuh +++ b/cpp/include/cudf/ast/detail/expression_evaluator.cuh @@ -719,7 +719,7 @@ struct expression_evaluator { typename ResultSubclass, typename T, bool result_has_nulls, - CUDF_ENABLE_IF(detail::is_valid_unary_op, + CUDF_ENABLE_IF(detail::is_valid_unary_op, possibly_null_value_t>)> __device__ inline void operator()( expression_result& output_object, @@ -730,19 +730,19 @@ struct expression_evaluator { { // The output data type is the same whether or not nulls are present, so // pull from the non-nullable operator. - using Out = cuda::std::invoke_result_t, Input>; + using Out = cuda::std::invoke_result_t, Input>; this->template resolve_output(output_object, output, output_row_index, thread_intermediate_storage, - detail::operator_functor{}(input)); + detail::operator_functor{}(input)); } template , + CUDF_ENABLE_IF(!detail::is_valid_unary_op, possibly_null_value_t>)> __device__ inline void operator()( expression_result& output_object, @@ -781,7 +781,7 @@ struct expression_evaluator { typename ResultSubclass, typename T, bool result_has_nulls, - CUDF_ENABLE_IF(detail::is_valid_binary_op, + CUDF_ENABLE_IF(detail::is_valid_binary_op, possibly_null_value_t, possibly_null_value_t>)> __device__ inline void operator()( @@ -794,19 +794,19 @@ struct expression_evaluator { { // The output data type is the same whether or not nulls are present, so // pull from the non-nullable operator. - using Out = cuda::std::invoke_result_t, LHS, RHS>; + using Out = cuda::std::invoke_result_t, LHS, RHS>; this->template resolve_output(output_object, output, output_row_index, thread_intermediate_storage, - detail::operator_functor{}(lhs, rhs)); + detail::operator_functor{}(lhs, rhs)); } template , + CUDF_ENABLE_IF(!detail::is_valid_binary_op, possibly_null_value_t, possibly_null_value_t>)> __device__ inline void operator()( diff --git a/cpp/include/cudf/ast/detail/operator_functor.cuh b/cpp/include/cudf/ast/detail/operator_functor.cuh index e245f0288be6..2d600247b41b 100644 --- a/cpp/include/cudf/ast/detail/operator_functor.cuh +++ b/cpp/include/cudf/ast/detail/operator_functor.cuh @@ -5,25 +5,91 @@ #pragma once #include -#include -#include -#include -#include +#include +#include #include #include -#include +#include +#include namespace CUDF_EXPORT cudf { namespace ast::detail { +template +struct operator_invoker; + +#define CUDF_AST_OPERATOR_MAP(OP, func_name, num_args) \ + template <> \ + struct operator_invoker { \ + static constexpr auto arity = num_args; \ + template \ + __device__ static inline auto eval(Args... a) -> decltype(cudf::detail::ops::func_name(a...)) \ + { \ + return cudf::detail::ops::func_name(a...); \ + } \ + }; + +CUDF_AST_OPERATOR_MAP(ADD, add, 2) +CUDF_AST_OPERATOR_MAP(SUB, sub, 2) +CUDF_AST_OPERATOR_MAP(MUL, mul, 2) +CUDF_AST_OPERATOR_MAP(DIV, div, 2) +CUDF_AST_OPERATOR_MAP(TRUE_DIV, true_div, 2) +CUDF_AST_OPERATOR_MAP(FLOOR_DIV, floor_div, 2) +CUDF_AST_OPERATOR_MAP(MOD, mod, 2) +CUDF_AST_OPERATOR_MAP(PYMOD, pymod, 2) +CUDF_AST_OPERATOR_MAP(POW, pow, 2) +CUDF_AST_OPERATOR_MAP(EQUAL, equal, 2) +CUDF_AST_OPERATOR_MAP(NOT_EQUAL, not_equal, 2) +CUDF_AST_OPERATOR_MAP(LESS, less, 2) +CUDF_AST_OPERATOR_MAP(GREATER, greater, 2) +CUDF_AST_OPERATOR_MAP(LESS_EQUAL, less_equal, 2) +CUDF_AST_OPERATOR_MAP(GREATER_EQUAL, greater_equal, 2) +CUDF_AST_OPERATOR_MAP(BITWISE_AND, bit_and, 2) +CUDF_AST_OPERATOR_MAP(BITWISE_OR, bit_or, 2) +CUDF_AST_OPERATOR_MAP(BITWISE_XOR, bit_xor, 2) +CUDF_AST_OPERATOR_MAP(LOGICAL_AND, logical_and, 2) +CUDF_AST_OPERATOR_MAP(LOGICAL_OR, logical_or, 2) +CUDF_AST_OPERATOR_MAP(IDENTITY, identity, 1) +CUDF_AST_OPERATOR_MAP(SIN, sin, 1) +CUDF_AST_OPERATOR_MAP(COS, cos, 1) +CUDF_AST_OPERATOR_MAP(TAN, tan, 1) +CUDF_AST_OPERATOR_MAP(ARCSIN, arcsin, 1) +CUDF_AST_OPERATOR_MAP(ARCCOS, arccos, 1) +CUDF_AST_OPERATOR_MAP(ARCTAN, arctan, 1) +CUDF_AST_OPERATOR_MAP(SINH, sinh, 1) +CUDF_AST_OPERATOR_MAP(COSH, cosh, 1) +CUDF_AST_OPERATOR_MAP(TANH, tanh, 1) +CUDF_AST_OPERATOR_MAP(ARCSINH, arcsinh, 1) +CUDF_AST_OPERATOR_MAP(ARCCOSH, arccosh, 1) +CUDF_AST_OPERATOR_MAP(ARCTANH, arctanh, 1) +CUDF_AST_OPERATOR_MAP(EXP, exp, 1) +CUDF_AST_OPERATOR_MAP(LOG, log, 1) +CUDF_AST_OPERATOR_MAP(SQRT, sqrt, 1) +CUDF_AST_OPERATOR_MAP(CBRT, cbrt, 1) +CUDF_AST_OPERATOR_MAP(CEIL, ceil, 1) +CUDF_AST_OPERATOR_MAP(FLOOR, floor, 1) +CUDF_AST_OPERATOR_MAP(ABS, abs, 1) +CUDF_AST_OPERATOR_MAP(RINT, rint, 1) +CUDF_AST_OPERATOR_MAP(BIT_INVERT, bit_invert, 1) +CUDF_AST_OPERATOR_MAP(NOT, logical_not, 1) +CUDF_AST_OPERATOR_MAP(CAST_TO_INT64, cast_to_i64, 1) +CUDF_AST_OPERATOR_MAP(CAST_TO_UINT64, cast_to_u64, 1) +CUDF_AST_OPERATOR_MAP(CAST_TO_FLOAT64, cast_to_f64, 1) +CUDF_AST_OPERATOR_MAP(IS_NULL, is_null, 1) +CUDF_AST_OPERATOR_MAP(NULL_EQUAL, null_equal, 2) +CUDF_AST_OPERATOR_MAP(NULL_LOGICAL_AND, null_logical_and, 2) +CUDF_AST_OPERATOR_MAP(NULL_LOGICAL_OR, null_logical_or, 2) + +#undef CUDF_AST_OPERATOR_MAP + /** * @brief Operator functor. * * This functor is templated on an `ast_operator`, with each template specialization defining a - * callable `operator()` that executes the operation. The functor specialization also has a member - * `arity` defining the number of operands that are accepted by the call to `operator()`. The - * `operator()` is templated on the types of its inputs (e.g. `typename LHS` and `typename RHS` for + * callable `eval` that executes the operation. The functor specialization also has a member + * `arity` defining the number of operands that are accepted by the call to `eval`. The + * `eval` is templated on the types of its inputs (e.g. `typename LHS` and `typename RHS` for * a binary operator). Trailing return types are defined as `decltype(result)` where `result` is * the returned value. The trailing return types allow SFINAE to only consider template * instantiations for valid combinations of types. This, in turn, allows the operator functors to be @@ -31,759 +97,84 @@ namespace ast::detail { * * @tparam op AST operator. */ -template -struct operator_functor {}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs + rhs) - { - return lhs + rhs; - } - - static constexpr int32_t fixed_point_result_scale(int32_t lhs, int32_t rhs) - { - return cuda::std::min(lhs, rhs); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs - rhs) - { - return lhs - rhs; - } - - static constexpr int32_t fixed_point_result_scale(int32_t lhs, int32_t rhs) - { - return cuda::std::min(lhs, rhs); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs * rhs) - { - return lhs * rhs; - } - - static constexpr int32_t fixed_point_result_scale(int32_t lhs, int32_t rhs) { return lhs + rhs; } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs / rhs) - { - return lhs / rhs; - } - - static constexpr int32_t fixed_point_result_scale(int32_t lhs, int32_t rhs) { return lhs - rhs; } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(static_cast(lhs) / static_cast(rhs)) - { - return static_cast(lhs) / static_cast(rhs); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; +template +struct operator_functor { + static constexpr auto arity = operator_invoker::arity; - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> CommonType - requires(cuda::std::is_integral_v) + template + __device__ inline auto operator()(T a) + requires(!cudf::detail::ops::nullable && requires { operator_invoker::eval(a); }) { - return cudf::detail::integral_floor_div(lhs, rhs); + return operator_invoker::eval(a); } - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> CommonType - requires(cuda::std::is_floating_point_v) + template + __device__ inline auto operator()(T a) + requires( + cudf::detail::ops::nullable && (requires { operator_invoker::eval(a); } || + requires { operator_invoker::eval(a.value()); }) + ) { - if constexpr (cuda::std::is_same_v) { - return cuda::std::floorf(static_cast(lhs) / static_cast(rhs)); + // If the operator is not defined for optional, but is defined for T then it is assumed to be + // null-propagating. + if constexpr (requires { operator_invoker::eval(a); }) { + return operator_invoker::eval(a); + } else { + using result_t = cuda::std::optional::eval(a.value()))>; + if (a.has_value()) { + return result_t{operator_invoker::eval(a.value())}; + } else { + return result_t{}; + } } - return cuda::std::floor(static_cast(lhs) / static_cast(rhs)); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(static_cast(lhs) % static_cast(rhs)) - requires(cuda::std::is_integral_v) - { - return static_cast(lhs) % static_cast(rhs); - } - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(fmodf(static_cast(lhs), static_cast(rhs))) - requires(cuda::std::is_same_v) - { - return fmodf(static_cast(lhs), static_cast(rhs)); - } - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(fmod(static_cast(lhs), static_cast(rhs))) - requires(cuda::std::is_same_v) - { - return fmod(static_cast(lhs), static_cast(rhs)); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(((static_cast(lhs) % static_cast(rhs)) + - static_cast(rhs)) % - static_cast(rhs)) - requires(cuda::std::is_integral_v) - { - return ((static_cast(lhs) % static_cast(rhs)) + - static_cast(rhs)) % - static_cast(rhs); - } - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(fmodf(fmodf(static_cast(lhs), static_cast(rhs)) + - static_cast(rhs), - static_cast(rhs))) - requires(cuda::std::is_same_v) - { - return fmodf(fmodf(static_cast(lhs), static_cast(rhs)) + - static_cast(rhs), - static_cast(rhs)); - } - - template > - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(fmod(fmod(static_cast(lhs), static_cast(rhs)) + - static_cast(rhs), - static_cast(rhs))) - requires(cuda::std::is_same_v) - { - return fmod(fmod(static_cast(lhs), static_cast(rhs)) + - static_cast(rhs), - static_cast(rhs)); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> LHS - requires(cuda::std::is_integral_v and cuda::std::is_integral_v) - { - return cudf::detail::integral_pow(lhs, rhs); - } - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept - -> decltype(cuda::std::pow(lhs, rhs)) - requires(not(cuda::std::is_integral_v and cuda::std::is_integral_v)) - { - return cuda::std::pow(lhs, rhs); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs == rhs) - { - return lhs == rhs; - } -}; - -// Alias NULL_EQUAL = EQUAL in the non-nullable case. -template <> -struct operator_functor - : public operator_functor {}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs != rhs) - { - return lhs != rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs < rhs) - { - return lhs < rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs > rhs) - { - return lhs > rhs; } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs <= rhs) - { - return lhs <= rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs >= rhs) - { - return lhs >= rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs & rhs) + template + __device__ inline auto operator()(T a, T b) + requires(!cudf::detail::ops::nullable && requires { operator_invoker::eval(a, b); }) { - return lhs & rhs; + return operator_invoker::eval(a, b); } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs | rhs) + template + __device__ inline auto operator()(T a, T b) + requires( + cudf::detail::ops::nullable && + (requires { operator_invoker::eval(a, b); } || + requires { operator_invoker::eval(a.value(), b.value()); })) { - return lhs | rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs ^ rhs) - { - return lhs ^ rhs; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs && rhs) - { - return lhs && rhs; - } -}; - -// Alias NULL_LOGICAL_AND = LOGICAL_AND in the non-nullable case. -template <> -struct operator_functor - : public operator_functor {}; - -template <> -struct operator_functor { - static constexpr auto arity{2}; - - template - __device__ inline auto operator()(LHS lhs, RHS rhs) const noexcept -> decltype(lhs || rhs) - { - return lhs || rhs; - } -}; - -// Alias NULL_LOGICAL_OR = LOGICAL_OR in the non-nullable case. -template <> -struct operator_functor - : public operator_functor {}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(input) - { - return input; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> bool - { - return false; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::sin(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::sin(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::cos(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::cos(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::tan(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::tan(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::asin(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::asin(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::acos(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::acos(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::atan(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::atan(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::sinh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::sinh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::cosh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::cosh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::tanh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::tanh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept - -> decltype(cuda::std::asinh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::asinh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept - -> decltype(cuda::std::acosh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::acosh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept - -> decltype(cuda::std::atanh(input)) - requires(cuda::std::is_floating_point_v) - { - return cuda::std::atanh(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::exp(input)) - { - return cuda::std::exp(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::log(input)) - { - return cuda::std::log(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::sqrt(input)) - { - return cuda::std::sqrt(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::cbrt(input)) - { - return cuda::std::cbrt(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::ceil(input)) - { - return cuda::std::ceil(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept - -> decltype(cuda::std::floor(input)) - { - return cuda::std::floor(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - // Only accept signed or unsigned types (both require is_arithmetic to be true) - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::abs(input)) - requires(cuda::std::is_signed_v) - { - return cuda::std::abs(input); - } - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(input) - requires(cuda::std::is_unsigned_v) - { - return input; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(cuda::std::rint(input)) - { - return cuda::std::rint(input); - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(~input) - { - return ~input; - } -}; - -template <> -struct operator_functor { - static constexpr auto arity{1}; - - template - __device__ inline auto operator()(InputT input) const noexcept -> decltype(!input) - { - return !input; - } -}; - -template -struct cast { - static constexpr auto arity{1}; - template - __device__ inline auto operator()(From f) const noexcept -> To - requires(is_fixed_point()) - { - if constexpr (cuda::std::is_floating_point_v) { - return convert_fixed_to_floating(f); + // If the operator is not defined for optional, but is defined for T then it is assumed to be + // null-propagating. + if constexpr (requires { operator_invoker::eval(a, b); }) { + return operator_invoker::eval(a, b); } else { - return static_cast(f); + using result_t = + cuda::std::optional::eval(a.value(), b.value()))>; + if (a.has_value() && b.has_value()) { + return result_t{operator_invoker::eval(a.value(), b.value())}; + } else { + return result_t{}; + } } } - template - __device__ inline auto operator()(From f) const noexcept -> decltype(static_cast(f)) - requires(!is_fixed_point()) - { - return static_cast(f); - } -}; - -template <> -struct operator_functor : cast {}; -template <> -struct operator_functor : cast {}; -template <> -struct operator_functor : cast {}; - -/* - * The default specialization of nullable operators is to fall back to the non-nullable - * implementation - */ -template -struct operator_functor { - using NonNullOperator = operator_functor; - static constexpr auto arity = NonNullOperator::arity; - - template - __device__ inline auto operator()(LHS const lhs, RHS const rhs) const noexcept - -> possibly_null_value_t - requires(arity_placeholder == 2) - { - using Out = possibly_null_value_t; - return (lhs.has_value() && rhs.has_value()) ? Out{NonNullOperator{}(*lhs, *rhs)} : Out{}; - } - - template - __device__ inline auto operator()(Input const input) const noexcept - -> possibly_null_value_t - requires(arity_placeholder == 1) - { - using Out = possibly_null_value_t; - return input.has_value() ? Out{NonNullOperator{}(*input)} : Out{}; - } -}; - -// IS_NULL(null) is true, IS_NULL(valid) is false -template <> -struct operator_functor { - using NonNullOperator = operator_functor; - static constexpr auto arity = NonNullOperator::arity; - - template - __device__ inline auto operator()(LHS const lhs) const noexcept -> bool - { - return !lhs.has_value(); - } -}; - -// NULL_EQUAL(null, null) is true, NULL_EQUAL(null, valid) is false, and NULL_EQUAL(valid, valid) == -// EQUAL(valid, valid) -template <> -struct operator_functor { - using NonNullOperator = operator_functor; - static constexpr auto arity = NonNullOperator::arity; - - template - __device__ inline auto operator()(LHS const lhs, RHS const rhs) const noexcept - -> possibly_null_value_t - { - // Case 1: Neither is null, so the output is given by the operation. - if (lhs.has_value() && rhs.has_value()) { return {NonNullOperator{}(*lhs, *rhs)}; } - // Case 2: Two nulls compare equal. - if (!lhs.has_value() && !rhs.has_value()) { return {true}; } - // Case 3: One value is null, while the other is not, so we return false. - return {false}; - } -}; - -///< NULL_LOGICAL_AND(null, null) is null, NULL_LOGICAL_AND(null, true) is null, -///< NULL_LOGICAL_AND(null, false) is false, and NULL_LOGICAL_AND(valid, valid) == -///< LOGICAL_AND(valid, valid) -template <> -struct operator_functor { - using NonNullOperator = operator_functor; - static constexpr auto arity = NonNullOperator::arity; - - template - __device__ inline auto operator()(LHS const lhs, RHS const rhs) const noexcept - -> possibly_null_value_t - { - // Case 1: Neither is null, so the output is given by the operation. - if (lhs.has_value() && rhs.has_value()) { return {NonNullOperator{}(*lhs, *rhs)}; } - // Case 2: Two nulls return null. - if (!lhs.has_value() && !rhs.has_value()) { return {}; } - // Case 3: One value is null, while the other is not. If it's true we return null, otherwise we - // return false. - auto const& valid_element = lhs.has_value() ? lhs : rhs; - if (*valid_element) { return {}; } - return {false}; - } -}; - -///< NULL_LOGICAL_OR(null, null) is null, NULL_LOGICAL_OR(null, true) is true, NULL_LOGICAL_OR(null, -///< false) is null, and NULL_LOGICAL_OR(valid, valid) == LOGICAL_OR(valid, valid) -template <> -struct operator_functor { - using NonNullOperator = operator_functor; - static constexpr auto arity = NonNullOperator::arity; - - template - __device__ inline auto operator()(LHS const lhs, RHS const rhs) const noexcept - -> possibly_null_value_t - { - // Case 1: Neither is null, so the output is given by the operation. - if (lhs.has_value() && rhs.has_value()) { return {NonNullOperator{}(*lhs, *rhs)}; } - // Case 2: Two nulls return null. - if (!lhs.has_value() && !rhs.has_value()) { return {}; } - // Case 3: One value is null, while the other is not. If it's true we return true, otherwise we - // return null. - auto const& valid_element = lhs.has_value() ? lhs : rhs; - if (*valid_element) { return {true}; } - return {}; + static constexpr int32_t fixed_point_result_scale(int32_t a, int32_t b) + requires(op == ast_operator::ADD || op == ast_operator::SUB || op == ast_operator::MUL || + op == ast_operator::DIV || op == ast_operator::MOD || op == ast_operator::PYMOD) + { + if constexpr (op == ast_operator::ADD || op == ast_operator::SUB) { + return cuda::std::min(a, b); + } else if constexpr (op == ast_operator::MUL) { + return a + b; + } else if constexpr (op == ast_operator::DIV) { + return a - b; + } else if constexpr (op == ast_operator::MOD) { + return cuda::std::min(a, b); + } else if constexpr (op == ast_operator::PYMOD) { + return cuda::std::min(a, b); + } } }; -constexpr bool predicate(possibly_null_value_t value) { return value; } - -constexpr bool predicate(possibly_null_value_t value) -{ - return value.has_value() && *value; -} - } // namespace ast::detail } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/ast/detail/operators.cuh b/cpp/include/cudf/ast/detail/operators.cuh index 292a26a56c26..a4879de95edc 100644 --- a/cpp/include/cudf/ast/detail/operators.cuh +++ b/cpp/include/cudf/ast/detail/operators.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -20,10 +20,10 @@ namespace ast::detail { // Traits for valid operator / type combinations template -constexpr bool is_valid_binary_op = cuda::std::is_invocable_v; +constexpr bool is_valid_binary_op = requires(Op op, LHS lhs, RHS rhs) { op(lhs, rhs); }; template -constexpr bool is_valid_unary_op = cuda::std::is_invocable_v; +constexpr bool is_valid_unary_op = requires(Op op, T value) { op(value); }; /** * @brief Operator dispatcher diff --git a/cpp/include/cudf/detail/operators/ansi_arithmetic.cuh b/cpp/include/cudf/detail/operators/ansi_arithmetic.cuh new file mode 100644 index 000000000000..1187eb4b44bf --- /dev/null +++ b/cpp/include/cudf/detail/operators/ansi_arithmetic.cuh @@ -0,0 +1,309 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Adds operands with overflow detection. + * + * @tparam T Value type. + * @param a Left operand. + * @param b Right operand. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected ansi_add(T a, T b) +{ + T r; + if (cuda::add_overflow(r, a, b).overflow) { return cuda::std::unexpected{errc::OVERFLOW}; } + return r; +} + +template +__device__ cuda::std::expected ansi_add(T a, T b) +{ + return a + b; +} + +/** + * @brief Adds operands with overflow detection. + * + * @tparam R Decimal representation type. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected, errc> ansi_add(numeric::decimal a, + numeric::decimal b) +{ + auto scale = cuda::std::min(a.scale(), b.scale()); + + if (numeric::addition_overflow(a.rescaled(scale).value(), b.rescaled(scale).value())) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + return numeric::decimal{numeric::scaled_integer{ + a.rescaled(scale).value() + b.rescaled(scale).value(), numeric::scale_type{scale}}}; +} + +template +__device__ cuda::std::expected ansi_sub(T a, T b) +{ + T r; + if (cuda::sub_overflow(r, a, b).overflow) { return cuda::std::unexpected{errc::OVERFLOW}; } + return r; +} + +/** + * @brief Subtracts operands with overflow detection. + * + * @tparam T Value type. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected ansi_sub(T a, T b) +{ + return a - b; +} + +template +__device__ cuda::std::expected, errc> ansi_sub(numeric::decimal a, + numeric::decimal b) +{ + auto scale = cuda::std::min(a.scale(), b.scale()); + + if (numeric::subtraction_overflow(a.rescaled(scale).value(), b.rescaled(scale).value())) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + return numeric::decimal{numeric::scaled_integer{ + a.rescaled(scale).value() - b.rescaled(scale).value(), numeric::scale_type{scale}}}; +} + +/** + * @brief Multiplies operands with overflow detection. + * + * @tparam T Value type. + * @param a Left operand. + * @param b Right operand. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected ansi_mul(T a, T b) +{ + T r; + if (cuda::mul_overflow(r, a, b).overflow) { return cuda::std::unexpected{errc::OVERFLOW}; } + return r; +} + +template +__device__ cuda::std::expected ansi_mul(T a, T b) +{ + return a * b; +} + +template +__device__ cuda::std::expected, errc> ansi_mul(numeric::decimal a, + numeric::decimal b) +{ + if (numeric::multiplication_overflow(a.value(), b.value())) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + return numeric::decimal{ + numeric::scaled_integer{a.value() * b.value(), numeric::scale_type{a.scale() + b.scale()}}}; +} + +/** + * @brief Divides operands with ANSI checks. + * + * @tparam T Value type. + * @param a Dividend. + * @param b Divisor. + * @return `errc::DIVISION_BY_ZERO` on zero divisor, `errc::OVERFLOW` on overflow, else + * `errc::SUCCESS`. + */ +template +__device__ cuda::std::expected ansi_div(T a, T b) +{ + if (b == 0) { return cuda::std::unexpected{errc::DIVISION_BY_ZERO}; } + T r; + if (cuda::div_overflow(r, a, b).overflow) { return cuda::std::unexpected{errc::OVERFLOW}; } + return r; +} + +template +__device__ cuda::std::expected ansi_div(T a, T b) +{ + return a / b; +} + +template +__device__ cuda::std::expected, errc> ansi_div(numeric::decimal a, + numeric::decimal b) +{ + if (b.value() == 0) { return cuda::std::unexpected{errc::DIVISION_BY_ZERO}; } + + if (numeric::division_overflow(a.value(), b.value())) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + return numeric::decimal{ + numeric::scaled_integer{a.value() / b.value(), numeric::scale_type{a.scale() - b.scale()}}}; +} + +/** + * @brief Computes modulus with ANSI checks. + * + * @tparam T Value type. + * @param a Dividend. + * @param b Divisor. + * @return `errc::DIVISION_BY_ZERO` on zero divisor, else the result. + */ +template +__device__ cuda::std::expected ansi_mod(T a, T b) +{ + if (b == 0) { return cuda::std::unexpected{errc::DIVISION_BY_ZERO}; } + + // avoid signed overflow UB / trap for minimum value divided by -1. + if (a == cuda::std::numeric_limits::min() && b == T{-1}) { return T{0}; } + + return a % b; +} + +template +__device__ cuda::std::expected ansi_mod(T a, T b) +{ + if (b == 0) { return cuda::std::unexpected{errc::DIVISION_BY_ZERO}; } + return a % b; +} + +template +__device__ cuda::std::expected ansi_mod(T a, T b) +{ + if (b == 0) { return cuda::std::unexpected{errc::DIVISION_BY_ZERO}; } + return a - b * cuda::std::floor(a / b); +} + +template +__device__ cuda::std::expected, errc> ansi_mod(numeric::decimal a, + numeric::decimal b) +{ + auto r = ansi_div(a, b); + if (r.has_error()) { return r.error(); } + return a - b * floor(r.value()); +} + +/** + * @brief Computes absolute value with ANSI overflow checks. + * + * @tparam T Value type. + * @param a Input value. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected ansi_abs(T a) +{ + if (a == cuda::std::numeric_limits::min()) { return cuda::std::unexpected{errc::OVERFLOW}; } + return (a < 0) ? -a : a; +} + +template +__device__ cuda::std::expected ansi_abs(T a) +{ + return a; +} + +template +__device__ cuda::std::expected ansi_abs(T a) +{ + return cuda::std::fabs(a); +} + +template +__device__ cuda::std::expected, errc> ansi_abs(numeric::decimal a) +{ + if (a.value() == cuda::std::numeric_limits::min()) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + auto rep = a.value() < 0 ? -a.value() : a.value(); + return numeric::decimal{numeric::scaled_integer{rep, numeric::scale_type{a.scale()}}}; +} + +/** + * @brief Computes unary negation with ANSI overflow checks. + * + * @tparam T Value type. + * @param a Input value. + * @return `errc::OVERFLOW` on overflow, else the result. + */ +template +__device__ cuda::std::expected ansi_neg(T a) +{ + if (a == cuda::std::numeric_limits::min()) { return cuda::std::unexpected{errc::OVERFLOW}; } + return -a; +} + +template +__device__ cuda::std::expected ansi_neg(T a) +{ + return -a; +} + +template +__device__ cuda::std::expected, errc> ansi_neg(numeric::decimal a) +{ + if (a.value() == cuda::std::numeric_limits::min()) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + auto rep = -a.value(); + return numeric::decimal{numeric::scaled_integer{rep, numeric::scale_type{a.scale()}}}; +} + +/** + * @brief Validates decimal precision against a target precision value. + * + * @tparam R Decimal representation type. + * @param a Input decimal value. + * @param precision Maximum allowed precision. + * @return `errc::OVERFLOW` when precision is invalid or exceeded, else the result. + */ +template +__device__ cuda::std::expected, errc> ansi_precision_check( + numeric::decimal a, int32_t precision) +{ + if (precision <= 0) { return cuda::std::unexpected{errc::OVERFLOW}; } + + auto value = a.value(); + if (value == cuda::std::numeric_limits::min()) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + auto abs_value = value < 0 ? -value : value; + + if (abs_value >= numeric::detail::ipow(precision)) { + return cuda::std::unexpected{errc::OVERFLOW}; + } + + return a; +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/arithmetic.cuh b/cpp/include/cudf/detail/operators/arithmetic.cuh new file mode 100644 index 000000000000..b335f6d7de4c --- /dev/null +++ b/cpp/include/cudf/detail/operators/arithmetic.cuh @@ -0,0 +1,197 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include + +#include +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Computes absolute value. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T abs(T a) + requires(cuda::std::is_signed_v) +{ + return cuda::std::abs(a); +} + +template +__device__ T abs(T a) +{ + return a; +} + +template +__device__ numeric::decimal abs(numeric::decimal a) +{ + auto rep = a.value() < 0 ? -a.value() : a.value(); + return numeric::decimal{numeric::scaled_integer{rep, a.scale()}}; +} + +/** + * @brief Computes sum of two values. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ auto add(A a, B b) -> decltype(a + b) +{ + return a + b; +} + +/** + * @brief Computes quotient of two values. + * + * @tparam A Dividend type. + * @tparam B Divisor type. + * @param a Dividend. + * @param b Divisor. + */ +template +__device__ auto div(A a, B b) -> decltype(a / b) +{ + return a / b; +} + +/** + * @brief Computes floor division of two values. + * + * @tparam A Dividend type. + * @tparam B Divisor type. + * @param a Dividend. + * @param b Divisor. + */ +template +__device__ auto floor_div(A a, B b) -> decltype(cudf::detail::integral_floor_div(a, b)) +{ + return cudf::detail::integral_floor_div(a, b); +} + +template +__device__ auto floor_div(A a, B b) -> decltype(cuda::std::floor(a / b)) +{ + return cuda::std::floor(a / b); +} + +/** + * @brief Computes remainder of two values. + * + * @tparam A Dividend type. + * @tparam B Divisor type. + * @param a Dividend. + * @param b Divisor. + */ +template +__device__ auto mod(A a, B b) -> decltype(a % b) +{ + return a % b; +} + +template +__device__ auto mod(A a, B b) -> decltype(cuda::std::fmod(a, b)) +{ + return cuda::std::fmod(a, b); +} + +/** + * @brief Computes Python-style modulus. + * + * @tparam A Dividend type. + * @tparam B Divisor type. + * @param a Dividend. + * @param b Divisor. + */ +template +__device__ auto pymod(A a, B b) -> decltype((a % b + b) % b) +{ + return (a % b + b) % b; +} + +template +__device__ auto pymod(A a, B b) -> decltype(cuda::std::fmod(cuda::std::fmod(a, b) + b, b)) +{ + return cuda::std::fmod(cuda::std::fmod(a, b) + b, b); +} + +/** + * @brief Computes product of two values. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ auto mul(A a, B b) -> decltype(a * b) +{ + return a * b; +} + +/** + * @brief Computes unary negation. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ auto neg(T a) -> decltype(-a) +{ + return -a; +} + +template +__device__ numeric::decimal neg(numeric::decimal a) +{ + auto rep = -a.value(); + return numeric::decimal{numeric::scaled_integer{rep, a.scale()}}; +} + +/** + * @brief Computes subtraction of two values. + * + * @tparam A Minuend type. + * @tparam B Subtrahend type. + * @param a Minuend. + * @param b Subtrahend. + */ +template +__device__ auto sub(A a, B b) -> decltype(a - b) +{ + return a - b; +} + +/** + * @brief Computes true division and returns a double. + * + * @tparam A Dividend type. + * @tparam B Divisor type. + * @param a Dividend. + * @param b Divisor. + */ +template +__device__ auto true_div(A a, B b) -> decltype(static_cast(a) / static_cast(b)) +{ + return static_cast(a) / static_cast(b); +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/bitwise.cuh b/cpp/include/cudf/detail/operators/bitwise.cuh new file mode 100644 index 000000000000..957c0feb2fdc --- /dev/null +++ b/cpp/include/cudf/detail/operators/bitwise.cuh @@ -0,0 +1,98 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Computes bitwise AND of two values. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ auto bit_and(A a, B b) -> decltype(a & b) +{ + return a & b; +} + +/** + * @brief Computes bitwise NOT of one value. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ auto bit_invert(T a) -> decltype(~a) +{ + return ~a; +} + +/** + * @brief Computes bitwise OR of two values. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ auto bit_or(A a, B b) -> decltype(a | b) +{ + return a | b; +} + +/** + * @brief Computes bitwise XOR of two values. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ auto bit_xor(A a, B b) -> decltype(a ^ b) +{ + return a ^ b; +} + +/** + * @brief Shifts a value left by a bit count. + * + * @tparam A Value type. + * @tparam B Shift count type. + * @param a Input value. + * @param b Shift count. + */ +template +__device__ auto bit_shift_left(A a, B b) -> decltype(a << b) +{ + return a << b; +} + +/** + * @brief Shifts a value right by a bit count. + * + * @tparam A Value type. + * @tparam B Shift count type. + * @param a Input value. + * @param b Shift count. + */ +template +__device__ auto bit_shift_right(A a, B b) -> decltype(a >> b) +{ + return a >> b; +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/casts.cuh b/cpp/include/cudf/detail/operators/casts.cuh new file mode 100644 index 000000000000..e6f24fa427c6 --- /dev/null +++ b/cpp/include/cudf/detail/operators/casts.cuh @@ -0,0 +1,242 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Casts input values to bool. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ bool cast_to_b8(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to int8_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ int8_t cast_to_i8(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to int16_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ int16_t cast_to_i16(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to int32_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ int32_t cast_to_i32(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to int64_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ int64_t cast_to_i64(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to uint8_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ uint8_t cast_to_u8(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to uint16_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ uint16_t cast_to_u16(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to uint32_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ uint32_t cast_to_u32(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to uint64_t. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ uint64_t cast_to_u64(T a) + requires(!nullable && cuda::std::convertible_to) +{ + return static_cast(a); +} + +/** + * @brief Casts input values to float. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ float cast_to_f32(T a) + requires(cuda::std::is_integral_v || cuda::std::is_floating_point_v) +{ + return static_cast(a); +} + +template +__device__ float cast_to_f32(numeric::decimal a) +{ + return convert_fixed_to_floating(a); +} + +/** + * @brief Casts input values to double. + * + * @tparam T Source type. + * @param a Input value. + */ +template +__device__ double cast_to_f64(T a) + requires(cuda::std::is_integral_v || floating_point || fixed_point) +{ + return static_cast(a); +} + +template +__device__ double cast_to_f64(numeric::decimal a) +{ + return convert_fixed_to_floating(a); +} + +namespace detail { + +/** + * @brief Casts one fixed-point decimal representation to another. + * + * @tparam To Destination representation type. + * @tparam From Source representation type. + * @param a Input value. + */ +template +__device__ numeric::decimal decimal_cast(numeric::decimal a) +{ + auto rep = static_cast(a.value()); + return numeric::decimal{numeric::scaled_integer{rep, a.scale()}}; +} + +} // namespace detail + +/** + * @brief Casts decimal input values to decimal32. + * + * @tparam R Source decimal representation type. + * @param a Input value. + */ +template +__device__ numeric::decimal32 cast_to_dec32(numeric::decimal a) +{ + return detail::decimal_cast(a); +} + +/** + * @brief Casts decimal input values to decimal64. + * + * @tparam R Source decimal representation type. + * @param a Input value. + */ +template +__device__ numeric::decimal64 cast_to_dec64(numeric::decimal a) +{ + return detail::decimal_cast(a); +} + +/** + * @brief Casts decimal input values to decimal128. + * + * @tparam R Source decimal representation type. + * @param a Input value. + */ +template +__device__ numeric::decimal128 cast_to_dec128(numeric::decimal a) +{ + return detail::decimal_cast(a); +} + +/** + * @brief Rescales decimal input values to a target scale. + * + * @tparam R Decimal representation type. + * @param a Input value. + * @param new_scale Target decimal scale. + */ +template +__device__ numeric::decimal rescale(numeric::decimal a, int32_t new_scale) +{ + return a.rescaled(numeric::scale_type{new_scale}); +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/comparison.cuh b/cpp/include/cudf/detail/operators/comparison.cuh new file mode 100644 index 000000000000..362f17039470 --- /dev/null +++ b/cpp/include/cudf/detail/operators/comparison.cuh @@ -0,0 +1,130 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Tests `a == b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool equal(A a, B b) + requires(!nullable && !nullable && requires { a == b; }) +{ + return a == b; +} + +/** + * @brief Tests `a != b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool not_equal(A a, B b) + requires(!nullable && !nullable && requires { a != b; }) +{ + return a != b; +} + +/** + * @brief Tests `a > b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool greater(A a, B b) + requires(!nullable && !nullable && requires { a > b; }) +{ + return a > b; +} + +/** + * @brief Tests `a >= b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool greater_equal(A a, B b) + requires(!nullable && !nullable && requires { a >= b; }) +{ + return a >= b; +} + +/** + * @brief Tests `a < b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool less(A a, B b) + requires(!nullable && !nullable && requires { a < b; }) +{ + return a < b; +} + +/** + * @brief Tests `a <= b`. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool less_equal(A a, B b) + requires(!nullable && !nullable && requires { a <= b; }) +{ + return a <= b; +} + +/** + * @brief Tests equality between two values for null-aware equality semantics. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool null_equal(A a, B b) + requires(!nullable && !nullable && requires { a == b; }) +{ + return a == b; +} + +template +__device__ bool null_equal(cuda::std::optional a, cuda::std::optional b) + requires(!nullable && !nullable && requires { a == b; }) +{ + return a == b; +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/concepts.cuh b/cpp/include/cudf/detail/operators/concepts.cuh new file mode 100644 index 000000000000..33b56d24067b --- /dev/null +++ b/cpp/include/cudf/detail/operators/concepts.cuh @@ -0,0 +1,45 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +template +concept integer = + cuda::std::is_integral_v && !cuda::std::is_same_v, bool>; + +template +concept signed_integer = integer && cuda::std::is_signed_v>; + +template +concept unsigned_integer = integer && cuda::std::is_unsigned_v>; + +template +concept fixed_point = cudf::is_fixed_point(); + +template +concept floating_point = cuda::std::is_floating_point_v; + +template +constexpr bool is_nullable = false; + +template +constexpr bool is_nullable> = true; + +template +concept nullable = is_nullable; + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/error.hpp b/cpp/include/cudf/detail/operators/error.hpp new file mode 100644 index 000000000000..d43a8f6302c9 --- /dev/null +++ b/cpp/include/cudf/detail/operators/error.hpp @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief An enumeration of error codes that can occur during operations. + */ +enum class errc : cuda::std::int8_t { SUCCESS = 0, OVERFLOW = 1, DIVISION_BY_ZERO = 2 }; + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/identity.cuh b/cpp/include/cudf/detail/operators/identity.cuh new file mode 100644 index 000000000000..b14c98b8bea9 --- /dev/null +++ b/cpp/include/cudf/detail/operators/identity.cuh @@ -0,0 +1,27 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Copies an input value to the output. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T identity(T a) +{ + return a; +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/logic.cuh b/cpp/include/cudf/detail/operators/logic.cuh new file mode 100644 index 000000000000..e9fa3890c55d --- /dev/null +++ b/cpp/include/cudf/detail/operators/logic.cuh @@ -0,0 +1,142 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Computes logical AND. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool logical_and(A a, B b) +{ + return a && b; +} + +/** + * @brief Computes logical OR. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool logical_or(A a, B b) +{ + return a || b; +} + +/** + * @brief Computes logical NOT. + * + * @tparam T Value type. + * @param a Input operand. + */ +template +__device__ bool logical_not(T a) +{ + return !a; +} + +/** + * @brief Computes logical AND with null-aware semantics. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool null_logical_and(A a, B b) +{ + return logical_and(a, b); +} + +template +__device__ cuda::std::optional null_logical_and(cuda::std::optional a, + cuda::std::optional b) +{ + if (a.has_value() && b.has_value()) { + return null_logical_and(a.value(), b.value()); + } else if (!a.has_value() && !b.has_value()) { + return {}; + } else { + if (a.has_value() ? *a : *b) { + return {}; + } else { + return false; + } + } +} + +/** + * @brief Computes logical OR with null-aware semantics. + * + * @tparam A Left operand type. + * @tparam B Right operand type. + * @param a Left operand. + * @param b Right operand. + */ +template +__device__ bool null_logical_or(A a, B b) +{ + return logical_or(a, b); +} + +template +__device__ cuda::std::optional null_logical_or(cuda::std::optional a, + cuda::std::optional b) +{ + if (a.has_value() && b.has_value()) { + return null_logical_or(a.value(), b.value()); + } else if (!a.has_value() && !b.has_value()) { + return {}; + } else { + if (a.has_value() ? *a : *b) { + return true; + } else { + return {}; + } + } +} + +/** + * @brief Selects one of two values based on a predicate. + * + * @tparam T Selected value type. + * @param true_value Value selected when @p pred is true. + * @param false_value Value selected when @p pred is false. + * @param pred Selection predicate. + */ +template +__device__ T if_else(T true_value, T false_value, bool pred) +{ + return pred ? true_value : false_value; +} + +template +__device__ cuda::std::optional if_else(cuda::std::optional true_value, + cuda::std::optional false_value, + cuda::std::optional pred) +{ + return pred.value_or(false) ? true_value : false_value; +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/math.cuh b/cpp/include/cudf/detail/operators/math.cuh new file mode 100644 index 000000000000..e33c6b022499 --- /dev/null +++ b/cpp/include/cudf/detail/operators/math.cuh @@ -0,0 +1,165 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include +#include + +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Computes cube root + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T cbrt(T a) +{ + return cuda::std::cbrt(a); +} + +/** + * @brief Computes ceiling. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T ceil(T a) +{ + return cuda::std::ceil(a); +} + +namespace detail { + +/** + * @brief Rounds a decimal value to an integral value. + * + * @tparam R Rep type of the decimal. + * @tparam ceil If true `ceil`s the value, otherwise `floor`s the value. + */ +template +__device__ numeric::decimal decimal_round(numeric::decimal a) +{ + if (a.scale() >= 0) { + return a; + } else { + auto factor = + numeric::detail::ipow(-static_cast(a.scale())); + auto div = a.value() / factor; + auto rem = a.value() % factor; + if (rem == 0) { + return a; + } else { + auto val = ceil ? (a.value() > 0 ? (div + 1) : div) : (a.value() > 0 ? div : (div - 1)); + return numeric::decimal{numeric::scaled_integer{val * factor, a.scale()}}; + } + } +} + +} // namespace detail + +template +__device__ numeric::decimal ceil(numeric::decimal a) +{ + return detail::decimal_round(a); +} + +/** + * @brief Computes natural exponential. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T exp(T a) +{ + return cuda::std::exp(a); +} + +/** + * @brief Computes floor of a value. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T floor(T a) +{ + return cuda::std::floor(a); +} + +template +__device__ numeric::decimal floor(numeric::decimal a) +{ + return detail::decimal_round(a); +} + +/** + * @brief Computes natural logarithm. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T log(T a) +{ + return cuda::std::log(a); +} + +/** + * @brief Computes exponentiation. + * + * @tparam A Base value type. + * @tparam B Exponent value type. + * @param a Base value. + * @param b Exponent value. + */ +template +__device__ auto pow(A a, B b) -> decltype(cuda::std::pow(a, b)) +{ + return cuda::std::pow(a, b); +} + +template +__device__ auto pow(A a, B b) -> decltype(cudf::detail::integral_pow(a, b)) +{ + return cudf::detail::integral_pow(a, b); +} + +/** + * @brief Rounds to integral value. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ T rint(T a) +{ + return cuda::std::rint(a); +} + +/** + * @brief Computes square root. + * + * @param a Input value. + */ +template +__device__ T sqrt(T a) +{ + return cuda::std::sqrt(a); +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/null_handling.cuh b/cpp/include/cudf/detail/operators/null_handling.cuh new file mode 100644 index 000000000000..e26331b72e1e --- /dev/null +++ b/cpp/include/cudf/detail/operators/null_handling.cuh @@ -0,0 +1,104 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include + +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Tests whether an input value is null. + * + * @tparam T Value type. + * @param a Input value. + */ +template +__device__ bool is_null(T a) + requires(!nullable) +{ + return false; +} + +template +__device__ bool is_null(T a) + requires(nullable) +{ + return !a.has_value(); +} + +/** + * @brief Sets the output to null when the condition is true. + * + * @tparam T Value type. + * @param a Input value. + * @param condition boolean condition. + */ +template +__device__ cuda::std::optional nullify_if(cuda::std::optional a, + cuda::std::optional condition) +{ + if (condition.has_value() && a.has_value()) { + if (condition.value()) { + return {}; + } else { + return a.value(); + } + } else { + return {}; + } +} + +/** + * @brief Returns the first non-null of two values. + * + * @tparam T Value type. + * @param a First value. + * @param b Second value. + */ +template +__device__ T coalesce(T a, T b) + requires(!nullable) +{ + return a; +} + +template +__device__ cuda::std::optional coalesce(cuda::std::optional a, cuda::std::optional b) +{ + if (a.has_value()) { + return a.value(); + } else if (b.has_value()) { + return b.value(); + } else { + return {}; + } +} + +/** + * @brief Converts an optional predicate to a non-nullable predicate. + * + * @param a Input boolean predicate. + */ +template T> +__device__ inline bool predicate(T a) +{ + return a; +} + +template T> +__device__ inline bool predicate(cuda::std::optional a) +{ + return a.value_or(false); +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/operators/operators.cuh b/cpp/include/cudf/detail/operators/operators.cuh new file mode 100644 index 000000000000..b52055d0a6b4 --- /dev/null +++ b/cpp/include/cudf/detail/operators/operators.cuh @@ -0,0 +1,18 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include diff --git a/cpp/include/cudf/detail/operators/trigonometric.cuh b/cpp/include/cudf/detail/operators/trigonometric.cuh new file mode 100644 index 000000000000..1d3274889227 --- /dev/null +++ b/cpp/include/cudf/detail/operators/trigonometric.cuh @@ -0,0 +1,163 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include +#include + +namespace CUDF_EXPORT cudf { +namespace detail { +namespace ops { + +/** + * @brief Computes inverse cosine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arccos(T a) +{ + return cuda::std::acos(a); +} + +/** + * @brief Computes inverse hyperbolic cosine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arccosh(T a) +{ + return cuda::std::acosh(a); +} + +/** + * @brief Computes inverse sine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arcsin(T a) +{ + return cuda::std::asin(a); +} + +/** + * @brief Computes inverse hyperbolic sine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arcsinh(T a) +{ + return cuda::std::asinh(a); +} + +/** + * @brief Computes inverse tangent. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arctan(T a) +{ + return cuda::std::atan(a); +} + +/** + * @brief Computes inverse hyperbolic tangent. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T arctanh(T a) +{ + return cuda::std::atanh(a); +} + +/** + * @brief Computes cosine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T cos(T a) +{ + return cuda::std::cos(a); +} + +/** + * @brief Computes hyperbolic cosine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T cosh(T a) +{ + return cuda::std::cosh(a); +} + +/** + * @brief Computes sine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T sin(T a) +{ + return cuda::std::sin(a); +} + +/** + * @brief Computes hyperbolic sine. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T sinh(T a) +{ + return cuda::std::sinh(a); +} + +/** + * @brief Computes tangent. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T tan(T a) +{ + return cuda::std::tan(a); +} + +/** + * @brief Computes hyperbolic tangent. + * + * @tparam T Value type + * @param a Input value. + */ +template +__device__ T tanh(T a) +{ + return cuda::std::tanh(a); +} + +} // namespace ops +} // namespace detail +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index da2632ea4095..987e8492c3ba 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -784,6 +784,9 @@ CUDF_HOST_DEVICE inline fixed_point operator%(fixed_point{scaled_integer{remainder, scale}}; } +template +using decimal = + fixed_point; ///< decimal fixed point with user defined representation type using decimal32 = fixed_point; ///< 32-bit decimal fixed point using decimal64 = fixed_point; ///< 64-bit decimal fixed point using decimal128 = fixed_point<__int128_t, Radix::BASE_10>; ///< 128-bit decimal fixed point diff --git a/cpp/src/ast/operators.cpp b/cpp/src/ast/operators.cpp index eaacabde0b74..35ea746dec9c 100644 --- a/cpp/src/ast/operators.cpp +++ b/cpp/src/ast/operators.cpp @@ -21,8 +21,7 @@ struct arity_functor { template void operator()(cudf::size_type& result) { - // Arity is not dependent on null handling, so just use the false implementation here. - result = operator_functor::arity; + result = operator_functor::arity; } }; @@ -156,7 +155,7 @@ struct type_dispatch_binary_op { type_dispatcher( lhs_type, // Always dispatch to the non-null operator for the purpose of type determination. - detail::single_dispatch_binary_operator_types>{}, + detail::single_dispatch_binary_operator_types>{}, std::forward(f), std::forward(args)...); } @@ -225,7 +224,7 @@ struct type_dispatch_unary_op { type_dispatcher( input_type, // Always dispatch to the non-null operator for the purpose of type determination. - detail::dispatch_unary_operator_types>{}, + detail::dispatch_unary_operator_types>{}, std::forward(f), std::forward(args)...); } diff --git a/cpp/src/jit/row_ir.cpp b/cpp/src/jit/row_ir.cpp index f2564e57028d..25861c7c6578 100644 --- a/cpp/src/jit/row_ir.cpp +++ b/cpp/src/jit/row_ir.cpp @@ -370,18 +370,17 @@ void node::emit_code(instance_context& instance, target_info const& info, code_s if (op_ == opcode::PREDICATE) { sink.emit(std::format( - R"***(bool {} = cudf::ast::detail::predicate({}); + R"***(bool {} = cudf::detail::ops::predicate({}); )***", id_, args_str)); } else { sink.emit(std::format( - R"***({} {} = cudf::ast::detail::operator_functor{{}}({}); + R"***({} {} = cudf::ast::detail::operator_functor{{}}({}); )***", type, id_, ast::detail::ast_operator_string(as_ast_op(op_)), - instance.has_nulls(), args_str)); } } break; diff --git a/cpp/tests/jit/row_ir.cpp b/cpp/tests/jit/row_ir.cpp index b1e5d1cd0a3b..a684077a2be2 100644 --- a/cpp/tests/jit/row_ir.cpp +++ b/cpp/tests/jit/row_ir.cpp @@ -136,7 +136,7 @@ TEST_F(RowIRCudaCodeGenTest, UnaryOperation) auto expected_code = R"***(int32_t tmp_0 = in_0; -int32_t tmp_1 = cudf::ast::detail::operator_functor{}(tmp_0); +int32_t tmp_1 = cudf::ast::detail::operator_functor{}(tmp_0); )***"; EXPECT_EQ(sink.get_code(), expected_code); @@ -155,7 +155,7 @@ int32_t tmp_1 = cudf::ast::detail::operator_functor{}(tmp_0); +numeric::decimal32 tmp_1 = cudf::ast::detail::operator_functor{}(tmp_0); )***"; EXPECT_EQ(sink.get_code(), expected_null_code); @@ -182,7 +182,7 @@ TEST_F(RowIRCudaCodeGenTest, BinaryOperation) auto expected_code = R"***(int32_t tmp_0 = in_0; int32_t tmp_1 = in_0; -int32_t tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); +int32_t tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); )***"; EXPECT_EQ(sink.get_code(), expected_code); @@ -204,7 +204,7 @@ int32_t tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); +numeric::decimal32 tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); )***"; EXPECT_EQ(sink.get_code(), expected_null_code); @@ -251,12 +251,12 @@ TEST_F(RowIRCudaCodeGenTest, VectorLengthOperation) auto expected_code = R"***(double tmp_0 = in_0; double tmp_1 = in_0; -double tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); +double tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); double tmp_3 = in_1; double tmp_4 = in_1; -double tmp_5 = cudf::ast::detail::operator_functor{}(tmp_3, tmp_4); -double tmp_6 = cudf::ast::detail::operator_functor{}(tmp_2, tmp_5); -double tmp_7 = cudf::ast::detail::operator_functor{}(tmp_6); +double tmp_5 = cudf::ast::detail::operator_functor{}(tmp_3, tmp_4); +double tmp_6 = cudf::ast::detail::operator_functor{}(tmp_2, tmp_5); +double tmp_7 = cudf::ast::detail::operator_functor{}(tmp_6); double tmp_8 = tmp_7; *out_0 = tmp_8; )***"; @@ -317,7 +317,7 @@ TEST_F(RowIRCudaCodeGenTest, AstConversionBasic) { int32_t tmp_0 = in_0; int32_t tmp_1 = in_1; -int32_t tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); +int32_t tmp_2 = cudf::ast::detail::operator_functor{}(tmp_0, tmp_1); int32_t tmp_3 = tmp_2; *out_0 = tmp_3; return; @@ -352,7 +352,7 @@ TEST_F(RowIRCudaCodeGenTest, FilterPredicate) filter_predicate.emit_code(ctx, target_info, sink); auto expected_code = R"***(bool tmp_0 = in_0; -bool tmp_1 = cudf::ast::detail::predicate(tmp_0); +bool tmp_1 = cudf::detail::ops::predicate(tmp_0); )***"; EXPECT_EQ(sink.get_code(), expected_code); @@ -370,7 +370,7 @@ bool tmp_1 = cudf::ast::detail::predicate(tmp_0); filter_predicate.emit_code(ctx, target_info, sink); auto expected_code = R"***(cuda::std::optional tmp_0 = in_0; -bool tmp_1 = cudf::ast::detail::predicate(tmp_0); +bool tmp_1 = cudf::detail::ops::predicate(tmp_0); )***"; EXPECT_EQ(sink.get_code(), expected_code);