diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 664b13dd971c..cd118a5c56b5 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1191,6 +1191,7 @@ add_library( src/text/vocabulary_tokenize.cu src/text/wordpiece_tokenize.cu src/transform/bools_to_mask.cu + src/transform/checked_arithmetic.cu src/transform/compute_column.cu src/transform/compute_column_kernel_complex.cu src/transform/compute_column_kernel_null_complex.cu diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 8a49d6f15c71..5d1c6b2d01f7 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -86,7 +86,12 @@ enum class binary_operator : int32_t { ///< operands are true, returns true; otherwise returns null NULL_LOGICAL_OR, ///< three-valued (Kleene) ||: if any operand is true, returns true; if both ///< operands are false, returns false; otherwise returns null - INVALID_BINARY ///< invalid operation + INVALID_BINARY, ///< invalid operation + ADD_OVERFLOW, ///< Addition with overflow detection + SUB_OVERFLOW, ///< Subtraction with overflow detection + MUL_OVERFLOW, ///< Multiplication with overflow detection + DIV_OVERFLOW, ///< Division with overflow and divide-by-zero detection + MOD_OVERFLOW ///< Modulo with divide-by-zero detection }; /// Binary operation common type default @@ -239,6 +244,91 @@ std::unique_ptr binary_operation( cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); +/** + * @brief Performs a checked binary operation between a scalar and a column. + * + * `PROPAGATE` throws `cudf::evaluation_error` when any row fails; `NULLIFY` makes failing rows + * null. + * + * @param lhs Left operand scalar + * @param rhs Right operand column + * @param op Checked binary operator + * @param output_type Desired output type + * @param policy Error handling policy + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column + * @return Output column + * @throws cudf::evaluation_error if @p policy is `error_policy::PROPAGATE` and any row fails + * @throws cudf::logic_error if @p op is not a checked arithmetic operator + * @throws cudf::data_type_error if the input and output types do not match, are not supported + * arithmetic or fixed-point types, or @p output_type has an invalid fixed-point scale + */ +std::unique_ptr binary_operation( + scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Performs a checked binary operation between a column and a scalar. + * + * `PROPAGATE` throws `cudf::evaluation_error` when any row fails; `NULLIFY` makes failing rows + * null. + * + * @param lhs Left operand column + * @param rhs Right operand scalar + * @param op Checked binary operator + * @param output_type Desired output type + * @param policy Error handling policy + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column + * @return Output column + * @throws cudf::evaluation_error if @p policy is `error_policy::PROPAGATE` and any row fails + * @throws cudf::logic_error if @p op is not a checked arithmetic operator + * @throws cudf::data_type_error if the input and output types do not match, are not supported + * arithmetic or fixed-point types, or @p output_type has an invalid fixed-point scale + */ +std::unique_ptr binary_operation( + column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Performs a checked binary operation between two columns. + * + * `PROPAGATE` throws `cudf::evaluation_error` when any row fails; `NULLIFY` makes failing rows + * null. + * + * @param lhs Left operand column + * @param rhs Right operand column + * @param op Checked binary operator + * @param output_type Desired output type + * @param policy Error handling policy + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column + * @return Output column + * @throws cudf::evaluation_error if @p policy is `error_policy::PROPAGATE` and any row fails + * @throws cudf::logic_error if @p op is not a checked arithmetic operator + * @throws cudf::data_type_error if the input and output types do not match, are not supported + * arithmetic or fixed-point types, or @p output_type has an invalid fixed-point scale + * @throws std::invalid_argument if @p lhs and @p rhs have different sizes + */ +std::unique_ptr binary_operation( + column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + /** * @brief Performs a binary operation between two columns using a * user-defined PTX function. diff --git a/cpp/include/cudf/unary.hpp b/cpp/include/cudf/unary.hpp index f0bcc8d343ec..1f4a130908b0 100644 --- a/cpp/include/cudf/unary.hpp +++ b/cpp/include/cudf/unary.hpp @@ -30,36 +30,40 @@ namespace CUDF_EXPORT cudf { * @brief Types of unary operations that can be performed on data. */ enum class unary_operator : int32_t { - SIN, ///< Trigonometric sine - COS, ///< Trigonometric cosine - TAN, ///< Trigonometric tangent - ARCSIN, ///< Trigonometric sine inverse - ARCCOS, ///< Trigonometric cosine inverse - ARCTAN, ///< Trigonometric tangent inverse - SINH, ///< Hyperbolic sine - COSH, ///< Hyperbolic cosine - TANH, ///< Hyperbolic tangent - ARCSINH, ///< Hyperbolic sine inverse - ARCCOSH, ///< Hyperbolic cosine inverse - ARCTANH, ///< Hyperbolic tangent inverse - EXP, ///< Exponential (base e, Euler number) - LOG, ///< Natural Logarithm (base e) - SQRT, ///< Square-root (x^0.5) - CBRT, ///< Cube-root (x^(1.0/3)) - CEIL, ///< Smallest integer value not less than arg - FLOOR, ///< largest integer value not greater than arg - ABS, ///< Absolute value - RINT, ///< Rounds the floating-point argument arg to an integer value - BIT_COUNT, ///< Count the number of bits set to 1 of an integer value - BIT_INVERT, ///< Bitwise Not (~) - NOT, ///< Logical Not (!) - NEGATE, ///< Unary negation (-), only for signed numeric and duration types. + SIN, ///< Trigonometric sine + COS, ///< Trigonometric cosine + TAN, ///< Trigonometric tangent + ARCSIN, ///< Trigonometric sine inverse + ARCCOS, ///< Trigonometric cosine inverse + ARCTAN, ///< Trigonometric tangent inverse + SINH, ///< Hyperbolic sine + COSH, ///< Hyperbolic cosine + TANH, ///< Hyperbolic tangent + ARCSINH, ///< Hyperbolic sine inverse + ARCCOSH, ///< Hyperbolic cosine inverse + ARCTANH, ///< Hyperbolic tangent inverse + EXP, ///< Exponential (base e, Euler number) + LOG, ///< Natural Logarithm (base e) + SQRT, ///< Square-root (x^0.5) + CBRT, ///< Cube-root (x^(1.0/3)) + CEIL, ///< Smallest integer value not less than arg + FLOOR, ///< largest integer value not greater than arg + ABS, ///< Absolute value + RINT, ///< Rounds the floating-point argument arg to an integer value + BIT_COUNT, ///< Count the number of bits set to 1 of an integer value + BIT_INVERT, ///< Bitwise Not (~) + NOT, ///< Logical Not (!) + NEGATE, ///< Unary negation (-), only for signed numeric and duration types. + NEG_OVERFLOW, ///< Negation with overflow detection + ABS_OVERFLOW ///< Absolute value with overflow detection }; /** * @brief Performs unary op on all values in column * * Note: For `decimal32` and `decimal64`, only `ABS`, `CEIL` and `FLOOR` are supported. + * Checked operators use error_policy::PROPAGATE; use the policy-aware overload to nullify + * errors. * * @param input A `column_view` as input * @param op operation to perform @@ -74,6 +78,29 @@ std::unique_ptr unary_operation( cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); +/** + * @brief Performs a checked unary operation on all values in a column. + * + * `PROPAGATE` throws `cudf::evaluation_error` when any row fails; `NULLIFY` makes failing rows + * null. + * + * @param input Input column + * @param op Checked unary operator + * @param policy Error handling policy + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column + * @return Output column + * @throws cudf::evaluation_error if @p policy is `error_policy::PROPAGATE` and any row fails + * @throws cudf::logic_error if @p op is not a checked arithmetic operator + * @throws cudf::data_type_error if @p input is not a supported arithmetic or fixed-point type + */ +std::unique_ptr unary_operation( + cudf::column_view const& input, + cudf::unary_operator op, + cudf::error_policy policy, + cuda::stream_ref stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + /** * @brief Creates a column of `type_id::BOOL8` elements where for every element in `input` `true` * indicates the value is null and `false` indicates the value is valid. diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index fd972555a642..66e086b1475b 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -32,6 +32,7 @@ #include #include #include +#include #include #include #include @@ -42,6 +43,8 @@ #include #include +#include + #include namespace cudf { @@ -49,6 +52,13 @@ namespace binops { bool is_supported_operation(data_type out, data_type lhs, data_type rhs, binary_operator op) { + if (cudf::detail::checked_arithmetic::is_checked(op)) { + if (out.id() != lhs.id() || lhs.id() != rhs.id()) { return false; } + if (is_fixed_point(lhs)) { + return out.scale() == binary_operation_fixed_point_scale(op, lhs.scale(), rhs.scale()); + } + return is_numeric(lhs) && lhs.id() != type_id::BOOL8; + } return cudf::binops::compiled::is_supported_operation(out, lhs, rhs, op); } @@ -89,15 +99,20 @@ inline bool is_null_dependent(binary_operator op) */ bool is_basic_arithmetic_binop(binary_operator op) { - return op == binary_operator::ADD or // operator + - op == binary_operator::SUB or // operator - - op == binary_operator::MUL or // operator * - op == binary_operator::DIV or // operator / using common type of lhs and rhs - op == binary_operator::NULL_MIN or // 2 null = null, 1 null = value, else min - op == binary_operator::NULL_MAX or // 2 null = null, 1 null = value, else max - op == binary_operator::MOD or // operator % - op == binary_operator::PMOD or // positive modulo operator - op == binary_operator::PYMOD; // operator % but following Python's negative sign rules + return op == binary_operator::ADD or // operator + + op == binary_operator::SUB or // operator - + op == binary_operator::MUL or // operator * + op == binary_operator::DIV or // operator / using common type of lhs and rhs + op == binary_operator::NULL_MIN or // 2 null = null, 1 null = value, else min + op == binary_operator::NULL_MAX or // 2 null = null, 1 null = value, else max + op == binary_operator::MOD or // operator % + op == binary_operator::PMOD or // positive modulo operator + op == binary_operator::PYMOD || // Python modulo + op == binary_operator::ADD_OVERFLOW || // checked addition + op == binary_operator::SUB_OVERFLOW || // checked subtraction + op == binary_operator::MUL_OVERFLOW || // checked multiplication + op == binary_operator::DIV_OVERFLOW || // checked division + op == binary_operator::MOD_OVERFLOW; // checked modulo } /** @@ -132,7 +147,8 @@ bool is_supported_fixed_point_binop(binary_operator op) */ bool is_same_scale_necessary(binary_operator op) { - return op != binary_operator::MUL && op != binary_operator::DIV; + return op != binary_operator::MUL && op != binary_operator::DIV && + op != binary_operator::MUL_OVERFLOW && op != binary_operator::DIV_OVERFLOW; } namespace jit { @@ -216,6 +232,11 @@ std::unique_ptr binary_operation(LhsType const& lhs, if constexpr (std::is_same_v and std::is_same_v) CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes don't match", std::invalid_argument); + if (cudf::detail::checked_arithmetic::is_checked(op)) { + return cudf::detail::checked_arithmetic::binary_operation( + lhs, rhs, op, output_type, error_policy::PROPAGATE, stream, mr); + } + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING and output_type.id() == type_id::STRING and (op == binary_operator::NULL_MAX or op == binary_operator::NULL_MIN)) @@ -407,8 +428,10 @@ int32_t binary_operation_fixed_point_scale(binary_operator op, { CUDF_EXPECTS(binops::is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation."); - if (op == binary_operator::MUL) return left_scale + right_scale; - if (op == binary_operator::DIV) return left_scale - right_scale; + if (op == binary_operator::MUL || op == binary_operator::MUL_OVERFLOW) + return left_scale + right_scale; + if (op == binary_operator::DIV || op == binary_operator::DIV_OVERFLOW) + return left_scale - right_scale; return std::min(left_scale, right_scale); } @@ -453,6 +476,45 @@ std::unique_ptr binary_operation(column_view const& lhs, return detail::binary_operation(lhs, rhs, op, output_type, stream, mr); } +std::unique_ptr binary_operation(scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::checked_arithmetic::binary_operation( + lhs, rhs, op, output_type, policy, stream, mr); +} + +std::unique_ptr binary_operation(column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::checked_arithmetic::binary_operation( + lhs, rhs, op, output_type, policy, stream, mr); +} + +std::unique_ptr binary_operation(column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::checked_arithmetic::binary_operation( + lhs, rhs, op, output_type, policy, stream, mr); +} + std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, std::string const& ptx, diff --git a/cpp/src/jit/transform_kernel.cuh b/cpp/src/jit/transform_kernel.cuh new file mode 100644 index 000000000000..6b1d2187f45b --- /dev/null +++ b/cpp/src/jit/transform_kernel.cuh @@ -0,0 +1,105 @@ +/* + * 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 +#include + +namespace cudf::detail { + +/** + * @brief Applies a row operation using transform input and output accessors. + * + * `operation` is invoked as `operation(row, arguments)`, where `arguments` is a tuple containing + * output pointers followed by input values. + */ +template +__device__ void transform_kernel(size_type row_size, + bitmask_type const* __restrict__ stencil, + column_device_view_core const* __restrict__ input_cols, + mutable_column_device_view_core const* __restrict__ output_cols, + int32_t* __restrict__ max_error, + RowOperation&& operation) +{ + auto const start = grid_1d::global_thread_id(); + auto const stride = grid_1d::grid_stride(); + auto thread_error = errc::SUCCESS; + + for (auto row = start;; row += stride) { + if constexpr (!IsNullAware) { + if (row >= row_size) { break; } + + if (stencil != nullptr && !bit_is_set(stencil, row)) { continue; } + + auto ins = InputAccessors::map( + [&]() { return cuda::std::tuple{A::element(input_cols, row)...}; }); + + auto outs = OutputAccessors::map( + [&]() { return cuda::std::tuple{A::output_arg(output_cols, row)...}; }); + + auto out_ptrs = + cuda::std::apply([&](auto&... args) { return cuda::std::tuple{&args...}; }, outs); + + auto const row_error = operation(row, cuda::std::tuple_cat(out_ptrs, ins)); + + OutputAccessors::map([&]() { + (A::assign(output_cols, row, cuda::std::get(outs)), ...); + }); + + thread_error = cuda::std::max(thread_error, row_error); + } else { + auto const active_mask = __ballot_sync(0xFFFF'FFFFu, row < row_size); + + // fully inactive warp, break the loop + if (active_mask == 0) { break; } + + // partially active warp, continue to next warp iteration if row is out of bounds + if (row >= row_size) { continue; } + + auto ins = InputAccessors::map( + [&]() { return cuda::std::tuple{A::nullable_element(input_cols, row)...}; }); + + auto outs = OutputAccessors::map( + [&]() { return cuda::std::tuple{A::null_output_arg(output_cols, row)...}; }); + + auto out_ptrs = + cuda::std::apply([&](auto&... args) { return cuda::std::tuple{&args...}; }, outs); + + auto const row_error = operation(row, cuda::std::tuple_cat(out_ptrs, ins)); + + OutputAccessors::map([&]() { + (A::assign(output_cols, row, *cuda::std::get(outs)), ...); + (jit::warp_compact_validity( + active_mask, output_cols, row, cuda::std::get(outs).has_value()), + ...); + }); + + thread_error = cuda::std::max(thread_error, row_error); + } + } + + if (thread_error == errc::SUCCESS) { return; } + + cuda::atomic_ref ref(*max_error); + ref.fetch_max(static_cast(thread_error), cuda::std::memory_order_relaxed); +} + +} // namespace cudf::detail diff --git a/cpp/src/transform/checked_arithmetic.cu b/cpp/src/transform/checked_arithmetic.cu new file mode 100644 index 000000000000..4c473949e047 --- /dev/null +++ b/cpp/src/transform/checked_arithmetic.cu @@ -0,0 +1,343 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "checked_arithmetic.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include +#include +#include + +#include +#include + +namespace cudf::detail::checked_arithmetic { +namespace { + +template +inline constexpr bool is_supported_type = + (is_numeric() && !std::is_same_v) || is_fixed_point(); + +struct checked_binary_row_operation { + binary_operator op; + + template + __device__ static errc assign(cuda::std::optional* out, cuda::std::expected result) + { + if (!result.has_value()) { + *out = cuda::std::nullopt; + return result.error(); + } + *out = result.value(); + return errc::SUCCESS; + } + + template + __device__ errc operator()(size_type, Args args) const + { + auto* out = cuda::std::get<0>(args); + auto const lhs = cuda::std::get<1>(args); + auto const rhs = cuda::std::get<2>(args); + + if (!lhs.has_value() || !rhs.has_value()) { + *out = cuda::std::nullopt; + return errc::SUCCESS; + } + + switch (op) { + case binary_operator::ADD_OVERFLOW: + return assign(out, cudf::detail::ops::add_overflow(*lhs, *rhs)); + case binary_operator::SUB_OVERFLOW: + return assign(out, cudf::detail::ops::sub_overflow(*lhs, *rhs)); + case binary_operator::MUL_OVERFLOW: + return assign(out, cudf::detail::ops::mul_overflow(*lhs, *rhs)); + case binary_operator::DIV_OVERFLOW: + return assign(out, cudf::detail::ops::div_overflow(*lhs, *rhs)); + case binary_operator::MOD_OVERFLOW: + return assign(out, cudf::detail::ops::mod_overflow(*lhs, *rhs)); + default: return errc::SUCCESS; + } + } +}; + +struct checked_unary_row_operation { + unary_operator op; + + template + __device__ static errc assign(cuda::std::optional* out, cuda::std::expected result) + { + if (!result.has_value()) { + *out = cuda::std::nullopt; + return result.error(); + } + *out = result.value(); + return errc::SUCCESS; + } + + template + __device__ errc operator()(size_type, Args args) const + { + auto* out = cuda::std::get<0>(args); + auto const input = cuda::std::get<1>(args); + + if (!input.has_value()) { + *out = cuda::std::nullopt; + return errc::SUCCESS; + } + + switch (op) { + case unary_operator::NEG_OVERFLOW: + return assign(out, cudf::detail::ops::neg_overflow(*input)); + case unary_operator::ABS_OVERFLOW: + return assign(out, cudf::detail::ops::abs_overflow(*input)); + default: return errc::SUCCESS; + } + } +}; + +template +CUDF_KERNEL void checked_binary_kernel(size_type row_size, + column_device_view lhs, + column_device_view rhs, + mutable_column_device_view out, + binary_operator op, + int32_t* max_error) +{ + using input_accessors = + jit::type_list, + jit::column_accessor<1, column_device_view_core, T, RhsIsScalar, 0>>; + using output_accessors = + jit::type_list>; + + column_device_view_core const inputs[] = {lhs, rhs}; + mutable_column_device_view_core const outputs[] = {out}; + cudf::detail::transform_kernel( + row_size, nullptr, inputs, outputs, max_error, checked_binary_row_operation{op}); +} + +template +CUDF_KERNEL void checked_unary_kernel(size_type row_size, + column_device_view input, + mutable_column_device_view out, + unary_operator op, + int32_t* max_error) +{ + using input_accessors = + jit::type_list>; + using output_accessors = + jit::type_list>; + + column_device_view_core const inputs[] = {input}; + mutable_column_device_view_core const outputs[] = {out}; + cudf::detail::transform_kernel( + row_size, nullptr, inputs, outputs, max_error, checked_unary_row_operation{op}); +} + +void throw_if_error(errc error, error_policy policy) +{ + if (error == errc::SUCCESS || policy == error_policy::NULLIFY) { return; } + throw evaluation_error( + error, + std::string{"Checked arithmetic evaluation failed with error `"} + to_string(error) + "`"); +} + +template +struct binary_launcher { + template + void operator()(column_view const& lhs, + column_view const& rhs, + mutable_column_view& out, + binary_operator op, + error_policy policy, + cuda::stream_ref stream) const + { + if constexpr (is_supported_type) { + auto lhs_device = column_device_view::create(lhs, stream); + auto rhs_device = column_device_view::create(rhs, stream); + auto out_device = mutable_column_device_view::create(out, stream); + cudf::detail::device_scalar max_error{ + static_cast(errc::SUCCESS), stream, cudf::get_current_device_resource_ref()}; + + cudf::detail::grid_1d config(out.size(), 256); + checked_binary_kernel + <<>>( + out.size(), *lhs_device, *rhs_device, *out_device, op, max_error.data()); + CUDF_CHECK_CUDA(stream.get()); + + throw_if_error(static_cast(max_error.value(stream)), policy); + } else { + CUDF_FAIL("Checked arithmetic requires matching arithmetic or fixed-point types", + cudf::data_type_error); + } + } +}; + +struct unary_launcher { + template + void operator()(column_view const& input, + mutable_column_view& out, + unary_operator op, + error_policy policy, + cuda::stream_ref stream) const + { + if constexpr (is_supported_type) { + auto input_device = column_device_view::create(input, stream); + auto out_device = mutable_column_device_view::create(out, stream); + cudf::detail::device_scalar max_error{ + static_cast(errc::SUCCESS), stream, cudf::get_current_device_resource_ref()}; + + cudf::detail::grid_1d config(out.size(), 256); + checked_unary_kernel<<>>( + out.size(), *input_device, *out_device, op, max_error.data()); + CUDF_CHECK_CUDA(stream.get()); + + throw_if_error(static_cast(max_error.value(stream)), policy); + } else { + CUDF_FAIL("Checked arithmetic requires an arithmetic or fixed-point type", + cudf::data_type_error); + } + } +}; + +void validate_binary(column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type) +{ + CUDF_EXPECTS(is_checked(op), + "Error policies are only supported for checked arithmetic operators"); + CUDF_EXPECTS(lhs.type().id() == rhs.type().id() && output_type.id() == lhs.type().id(), + "Checked arithmetic requires matching input and output storage types", + cudf::data_type_error); + CUDF_EXPECTS( + (is_numeric(lhs.type()) && lhs.type().id() != type_id::BOOL8) || is_fixed_point(lhs.type()), + "Checked arithmetic requires arithmetic or fixed-point inputs", + cudf::data_type_error); + + if (is_fixed_point(lhs.type())) { + auto const expected_scale = + cudf::binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); + CUDF_EXPECTS(output_type.scale() == expected_scale, + "Checked fixed-point output has an invalid scale", + cudf::data_type_error); + } else { + CUDF_EXPECTS(lhs.type() == rhs.type() && output_type == lhs.type(), + "Checked arithmetic requires identical input and output types", + cudf::data_type_error); + } +} + +template +std::unique_ptr binary_operation_impl(column_view const& lhs, + column_view const& rhs, + size_type size, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + validate_binary(lhs, rhs, op, output_type); + + auto result = make_fixed_width_column(output_type, size, mask_state::ALL_VALID, stream, mr); + if (size == 0) { return result; } + + auto result_view = result->mutable_view(); + cudf::type_dispatcher(lhs.type(), + binary_launcher{}, + lhs, + rhs, + result_view, + op, + policy, + stream); + result->set_null_count( + cudf::detail::null_count(result_view.null_mask(), 0, result_view.size(), stream)); + return result; +} + +} // namespace + +std::unique_ptr binary_operation(scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto lhs_column = + make_column_from_scalar(lhs, 1, stream, cudf::get_current_device_resource_ref()); + return binary_operation_impl( + lhs_column->view(), rhs, rhs.size(), op, output_type, policy, stream, mr); +} + +std::unique_ptr binary_operation(column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + auto rhs_column = + make_column_from_scalar(rhs, 1, stream, cudf::get_current_device_resource_ref()); + return binary_operation_impl( + lhs, rhs_column->view(), lhs.size(), op, output_type, policy, stream, mr); +} + +std::unique_ptr binary_operation(column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes do not match", std::invalid_argument); + return binary_operation_impl( + lhs, rhs, lhs.size(), op, output_type, policy, stream, mr); +} + +std::unique_ptr unary_operation(column_view const& input, + unary_operator op, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(is_checked(op), + "Error policies are only supported for checked arithmetic operators"); + CUDF_EXPECTS((is_numeric(input.type()) && input.type().id() != type_id::BOOL8) || + is_fixed_point(input.type()), + "Checked arithmetic requires an arithmetic or fixed-point input", + cudf::data_type_error); + + auto result = + make_fixed_width_column(input.type(), input.size(), mask_state::ALL_VALID, stream, mr); + if (input.is_empty()) { return result; } + + auto result_view = result->mutable_view(); + cudf::type_dispatcher(input.type(), unary_launcher{}, input, result_view, op, policy, stream); + result->set_null_count( + cudf::detail::null_count(result_view.null_mask(), 0, result_view.size(), stream)); + return result; +} + +} // namespace cudf::detail::checked_arithmetic diff --git a/cpp/src/transform/checked_arithmetic.hpp b/cpp/src/transform/checked_arithmetic.hpp new file mode 100644 index 000000000000..9b033f5583fb --- /dev/null +++ b/cpp/src/transform/checked_arithmetic.hpp @@ -0,0 +1,55 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include + +namespace cudf::detail::checked_arithmetic { + +[[nodiscard]] constexpr bool is_checked(binary_operator op) +{ + return op == binary_operator::ADD_OVERFLOW || op == binary_operator::SUB_OVERFLOW || + op == binary_operator::MUL_OVERFLOW || op == binary_operator::DIV_OVERFLOW || + op == binary_operator::MOD_OVERFLOW; +} + +[[nodiscard]] constexpr bool is_checked(unary_operator op) +{ + return op == unary_operator::NEG_OVERFLOW || op == unary_operator::ABS_OVERFLOW; +} + +std::unique_ptr binary_operation(scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + +std::unique_ptr binary_operation(column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + +std::unique_ptr binary_operation(column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + +std::unique_ptr unary_operation(column_view const& input, + unary_operator op, + error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr); + +} // namespace cudf::detail::checked_arithmetic diff --git a/cpp/src/transform/jit/kernel.cu b/cpp/src/transform/jit/kernel.cu index 1ec0c3c3752b..def6e24facb6 100644 --- a/cpp/src/transform/jit/kernel.cu +++ b/cpp/src/transform/jit/kernel.cu @@ -5,23 +5,19 @@ #include #include -#include #include #include #include -#include #include #include -#include #include #include +#include #include -#include #include -#include -#include +#include #pragma nv_hdrstop // The above headers are used by the kernel below and need to be included before // it. Each UDF will have a different operation_udf.cuh generated for it, so we @@ -64,79 +60,25 @@ __device__ void transform_kernel(size_type row_size, mutable_column_device_view_core const* __restrict__ output_cols, int32_t* __restrict__ max_error) { - auto start = detail::grid_1d::global_thread_id(); - auto stride = detail::grid_1d::grid_stride(); - auto thread_error = errc::SUCCESS; - - for (auto row = start; row < row_size; row += stride) { - auto operation = [&](Args args) { - // TODO: static assert invocable - auto func = [&](auto... a) { - if constexpr (!cuda::std::is_void_v) { - return static_cast(GENERIC_TRANSFORM_OP(a...)); - } else { - (void)GENERIC_TRANSFORM_OP(a...); - return errc::SUCCESS; - } - }; - - if constexpr (has_user_data) { - return cuda::std::apply(func, cuda::std::tuple_cat(cuda::std::tuple{user_data, row}, args)); + auto operation = [&](size_type row, Args args) { + auto func = [&](auto... a) { + if constexpr (!cuda::std::is_void_v) { + return static_cast(GENERIC_TRANSFORM_OP(a...)); } else { - return cuda::std::apply(func, args); + (void)GENERIC_TRANSFORM_OP(a...); + return errc::SUCCESS; } }; - if constexpr (!is_null_aware) { - if (stencil != nullptr && !bit_is_set(stencil, row)) { continue; } - - auto ins = InputAccessors::map( - [&]() { return cuda::std::tuple{A::element(input_cols, row)...}; }); - - auto outs = OutputAccessors::map( - [&]() { return cuda::std::tuple{A::output_arg(output_cols, row)...}; }); - - auto out_ptrs = - cuda::std::apply([&](auto&... args) { return cuda::std::tuple{&args...}; }, outs); - - auto row_error = operation(cuda::std::tuple_cat(out_ptrs, ins)); - - OutputAccessors::map([&]() { - (A::assign(output_cols, row, cuda::std::get(outs)), ...); - }); - - thread_error = cuda::std::max(thread_error, row_error); - + if constexpr (has_user_data) { + return cuda::std::apply(func, cuda::std::tuple_cat(cuda::std::tuple{user_data, row}, args)); } else { - auto active_mask = __ballot_sync(__activemask(), row < row_size); - - auto ins = InputAccessors::map( - [&]() { return cuda::std::tuple{A::nullable_element(input_cols, row)...}; }); - - auto outs = OutputAccessors::map( - [&]() { return cuda::std::tuple{A::null_output_arg(output_cols, row)...}; }); - - auto out_ptrs = - cuda::std::apply([&](auto&... args) { return cuda::std::tuple{&args...}; }, outs); - - auto row_error = operation(cuda::std::tuple_cat(out_ptrs, ins)); - - OutputAccessors::map([&]() { - (A::assign(output_cols, row, *cuda::std::get(outs)), ...); - (warp_compact_validity( - active_mask, output_cols, row, cuda::std::get(outs).has_value()), - ...); - }); - - thread_error = cuda::std::max(thread_error, row_error); + return cuda::std::apply(func, args); } - } - - // early exit if no error occurred - if (thread_error == errc::SUCCESS) { return; } + }; - cuda::atomic_ref ref(*max_error); - ref.fetch_max(static_cast(thread_error), cuda::std::memory_order_relaxed); + detail::transform_kernel( + row_size, stencil, input_cols, output_cols, max_error, operation); } } // namespace jit diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index e294e9b00715..2d4b42e52733 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -22,6 +22,8 @@ #include #include +#include + namespace cudf { namespace detail { namespace { @@ -562,6 +564,11 @@ std::unique_ptr unary_operation(cudf::column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { + if (checked_arithmetic::is_checked(op)) { + // Omitting an error policy for a checked operator means propagate any row error. + return checked_arithmetic::unary_operation(input, op, error_policy::PROPAGATE, stream, mr); + } + if (cudf::is_fixed_point(input.type())) return type_dispatcher(input.type(), detail::FixedPointOpDispatcher{}, input, op, stream, mr); @@ -664,4 +671,14 @@ std::unique_ptr unary_operation(cudf::column_view const& input, return detail::unary_operation(input, op, stream, mr); } +std::unique_ptr unary_operation(cudf::column_view const& input, + cudf::unary_operator op, + cudf::error_policy policy, + cuda::stream_ref stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::checked_arithmetic::unary_operation(input, op, policy, stream, mr); +} + } // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 29cfd77d67ea..19480ee667c9 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -264,7 +264,10 @@ ConfigureTest(FIXED_POINT_TEST fixed_point/fixed_point_tests.cpp fixed_point/fix # ################################################################################################## # * unary tests ----------------------------------------------------------------------------------- -ConfigureTest(UNARY_TEST unary/math_ops_test.cpp unary/unary_ops_test.cpp unary/cast_tests.cpp) +ConfigureTest( + UNARY_TEST unary/math_ops_test.cpp unary/unary_ops_test.cpp unary/cast_tests.cpp + unary/operator_parity_test.cpp +) # ################################################################################################## # * round tests ----------------------------------------------------------------------------------- @@ -279,6 +282,7 @@ ConfigureTest( binaryop/binop-compiled-test.cpp binaryop/binop-compiled-fixed_point-test.cpp binaryop/binop-generic-ptx-test.cpp + binaryop/operator_parity_test.cpp ) # ################################################################################################## diff --git a/cpp/tests/binaryop/operator_parity_test.cpp b/cpp/tests/binaryop/operator_parity_test.cpp new file mode 100644 index 000000000000..78374d5554a9 --- /dev/null +++ b/cpp/tests/binaryop/operator_parity_test.cpp @@ -0,0 +1,160 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +#include +#include +#include + +#include + +struct BinaryOperatorParityTest : public cudf::test::BaseFixture {}; + +static_assert(static_cast(cudf::binary_operator::INVALID_BINARY) == 34); + +TEST_F(BinaryOperatorParityTest, CheckedDecimalSupportRequiresOutputScale) +{ + auto const lhs_type = cudf::data_type{cudf::type_id::DECIMAL32, -2}; + auto const rhs_type = cudf::data_type{cudf::type_id::DECIMAL32, -1}; + + auto expect_scale = [&](cudf::binary_operator op, int32_t expected_scale) { + EXPECT_TRUE(cudf::binops::is_supported_operation( + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, lhs_type, rhs_type, op)); + EXPECT_FALSE(cudf::binops::is_supported_operation( + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale + 1}, lhs_type, rhs_type, op)); + EXPECT_FALSE(cudf::binops::is_supported_operation( + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale - 1}, lhs_type, rhs_type, op)); + }; + + expect_scale(cudf::binary_operator::ADD_OVERFLOW, -2); + expect_scale(cudf::binary_operator::SUB_OVERFLOW, -2); + expect_scale(cudf::binary_operator::MUL_OVERFLOW, -3); + expect_scale(cudf::binary_operator::DIV_OVERFLOW, -1); + expect_scale(cudf::binary_operator::MOD_OVERFLOW, -2); +} + +TEST_F(BinaryOperatorParityTest, CheckedArithmeticPropagates) +{ + auto max = std::numeric_limits::max(); + auto min = std::numeric_limits::min(); + auto lhs = cudf::test::fixed_width_column_wrapper{max, min, max, 1, min}; + auto rhs = cudf::test::fixed_width_column_wrapper{1, 1, 2, 0, -1}; + auto type = cudf::data_type{cudf::type_id::INT32}; + + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD_OVERFLOW, type), + cudf::evaluation_error); + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::SUB_OVERFLOW, type), + cudf::evaluation_error); + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL_OVERFLOW, type), + cudf::evaluation_error); + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV_OVERFLOW, type), + cudf::evaluation_error); + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::MOD_OVERFLOW, type), + cudf::evaluation_error); +} + +TEST_F(BinaryOperatorParityTest, CheckedArithmeticNullifies) +{ + auto max = std::numeric_limits::max(); + auto min = std::numeric_limits::min(); + auto type = cudf::data_type{cudf::type_id::INT32}; + + auto add_lhs = cudf::test::fixed_width_column_wrapper{max, 4}; + auto add_rhs = cudf::test::fixed_width_column_wrapper{1, 5}; + auto add_expected = cudf::test::fixed_width_column_wrapper{{0, 9}, {false, true}}; + auto add_result = cudf::binary_operation( + add_lhs, add_rhs, cudf::binary_operator::ADD_OVERFLOW, type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(add_expected, add_result->view()); + + auto div_lhs = cudf::test::fixed_width_column_wrapper{min, 4, 6}; + auto div_rhs = cudf::test::fixed_width_column_wrapper{-1, 0, 3}; + auto div_expected = + cudf::test::fixed_width_column_wrapper{{0, 0, 2}, {false, false, true}}; + auto div_result = cudf::binary_operation( + div_lhs, div_rhs, cudf::binary_operator::DIV_OVERFLOW, type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(div_expected, div_result->view()); + + auto mod_expected = + cudf::test::fixed_width_column_wrapper{{0, 0, 0}, {true, false, true}}; + auto mod_result = cudf::binary_operation( + div_lhs, div_rhs, cudf::binary_operator::MOD_OVERFLOW, type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(mod_expected, mod_result->view()); +} + +TEST_F(BinaryOperatorParityTest, CheckedArithmeticSkipsNullInputs) +{ + auto lhs = cudf::test::fixed_width_column_wrapper{{1, 8}, {false, true}}; + auto rhs = cudf::test::fixed_width_column_wrapper{0, 2}; + auto expected = cudf::test::fixed_width_column_wrapper{{0, 4}, {false, true}}; + + auto result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::DIV_OVERFLOW, cudf::data_type{cudf::type_id::INT32}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TEST_F(BinaryOperatorParityTest, CheckedArithmeticScalarOperands) +{ + auto lhs = cudf::numeric_scalar{20}; + auto rhs = cudf::test::fixed_width_column_wrapper{2, 4}; + auto expected = cudf::test::fixed_width_column_wrapper{{10, 5}, {true, true}}; + + auto left_result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::DIV_OVERFLOW, lhs.type(), cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, left_result->view()); + + auto divisor = cudf::numeric_scalar{3}; + auto values = cudf::test::fixed_width_column_wrapper{7, 8}; + auto mod_expected = cudf::test::fixed_width_column_wrapper{{1, 2}, {true, true}}; + auto right_result = cudf::binary_operation(values, + divisor, + cudf::binary_operator::MOD_OVERFLOW, + divisor.type(), + cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(mod_expected, right_result->view()); +} + +TEST_F(BinaryOperatorParityTest, CheckedDecimalScaleRules) +{ + auto lhs = cudf::test::fixed_point_column_wrapper{{120, 250}, numeric::scale_type{-2}}; + auto rhs = cudf::test::fixed_point_column_wrapper{{3, 5}, numeric::scale_type{-1}}; + + auto add_expected = cudf::test::fixed_point_column_wrapper{ + {150, 300}, {true, true}, numeric::scale_type{-2}}; + auto add_type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD_OVERFLOW, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto add_result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::ADD_OVERFLOW, add_type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(add_expected, add_result->view()); + + auto mul_expected = cudf::test::fixed_point_column_wrapper{ + {360, 1250}, {true, true}, numeric::scale_type{-3}}; + auto mul_type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL_OVERFLOW, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto mul_result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::MUL_OVERFLOW, mul_type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(mul_expected, mul_result->view()); + + auto div_expected = cudf::test::fixed_point_column_wrapper{ + {40, 50}, {true, true}, numeric::scale_type{-1}}; + auto div_type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::DIV_OVERFLOW, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto div_result = cudf::binary_operation( + lhs, rhs, cudf::binary_operator::DIV_OVERFLOW, div_type, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(div_expected, div_result->view()); + + EXPECT_THROW( + cudf::binary_operation( + lhs, rhs, cudf::binary_operator::ADD_OVERFLOW, cudf::data_type{cudf::type_id::DECIMAL32, -1}), + cudf::data_type_error); +} diff --git a/cpp/tests/unary/operator_parity_test.cpp b/cpp/tests/unary/operator_parity_test.cpp new file mode 100644 index 000000000000..6a55d50b7223 --- /dev/null +++ b/cpp/tests/unary/operator_parity_test.cpp @@ -0,0 +1,69 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include + +#include +#include + +#include + +struct UnaryOperatorParityTest : public cudf::test::BaseFixture {}; + +static_assert(static_cast(cudf::unary_operator::NEGATE) == 23); + +TEST_F(UnaryOperatorParityTest, CheckedUnaryPropagates) +{ + auto min = std::numeric_limits::min(); + auto input = cudf::test::fixed_width_column_wrapper{min, -2, 3}; + + EXPECT_THROW(cudf::unary_operation(input, cudf::unary_operator::NEG_OVERFLOW), + cudf::evaluation_error); + EXPECT_THROW(cudf::unary_operation(input, cudf::unary_operator::ABS_OVERFLOW), + cudf::evaluation_error); +} + +TEST_F(UnaryOperatorParityTest, CheckedUnaryNullifies) +{ + auto min = std::numeric_limits::min(); + auto input = cudf::test::fixed_width_column_wrapper{min, -2, 3}; + + auto neg_expected = + cudf::test::fixed_width_column_wrapper{{0, 2, -3}, {false, true, true}}; + auto neg_result = + cudf::unary_operation(input, cudf::unary_operator::NEG_OVERFLOW, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(neg_expected, neg_result->view()); + + auto abs_expected = + cudf::test::fixed_width_column_wrapper{{0, 2, 3}, {false, true, true}}; + auto abs_result = + cudf::unary_operation(input, cudf::unary_operator::ABS_OVERFLOW, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(abs_expected, abs_result->view()); +} + +TEST_F(UnaryOperatorParityTest, CheckedUnarySkipsNullInputs) +{ + auto min = std::numeric_limits::min(); + auto input = cudf::test::fixed_width_column_wrapper{{min, 5}, {false, true}}; + auto expected = cudf::test::fixed_width_column_wrapper{{0, -5}, {false, true}}; + + auto result = cudf::unary_operation(input, cudf::unary_operator::NEG_OVERFLOW); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TEST_F(UnaryOperatorParityTest, CheckedUnaryDecimal) +{ + auto min = std::numeric_limits::min(); + auto input = cudf::test::fixed_point_column_wrapper{{min, -25}, numeric::scale_type{-2}}; + auto expected = cudf::test::fixed_point_column_wrapper{ + {0, 25}, {false, true}, numeric::scale_type{-2}}; + + auto result = + cudf::unary_operation(input, cudf::unary_operator::ABS_OVERFLOW, cudf::error_policy::NULLIFY); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +}