From fe6c79b0fc78745ec16cd9c4d4fd3bc822b73f86 Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Wed, 2 Sep 2026 17:28:16 +0100 Subject: [PATCH 1/9] Add precompiled checked arithmetic operators --- cpp/CMakeLists.txt | 1 + cpp/include/cudf/binaryop.hpp | 55 +++- cpp/include/cudf/unary.hpp | 24 ++ cpp/src/binaryop/binaryop.cpp | 83 ++++- cpp/src/jit/transform_kernel.cuh | 97 ++++++ cpp/src/transform/checked_arithmetic.cu | 343 ++++++++++++++++++++ cpp/src/transform/checked_arithmetic.hpp | 55 ++++ cpp/src/transform/jit/kernel.cu | 86 +---- cpp/src/unary/math_ops.cu | 17 + cpp/tests/CMakeLists.txt | 3 +- cpp/tests/binaryop/operator_parity_test.cpp | 141 ++++++++ cpp/tests/unary/operator_parity_test.cpp | 71 ++++ 12 files changed, 890 insertions(+), 86 deletions(-) create mode 100644 cpp/src/jit/transform_kernel.cuh create mode 100644 cpp/src/transform/checked_arithmetic.cu create mode 100644 cpp/src/transform/checked_arithmetic.hpp create mode 100644 cpp/tests/binaryop/operator_parity_test.cpp create mode 100644 cpp/tests/unary/operator_parity_test.cpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 193aa94409a7..cd781d1c287b 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1130,6 +1130,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 e5c10ed5e481..626537889f34 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 @@ -236,6 +241,54 @@ 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 + */ +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()); + +/** @copydoc binary_operation(scalar const&, column_view const&, binary_operator, data_type, + * error_policy, cuda::stream_ref, rmm::device_async_resource_ref) + */ +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()); + +/** @copydoc binary_operation(scalar const&, column_view const&, binary_operator, data_type, + * error_policy, cuda::stream_ref, rmm::device_async_resource_ref) + */ +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..762f4dfd04d7 100644 --- a/cpp/include/cudf/unary.hpp +++ b/cpp/include/cudf/unary.hpp @@ -54,12 +54,16 @@ enum class unary_operator : int32_t { 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,26 @@ 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 + */ +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 c4d39655eb78..e7107253fad2 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,10 @@ 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)) { + return out.id() == lhs.id() && lhs.id() == rhs.id() && + ((is_numeric(lhs) && lhs.id() != type_id::BOOL8) || is_fixed_point(lhs)); + } return cudf::binops::compiled::is_supported_operation(out, lhs, rhs, op); } @@ -89,15 +96,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 +144,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 { @@ -209,6 +222,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)) @@ -400,8 +418,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); } @@ -446,6 +466,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..923c98db193f --- /dev/null +++ b/cpp/src/jit/transform_kernel.cuh @@ -0,0 +1,97 @@ +/* + * 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 < row_size; row += stride) { + if constexpr (!IsNullAware) { + 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(__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 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 9153e1cbb82d..c25da8ac7f70 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -264,7 +264,7 @@ 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 +279,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..90e1ab7c27a7 --- /dev/null +++ b/cpp/tests/binaryop/operator_parity_test.cpp @@ -0,0 +1,141 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +#include +#include +#include + +#include + +namespace { + +struct BinaryOperatorParityTest : public cudf::test::BaseFixture {}; + +static_assert(static_cast(cudf::binary_operator::INVALID_BINARY) == 34); + +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); +} +} // namespace diff --git a/cpp/tests/unary/operator_parity_test.cpp b/cpp/tests/unary/operator_parity_test.cpp new file mode 100644 index 000000000000..acdaab0c7fe6 --- /dev/null +++ b/cpp/tests/unary/operator_parity_test.cpp @@ -0,0 +1,71 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +#include +#include + +#include + +namespace { + +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()); +} +} // namespace From 414020127e4a53589fb1a7e65610d5d318587682 Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Wed, 9 Sep 2026 17:09:11 +0000 Subject: [PATCH 2/9] Enhance error handling and add tests for checked arithmetic operations --- cpp/include/cudf/binaryop.hpp | 6 +++ cpp/include/cudf/unary.hpp | 55 +++++++++++---------- cpp/src/binaryop/binaryop.cpp | 7 ++- cpp/tests/CMakeLists.txt | 5 +- cpp/tests/binaryop/operator_parity_test.cpp | 23 +++++++-- cpp/tests/unary/operator_parity_test.cpp | 4 +- 6 files changed, 65 insertions(+), 35 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 626537889f34..117bd3c2d756 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -255,6 +255,10 @@ std::unique_ptr binary_operation( * @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, @@ -279,6 +283,8 @@ std::unique_ptr binary_operation( /** @copydoc binary_operation(scalar const&, column_view const&, binary_operator, data_type, * error_policy, cuda::stream_ref, rmm::device_async_resource_ref) + * + * @throws std::invalid_argument if @p lhs and @p rhs have different sizes */ std::unique_ptr binary_operation( column_view const& lhs, diff --git a/cpp/include/cudf/unary.hpp b/cpp/include/cudf/unary.hpp index 762f4dfd04d7..1f4a130908b0 100644 --- a/cpp/include/cudf/unary.hpp +++ b/cpp/include/cudf/unary.hpp @@ -30,32 +30,32 @@ 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. - NEG_OVERFLOW, ///< Negation with overflow detection - ABS_OVERFLOW ///< Absolute value with overflow detection + 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 }; /** @@ -90,6 +90,9 @@ std::unique_ptr unary_operation( * @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, diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index e7107253fad2..b097d6538a7b 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -53,8 +53,11 @@ 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)) { - return out.id() == lhs.id() && lhs.id() == rhs.id() && - ((is_numeric(lhs) && lhs.id() != type_id::BOOL8) || is_fixed_point(lhs)); + 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); } diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index c25da8ac7f70..c6e491a3cdeb 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 unary/operator_parity_test.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 ----------------------------------------------------------------------------------- diff --git a/cpp/tests/binaryop/operator_parity_test.cpp b/cpp/tests/binaryop/operator_parity_test.cpp index 90e1ab7c27a7..bdb63dcb149b 100644 --- a/cpp/tests/binaryop/operator_parity_test.cpp +++ b/cpp/tests/binaryop/operator_parity_test.cpp @@ -5,6 +5,7 @@ #include #include +#include #include #include @@ -12,12 +13,29 @@ #include -namespace { - 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_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(); @@ -138,4 +156,3 @@ TEST_F(BinaryOperatorParityTest, CheckedDecimalScaleRules) lhs, rhs, cudf::binary_operator::ADD_OVERFLOW, cudf::data_type{cudf::type_id::DECIMAL32, -1}), cudf::data_type_error); } -} // namespace diff --git a/cpp/tests/unary/operator_parity_test.cpp b/cpp/tests/unary/operator_parity_test.cpp index acdaab0c7fe6..6a55d50b7223 100644 --- a/cpp/tests/unary/operator_parity_test.cpp +++ b/cpp/tests/unary/operator_parity_test.cpp @@ -6,14 +6,13 @@ #include #include #include +#include #include #include #include -namespace { - struct UnaryOperatorParityTest : public cudf::test::BaseFixture {}; static_assert(static_cast(cudf::unary_operator::NEGATE) == 23); @@ -68,4 +67,3 @@ TEST_F(UnaryOperatorParityTest, CheckedUnaryDecimal) cudf::unary_operation(input, cudf::unary_operator::ABS_OVERFLOW, cudf::error_policy::NULLIFY); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -} // namespace From 3c587f8c6fb941a7530a31c484de936011b2c4ac Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Thu, 10 Sep 2026 01:44:57 +0100 Subject: [PATCH 3/9] Test lower invalid checked decimal scale --- cpp/tests/binaryop/operator_parity_test.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/tests/binaryop/operator_parity_test.cpp b/cpp/tests/binaryop/operator_parity_test.cpp index bdb63dcb149b..78374d5554a9 100644 --- a/cpp/tests/binaryop/operator_parity_test.cpp +++ b/cpp/tests/binaryop/operator_parity_test.cpp @@ -27,6 +27,8 @@ TEST_F(BinaryOperatorParityTest, CheckedDecimalSupportRequiresOutputScale) 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); From 5e6ac7feaf4314367c4e47b2168ed8acd650d776 Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Thu, 10 Sep 2026 01:54:45 +0100 Subject: [PATCH 4/9] Fix partial-warp transform validity ballot --- cpp/src/jit/transform_kernel.cuh | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/cpp/src/jit/transform_kernel.cuh b/cpp/src/jit/transform_kernel.cuh index 923c98db193f..e85db7f5edfc 100644 --- a/cpp/src/jit/transform_kernel.cuh +++ b/cpp/src/jit/transform_kernel.cuh @@ -43,8 +43,10 @@ __device__ void transform_kernel(size_type row_size, auto const stride = grid_1d::grid_stride(); auto thread_error = errc::SUCCESS; - for (auto row = start; row < row_size; row += stride) { + 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( @@ -64,7 +66,10 @@ __device__ void transform_kernel(size_type row_size, thread_error = cuda::std::max(thread_error, row_error); } else { - auto const active_mask = __ballot_sync(__activemask(), row < row_size); + auto const active_mask = __ballot_sync(0xFFFF'FFFFu, row < row_size); + + if (active_mask == 0) { break; } + if (row >= row_size) { continue; } auto ins = InputAccessors::map( [&]() { return cuda::std::tuple{A::nullable_element(input_cols, row)...}; }); From cd35c43c1fea21e4ccd842b0f3293b3005f278d5 Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Thu, 10 Sep 2026 00:59:05 +0000 Subject: [PATCH 5/9] Improve warp handling in transform_kernel to break on fully inactive warps and continue on partially active ones --- cpp/src/jit/transform_kernel.cuh | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/src/jit/transform_kernel.cuh b/cpp/src/jit/transform_kernel.cuh index e85db7f5edfc..6b1d2187f45b 100644 --- a/cpp/src/jit/transform_kernel.cuh +++ b/cpp/src/jit/transform_kernel.cuh @@ -68,7 +68,10 @@ __device__ void transform_kernel(size_type row_size, } 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( From eba57e1a3fc87478676354bf0a370091ff9e10f0 Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Thu, 10 Sep 2026 02:00:46 +0100 Subject: [PATCH 6/9] Fix checked binary operation documentation --- cpp/include/cudf/binaryop.hpp | 39 +++++++++++++++++++++++++++++++---- 1 file changed, 35 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 9b0ce8d09d4e..5d1c6b2d01f7 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -272,8 +272,24 @@ 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()); -/** @copydoc binary_operation(scalar const&, column_view const&, binary_operator, data_type, - * error_policy, cuda::stream_ref, rmm::device_async_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, @@ -284,9 +300,24 @@ 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()); -/** @copydoc binary_operation(scalar const&, column_view const&, binary_operator, data_type, - * error_policy, cuda::stream_ref, rmm::device_async_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( From e59614e1b9c6d754169fbc16ef1d21d77e47f27d Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Wed, 23 Sep 2026 12:35:41 +0000 Subject: [PATCH 7/9] REMOVE: Static assert for invalid binary operator in operator parity tests --- cpp/tests/binaryop/operator_parity_test.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/tests/binaryop/operator_parity_test.cpp b/cpp/tests/binaryop/operator_parity_test.cpp index 78374d5554a9..b9cd3ab0cba0 100644 --- a/cpp/tests/binaryop/operator_parity_test.cpp +++ b/cpp/tests/binaryop/operator_parity_test.cpp @@ -15,8 +15,6 @@ 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}; From ae2fe8629dd2c04827558bc7e27e63ba2b72747e Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Wed, 23 Sep 2026 12:44:23 +0000 Subject: [PATCH 8/9] update binop and unop enums --- cpp/include/cudf/binaryop.hpp | 4 ++-- java/src/main/java/ai/rapids/cudf/BinaryOp.java | 10 ++++++++-- java/src/main/java/ai/rapids/cudf/UnaryOp.java | 7 +++++-- python/pylibcudf/pylibcudf/binaryop.pyi | 7 ++++++- python/pylibcudf/pylibcudf/libcudf/binaryop.pxd | 7 ++++++- python/pylibcudf/pylibcudf/libcudf/unary.pxd | 2 ++ python/pylibcudf/pylibcudf/unary.pyi | 2 ++ 7 files changed, 31 insertions(+), 8 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 5d1c6b2d01f7..60a65cf2cc4d 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -86,12 +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 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 + MOD_OVERFLOW, ///< Modulo with divide-by-zero detection + INVALID_BINARY ///< invalid operation }; /// Binary operation common type default diff --git a/java/src/main/java/ai/rapids/cudf/BinaryOp.java b/java/src/main/java/ai/rapids/cudf/BinaryOp.java index 7dec9072d0e5..fd9b19885f92 100644 --- a/java/src/main/java/ai/rapids/cudf/BinaryOp.java +++ b/java/src/main/java/ai/rapids/cudf/BinaryOp.java @@ -46,11 +46,17 @@ public enum BinaryOp { NULL_NOT_EQUALS(28), // negation of NULL_EQUALS NULL_MAX(29), // MAX but NULL < not NULL NULL_MIN(30), // MIN but NULL > not NULL - //NOT IMPLEMENTED YET GENERIC_BINARY(30); + //NOT IMPLEMENTED YET GENERIC_BINARY(31); NULL_LOGICAL_AND(32), // three-valued (Kleene) &&: if any operand is false, returns false; if both // operands are true, returns true; otherwise returns null - NULL_LOGICAL_OR(33); // three-valued (Kleene) ||: if any operand is true, returns true; if both + NULL_LOGICAL_OR(33), // three-valued (Kleene) ||: if any operand is true, returns true; if both // operands are false, returns false; otherwise returns null + ADD_OVERFLOW(34), + SUB_OVERFLOW(35), + MUL_OVERFLOW(36), + DIV_OVERFLOW(37), + MOD_OVERFLOW(38), + INVALID_BINARY(39); static final EnumSet COMPARISON = EnumSet.of( diff --git a/java/src/main/java/ai/rapids/cudf/UnaryOp.java b/java/src/main/java/ai/rapids/cudf/UnaryOp.java index 6c09e8957456..6154a5fb36d9 100644 --- a/java/src/main/java/ai/rapids/cudf/UnaryOp.java +++ b/java/src/main/java/ai/rapids/cudf/UnaryOp.java @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ package ai.rapids.cudf; @@ -30,7 +30,10 @@ public enum UnaryOp { RINT(19), BIT_COUNT(20), BIT_INVERT(21), - NOT(22); + NOT(22), + NEGATE(23), + NEG_OVERFLOW(24), + ABS_OVERFLOW(25); private static final UnaryOp[] OPS = UnaryOp.values(); final int nativeId; diff --git a/python/pylibcudf/pylibcudf/binaryop.pyi b/python/pylibcudf/pylibcudf/binaryop.pyi index 1f3c9a2cb64f..9c1efa2bed9d 100644 --- a/python/pylibcudf/pylibcudf/binaryop.pyi +++ b/python/pylibcudf/pylibcudf/binaryop.pyi @@ -39,12 +39,17 @@ class BinaryOperator(IntEnum): LESS_EQUAL = ... GREATER_EQUAL = ... NULL_EQUALS = ... + NULL_NOT_EQUALS = ... NULL_MAX = ... NULL_MIN = ... - NULL_NOT_EQUALS = ... GENERIC_BINARY = ... NULL_LOGICAL_AND = ... NULL_LOGICAL_OR = ... + ADD_OVERFLOW = ... + SUB_OVERFLOW = ... + MUL_OVERFLOW = ... + DIV_OVERFLOW = ... + MOD_OVERFLOW = ... INVALID_BINARY = ... def binary_operation( diff --git a/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd b/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd index 303b112f71ee..a5c128560537 100644 --- a/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd @@ -44,12 +44,17 @@ cdef extern from "cudf/binaryop.hpp" namespace "cudf" nogil: LESS_EQUAL GREATER_EQUAL NULL_EQUALS + NULL_NOT_EQUALS NULL_MAX NULL_MIN - NULL_NOT_EQUALS GENERIC_BINARY NULL_LOGICAL_AND NULL_LOGICAL_OR + ADD_OVERFLOW + SUB_OVERFLOW + MUL_OVERFLOW + DIV_OVERFLOW + MOD_OVERFLOW INVALID_BINARY cdef unique_ptr[column] binary_operation ( diff --git a/python/pylibcudf/pylibcudf/libcudf/unary.pxd b/python/pylibcudf/pylibcudf/libcudf/unary.pxd index 6f59ff8d5e0f..319640bfb5d8 100644 --- a/python/pylibcudf/pylibcudf/libcudf/unary.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/unary.pxd @@ -38,6 +38,8 @@ cdef extern from "cudf/unary.hpp" namespace "cudf" nogil: BIT_INVERT NOT NEGATE + NEG_OVERFLOW + ABS_OVERFLOW cdef extern unique_ptr[column] unary_operation( column_view input, diff --git a/python/pylibcudf/pylibcudf/unary.pyi b/python/pylibcudf/pylibcudf/unary.pyi index 821254912db5..71c04338cfc3 100644 --- a/python/pylibcudf/pylibcudf/unary.pyi +++ b/python/pylibcudf/pylibcudf/unary.pyi @@ -34,6 +34,8 @@ class UnaryOperator(IntEnum): BIT_INVERT = ... NOT = ... NEGATE = ... + NEG_OVERFLOW = ... + ABS_OVERFLOW = ... def unary_operation( input: Column, From 5637093fc8dcdafcfbda0c854fb2569e656d012d Mon Sep 17 00:00:00 2001 From: Basit Ayantunde Date: Wed, 23 Sep 2026 12:57:36 +0000 Subject: [PATCH 9/9] Update copyright notices to include "AFFILIATES" in multiple files --- java/src/main/java/ai/rapids/cudf/UnaryOp.java | 2 +- python/pylibcudf/pylibcudf/binaryop.pyi | 2 +- python/pylibcudf/pylibcudf/libcudf/binaryop.pxd | 2 +- python/pylibcudf/pylibcudf/libcudf/unary.pxd | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/java/src/main/java/ai/rapids/cudf/UnaryOp.java b/java/src/main/java/ai/rapids/cudf/UnaryOp.java index 6154a5fb36d9..abc3af090a15 100644 --- a/java/src/main/java/ai/rapids/cudf/UnaryOp.java +++ b/java/src/main/java/ai/rapids/cudf/UnaryOp.java @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ package ai.rapids.cudf; diff --git a/python/pylibcudf/pylibcudf/binaryop.pyi b/python/pylibcudf/pylibcudf/binaryop.pyi index 9c1efa2bed9d..2d11baa3a061 100644 --- a/python/pylibcudf/pylibcudf/binaryop.pyi +++ b/python/pylibcudf/pylibcudf/binaryop.pyi @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from enum import IntEnum diff --git a/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd b/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd index a5c128560537..10331636f5cb 100644 --- a/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/binaryop.pxd @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from libc.stdint cimport int32_t from libcpp cimport bool diff --git a/python/pylibcudf/pylibcudf/libcudf/unary.pxd b/python/pylibcudf/pylibcudf/libcudf/unary.pxd index 319640bfb5d8..86d9cb598a02 100644 --- a/python/pylibcudf/pylibcudf/libcudf/unary.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/unary.pxd @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from libc.stdint cimport int32_t from libcpp cimport bool