diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 490a69f3c2a7..58b836ca6703 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -99,7 +99,6 @@ mark_as_advanced(CUDF_BUILD_STREAMS_TEST_UTIL) option(CUDF_CLANG_TIDY "Enable clang-tidy during compilation" OFF) option(CUDF_IWYU "Enable IWYU during compilation" OFF) option(CUDF_CLANG_TIDY_AUTOFIX "Enable clang-tidy autofixes" OFF) - option( CUDF_KVIKIO_REMOTE_IO "Enable remote IO (e.g. AWS S3) support through KvikIO. If disabled, cudf-python will still be able to do remote IO through fsspec." @@ -452,6 +451,7 @@ add_library( src/binaryop/compiled/Sub.cu src/binaryop/compiled/TrueDiv.cu src/binaryop/compiled/binary_ops.cu + src/binaryop/compiled/binary_ops_safe.cu src/binaryop/compiled/equality_ops.cu src/binaryop/compiled/util.cpp src/labeling/label_bins.cu diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index e745bfdab891..969d07bb4909 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -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 */ @@ -10,7 +10,9 @@ #include #include +#include #include +#include namespace CUDF_EXPORT cudf { @@ -221,6 +223,130 @@ std::unique_ptr binary_operation( rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); +/** + * @brief Decimal fixed-point binary operation between a scalar and a column, + * with a per-row overflow column. + * + * The result column contains `op(lhs, rhs[i])` for all `0 <= i < rhs.size()`, + * matching the semantics of `binary_operation` for the seven supported decimal + * arithmetic operators (ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD). + * + * Additionally returns a `BOOL8` column of the same length as the result. + * Element `i` is `true` iff row `i` is an active (non-null on both sides) row + * whose arithmetic or rescale to @p output_type overflowed (or, for DIV / MOD + * / PMOD / PYMOD, divided by zero). Null rows and clean rows hold `false`. + * The overflow column has no null mask. + * + * @p lhs, @p rhs, and @p output_type must share the same decimal storage type + * (e.g. all `DECIMAL64`); mixing decimal widths or pairing decimal with a + * non-decimal operand is not supported on this path. + * + * @param lhs The left operand decimal scalar + * @param rhs The right operand decimal column + * @param op The binary operator + * @param output_type The desired data type of the result column (must be a base-10 decimal) + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned columns' device memory + * @return A pair `{result, overflow}` where `result` is the arithmetic result column and + * `overflow` is the per-row `BOOL8` overflow column described above. + * @throw cudf::logic_error if @p lhs or @p rhs is not a fixed-point type + * @throw cudf::logic_error if @p op is not one of ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD + * @throw cudf::logic_error if @p lhs, @p rhs, and @p output_type do not share the same + * decimal storage type + * @throw cudf::data_type_error if the operation is not supported for the types of + * @p lhs and @p rhs + */ +std::pair, std::unique_ptr> binary_operation_safe( + scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Decimal fixed-point binary operation between a column and a scalar, + * with a per-row overflow column. + * + * The result column contains `op(lhs[i], rhs)` for all `0 <= i < lhs.size()`, + * matching the semantics of `binary_operation` for the seven supported decimal + * arithmetic operators (ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD). + * + * Additionally returns a `BOOL8` column of the same length as the result. + * Element `i` is `true` iff row `i` is an active (non-null on both sides) row + * whose arithmetic or rescale to @p output_type overflowed (or, for DIV / MOD + * / PMOD / PYMOD, divided by zero). Null rows and clean rows hold `false`. + * The overflow column has no null mask. + * + * @p lhs, @p rhs, and @p output_type must share the same decimal storage type + * (e.g. all `DECIMAL64`); mixing decimal widths or pairing decimal with a + * non-decimal operand is not supported on this path. + * + * @param lhs The left operand decimal column + * @param rhs The right operand decimal scalar + * @param op The binary operator + * @param output_type The desired data type of the result column (must be a base-10 decimal) + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned columns' device memory + * @return A pair `{result, overflow}` where `result` is the arithmetic result column and + * `overflow` is the per-row `BOOL8` overflow column described above. + * @throw cudf::logic_error if @p lhs or @p rhs is not a fixed-point type + * @throw cudf::logic_error if @p op is not one of ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD + * @throw cudf::logic_error if @p lhs, @p rhs, and @p output_type do not share the same + * decimal storage type + * @throw cudf::data_type_error if the operation is not supported for the types of + * @p lhs and @p rhs + */ +std::pair, std::unique_ptr> binary_operation_safe( + column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Decimal fixed-point binary operation between two columns, + * with a per-row overflow column. + * + * The result column contains `op(lhs[i], rhs[i])` for all `0 <= i < lhs.size()`, + * matching the semantics of `binary_operation` for the seven supported decimal + * arithmetic operators (ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD). + * + * Additionally returns a `BOOL8` column of the same length as the result. + * Element `i` is `true` iff row `i` is an active (non-null on both sides) row + * whose arithmetic or rescale to @p output_type overflowed (or, for DIV / MOD + * / PMOD / PYMOD, divided by zero). Null rows and clean rows hold `false`. + * The overflow column has no null mask. + * + * @p lhs, @p rhs, and @p output_type must share the same decimal storage type + * (e.g. all `DECIMAL64`); mixing decimal widths or pairing decimal with a + * non-decimal operand is not supported on this path. + * + * @param lhs The left operand decimal column + * @param rhs The right operand decimal column + * @param op The binary operator + * @param output_type The desired data type of the result column (must be a base-10 decimal) + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned columns' device memory + * @return A pair `{result, overflow}` where `result` is the arithmetic result column and + * `overflow` is the per-row `BOOL8` overflow column described above. + * @throw cudf::logic_error if @p lhs and @p rhs are different sizes + * @throw cudf::logic_error if @p lhs or @p rhs is not a fixed-point type + * @throw cudf::logic_error if @p op is not one of ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD + * @throw cudf::logic_error if @p lhs, @p rhs, and @p output_type do not share the same + * decimal storage type + * @throw cudf::data_type_error if the operation is not supported for the types of + * @p lhs and @p rhs + */ +std::pair, std::unique_ptr> binary_operation_safe( + column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view 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/fixed_point/detail/floating_conversion.hpp b/cpp/include/cudf/fixed_point/detail/floating_conversion.hpp index fecc1b6b330d..56771781948e 100644 --- a/cpp/include/cudf/fixed_point/detail/floating_conversion.hpp +++ b/cpp/include/cudf/fixed_point/detail/floating_conversion.hpp @@ -508,6 +508,51 @@ CUDF_HOST_DEVICE inline IntegerType guarded_left_shift(IntegerType value, int bi : cuda::std::numeric_limits::max(); } +/** + * @brief Perform a bit-shift left, optionally detecting overflow (saturating) + * + * @tparam CheckOverflow Whether to detect overflow in addition to guarding undefined behavior + * @tparam IntegerType Type of input unsigned integer value + * @param value The integer whose bits are being shifted + * @param bit_shift The number of bits to shift left + * @return `{value, overflow}` where `overflow` is true only when `CheckOverflow` is true and the + * shift overflows + */ +template )> +CUDF_HOST_DEVICE inline cuda::std::pair checked_left_shift(IntegerType value, + int bit_shift) +{ + if constexpr (!CheckOverflow) { + // Mirror the negative-shift guard used in the `CheckOverflow` branch so we + // never forward `value << bit_shift` with a negative shift count + // (undefined behavior). A negative shift is not expected from any current + // caller; return zero in line with `guarded_right_shift`'s underflow case. + if (bit_shift < 0) { return {IntegerType{0}, false}; } + return {guarded_left_shift(value, bit_shift), false}; + } else { + constexpr int digits = cuda::std::numeric_limits::digits; + constexpr int max_safe_bit_shift = digits - 1; + + if (bit_shift < 0) { + // Not expected for callers; treat as overflow (would be a right-shift). + return {cuda::std::numeric_limits::max(), true}; + } + if (bit_shift > max_safe_bit_shift) { + return {cuda::std::numeric_limits::max(), true}; + } + if (bit_shift == 0) { return {value, false}; } + + // Detect whether any bits would be shifted out. + auto const max_value_before_shift = cuda::std::numeric_limits::max() >> bit_shift; + if (value > max_value_before_shift) { + return {cuda::std::numeric_limits::max(), true}; + } + return {static_cast(value << bit_shift), false}; + } +} + /** * @brief Perform a bit-shift right, guarding against undefined behavior * @@ -524,6 +569,44 @@ CUDF_HOST_DEVICE inline IntegerType guarded_right_shift(IntegerType value, int b return (bit_shift <= max_safe_bit_shift) ? value >> bit_shift : 0; } +/** + * @brief Cast `value` to a narrower unsigned type, optionally detecting overflow (saturating) + */ +template && cuda::std::is_unsigned_v)> +CUDF_HOST_DEVICE inline cuda::std::pair checked_narrow_cast(From value) +{ + if constexpr (CheckOverflow) { + if (value > static_cast(cuda::std::numeric_limits::max())) { + return {cuda::std::numeric_limits::max(), true}; + } + } + return {static_cast(value), false}; +} + +/** + * @brief Multiply by 10^pow10 with overflow detection (saturating) + * + * Reuses the bit-size-specialized `divide_power10` / `multiply_power10` helpers above so the + * check and the multiply are both single table-style ops, independent of `pow10`. + */ +template )> +CUDF_HOST_DEVICE inline cuda::std::pair multiply_power10_saturating(T value, int pow10) +{ + if (pow10 <= 0) { return {value, false}; } + if constexpr (!CheckOverflow) { return {multiply_power10(value, pow10), false}; } + + // value * 10^pow10 fits in T iff value <= floor(max_v / 10^pow10). When 10^pow10 itself + // overflows T, divide_power10 returns 0, so the threshold becomes 0 and any nonzero value + // correctly saturates -- matching the per-digit loop's behavior. + auto const max_v = cuda::std::numeric_limits::max(); + auto const value_max = divide_power10(max_v, pow10); + if (value > value_max) { return {max_v, true}; } + return {multiply_power10(value, pow10), false}; +} + /** * @brief Helper struct with common constants needed by the floating <--> decimal conversions */ @@ -670,16 +753,20 @@ add_half_if_truncates(FloatingType floating, * * @tparam Rep The type of the storage for the decimal value * @tparam FloatingType The type of the original floating-point value we are converting from + * @tparam CheckOverflow Whether to detect overflow while narrowing or shifting * @param base2_value The base-2 fixed-point value we are converting from * @param pow2 The number of powers of 2 to apply to convert from base-2 * @param pow10 The number of powers of 10 to apply to reach the desired scale factor - * @return Magnitude of the converted-to decimal integer + * @return `{magnitude, overflow}` for the converted-to decimal integer */ template )> -CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( - typename shifting_constants::IntegerRep const base2_value, int pow2, int pow10) +CUDF_HOST_DEVICE inline cuda::std::pair, bool> +shift_to_decimal_pospow(typename shifting_constants::IntegerRep const base2_value, + int pow2, + int pow10) { // To convert to decimal, we need to apply the input powers of 2 and 10 // The result will be (integer) base2_value * (2^pow2) / (10^pow10) @@ -694,6 +781,7 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( // So we'll cast up to ShiftingRep: uint64 for float's, __uint128_t for double's using Constants = shifting_constants; using ShiftingRep = typename Constants::ShiftingRep; + using UnsignedRep = cuda::std::make_unsigned_t; auto shifting_rep = static_cast(base2_value); // We want to start with our significand bits at the top of the shifting range, @@ -706,12 +794,10 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( static constexpr int max_init_shift = shift_up_to - shift_from; // If our total bit shift is less than this, we don't need to iterate - using UnsignedRep = cuda::std::make_unsigned_t; if (pow2 <= max_init_shift) { // Shift bits left, divide by 10s to apply the scale factor, and we're done. shifting_rep = divide_power10(shifting_rep << pow2, pow10); - // NOTE: Cast can overflow! - return static_cast(shifting_rep); + return checked_narrow_cast(shifting_rep); } // We need to iterate. Do the combined initial shift @@ -728,9 +814,7 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( if (pow2 <= Constants::max_bits_shift) { // Shift bits left, divide by 10s to apply the scale factor, and we're done. shifting_rep = divide_power10(shifting_rep << pow2, pow10); - - // NOTE: Cast can overflow! - return static_cast(shifting_rep); + return checked_narrow_cast(shifting_rep); } // Shift the max number of bits left again @@ -748,8 +832,10 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( } // Final bit shift: Shift may be large, guard against UB - // NOTE: This can overflow (both cast and shift)! - return guarded_left_shift(static_cast(shifting_rep), pow2); + auto const [narrowed, narrow_overflow] = + checked_narrow_cast(shifting_rep); + auto const [shifted, shift_overflow] = checked_left_shift(narrowed, pow2); + return {shifted, narrow_overflow || shift_overflow}; } /** @@ -757,16 +843,20 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_pospow( * * @tparam Rep The type of the storage for the decimal value * @tparam FloatingType The type of the original floating-point value we are converting from + * @tparam CheckOverflow Whether to detect overflow while narrowing, shifting, or multiplying * @param base2_value The base-2 fixed-point value we are converting from * @param pow2 The number of powers of 2 to apply to convert from base-2 * @param pow10 The number of powers of 10 to apply to reach the desired scale factor - * @return Magnitude of the converted-to decimal integer + * @return `{magnitude, overflow}` for the converted-to decimal integer */ template )> -CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( - typename shifting_constants::IntegerRep base2_value, int pow2, int pow10) +CUDF_HOST_DEVICE inline cuda::std::pair, bool> +shift_to_decimal_negpow(typename shifting_constants::IntegerRep base2_value, + int pow2, + int pow10) { // This is similar to shift_to_decimal_pospow(), except pow10 < 0 & pow2 < 0 // See comments in that function for details. @@ -775,6 +865,7 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( // ShiftingRep: uint64 for float's, __uint128_t for double's using Constants = shifting_constants; using ShiftingRep = typename Constants::ShiftingRep; + using UnsignedRep = cuda::std::make_unsigned_t; auto shifting_rep = static_cast(base2_value); // Convert to using positive values so we don't have keep negating @@ -782,7 +873,6 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( int pow2_mag = -pow2; // For performing final 10s-shift - using UnsignedRep = cuda::std::make_unsigned_t; auto final_shifts_low10s = [&]() { // Last 10s-shift: multiply all remaining decimal places, shift all remaining bits, then bail // The multiplier is less than the max-shift, and thus fits within 64 / 32 bits @@ -792,8 +882,9 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( shifting_rep = multiply_power10_32bit(shifting_rep, pow10_mag); } - // Final bit shifting: Shift may be large, guard against UB - return static_cast(guarded_right_shift(shifting_rep, pow2_mag)); + // guarded_right_shift does not overflow the representable range; it may drop to 0 on UB. + return checked_narrow_cast( + guarded_right_shift(shifting_rep, pow2_mag)); }; // If our total decimal shift is less than the max, we don't need to iterate @@ -825,8 +916,11 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( // We need to convert to the output rep for the final scale-factor multiply, because if (e.g.) // float -> dec128 and some large pow10_mag, it might overflow the 64bit shifting rep. // It's not needed for pow10 > 0 because we're dividing by 10s there instead of multiplying. - // NOTE: This can overflow! (Both multiply and cast) - return multiply_power10(static_cast(shifting_rep), pow10_mag); + auto const [narrowed, narrow_overflow] = + checked_narrow_cast(shifting_rep); + auto const [scaled, multiply_overflow] = + multiply_power10_saturating(narrowed, pow10_mag); + return {scaled, narrow_overflow || multiply_overflow}; } // More bits to shift than we have room: Shift the max number of 2s @@ -838,20 +932,34 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t shift_to_decimal_negpow( return final_shifts_low10s(); } +template +CUDF_HOST_DEVICE inline auto maybe_with_overflow(T value, bool overflow) +{ + if constexpr (CheckOverflow) { + return cuda::std::make_pair(value, overflow); + } else { + return value; + } +} + /** * @brief Perform base-2 -> base-10 fixed-point conversion * * @tparam Rep The type of integer we are converting to, to store the decimal value * @tparam FloatingType The type of floating-point object we are converting from + * @tparam CheckOverflow Whether to detect overflow while applying the scale * @param base2_value The base-2 fixed-point value we are converting from * @param pow2 The number of powers of 2 to apply to convert from base-2 * @param pow10 The number of powers of 10 to apply to reach the desired scale factor - * @return Integer representation of the floating-point value, given the desired scale + * @return `{magnitude, overflow}` for the converted-to decimal integer when `CheckOverflow` is + * true, otherwise just the magnitude (legacy unchecked behavior, preserved for downstream callers + * such as spark-rapids-jni). */ template )> -CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t convert_floating_to_integral_shifting( +CUDF_HOST_DEVICE inline auto convert_floating_to_integral_shifting( typename floating_converter::IntegralType base2_value, int pow10, int pow2) { // Apply the powers of 2 and 10 to convert to decimal. @@ -862,54 +970,71 @@ CUDF_HOST_DEVICE inline cuda::std::make_unsigned_t convert_floating_to_inte // Also data within a column tends to be similar, so they will often take the // same branches on pow2 as well. - // NOTE: some returns here can overflow (e.g. ShiftingRep -> UnsignedRep) using UnsignedRep = cuda::std::make_unsigned_t; + if (pow10 == 0) { - // NOTE: Left Bit-shift can overflow! As can cast! (e.g. double -> decimal32) - // Bit shifts may be large, guard against UB if (pow2 >= 0) { - return guarded_left_shift(static_cast(base2_value), pow2); - } else { - return static_cast(guarded_right_shift(base2_value, -pow2)); + auto const [narrowed, narrow_overflow] = + checked_narrow_cast(base2_value); + auto const [shifted, shift_overflow] = checked_left_shift(narrowed, pow2); + return maybe_with_overflow(shifted, narrow_overflow || shift_overflow); } - } else if (pow10 > 0) { + auto const [v, o] = + checked_narrow_cast(guarded_right_shift(base2_value, -pow2)); + return maybe_with_overflow(v, o); + } + + if (pow10 > 0) { if (pow2 <= 0) { // Power-2/10 shifts both downward: order doesn't matter, apply and bail. - // Guard against shift being undefined behavior auto const shifted = guarded_right_shift(base2_value, -pow2); - return static_cast(divide_power10(shifted, pow10)); + auto const divided = divide_power10(shifted, pow10); + auto const [v, o] = checked_narrow_cast(divided); + return maybe_with_overflow(v, o); } - return shift_to_decimal_pospow(base2_value, pow2, pow10); - } else { // pow10 < 0 - if (pow2 >= 0) { - // Power-2/10 shifts both upward: order doesn't matter, apply and bail. - // NOTE: Either shift, multiply, or cast (e.g. double -> decimal32) can overflow! - auto const shifted = guarded_left_shift(static_cast(base2_value), pow2); - return multiply_power10(shifted, -pow10); - } - return shift_to_decimal_negpow(base2_value, pow2, pow10); + auto const [v, o] = + shift_to_decimal_pospow(base2_value, pow2, pow10); + return maybe_with_overflow(v, o); } + + // pow10 < 0 + if (pow2 >= 0) { + // Power-2/10 shifts both upward: order doesn't matter, apply and bail. + auto const [narrowed, narrow_overflow] = + checked_narrow_cast(base2_value); + auto const [shifted, shift_overflow] = checked_left_shift(narrowed, pow2); + auto const [scaled, multiply_overflow] = + multiply_power10_saturating(shifted, -pow10); + return maybe_with_overflow( + scaled, narrow_overflow || shift_overflow || multiply_overflow); + } + auto const [v, o] = + shift_to_decimal_negpow(base2_value, pow2, pow10); + return maybe_with_overflow(v, o); } /** * @brief Perform floating-point -> integer decimal conversion * * @tparam Rep The type of integer we are converting to, to store the decimal value + * @tparam CheckOverflow Whether to return overflow detection with the converted value * @tparam FloatingType The type of floating-point object we are converting from * @param floating The floating point value to convert * @param scale The desired base-10 scale factor: decimal value = returned value * 10^scale - * @return Integer representation of the floating-point value, given the desired scale + * @return Integer representation of the floating-point value, or `{value, overflow}` when + * `CheckOverflow` is true */ template )> -CUDF_HOST_DEVICE inline Rep convert_floating_to_integral(FloatingType const& floating, - scale_type const& scale) +CUDF_HOST_DEVICE inline auto convert_floating_to_integral(FloatingType const& floating, + scale_type const& scale) { // Extract components of the floating point number using converter = floating_converter; auto const integer_rep = converter::bit_cast_to_integer(floating); - if (converter::is_zero(integer_rep)) { return 0; } + if (converter::is_zero(integer_rep)) { return maybe_with_overflow(Rep{0}, false); } // Note that the significand here is an unsigned integer with sizeof(FloatingType) auto const is_negative = converter::get_is_negative(integer_rep); @@ -921,13 +1046,52 @@ CUDF_HOST_DEVICE inline Rep convert_floating_to_integral(FloatingType const& flo add_half_if_truncates(floating, significand, floating_pow2, pow10); // Apply the powers of 2 and 10 to convert to decimal. - auto const magnitude = - convert_floating_to_integral_shifting(base2_value, pow10, pow2); + auto const [magnitude_u, overflow] = [&] { + if constexpr (CheckOverflow) { + return convert_floating_to_integral_shifting( + base2_value, pow10, pow2); + } else { + auto const v = + convert_floating_to_integral_shifting(base2_value, pow10, pow2); + return cuda::std::pair{v, false}; + } + }(); + + // Reapply the sign. Negative range has one extra representable value for two's + // complement: magnitude == max+1 maps to min, and we can't get there by casting + // then negating (that would invoke signed overflow / undefined behavior). + using UnsignedRep = cuda::std::make_unsigned_t; + auto const umax = static_cast(cuda::std::numeric_limits::max()); + auto const umin_mag = umax + UnsignedRep{1}; + + if constexpr (!CheckOverflow) { + // Reapply sign with saturation on representational overflow. + if (!is_negative) { + if (magnitude_u > umax) { return cuda::std::numeric_limits::max(); } + return static_cast(magnitude_u); + } + + if (magnitude_u >= umin_mag) { return cuda::std::numeric_limits::min(); } - // Reapply the sign and return - // NOTE: Cast can overflow! - auto const signed_magnitude = static_cast(magnitude); - return is_negative ? -signed_magnitude : signed_magnitude; + return -static_cast(magnitude_u); + } else { + // Reapply sign with saturation on representational overflow. + if (!is_negative) { + if (magnitude_u > umax) { + return cuda::std::make_pair(cuda::std::numeric_limits::max(), true); + } + return cuda::std::make_pair(static_cast(magnitude_u), overflow); + } + + if (magnitude_u > umin_mag) { + return cuda::std::make_pair(cuda::std::numeric_limits::min(), true); + } + if (magnitude_u == umin_mag) { + return cuda::std::make_pair(cuda::std::numeric_limits::min(), overflow); + } + + return cuda::std::make_pair(-static_cast(magnitude_u), overflow); + } } /** diff --git a/cpp/include/cudf/fixed_point/detail/safe_arithmetic.hpp b/cpp/include/cudf/fixed_point/detail/safe_arithmetic.hpp new file mode 100644 index 000000000000..1bd02197e1ad --- /dev/null +++ b/cpp/include/cudf/fixed_point/detail/safe_arithmetic.hpp @@ -0,0 +1,299 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include +#include + +#include +#include + +/** + * @file safe_arithmetic.hpp + * @brief Overflow-aware free functions for `numeric::fixed_point`. + * + * These free functions intentionally live outside of `fixed_point` itself + * (mirroring how `floating_conversion.hpp` lives outside `fixed_point` for the + * float<->decimal conversion path). They take regular, non-tracking + * `fixed_point` operands and return a `safe_result` that pairs the computed + * value with an `overflow` flag describing whether the operation -- including + * any internal rescale -- would have wrapped the underlying integer storage. + * + * When overflow is detected the actual arithmetic is skipped: computing the + * wrapped result would be invalid (callers are not expected to consume a value + * whose `overflow` flag is set) and, for signed integer types, performing it + * would invoke undefined behavior. In that case `value` holds a defined + * placeholder (zero) and `overflow` is `true`. + * + * The intent is that overflow-aware code paths (the `binary_operation_safe` + * kernel today; specialized reduce/groupby aggregations in the future) compose + * these primitives explicitly, instead of relying on a sticky bit baked into + * the value type. + */ + +namespace CUDF_EXPORT numeric { +namespace detail { + +/** + * @brief Result of an overflow-checked `fixed_point` operation. + * + * @tparam Rep Storage type of the wrapped `fixed_point` value + * @tparam Rad Radix of the wrapped `fixed_point` value + */ +template +struct safe_result { + fixed_point value; ///< Computed `fixed_point` value + bool overflow; ///< Whether the producing operation overflowed +}; + +/** + * @brief Whether `shift(val, scale)` would incur signed-integer overflow + * + * Mirrors the overflow conditions of `multiplication_overflow` / + * `division_overflow` on the intermediate scale factor. + * + * @tparam Rep Representation type + * @tparam Rad Radix + * @tparam T Type of the value being shifted (typically `Rep`) + * @param val The value being shifted + * @param scale The amount to shift the value by + * @return true if the shift would overflow `Rep`, false otherwise + */ +template +CUDF_HOST_DEVICE inline constexpr bool shift_overflows(T const& val, scale_type const& scale) +{ + auto const v = static_cast(val); + if (scale == 0) { return false; } + if (scale > 0) { + Rep const divisor = ipow(static_cast(scale)); + return division_overflow(v, divisor); + } + Rep const multiplier = ipow(static_cast(-scale)); + return multiplication_overflow(v, multiplier); +} + +/** + * @brief Rescale a `fixed_point` value, reporting whether the underlying shift overflows + * + * Equivalent to `x.rescaled(new_scale)` but surfaces the shift overflow flag + * instead of dropping it. + * + * @tparam Rep Storage type + * @tparam Rad Radix + * @param x The value to rescale + * @param new_scale The target scale + * @return `{rescaled_value, overflow}`; `rescaled_value` is zero on overflow + */ +template +CUDF_HOST_DEVICE inline safe_result safe_rescaled(fixed_point x, + scale_type new_scale) +{ + if (new_scale == x.scale()) { return safe_result{x, false}; } + auto const scale_delta = scale_type{new_scale - x.scale()}; + // Skip the shift when it would overflow: the rescaled value is meaningless + // and performing the multiply would be signed-integer-overflow UB. + if (shift_overflows(x.value(), scale_delta)) { + return safe_result{fixed_point{scaled_integer{Rep{0}, new_scale}}, + true}; + } + Rep const value = shift(x.value(), scale_delta); + return safe_result{fixed_point{scaled_integer{value, new_scale}}, false}; +} + +/** + * @brief Overflow-checked addition of two `fixed_point` values + * + * Rescales both operands to the smaller of their two scales (matching + * `operator+`), then performs the add. The returned `overflow` flag is the + * disjunction of any rescale overflow and the integer add overflow. When + * overflow is detected the add is skipped and `value` is zero. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_add(fixed_point lhs, + fixed_point rhs) +{ + auto const common_scale = cuda::std::min(lhs.scale(), rhs.scale()); + auto const lhs_r = safe_rescaled(lhs, common_scale); + auto const rhs_r = safe_rescaled(rhs, common_scale); + Rep const lv = lhs_r.value.value(); + Rep const rv = rhs_r.value.value(); + bool const overflow = lhs_r.overflow || rhs_r.overflow || addition_overflow(lv, rv); + Rep const sum = overflow ? Rep{0} : lv + rv; + return safe_result{fixed_point{scaled_integer{sum, common_scale}}, + overflow}; +} + +/** + * @brief Overflow-checked subtraction of two `fixed_point` values + * + * When overflow is detected the subtract is skipped and `value` is zero. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_sub(fixed_point lhs, + fixed_point rhs) +{ + auto const common_scale = cuda::std::min(lhs.scale(), rhs.scale()); + auto const lhs_r = safe_rescaled(lhs, common_scale); + auto const rhs_r = safe_rescaled(rhs, common_scale); + Rep const lv = lhs_r.value.value(); + Rep const rv = rhs_r.value.value(); + bool const overflow = lhs_r.overflow || rhs_r.overflow || subtraction_overflow(lv, rv); + Rep const diff = overflow ? Rep{0} : lv - rv; + return safe_result{fixed_point{scaled_integer{diff, common_scale}}, + overflow}; +} + +/** + * @brief Overflow-checked multiplication of two `fixed_point` values + * + * No rescale is needed -- the result scale is `lhs.scale() + rhs.scale()`. + * When overflow is detected the multiply is skipped and `value` is zero. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_mul(fixed_point lhs, + fixed_point rhs) +{ + Rep const lv = lhs.value(); + Rep const rv = rhs.value(); + bool const overflow = multiplication_overflow(lv, rv); + Rep const prod = overflow ? Rep{0} : lv * rv; + scale_type const out_scale{lhs.scale() + rhs.scale()}; + return safe_result{fixed_point{scaled_integer{prod, out_scale}}, + overflow}; +} + +/** + * @brief Overflow-checked division of two `fixed_point` values + * + * Two failure modes are reported as overflow: `INT_MIN / -1` (caught by + * `division_overflow`) and division by zero. In both cases the divide is + * skipped -- so we never invoke signed-integer divide overflow or + * divide-by-zero UB -- and a zero-valued result is returned. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_div(fixed_point lhs, + fixed_point rhs) +{ + Rep const lv = lhs.value(); + Rep const rv = rhs.value(); + scale_type const out_scale{lhs.scale() - rhs.scale()}; + // Short-circuit on a zero divisor before touching `division_overflow` (which + // would itself divide) or the divide below. + bool const overflow = (rv == Rep{0}) || division_overflow(lv, rv); + Rep const quot = overflow ? Rep{0} : lv / rv; + return safe_result{fixed_point{scaled_integer{quot, out_scale}}, + overflow}; +} + +/** + * @brief Overflow-checked modulo of two `fixed_point` values + * + * The op-level failure modes are divide-by-zero and the `INT_MIN % -1` + * signed-overflow boundary; the rescale to the common scale can also overflow. + * All are OR'd into the returned flag, and whenever any of them is set the `%` + * is skipped and a zero-valued result is returned (so we never invoke + * `%`-by-zero or signed-overflow UB). + */ +template +CUDF_HOST_DEVICE inline safe_result safe_mod(fixed_point lhs, + fixed_point rhs) +{ + auto const common_scale = cuda::std::min(lhs.scale(), rhs.scale()); + auto const lhs_r = safe_rescaled(lhs, common_scale); + auto const rhs_r = safe_rescaled(rhs, common_scale); + Rep const lv = lhs_r.value.value(); + Rep const rv = rhs_r.value.value(); + bool const overflow = + lhs_r.overflow || rhs_r.overflow || (rv == Rep{0}) || division_overflow(lv, rv); + Rep const remainder = overflow ? Rep{0} : lv % rv; + return safe_result{fixed_point{scaled_integer{remainder, common_scale}}, + overflow}; +} + +/** + * @brief Overflow-checked positive modulo, matching `ops::PMod` semantics for decimals. + * + * Implements `rem = x % y; (rem < 0) ? (rem + y) % y : rem`. The intermediate + * `rem + y` can overflow even though `%` itself cannot. If the base modulo + * already overflowed (incl. divide-by-zero), or the correcting add overflows, + * the correction is skipped and a zero-valued result is returned. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_pmod(fixed_point lhs, + fixed_point rhs) +{ + auto const m = safe_mod(lhs, rhs); + if (m.overflow || !(m.value.value() < Rep{0})) { return m; } + + // `m.overflow` is false here, so `safe_mod` saw a non-zero divisor: `rv != 0`. + auto const rhs_r = safe_rescaled(rhs, m.value.scale()); + Rep const mv = m.value.value(); + Rep const rv = rhs_r.value.value(); + bool const overflow = rhs_r.overflow || addition_overflow(mv, rv); + Rep const corrected = overflow ? Rep{0} : (mv + rv) % rv; + return safe_result{ + fixed_point{scaled_integer{corrected, m.value.scale()}}, overflow}; +} + +/** + * @brief Overflow-checked Python-style modulo: `((x % y) + y) % y`. + * + * The intermediate add can overflow; the final `%` cannot. If the base modulo + * already overflowed (incl. divide-by-zero), or the correcting add overflows, + * the correction is skipped and a zero-valued result is returned, so we never + * invoke `%`-by-zero or signed-overflow UB. + */ +template +CUDF_HOST_DEVICE inline safe_result safe_pymod(fixed_point lhs, + fixed_point rhs) +{ + auto const m = safe_mod(lhs, rhs); + if (m.overflow) { return m; } + + // `m.overflow` is false here, so `safe_mod` saw a non-zero divisor: `rv != 0`. + auto const rhs_r = safe_rescaled(rhs, m.value.scale()); + Rep const mv = m.value.value(); + Rep const rv = rhs_r.value.value(); + bool const overflow = rhs_r.overflow || addition_overflow(mv, rv); + Rep const corrected = overflow ? Rep{0} : (mv + rv) % rv; + return safe_result{ + fixed_point{scaled_integer{corrected, m.value.scale()}}, overflow}; +} + +/** + * @brief Overflow-checked floating-point -> `fixed_point` conversion + * + * Uses `convert_floating_to_integral` (in + * `floating_conversion.hpp`) for base-10 decimals, which saturates and reports + * overflow. For base-2 radixes there is no checked path today, so `overflow` + * is always `false`. + * + * @tparam Fixed Target `fixed_point` instantiation + * @tparam Floating Source floating-point type + * @param floating The floating-point value to convert + * @param scale The desired scale of the result + * @return `{fixed_point_value, overflow}` + */ +template && cudf::is_fixed_point())> +CUDF_HOST_DEVICE inline safe_result safe_convert_floating_to_fixed( + Floating floating, scale_type scale) +{ + using Rep = typename Fixed::rep; + if constexpr (Fixed::rad == Radix::BASE_10) { + auto const [value, overflow] = convert_floating_to_integral(floating, scale); + return safe_result{Fixed{scaled_integer{value, scale}}, overflow}; + } else { + Rep const value = static_cast(shift(floating, scale)); + return safe_result{Fixed{scaled_integer{value, scale}}, false}; + } +} + +} // namespace detail +} // namespace CUDF_EXPORT numeric diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 987e8492c3ba..afabf3f6e552 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -663,10 +663,8 @@ CUDF_HOST_DEVICE inline fixed_point operator+(fixed_point(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value) && "fixed_point overflow"); - #endif return fixed_point{scaled_integer{sum, scale}}; @@ -681,10 +679,8 @@ CUDF_HOST_DEVICE inline fixed_point operator-(fixed_point(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value) && "fixed_point overflow"); - #endif return fixed_point{scaled_integer{diff, scale}}; @@ -696,9 +692,7 @@ CUDF_HOST_DEVICE inline fixed_point operator*(fixed_point const& rhs) { #if defined(__CUDACC_DEBUG__) - assert(!multiplication_overflow(lhs._value, rhs._value) && "fixed_point overflow"); - #endif return fixed_point{ @@ -711,9 +705,7 @@ CUDF_HOST_DEVICE inline fixed_point operator/(fixed_point const& rhs) { #if defined(__CUDACC_DEBUG__) - assert(!division_overflow(lhs._value, rhs._value) && "fixed_point overflow"); - #endif return fixed_point{ diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index bb4a48629ec5..a8a3c0eeb673 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -46,6 +46,7 @@ #include #include +#include namespace cudf { namespace binops { @@ -232,6 +233,82 @@ std::unique_ptr binary_operation(LhsType const& lhs, out->set_null_count(cudf::detail::null_count(out_view.null_mask(), 0, out->size(), stream)); return out; } + +namespace { +[[nodiscard]] constexpr bool is_decimal_safe_overflow_binary_operator(binary_operator op) +{ + switch (op) { + case binary_operator::ADD: + case binary_operator::SUB: + case binary_operator::MUL: + case binary_operator::DIV: + case binary_operator::MOD: + case binary_operator::PMOD: + case binary_operator::PYMOD: return true; + default: return false; + } +} +} // namespace + +/** + * @brief Like `binary_operation` for fixed-point decimal operands, but additionally returns a + * per-row BOOL8 overflow column. Element `i` of the overflow column is `true` iff row + * `i` is an active (non-null) row whose arithmetic or rescale overflowed. + */ +template +std::pair, std::unique_ptr> binary_operation_safe( + LhsType const& lhs, + RhsType const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(cudf::is_fixed_point(lhs.type()) && cudf::is_fixed_point(rhs.type()), + "binary_operation_safe requires both operands to use fixed-point types."); + CUDF_EXPECTS(is_decimal_safe_overflow_binary_operator(op), + "binary_operation_safe only supports ADD, SUB, MUL, DIV, MOD, PMOD, and PYMOD."); + + 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 (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)) + CUDF_FAIL("binary_operation_safe does not support string null min/max", cudf::logic_error); + + if (not cudf::binops::compiled::is_supported_operation(output_type, lhs.type(), rhs.type(), op)) + CUDF_FAIL("Unsupported operator for these types", cudf::data_type_error); + + cudf::binops::compiled::fixed_point_binary_operation_validation( + op, lhs.type(), rhs.type(), output_type); + + auto const common_intermediate = + cudf::binops::compiled::get_common_type(output_type, lhs.type(), rhs.type()); + CUDF_EXPECTS(common_intermediate.has_value(), + "binary_operation_safe could not determine a common intermediate type."); + CUDF_EXPECTS( + lhs.type().id() == rhs.type().id() && lhs.type().id() == common_intermediate->id(), + "binary_operation_safe requires lhs, rhs, and their common intermediate type to share the " + "same decimal storage type (e.g. DECIMAL64 operands with DECIMAL64 output)."); + + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); + auto overflow_col = make_fixed_width_column( + data_type{type_id::BOOL8}, out->size(), mask_state::UNALLOCATED, stream, mr); + + if constexpr (std::is_same_v) + if (lhs.is_empty()) return {std::move(out), std::move(overflow_col)}; + if constexpr (std::is_same_v) + if (rhs.is_empty()) return {std::move(out), std::move(overflow_col)}; + + auto out_view = out->mutable_view(); + mutable_column_view overflow_mv = overflow_col->mutable_view(); + cudf::binops::compiled::binary_operation_safe( + out_view, lhs, rhs, op, overflow_mv.head(), stream, mr); + out->set_null_count(cudf::detail::null_count(out_view.null_mask(), 0, out->size(), stream)); + return {std::move(out), std::move(overflow_col)}; +} } // namespace compiled } // namespace binops @@ -443,6 +520,45 @@ std::unique_ptr binary_operation(column_view const& lhs, return detail::binary_operation(lhs, rhs, op, output_type, stream, mr); } +std::pair, std::unique_ptr> binary_operation_safe( + scalar const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return cudf::binops::compiled::binary_operation_safe( + lhs, rhs, op, output_type, stream, mr); +} + +std::pair, std::unique_ptr> binary_operation_safe( + column_view const& lhs, + scalar const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return cudf::binops::compiled::binary_operation_safe( + lhs, rhs, op, output_type, stream, mr); +} + +std::pair, std::unique_ptr> binary_operation_safe( + column_view const& lhs, + column_view const& rhs, + binary_operator op, + data_type output_type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return cudf::binops::compiled::binary_operation_safe( + lhs, rhs, op, output_type, stream, mr); +} + std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, std::string const& ptx, diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index b6a5d342585d..10059e567b93 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -402,6 +402,42 @@ void binary_operation(mutable_column_view& out, operator_dispatcher(out, lhs, rhsv, false, true, op, stream); } +// Safe (decimal-overflow-tracking) overloads. The heavy lifting lives in +// `apply_binary_op_safe` (see `binary_ops_safe.cu`); these three thin wrappers only +// convert any scalar operand to a single-element `column_view` before forwarding. +void binary_operation_safe(mutable_column_view& out, + column_view const& lhs, + column_view const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + apply_binary_op_safe(out, lhs, rhs, false, false, op, d_overflow_per_row, stream, mr); +} +void binary_operation_safe(mutable_column_view& out, + scalar const& lhs, + column_view const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto [lhsv, aux] = scalar_to_column_view(lhs, stream); + apply_binary_op_safe(out, lhsv, rhs, true, false, op, d_overflow_per_row, stream, mr); +} +void binary_operation_safe(mutable_column_view& out, + column_view const& lhs, + scalar const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto [rhsv, aux] = scalar_to_column_view(rhs, stream); + apply_binary_op_safe(out, lhs, rhsv, false, true, op, d_overflow_per_row, stream, mr); +} + namespace detail { void apply_sorting_struct_binary_op(mutable_column_view& out, column_view const& lhs, diff --git a/cpp/src/binaryop/compiled/binary_ops.hpp b/cpp/src/binaryop/compiled/binary_ops.hpp index e1779cf5f967..fed842675b8c 100644 --- a/cpp/src/binaryop/compiled/binary_ops.hpp +++ b/cpp/src/binaryop/compiled/binary_ops.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -136,6 +136,56 @@ void binary_operation(mutable_column_view& out, binary_operator op, rmm::cuda_stream_view stream); +/** + * @brief Decimal fixed-point binary operation that records per-row overflow. + * + * Like `binary_operation` but only supports the seven base-10 decimal arithmetic operators + * (ADD, SUB, MUL, DIV, MOD, PMOD, PYMOD). For each row, writes `true` to + * @p d_overflow_per_row[i] iff row `i` is an active (non-null on both sides) row whose + * arithmetic or rescale to the output scale overflows; otherwise writes `false`. The + * @p d_overflow_per_row buffer is expected to point at @p out.size() bytes of device + * memory and is fully overwritten by the kernel. + * + * Both operands must be base-10 decimals of the same storage type as @p out. + */ +void binary_operation_safe(mutable_column_view& out, + scalar const& lhs, + column_view const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +void binary_operation_safe(mutable_column_view& out, + column_view const& lhs, + scalar const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); +void binary_operation_safe(mutable_column_view& out, + column_view const& lhs, + column_view const& rhs, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + +/** + * @brief Core decimal-safe column-column kernel. + * + * Defined in `binary_ops_safe.cu`. The three `binary_operation_safe` overloads above + * route here after converting any `scalar` operand to a single-element `column_view`. + */ +void apply_binary_op_safe(mutable_column_view& out, + column_view const& lhs, + column_view const& rhs, + bool is_lhs_scalar, + bool is_rhs_scalar, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr); + // Defined in util.cpp /** * @brief Get the common type among all input types. diff --git a/cpp/src/binaryop/compiled/binary_ops_safe.cu b/cpp/src/binaryop/compiled/binary_ops_safe.cu new file mode 100644 index 000000000000..ab31e9bac31d --- /dev/null +++ b/cpp/src/binaryop/compiled/binary_ops_safe.cu @@ -0,0 +1,260 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "binary_ops.hpp" + +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include + +namespace cudf { +namespace binops { +namespace compiled { +namespace { + +// One thin overflow-aware functor per supported binary operator. Each functor +// wraps the matching free function in `cudf/fixed_point/detail/safe_arithmetic.hpp` +// and returns the `safe_result` directly, so the kernel only needs to OR the +// op-level overflow with the rescale-to-output overflow before recording the +// global flag. +struct SafeAdd { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_add(lhs, rhs); + } +}; +struct SafeSub { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_sub(lhs, rhs); + } +}; +struct SafeMul { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_mul(lhs, rhs); + } +}; +struct SafeDiv { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_div(lhs, rhs); + } +}; +struct SafeMod { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_mod(lhs, rhs); + } +}; +struct SafePMod { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_pmod(lhs, rhs); + } +}; +struct SafePyMod { + template + __device__ __forceinline__ auto operator()(Decimal lhs, Decimal rhs) const + { + return numeric::detail::safe_pymod(lhs, rhs); + } +}; + +// Per-row functor that performs an overflow-checked decimal binary op. +// +// Both operands and the output share the same base-10 decimal storage type +// (caller-validated). Each thread loads the raw integer rep, wraps it in a +// regular `fixed_point` value (no sticky-flag layer), calls the matching +// `safe_*` free function, rescales to the output scale via `safe_rescaled`, +// and writes a per-row overflow bit (`true` iff active row && (op or rescale) +// overflowed) to the `d_overflow_per_row` buffer. +template +struct decimal_safe_op_kernel { + using Rep = typename Decimal::rep; + + mutable_column_device_view out; + column_device_view lhs; + column_device_view rhs; + bool is_lhs_scalar; + bool is_rhs_scalar; + numeric::scale_type out_scale; + bool* d_overflow_per_row; + + __device__ __forceinline__ void operator()(size_type i) const + { + auto const li = is_lhs_scalar ? 0 : i; + auto const ri = is_rhs_scalar ? 0 : i; + // Per cuDF convention, null payloads of fixed-width columns are unspecified + // and must not be read. Short-circuit so we never call `element` on a + // null row, and write defined neutral values into the output's null slots + // (the result column's null mask, set by the caller, is what makes them null). + if (!(lhs.is_valid(li) && rhs.is_valid(ri))) { + d_overflow_per_row[i] = false; + out.data()[i] = Rep{}; + return; + } + + auto const lscale = numeric::scale_type{lhs.type().scale()}; + auto const rscale = numeric::scale_type{rhs.type().scale()}; + + Decimal const x{numeric::scaled_integer{lhs.element(li), lscale}}; + Decimal const y{numeric::scaled_integer{rhs.element(ri), rscale}}; + + auto const op_res = SafeOp{}(x, y); + auto const rescaled_res = numeric::detail::safe_rescaled(op_res.value, out_scale); + + d_overflow_per_row[i] = op_res.overflow || rescaled_res.overflow; + out.data()[i] = rescaled_res.value.value(); + } +}; + +template +void launch_decimal_safe_kernel(mutable_column_device_view& outd, + column_device_view const& lhsd, + column_device_view const& rhsd, + bool is_lhs_scalar, + bool is_rhs_scalar, + size_type n, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto const out_scale = numeric::scale_type{outd.type().scale()}; + decimal_safe_op_kernel kern{ + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, out_scale, d_overflow_per_row}; + thrust::for_each_n( + rmm::exec_policy_nosync(stream, mr), cuda::counting_iterator{0}, n, kern); +} + +template +void dispatch_op_and_run(mutable_column_device_view& outd, + column_device_view const& lhsd, + column_device_view const& rhsd, + bool is_lhs_scalar, + bool is_rhs_scalar, + size_type n, + bool* d_overflow_per_row, + binary_operator op, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + switch (op) { + case binary_operator::ADD: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::SUB: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::MUL: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::DIV: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::MOD: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::PMOD: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + case binary_operator::PYMOD: + launch_decimal_safe_kernel( + outd, lhsd, rhsd, is_lhs_scalar, is_rhs_scalar, n, d_overflow_per_row, stream, mr); + break; + default: + CUDF_FAIL("binary_operation_safe only supports ADD, SUB, MUL, DIV, MOD, PMOD, and PYMOD."); + } +} + +} // namespace + +void apply_binary_op_safe(mutable_column_view& out, + column_view const& lhs, + column_view const& rhs, + bool is_lhs_scalar, + bool is_rhs_scalar, + binary_operator op, + bool* d_overflow_per_row, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(d_overflow_per_row != nullptr, + "binary_operation_safe requires a non-null device per-row overflow buffer."); + CUDF_EXPECTS( + lhs.type().id() == rhs.type().id() && lhs.type().id() == out.type().id(), + "binary_operation_safe requires lhs/rhs/out to share the same decimal storage type."); + + if (out.size() == 0) { return; } + + auto lhsd = column_device_view::create(lhs, stream); + auto rhsd = column_device_view::create(rhs, stream); + auto outd = mutable_column_device_view::create(out, stream); + + switch (out.type().id()) { + case type_id::DECIMAL32: + dispatch_op_and_run(*outd, + *lhsd, + *rhsd, + is_lhs_scalar, + is_rhs_scalar, + out.size(), + d_overflow_per_row, + op, + stream, + mr); + break; + case type_id::DECIMAL64: + dispatch_op_and_run(*outd, + *lhsd, + *rhsd, + is_lhs_scalar, + is_rhs_scalar, + out.size(), + d_overflow_per_row, + op, + stream, + mr); + break; + case type_id::DECIMAL128: + dispatch_op_and_run(*outd, + *lhsd, + *rhsd, + is_lhs_scalar, + is_rhs_scalar, + out.size(), + d_overflow_per_row, + op, + stream, + mr); + break; + default: CUDF_FAIL("binary_operation_safe requires a base-10 decimal output type."); + } +} + +} // namespace compiled +} // namespace binops +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index e709a52ace20..3c03ecfb0526 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -261,6 +261,7 @@ ConfigureTest(CLAMP_TEST replace/clamp_test.cpp) # ################################################################################################## # * fixed_point tests ----------------------------------------------------------------------------- ConfigureTest(FIXED_POINT_TEST fixed_point/fixed_point_tests.cpp fixed_point/fixed_point_tests.cu) +ConfigureTest(SAFE_ARITHMETIC_TEST fixed_point/safe_arithmetic_tests.cpp) # ################################################################################################## # * unary tests ----------------------------------------------------------------------------------- diff --git a/cpp/tests/binaryop/binop-compiled-fixed_point-test.cpp b/cpp/tests/binaryop/binop-compiled-fixed_point-test.cpp index 3da2aff0157a..cba52e645216 100644 --- a/cpp/tests/binaryop/binop-compiled-fixed_point-test.cpp +++ b/cpp/tests/binaryop/binop-compiled-fixed_point-test.cpp @@ -9,6 +9,7 @@ #include #include +#include #include #include #include @@ -17,6 +18,10 @@ #include +#include +#include +#include + template struct FixedPointCompiledTest : public cudf::test::BaseFixture {}; @@ -926,3 +931,192 @@ TYPED_TEST(FixedPointCompiledTest, FixedPointWithFloating) test_fixed_floating(cudf::binary_operator::NULL_MAX, 4.0, 20, -1, decimal_result); test_fixed_floating(cudf::binary_operator::NULL_MIN, 4.0, 200, -1, decimal_result); } + +TEST(BinaryOperationSafeDecimal, mulNoOverflowFlagZero) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + auto const lhs = fp_wrapper{{11, 22}, scale_type{-1}}; + auto const rhs = fp_wrapper{{10, 10}, scale_type{0}}; + auto const expected = fp_wrapper{{110, 220}, scale_type{-1}}; + auto const expected_overflow = wrapper{{false, false}}; + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::MUL, type, stream, mr); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} + +TEST(BinaryOperationSafeDecimal, mulOverflowSetsGlobalFlag) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + // Product of scaled integers overflows int64; sticky overflow path should mark the row. + Rep const a = (std::numeric_limits::max() / 2) + 1; + Rep const b = 3; + auto const lhs = fp_wrapper{{a}, scale_type{0}}; + auto const rhs = fp_wrapper{{b}, scale_type{0}}; + auto const expected_overflow = wrapper{{true}}; + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::MUL, type, stream, mr); + + EXPECT_EQ(1, result->size()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} + +TEST(BinaryOperationSafeDecimal, mulEmptyInput) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + auto const lhs = fp_wrapper{{}, scale_type{0}}; + auto const rhs = fp_wrapper{{}, scale_type{0}}; + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::MUL, type, stream, mr); + + EXPECT_EQ(0, result->size()); + EXPECT_EQ(0, overflow->size()); + EXPECT_EQ(cudf::type_id::BOOL8, overflow->type().id()); +} + +TEST(BinaryOperationSafeDecimal, mulNullRowsKeepOverflowFalse) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + Rep const big = (std::numeric_limits::max() / 2) + 1; + // Row 0: valid + valid, overflowing (big * 3 wraps int64) -> overflow=true + // Row 1: lhs null, rhs valid, would have overflowed -> overflow=false + // Row 2: valid + valid, overflowing -> overflow=true + // Row 3: lhs valid, rhs null, would have overflowed -> overflow=false + auto const lhs = fp_wrapper{{big, big, big, big}, {1, 0, 1, 1}, scale_type{0}}; + auto const rhs = fp_wrapper{{3, 3, 3, 3}, {1, 1, 1, 0}, scale_type{0}}; + auto const expected_overflow = wrapper{{true, false, true, false}}; + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::MUL, type, stream, mr); + + EXPECT_EQ(4, result->size()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} + +TEST(BinaryOperationSafeDecimal, mulSlicedColumns) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + Rep const big = (std::numeric_limits::max() / 2) + 1; + // Build 6-row columns and slice [1, 4) so that only the slice is fed to the kernel. + // Row layout (full): [1*1, big*3, 2*2, big*3, 3*3, big*3] + // Slice rows 1..3: [ big*3, 2*2, big*3 ] -> overflow=[true,false,true] + auto const full_lhs = fp_wrapper{{Rep{1}, big, Rep{2}, big, Rep{3}, big}, scale_type{0}}; + auto const full_rhs = + fp_wrapper{{Rep{1}, Rep{3}, Rep{2}, Rep{3}, Rep{3}, Rep{3}}, scale_type{0}}; + auto const lhs_slice = + cudf::slice(static_cast(full_lhs), std::vector{1, 4})[0]; + auto const rhs_slice = + cudf::slice(static_cast(full_rhs), std::vector{1, 4})[0]; + + auto const expected_overflow = wrapper{{true, false, true}}; + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::MUL, lhs_slice.type(), rhs_slice.type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs_slice, rhs_slice, cudf::binary_operator::MUL, type, stream, mr); + + EXPECT_EQ(3, result->size()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} + +TEST(BinaryOperationSafeDecimal, mulMultiBlockSize) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + Rep const big = (std::numeric_limits::max() / 2) + 1; + auto const sz = std::size_t{10'000}; // spans many CUDA blocks + auto const stride = std::size_t{137}; // sparse overflow pattern + + std::vector lhs_data(sz, Rep{1}); + std::vector rhs_data(sz, Rep{1}); + std::vector expected_overflow_data(sz, false); + for (std::size_t i = 0; i < sz; i += stride) { + lhs_data[i] = big; + rhs_data[i] = Rep{3}; + expected_overflow_data[i] = true; + } + + auto const lhs = fp_wrapper(lhs_data.begin(), lhs_data.end(), scale_type{0}); + auto const rhs = fp_wrapper(rhs_data.begin(), rhs_data.end(), scale_type{0}); + auto const expected_overflow = + wrapper(expected_overflow_data.begin(), expected_overflow_data.end()); + + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::MUL, type, stream, mr); + + EXPECT_EQ(static_cast(sz), result->size()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} + +TEST(BinaryOperationSafeDecimal, divByZeroSetsOverflowAndAvoidsUB) +{ + using namespace numeric; + using Rep = cudf::device_storage_type_t; + // Mix in a zero divisor; CodeRabbit's critical finding requires we never + // execute `/ 0` and that the row is flagged. + auto const lhs = fp_wrapper{{Rep{10}, Rep{20}, Rep{30}}, scale_type{0}}; + auto const rhs = fp_wrapper{{Rep{2}, Rep{0}, Rep{5}}, scale_type{0}}; + auto const expected_overflow = wrapper{{false, true, false}}; + + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::DIV, + static_cast(lhs).type(), + static_cast(rhs).type()); + + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const [result, overflow] = + cudf::binary_operation_safe(lhs, rhs, cudf::binary_operator::DIV, type, stream, mr); + + EXPECT_EQ(3, result->size()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_overflow, overflow->view()); +} diff --git a/cpp/tests/fixed_point/safe_arithmetic_tests.cpp b/cpp/tests/fixed_point/safe_arithmetic_tests.cpp new file mode 100644 index 000000000000..f75b653f4a04 --- /dev/null +++ b/cpp/tests/fixed_point/safe_arithmetic_tests.cpp @@ -0,0 +1,669 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +using namespace numeric; + +struct SafeArithmeticTest : public cudf::test::BaseFixture {}; +struct FloatingConversionOverflowTest : public cudf::test::BaseFixture {}; + +// --------------------------------------------------------------------------- +// safe_add +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, AddNoOverflow) +{ + decimal64 const a{scaled_integer{1, scale_type{0}}}; + decimal64 const b{scaled_integer{2, scale_type{0}}}; + auto const r = detail::safe_add(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{3}); +} + +TEST_F(SafeArithmeticTest, AddOverflowSetsFlag) +{ + auto constexpr near_max = std::numeric_limits::max() - 100; + decimal64 const a{scaled_integer{near_max, scale_type{0}}}; + decimal64 const b{scaled_integer{200, scale_type{0}}}; + auto const r = detail::safe_add(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, AddRescaleOverflowSetsFlag) +{ + // Rescaling lhs from scale 0 down to scale -3 multiplies by 10^3; near-max value overflows. + decimal64 const a{ + scaled_integer{std::numeric_limits::max() / 2, scale_type{0}}}; + decimal64 const b{scaled_integer{0, scale_type{-3}}}; + auto const r = detail::safe_add(a, b); + EXPECT_TRUE(r.overflow); +} + +// --------------------------------------------------------------------------- +// safe_sub +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, SubNoOverflow) +{ + decimal64 const a{scaled_integer{10, scale_type{0}}}; + decimal64 const b{scaled_integer{3, scale_type{0}}}; + auto const r = detail::safe_sub(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{7}); +} + +TEST_F(SafeArithmeticTest, SubOverflowSetsFlag) +{ + auto constexpr near_min = std::numeric_limits::min() + 100; + decimal64 const a{scaled_integer{near_min, scale_type{0}}}; + decimal64 const b{scaled_integer{200, scale_type{0}}}; + auto const r = detail::safe_sub(a, b); + EXPECT_TRUE(r.overflow); +} + +// --------------------------------------------------------------------------- +// safe_mul +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, MulNoOverflow) +{ + decimal64 const a{scaled_integer{2, scale_type{0}}}; + decimal64 const b{scaled_integer{3, scale_type{0}}}; + auto const r = detail::safe_mul(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{6}); +} + +TEST_F(SafeArithmeticTest, MulOverflowSetsFlag) +{ + decimal64 const a{scaled_integer{1'000'000'000'000LL, scale_type{0}}}; + decimal64 const b{scaled_integer{1'000'000'000'000LL, scale_type{0}}}; + auto const r = detail::safe_mul(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Mul32OverflowSetsFlag) +{ + decimal32 const a{scaled_integer{1'000'000, scale_type{0}}}; + decimal32 const b{scaled_integer{1'000'000, scale_type{0}}}; + auto const r = detail::safe_mul(a, b); + EXPECT_TRUE(r.overflow); +} + +// --------------------------------------------------------------------------- +// safe_div +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, DivNoOverflow) +{ + decimal64 const a{scaled_integer{10, scale_type{0}}}; + decimal64 const b{scaled_integer{2, scale_type{0}}}; + auto const r = detail::safe_div(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{5}); +} + +TEST_F(SafeArithmeticTest, DivIntMinByNegativeOneOverflows) +{ + // INT64_MIN / -1 is the canonical signed-integer division overflow. + decimal64 const a{scaled_integer{std::numeric_limits::min(), scale_type{0}}}; + decimal64 const b{scaled_integer{-1, scale_type{0}}}; + auto const r = detail::safe_div(a, b); + EXPECT_TRUE(r.overflow); +} + +// --------------------------------------------------------------------------- +// safe_mod / safe_pmod / safe_pymod +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, ModNoOverflow) +{ + decimal64 const a{scaled_integer{7, scale_type{0}}}; + decimal64 const b{scaled_integer{3, scale_type{0}}}; + auto const r = detail::safe_mod(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{1}); +} + +TEST_F(SafeArithmeticTest, PModMatchesOpsPModSemantics) +{ + // -7 % 3 should be -1 from %, then pmod produces 2. + decimal64 const a{scaled_integer{-7, scale_type{0}}}; + decimal64 const b{scaled_integer{3, scale_type{0}}}; + auto const r = detail::safe_pmod(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{2}); +} + +TEST_F(SafeArithmeticTest, PyModMatchesOpsPyModSemantics) +{ + // ((x % y) + y) % y == 2 for x = -7, y = 3 + decimal64 const a{scaled_integer{-7, scale_type{0}}}; + decimal64 const b{scaled_integer{3, scale_type{0}}}; + auto const r = detail::safe_pymod(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{2}); +} + +// --------------------------------------------------------------------------- +// safe_rescaled +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, RescaledNoOpDoesNotSetFlag) +{ + decimal64 const a{ + scaled_integer{std::numeric_limits::max() / 2, scale_type{0}}}; + auto const r = detail::safe_rescaled(a, scale_type{0}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::max() / 2); +} + +TEST_F(SafeArithmeticTest, RescaledShiftOverflowSetsFlag) +{ + decimal64 const a{ + scaled_integer{std::numeric_limits::max() / 2, scale_type{0}}}; + // Rescaling to a more negative scale multiplies by a power of 10 and overflows. + auto const r = detail::safe_rescaled(a, scale_type{-3}); + EXPECT_TRUE(r.overflow); +} + +// --------------------------------------------------------------------------- +// safe_convert_floating_to_fixed +// --------------------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal32DetectsPositiveOverflow) +{ + auto const r = detail::safe_convert_floating_to_fixed(1e20, scale_type{0}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::max()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal32DetectsNegativeOverflow) +{ + auto const r = detail::safe_convert_floating_to_fixed(-1e20, scale_type{0}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::min()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal64DetectsPositiveOverflowViaScale) +{ + // Overflow via scale factor multiplication even for a "moderate" input. + // scale -19 implies multiplying by 10^19 in the decimal rep. + auto const r = detail::safe_convert_floating_to_fixed(1.0, scale_type{-19}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::max()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal64NoOverflow) +{ + auto const r = detail::safe_convert_floating_to_fixed(123.456, scale_type{-3}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{123456}); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal64DetectsNegativeOverflowViaScale) +{ + // Mirrors the positive-overflow-via-scale test for the negative branch: + // -1.0 at scale -19 implies multiplying by 10^19, which exceeds INT64 range. + auto const r = detail::safe_convert_floating_to_fixed(-1.0, scale_type{-19}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::min()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatToDecimal32NoOverflow) +{ + // 1.5f is exactly representable in float, and 1.5 at scale -1 is the integer 15. + auto const r = detail::safe_convert_floating_to_fixed(1.5f, scale_type{-1}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int32_t{15}); +} + +TEST_F(SafeArithmeticTest, ConvertFloatToDecimal32DetectsOverflow) +{ + // Float overflow path exercises the 32-bit ShiftingRep branch. + auto const r = detail::safe_convert_floating_to_fixed(1e20f, scale_type{0}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::max()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal128DetectsOverflow) +{ + // INT128_MAX is roughly 1.7e38; 1.0 at scale -39 implies 10^39 and overflows. + auto const r = detail::safe_convert_floating_to_fixed(1.0, scale_type{-39}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits<__int128_t>::max()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal128NoOverflow) +{ + // 1.0 at scale -30 is 10^30, which fits in INT128. + auto const r = detail::safe_convert_floating_to_fixed(1.0, scale_type{-30}); + EXPECT_FALSE(r.overflow); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal64PositiveScaleNoOverflow) +{ + // pow10 > 0 path: 12300.0 at scale +2 is 12300 / 10^2 = 123. + auto const r = detail::safe_convert_floating_to_fixed(12300.0, scale_type{2}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{123}); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingToDecimal64PositiveScaleDetectsOverflow) +{ + // pow10 > 0 path with input far above INT64*10^pow10: + // 1e25 / 10^5 = 1e20 > INT64_MAX (~9.2e18). + auto const r = detail::safe_convert_floating_to_fixed(1e25, scale_type{5}); + EXPECT_TRUE(r.overflow); + EXPECT_EQ(r.value.value(), std::numeric_limits::max()); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingZeroIsExact) +{ + auto const r = detail::safe_convert_floating_to_fixed(0.0, scale_type{-3}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{0}); +} + +TEST_F(SafeArithmeticTest, ConvertFloatingNegativeZeroIsExact) +{ + auto const r = detail::safe_convert_floating_to_fixed(-0.0, scale_type{-3}); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), int64_t{0}); +} + +// --------------------------------------------------------------------------- +// Composability: callers OR the overflow flags from a chain explicitly. +// --------------------------------------------------------------------------- + +// --------------------------------------------------------------------------- +// Cross-width overflow coverage: decimal32 / decimal64 / decimal128 +// +// `safe_*` operations preserve the input `Rep`; they do not upcast the +// representation type. These compile-time assertions verify that contract, +// and the runtime tests exercise overflow detection on each width. +// --------------------------------------------------------------------------- + +namespace { + +// `safe_result` returned by each operation must keep the operand `Rep`/`Rad`. +// Use `declval` to keep the check in an unevaluated context (the `fixed_point` +// default constructor is not `constexpr`). +template +constexpr bool safe_ops_preserve_rep() +{ + using Rep = typename Fixed::rep; + constexpr auto Rad = Fixed::rad; + using Expected = numeric::detail::safe_result; + return std::is_same_v(), std::declval())), + Expected> && + std::is_same_v(), std::declval())), + Expected> && + std::is_same_v(), std::declval())), + Expected> && + std::is_same_v(), std::declval())), + Expected>; +} + +static_assert(safe_ops_preserve_rep(), + "safe_* over decimal32 must preserve int32_t representation"); +static_assert(safe_ops_preserve_rep(), + "safe_* over decimal64 must preserve int64_t representation"); +static_assert(safe_ops_preserve_rep(), + "safe_* over decimal128 must preserve __int128_t representation"); + +} // namespace + +// --- decimal32 ------------------------------------------------------------- + +TEST_F(SafeArithmeticTest, Add32OverflowSetsFlag) +{ + auto constexpr near_max = std::numeric_limits::max() - 100; + decimal32 const a{scaled_integer{near_max, scale_type{0}}}; + decimal32 const b{scaled_integer{200, scale_type{0}}}; + auto const r = detail::safe_add(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Sub32OverflowSetsFlag) +{ + auto constexpr near_min = std::numeric_limits::min() + 100; + decimal32 const a{scaled_integer{near_min, scale_type{0}}}; + decimal32 const b{scaled_integer{200, scale_type{0}}}; + auto const r = detail::safe_sub(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Div32IntMinByNegativeOneOverflows) +{ + decimal32 const a{scaled_integer{std::numeric_limits::min(), scale_type{0}}}; + decimal32 const b{scaled_integer{-1, scale_type{0}}}; + auto const r = detail::safe_div(a, b); + EXPECT_TRUE(r.overflow); +} + +// --- decimal128 ------------------------------------------------------------ + +TEST_F(SafeArithmeticTest, Add128OverflowSetsFlag) +{ + auto const near_max = std::numeric_limits<__int128_t>::max() - __int128_t{100}; + decimal128 const a{scaled_integer<__int128_t>{near_max, scale_type{0}}}; + decimal128 const b{scaled_integer<__int128_t>{200, scale_type{0}}}; + auto const r = detail::safe_add(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Sub128OverflowSetsFlag) +{ + auto const near_min = std::numeric_limits<__int128_t>::min() + __int128_t{100}; + decimal128 const a{scaled_integer<__int128_t>{near_min, scale_type{0}}}; + decimal128 const b{scaled_integer<__int128_t>{200, scale_type{0}}}; + auto const r = detail::safe_sub(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Mul128OverflowSetsFlag) +{ + // 10^20 * 10^20 = 10^40, which exceeds INT128_MAX (~1.7e38). + auto const ten_pow_20 = static_cast<__int128_t>(1'000'000'000'000'000'000LL) * __int128_t{100}; + decimal128 const a{scaled_integer<__int128_t>{ten_pow_20, scale_type{0}}}; + decimal128 const b{scaled_integer<__int128_t>{ten_pow_20, scale_type{0}}}; + auto const r = detail::safe_mul(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, Mul128NoOverflow) +{ + decimal128 const a{scaled_integer<__int128_t>{__int128_t{12345}, scale_type{0}}}; + decimal128 const b{scaled_integer<__int128_t>{__int128_t{6789}, scale_type{0}}}; + auto const r = detail::safe_mul(a, b); + EXPECT_FALSE(r.overflow); + EXPECT_EQ(r.value.value(), __int128_t{12345} * __int128_t{6789}); +} + +TEST_F(SafeArithmeticTest, Div128IntMinByNegativeOneOverflows) +{ + decimal128 const a{ + scaled_integer<__int128_t>{std::numeric_limits<__int128_t>::min(), scale_type{0}}}; + decimal128 const b{scaled_integer<__int128_t>{-1, scale_type{0}}}; + auto const r = detail::safe_div(a, b); + EXPECT_TRUE(r.overflow); +} + +TEST_F(SafeArithmeticTest, CallerComposesOverflowAcrossChain) +{ + // (a * b) overflows; the caller is responsible for OR'ing flags as the + // chain progresses. + decimal64 const a{scaled_integer{1'000'000'000'000LL, scale_type{0}}}; + decimal64 const b{scaled_integer{1'000'000'000'000LL, scale_type{0}}}; + decimal64 const c{scaled_integer{0, scale_type{0}}}; + + auto const m1 = detail::safe_mul(a, b); // overflow here + auto const m2 = detail::safe_add(m1.value, c); // add itself doesn't overflow + bool const composed_overflow = m1.overflow || m2.overflow; + EXPECT_TRUE(composed_overflow); +} + +// --------------------------------------------------------------------------- +// Floating <-> decimal conversion overflow primitives +// +// `safe_convert_floating_to_fixed` (above) exercises the conversion path +// end-to-end, but the lower-level `CheckOverflow=true` primitives in +// `floating_conversion.hpp` (`checked_left_shift`, `multiply_power10_saturating`, +// `checked_narrow_cast`) and the top-level `convert_floating_to_integral` +// gain little direct test coverage from that. The tests below pin down each +// primitive's contract: on overflow, return the saturated max with the flag +// set; otherwise return the exact result with the flag clear. +// --------------------------------------------------------------------------- + +// --- checked_left_shift(value, bit_shift) ---------------------------- + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftZeroIsIdentity) +{ + auto const [v, ovf] = numeric::detail::checked_left_shift(uint64_t{42}, 0); + EXPECT_EQ(v, uint64_t{42}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftWithinRange) +{ + auto const [v, ovf] = numeric::detail::checked_left_shift(uint64_t{0x1234}, 4); + EXPECT_EQ(v, uint64_t{0x12340}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftAtBoundaryDoesNotOverflow) +{ + // value == max >> bit_shift is the largest input that still fits after the shift. + constexpr int bit_shift = 8; + auto const at_boundary = std::numeric_limits::max() >> bit_shift; + auto const [v, ovf] = numeric::detail::checked_left_shift(at_boundary, bit_shift); + EXPECT_EQ(v, static_cast(at_boundary << bit_shift)); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftJustOverBoundarySaturates) +{ + constexpr int bit_shift = 8; + auto const just_over = (std::numeric_limits::max() >> bit_shift) + 1u; + auto const [v, ovf] = numeric::detail::checked_left_shift(just_over, bit_shift); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftTooLargeBitShiftSaturates) +{ + // Shifting by >= digits is UB on the underlying operator and saturates here. + auto const [v, ovf] = + numeric::detail::checked_left_shift(uint64_t{1}, std::numeric_limits::digits); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftNegativeBitShiftSaturates) +{ + // Negative shifts aren't expected from callers; the checked path treats them + // as overflow rather than falling into the underlying UB path. + auto const [v, ovf] = numeric::detail::checked_left_shift(uint64_t{1}, -1); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedLeftShiftUncheckedNeverOverflows) +{ + // CheckOverflow=false (the default) must always report overflow=false. + auto const [v, ovf] = numeric::detail::checked_left_shift(uint64_t{0xFFFF}, 60); + EXPECT_FALSE(ovf); + // value is whatever guarded_left_shift returns; only the flag is tested here. + (void)v; +} + +// --- multiply_power10_saturating(value, pow10) -------------------- + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10ZeroPow10IsIdentity) +{ + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(123, 0); + EXPECT_EQ(v, uint64_t{123}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10NegativePow10IsIdentity) +{ + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(123, -5); + EXPECT_EQ(v, uint64_t{123}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10WithinRange) +{ + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(123, 6); + EXPECT_EQ(v, uint64_t{123'000'000}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10AtBoundaryDoesNotOverflow) +{ + // value == max / 10^pow10 is the largest input that fits after the multiply. + // For uint32_t and pow10=3: 4294967295 / 1000 == 4294967. + uint32_t const at_boundary = std::numeric_limits::max() / 1000u; + auto const [v, ovf] = + numeric::detail::multiply_power10_saturating(at_boundary, 3); + EXPECT_EQ(v, at_boundary * 1000u); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10JustOverBoundarySaturates) +{ + uint32_t const just_over = (std::numeric_limits::max() / 1000u) + 1u; + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(just_over, 3); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10HugePow10SaturatesNonzero) +{ + // 10^pow10 itself doesn't fit in uint32_t (10^10 > 2^32); divide_power10 returns 0, + // so the threshold is 0 and any nonzero value saturates. + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(1, 12); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, MultiplyPower10ZeroValueNeverOverflows) +{ + auto const [v, ovf] = numeric::detail::multiply_power10_saturating(0, 12); + EXPECT_EQ(v, uint32_t{0}); + EXPECT_FALSE(ovf); +} + +// --- checked_narrow_cast(value) ---------------------------------- + +TEST_F(FloatingConversionOverflowTest, CheckedNarrowCastFits) +{ + auto const [v, ovf] = numeric::detail::checked_narrow_cast(uint64_t{12345}); + EXPECT_EQ(v, uint32_t{12345}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedNarrowCastAtBoundaryDoesNotOverflow) +{ + auto const at_max = static_cast(std::numeric_limits::max()); + auto const [v, ovf] = numeric::detail::checked_narrow_cast(at_max); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedNarrowCastOverflowSaturates) +{ + auto const just_over = static_cast(std::numeric_limits::max()) + 1u; + auto const [v, ovf] = numeric::detail::checked_narrow_cast(just_over); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, CheckedNarrowCastUncheckedNeverOverflows) +{ + // CheckOverflow=false truncates without ever setting the flag. + auto const big = uint64_t{0x1'0000'0000ULL}; + auto const [v, ovf] = numeric::detail::checked_narrow_cast(big); + EXPECT_EQ(v, uint32_t{0}); // low 32 bits + EXPECT_FALSE(ovf); +} + +// --- convert_floating_to_integral(floating, scale) -------------- +// +// The end-to-end conversion entry point. These tests target paths that the +// `safe_convert_floating_to_fixed` wrapper tests above don't exercise directly, +// in particular the sign-reapply branches that distinguish INT_MIN-exact from +// representational overflow. + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralZero) +{ + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(0.0, scale_type{-3}); + EXPECT_EQ(v, int64_t{0}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralPositiveFits) +{ + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(123.456, scale_type{-3}); + EXPECT_EQ(v, int64_t{123456}); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralPositiveSaturatesToMax) +{ + // |1e20| >> INT32_MAX (~2.15e9) so the sign-reapply path saturates to max with overflow=true. + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(1e20, scale_type{0}); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralNegativeIntMinExactDoesNotOverflow) +{ + // INT32_MIN == -2^31 = -2147483648 is exactly representable in IEEE 754 double. + // magnitude_u == umin_mag, which the checked path maps to {INT32_MIN, overflow=shifting_ovf}. + // No shifting overflow occurs for this input, so the flag should remain false. + auto const int32_min_as_double = static_cast(std::numeric_limits::min()); + auto const [v, ovf] = numeric::detail::convert_floating_to_integral( + int32_min_as_double, scale_type{0}); + EXPECT_EQ(v, std::numeric_limits::min()); + EXPECT_FALSE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralNegativeJustBelowIntMinSaturates) +{ + // One step past INT64_MIN: magnitude_u > umin_mag -> {INT64_MIN, true}. + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(-1e20, scale_type{0}); + EXPECT_EQ(v, std::numeric_limits::min()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralOverflowViaNegativeScale) +{ + // Scale -19 multiplies by 10^19 in the decimal rep, which exceeds INT64 range. + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(1.0, scale_type{-19}); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralOverflowViaPositiveScale) +{ + // pow10 > 0 path: 1e25 / 10^5 == 1e20 > INT64_MAX (~9.2e18). + auto const [v, ovf] = + numeric::detail::convert_floating_to_integral(1e25, scale_type{5}); + EXPECT_EQ(v, std::numeric_limits::max()); + EXPECT_TRUE(ovf); +} + +TEST_F(FloatingConversionOverflowTest, ConvertIntegralUncheckedSameValueAsChecked) +{ + // The unchecked path must produce the same Rep value as the checked path's + // first element for inputs that don't overflow (no regression in the + // unchecked code path after the refactor). + double const x = 123.456; + auto const checked = + numeric::detail::convert_floating_to_integral(x, scale_type{-3}); + auto const unchecked = numeric::detail::convert_floating_to_integral(x, scale_type{-3}); + EXPECT_EQ(checked.first, unchecked); +} + +CUDF_TEST_PROGRAM_MAIN()