Add decimal division functionality with scale preservation - #22261
Open
simoneves wants to merge 4 commits into
Open
Add decimal division functionality with scale preservation#22261simoneves wants to merge 4 commits into
simoneves wants to merge 4 commits into
Conversation
simoneves
requested review from
galipremsagar,
karthikeyann,
mroeschke and
ttnghia
April 22, 2026 23:48
github-actions
Bot
added
libcudf
simoneves
force-pushed
the
simoneves/a-hirota/decimal-division-operations-rebased
branch
from
April 23, 2026 00:53
aefd0bd to
e3bbe5b
Compare
Contributor
3 tasks
mhaseeb123
changed the title
mhaseeb123
requested changes
May 18, 2026
Comment on lines
+1
to
+15
| /* | ||
| * Copyright (c) 2025, NVIDIA CORPORATION. | ||
| * | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ |
Contributor
There was a problem hiding this comment.
Suggested change
| /* | |
| * Copyright (c) 2025, NVIDIA CORPORATION. | |
| * | |
| * Licensed under the Apache License, Version 2.0 (the "License"); | |
| * you may not use this file except in compliance with the License. | |
| * You may obtain a copy of the License at | |
| * | |
| * http://www.apache.org/licenses/LICENSE-2.0 | |
| * | |
| * Unless required by applicable law or agreed to in writing, software | |
| * distributed under the License is distributed on an "AS IS" BASIS, | |
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | |
| * See the License for the specific language governing permissions and | |
| * limitations under the License. | |
| */ | |
| /* | |
| * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. | |
| * SPDX-License-Identifier: Apache-2.0 | |
| */ |
| namespace CUDF_EXPORT cudf { | ||
|
|
||
| /** | ||
| * @addtogroup transformation_decimalops |
Contributor
There was a problem hiding this comment.
I think this doxygen group does not already exist and we need to add it. Might need to rename it to match existing patterns if any
Comment on lines
+42
to
+49
| * @param lhs The left operand column | ||
| * @param rhs The right operand column | ||
| * @param rounding_mode The rounding mode to use | ||
| * @param stream CUDA stream used for device memory operations and kernel launches | ||
| * @param mr Device memory resource used to allocate the returned column's device memory | ||
| * @return Output column containing the result of the decimal division | ||
| * @throw cudf::logic_error if @p lhs and @p rhs are different sizes | ||
| * @throw cudf::logic_error if @p lhs and @p rhs are not decimal types |
Contributor
There was a problem hiding this comment.
clang-format likely needed here
| /** | ||
| * @brief Performs decimal division between two columns with scale preservation. | ||
| * | ||
| * The output contains the result of `divide_decimal(lhs[i], rhs[i])` for all `0 <= i < lhs.size()` |
Contributor
There was a problem hiding this comment.
The output contains the result of
divide_decimal(lhs[i], rhs[i])for all0 <= i < lhs.size()
Result of divide_decimal(lhs[i], rhs[i]) isn't yet defined as this is the first instance of this API overload
| @@ -11,6 +11,7 @@ | |||
| import numpy as np | |||
Contributor
There was a problem hiding this comment.
Please separate out python bindings and pytests into a separate follow up PR if possible to limit the scope of this PR
| * @tparam Rad1 The radix of the fixed-point numbers | ||
| * @param lhs The dividend (left-hand side of division) | ||
| * @param rhs The divisor (right-hand side of division) | ||
| * @param rounding_mode The rounding mode to use (default: HALF_UP) |
Contributor
There was a problem hiding this comment.
No need to add default value here as clear from the param below. Also less future maintenance
Suggested change
| * @param rounding_mode The rounding mode to use (default: HALF_UP) | |
| * @param rounding_mode The rounding mode to use |
| // We need to compensate for the scale difference to preserve the dividend's scale | ||
| // Standard division would give us scale = lhs.scale() - rhs.scale() | ||
| // To preserve lhs.scale(), we need to scale up by 10^(-rhs.scale()) | ||
| auto const scale_factor = detail::ipow<Rep1, Rad1>(-static_cast<int>(rhs.scale())); |
Contributor
There was a problem hiding this comment.
We should look into reusing utils from src/round/round.cu here if possible to avoid code duplication
| * @return A fixed-point number with the same scale as the dividend | ||
| */ | ||
| template <typename Rep1, Radix Rad1> | ||
| CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> divide_decimal( |
Contributor
There was a problem hiding this comment.
Is the CUDF_HOST_DEVICE needed or could this just be a device only util since i only see it being called from a device function below
| namespace detail { | ||
|
|
||
| template <typename DecimalType> | ||
| struct divide_decimal_functor { |
Contributor
There was a problem hiding this comment.
Looks like unused functor?
| } | ||
| }; | ||
|
|
||
| std::unique_ptr<column> divide_decimal_impl(column_view const& lhs, |
Contributor
There was a problem hiding this comment.
Would recommend restructuring the code in this file as follows:
namespace CUDF_EXPORT cudf {
namespace detail {
// Provided two fixed point iterators here and computes decimal_divide from `fixed_point.hpp` using a transform on them.
template <LHS, RHS>
std::unique_ptr<column> divide_decimal(LHS lhs, RHS rhs, size_type size ...)
{}
// Detail APIs for decimal_divide(col, col), decimal_divide(col, scalar), decimal_divide(scalar, col) overloads that call the above one with combination of regular (columns) or `cuda::constant_iterator` (scalars)
} // namespace detail
// public APIs here that just do this
std::unique_ptr<column> divide_decimal(..)
{
CUDF_FUNC_RANGE();
return detail::divide_decimal(..)
}
} // namespace cudf
mhaseeb123
requested changes
May 18, 2026
mhaseeb123
reviewed
May 18, 2026
| @@ -0,0 +1,103 @@ | |||
| /* | |||
Contributor
There was a problem hiding this comment.
Nit: Reconsider renaming the file. Here decimal/decimal_ops.hpp the word decimal seems to be repeating, also adds a new folder. Perhaps we could move this to cudf/fixed_point/math_ops.hpp.
Open to suggestions
Contributor
|
@simoneves I hope you can pick this back up! |
Author
It's on the list. I wasn't planning to pick up any new Velox work before HeavyDB OSS is done (other than getting Decimal Part 4 landed) but I'm sure that's negotiable! |
lamarrr
reviewed
Jun 26, 2026
| } | ||
|
|
||
| // Fallback to __int128_t for severe overflow cases | ||
| using WidestRep = __int128_t; |
Contributor
There was a problem hiding this comment.
what if the existing rep is __int128_t?
lamarrr
requested changes
Jun 26, 2026
lamarrr
left a comment
lamarrr
left a comment
Contributor
There was a problem hiding this comment.
Initial round of reviews.
This is really great work
Comment on lines
+838
to
+860
| // Fallback to __int128_t for severe overflow cases | ||
| using WidestRep = __int128_t; | ||
| WidestRep wide_scale = static_cast<WidestRep>(scale_factor); | ||
| WidestRep scaled_dividend = static_cast<WidestRep>(lhs.value()) * wide_scale; | ||
| WidestRep wide_divisor = static_cast<WidestRep>(rhs.value()); | ||
| WidestRep quotient = scaled_dividend / wide_divisor; | ||
| WidestRep remainder = scaled_dividend % wide_divisor; | ||
|
|
||
| // Apply rounding | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_rem * 2 >= abs_div) { quotient += (quotient >= 0) ? 1 : -1; } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| WidestRep abs_rem_2 = abs_rem * 2; | ||
| if (abs_rem_2 > abs_div || (abs_rem_2 == abs_div && (quotient % 2) != 0)) { | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; |
Contributor
There was a problem hiding this comment.
Suggested change
| // Fallback to __int128_t for severe overflow cases | |
| using WidestRep = __int128_t; | |
| WidestRep wide_scale = static_cast<WidestRep>(scale_factor); | |
| WidestRep scaled_dividend = static_cast<WidestRep>(lhs.value()) * wide_scale; | |
| WidestRep wide_divisor = static_cast<WidestRep>(rhs.value()); | |
| WidestRep quotient = scaled_dividend / wide_divisor; | |
| WidestRep remainder = scaled_dividend % wide_divisor; | |
| // Apply rounding | |
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | |
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | |
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | |
| if (abs_rem * 2 >= abs_div) { quotient += (quotient >= 0) ? 1 : -1; } | |
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | |
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | |
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | |
| WidestRep abs_rem_2 = abs_rem * 2; | |
| if (abs_rem_2 > abs_div || (abs_rem_2 == abs_div && (quotient % 2) != 0)) { | |
| quotient += (quotient >= 0) ? 1 : -1; | |
| } | |
| } | |
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; | |
| // Fallback to __int128_t for severe overflow cases | |
| using widest_rep = __int128_t; | |
| widest_rep wide_scale = static_cast<widest_rep>(scale_factor); | |
| widest_rep scaled_dividend = static_cast<widest_rep>(lhs.value()) * wide_scale; | |
| widest_rep wide_divisor = static_cast<widest_rep>(rhs.value()); | |
| widest_rep quotient = scaled_dividend / wide_divisor; | |
| widest_rep remainder = scaled_dividend % wide_divisor; | |
| // Apply rounding | |
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | |
| widest_rep abs_rem = (remainder < 0) ? -remainder : remainder; | |
| widest_rep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | |
| if (abs_rem * 2 >= abs_div) { quotient += (quotient >= 0) ? 1 : -1; } | |
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | |
| widest_rep abs_rem = (remainder < 0) ? -remainder : remainder; | |
| widest_rep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | |
| widest_rep abs_rem_2 = abs_rem * 2; | |
| if (abs_rem_2 > abs_div || (abs_rem_2 == abs_div && (quotient % 2) != 0)) { | |
| quotient += (quotient >= 0) ? 1 : -1; | |
| } | |
| } | |
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; |
For function-local aliases, we typically use lower-case
Comment on lines
+73
to
+78
| std::make_unique<column>(lhs_type, | ||
| size, | ||
| rmm::device_buffer{size * cudf::size_of(lhs_type), stream, mr}, | ||
| rmm::device_buffer{}, // Empty buffer = non-nullable | ||
| 0, // null_count = 0 | ||
| std::vector<std::unique_ptr<column>>{}); |
Contributor
There was a problem hiding this comment.
Suggested change
| std::make_unique<column>(lhs_type, | |
| size, | |
| rmm::device_buffer{size * cudf::size_of(lhs_type), stream, mr}, | |
| rmm::device_buffer{}, // Empty buffer = non-nullable | |
| 0, // null_count = 0 | |
| std::vector<std::unique_ptr<column>>{}); | |
| make_fixed_width_column(lhs_type, size, mask_state::UNALLOCATED, stream, mr); |
Comment on lines
+88
to
+142
| if (lhs_type.id() == type_id::DECIMAL32) { | ||
| using Type = int32_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (lhs_type.id() == type_id::DECIMAL64) { | ||
| using Type = int64_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (lhs_type.id() == type_id::DECIMAL128) { | ||
| using Type = __int128_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } |
Contributor
There was a problem hiding this comment.
giant type-id switch here. it is also non-exhaustive and would fail if it is none of the existing decimal type-ids.
we should split it up to use the existing type-dispatch pattern in CUDF.
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| using namespace numeric; |
Contributor
There was a problem hiding this comment.
we shouldn't do this. this would cause collisions and ambiguous resolution across namespaces.
Comment on lines
+188
to
+194
| result = | ||
| std::make_unique<column>(lhs_type, | ||
| size, | ||
| rmm::device_buffer{size * cudf::size_of(lhs_type), stream, mr}, | ||
| rmm::device_buffer{}, // Empty buffer = non-nullable | ||
| 0, // null_count = 0 | ||
| std::vector<std::unique_ptr<column>>{}); |
Contributor
There was a problem hiding this comment.
same as above, use make_fixed_width_column
|
|
||
| if (lhs.is_empty()) { return make_empty_column(lhs.type()); } | ||
|
|
||
| using namespace numeric; |
Contributor
There was a problem hiding this comment.
same as above: no using namespace
| CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL32 || lhs.type().id() == type_id::DECIMAL64 || | ||
| lhs.type().id() == type_id::DECIMAL128, | ||
| "Column must be decimal type"); | ||
| CUDF_EXPECTS(rhs.type() == lhs.type(), "Scalar type (" + std::to_string(int(rhs.type().id())) + ") must match column type (" + std::to_string(int(lhs.type().id())) + ")"); |
Contributor
There was a problem hiding this comment.
use std::format instead of +-concatenation
Comment on lines
+758
to
+860
| // Check for division by zero | ||
| // In CUDA device code, we cannot throw exceptions, so we assert | ||
| // In host code, this will cause undefined behavior (same as standard division) | ||
| #if defined(__CUDACC_DEBUG__) | ||
| assert(rhs.value() != 0 && "division by zero"); | ||
| assert(!detail::division_overflow<Rep1>(lhs.value(), rhs.value()) && "fixed_point overflow"); | ||
| #endif | ||
|
|
||
| // Scale up the dividend to maintain precision | ||
| // Result will have scale = lhs.scale() | ||
| // We need to compensate for the scale difference to preserve the dividend's scale | ||
| // Standard division would give us scale = lhs.scale() - rhs.scale() | ||
| // To preserve lhs.scale(), we need to scale up by 10^(-rhs.scale()) | ||
| auto const scale_factor = detail::ipow<Rep1, Rad1>(-static_cast<int>(rhs.scale())); | ||
|
|
||
| // Check for potential overflow when scaling | ||
| bool overflow = multiplication_overflow<Rep1>(lhs.value(), scale_factor); | ||
|
|
||
| if (!overflow) { | ||
| // Standard calculation without overflow | ||
| Rep1 scaled_dividend = lhs.value() * scale_factor; | ||
| Rep1 quotient = scaled_dividend / rhs.value(); | ||
| Rep1 remainder = scaled_dividend % rhs.value(); | ||
|
|
||
| // Apply rounding based on remainder | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| // Round half away from zero | ||
| // Avoid abs() ambiguity for __int128 by using conditional | ||
| auto abs_remainder = (remainder < 0) ? -remainder : remainder; | ||
| auto abs_divisor = (rhs.value() < 0) ? -rhs.value() : rhs.value(); | ||
| if (abs_remainder * 2 >= abs_divisor) { | ||
| // Round away from zero: if quotient is positive, add 1; if negative, subtract 1 | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| // Banker's rounding | ||
| // Avoid abs() ambiguity for __int128 by using conditional | ||
| auto abs_rem_2 = ((remainder < 0) ? -remainder : remainder) * 2; | ||
| auto abs_divisor = (rhs.value() < 0) ? -rhs.value() : rhs.value(); | ||
| if (abs_rem_2 > abs_divisor || (abs_rem_2 == abs_divisor && (quotient % 2) != 0)) { | ||
| // Round to nearest even: direction depends on quotient sign | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{quotient, lhs.scale()}}; | ||
| } | ||
|
|
||
| // Handle overflow cases with type promotion | ||
| if constexpr (cuda::std::is_same_v<Rep1, int32_t>) { | ||
| // Try int64_t first | ||
| using WiderRep = int64_t; | ||
| WiderRep wide_scale = static_cast<WiderRep>(scale_factor); | ||
| bool overflow_in_int64 = | ||
| multiplication_overflow<WiderRep>(static_cast<WiderRep>(lhs.value()), wide_scale); | ||
|
|
||
| if (!overflow_in_int64) { | ||
| WiderRep scaled_dividend = static_cast<WiderRep>(lhs.value()) * wide_scale; | ||
| WiderRep wide_divisor = static_cast<WiderRep>(rhs.value()); | ||
| WiderRep quotient = scaled_dividend / wide_divisor; | ||
| WiderRep remainder = scaled_dividend % wide_divisor; | ||
|
|
||
| // Apply rounding | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| auto abs_remainder = (remainder < 0) ? -remainder : remainder; | ||
| auto abs_divisor = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_remainder * 2 >= abs_divisor) { quotient += (quotient >= 0) ? 1 : -1; } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| auto abs_rem_2 = ((remainder < 0) ? -remainder : remainder) * 2; | ||
| auto abs_divisor = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_rem_2 > abs_divisor || (abs_rem_2 == abs_divisor && (quotient % 2) != 0)) { | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{ | ||
| scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; | ||
| } | ||
| } | ||
|
|
||
| // Fallback to __int128_t for severe overflow cases | ||
| using WidestRep = __int128_t; | ||
| WidestRep wide_scale = static_cast<WidestRep>(scale_factor); | ||
| WidestRep scaled_dividend = static_cast<WidestRep>(lhs.value()) * wide_scale; | ||
| WidestRep wide_divisor = static_cast<WidestRep>(rhs.value()); | ||
| WidestRep quotient = scaled_dividend / wide_divisor; | ||
| WidestRep remainder = scaled_dividend % wide_divisor; | ||
|
|
||
| // Apply rounding | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_rem * 2 >= abs_div) { quotient += (quotient >= 0) ? 1 : -1; } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| WidestRep abs_rem_2 = abs_rem * 2; | ||
| if (abs_rem_2 > abs_div || (abs_rem_2 == abs_div && (quotient % 2) != 0)) { | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; |
Contributor
There was a problem hiding this comment.
The approach to handling overflow here, while correct, would be difficult to debug and make changes to.
There are a lot of data-dependent cases and logic around the overflow checking, e.g., having 2 fallbacks for a single data type.
Perhaps we could simplify this code and make it less data-dependent, maybe by pre-promoting some data types before performing overflow checks.
Comment on lines
+327
to
+383
| if (rhs_type.id() == type_id::DECIMAL32) { | ||
| using Type = int32_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (rhs_type.id() == type_id::DECIMAL64) { | ||
| using Type = int64_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else { | ||
| using Type = __int128_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); |
Contributor
There was a problem hiding this comment.
Suggested change
| if (rhs_type.id() == type_id::DECIMAL32) { | |
| using Type = int32_t; | |
| using DecType = fixed_point<Type, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Type lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Type>(), | |
| rhs.end<Type>(), | |
| result_view.begin<Type>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); | |
| } else if (rhs_type.id() == type_id::DECIMAL64) { | |
| using Type = int64_t; | |
| using DecType = fixed_point<Type, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Type lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Type>(), | |
| rhs.end<Type>(), | |
| result_view.begin<Type>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); | |
| } else { | |
| using Type = __int128_t; | |
| using DecType = fixed_point<Type, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Type lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Type>(), | |
| rhs.end<Type>(), | |
| result_view.begin<Type>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); | |
| if (rhs_type.id() == type_id::DECIMAL32) { | |
| using Rep = int32_t; | |
| using DecType = fixed_point<Rep, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Rep lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Rep>(), | |
| rhs.end<Rep>(), | |
| result_view.begin<Rep>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Rep rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); | |
| } else if (rhs_type.id() == type_id::DECIMAL64) { | |
| using Rep = int64_t; | |
| using DecType = fixed_point<Rep, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Rep lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Rep>(), | |
| rhs.end<Rep>(), | |
| result_view.begin<Rep>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Rep rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Rep>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Rep>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); | |
| } else { | |
| using Rep = __int128_t; | |
| using DecType = fixed_point<Rep, Radix::BASE_10>; | |
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | |
| DecType lhs_fp = decimal_scalar.value(stream); | |
| Rep lhs_val = lhs_fp.value(); | |
| thrust::transform( | |
| rmm::exec_policy(stream), | |
| rhs.begin<Rep>(), | |
| rhs.end<Rep>(), | |
| result_view.begin<Rep>(), | |
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Rep rhs_val) { | |
| DecType lhs_fp{numeric::scaled_integer<Rep>{lhs_val, numeric::scale_type{lhs_scale}}}; | |
| DecType rhs_fp{numeric::scaled_integer<Rep>{rhs_val, numeric::scale_type{rhs_scale}}}; | |
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | |
| return result_fp.value(); | |
| }); |
| "Columns must be decimal type"); | ||
|
|
||
| // Note: Zero division check is handled in the kernel | ||
| // GPU kernels cannot throw exceptions, so they produce special values or assert in debug mode |
Contributor
There was a problem hiding this comment.
i think it is worth clarifying that the implementation doesn't throw any special value presently.
davidwendt
reviewed
Jun 26, 2026
| fixed_point<Rep1, Rad1> const& lhs, | ||
| fixed_point<Rep1, Rad1> const& rhs, | ||
| decimal_rounding_mode rounding_mode = decimal_rounding_mode::HALF_UP) | ||
| { |
Contributor
There was a problem hiding this comment.
Would it be possible to not put any of this into fixed_point.hpp?
This header is included everywhere and this new code is only needed in decimal_ops.cu I believe.
Perhaps decimal_ops.cuh since I agree that CUDF_HOST_DEVICE is not needed.
3 tasks
Author
|
Apologies @mhaseeb123 and @lamarrr I haven't been keeping an eye on this one. Thank you for your reviews. I will get back to this very soon. |
Contributor
|
@simoneves, this looks close - can you target 26.08? |
Implements divide_decimal for fixed-point decimal columns that preserves the dividend's scale, similar to Java BigDecimal.divide() with rounding mode. - Add divide_decimal C++ implementation with HALF_UP and HALF_EVEN rounding - Add Python bindings via pylibcudf - Add divide_decimal method to DecimalBaseColumn - Add comprehensive tests for various scales and rounding modes - Fix critical bug in C++ implementation for different scales - Add Doxygen documentation for new functions Addresses issue NVIDIA#17448 Co-authored-by: Akihiro Hirota <akihiro-hirota@gmobiz.com> (cherry picked from commit f8b1ff34ef9490b486ab76f2071b86c6d4d3423c)
simoneves
force-pushed
the
simoneves/a-hirota/decimal-division-operations-rebased
branch
from
July 14, 2026 03:45
41c2ab1 to
de0aea1
Compare
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesDecimal division
Estimated code review effort: 4 (Complex) | ~60 minutes Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 11
🧹 Nitpick comments (3)
python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd (1)
14-30: 🗄️ Data Integrity & Integration | 🔵 Trivial | ⚖️ Poor tradeoffStream/memory-resource parameters are dropped from the pylibcudf decimal-division API. The upstream C++
divide_decimaloverloads acceptstreamandmrwith defaults; the pxd declarations here truncate the signature to 3 arguments, so none of the.pyxwrappers can forward a caller-supplied stream or memory resource, unlike the rest of pylibcudf's public surface.
python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd#L14-L30: addstream(Stream) andmrparameters to all three extern declarations matching the C++ defaults.python/pylibcudf/pylibcudf/decimal_division.pyx#L42-L115: thread an optionalStream stream=None, DeviceMemoryResource mr=Nonethroughdivide_decimaland forward tocpp_divide_decimal.python/pylibcudf/pylibcudf/decimal_division.pyx#L118-L192: same fordivide_decimal_column_scalar.python/pylibcudf/pylibcudf/decimal_division.pyx#L195-L269: same fordivide_decimal_scalar_column.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd` around lines 14 - 30, Expose stream and memory-resource forwarding across the decimal division API. In python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd lines 14-30, add Stream and mr parameters with the upstream C++ defaults to all three divide_decimal declarations; in python/pylibcudf/pylibcudf/decimal_division.pyx lines 42-115, 118-192, and 195-269, add optional Stream stream=None and DeviceMemoryResource mr=None parameters to divide_decimal, divide_decimal_column_scalar, and divide_decimal_scalar_column respectively, forwarding both values to cpp_divide_decimal.python/cudf/cudf/tests/test_decimal_division.py (1)
156-170: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winStub test and missing edge-case coverage.
test_different_decimal_typesis a no-op (pass) despite its docstring claiming Decimal32/64/128 coverage — the C++ layer has explicitTYPED_TESTs across all three widths, but nothing exercises this at the Python level. Also missing: empty Series, fully-null Series, and single-element Series cases.As per coding guidelines: "Ensure test files provide comprehensive edge case coverage (empty, all-null, single-element, mixed types) and do not depend on external datasets."
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudf/cudf/tests/test_decimal_division.py` around lines 156 - 170, Implement test_different_decimal_types with Python-level coverage for Decimal32, Decimal64, and Decimal128, replacing the pass stub and validating decimal division results. Add focused tests in the same test module for empty Series, fully-null Series, and single-element Series, using only inline data and asserting the expected values or null behavior.Source: Coding guidelines
python/pylibcudf/pylibcudf/decimal_division.pyx (1)
85-104: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueFactor out shared decimal validation
The rounding-mode conversion and decimal-family checks are repeated across all three overloads; a small helper would keep them aligned.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/pylibcudf/pylibcudf/decimal_division.pyx` around lines 85 - 104, Create a shared helper for rounding-mode conversion and decimal-family validation, then update all three overloads to call it instead of duplicating the logic. Reuse the existing DecimalRoundingMode, cpp_decimal_rounding_mode, and type_id checks, preserving the current ValueError and TypeError behavior and messages.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/include/cudf/fixed_point/fixed_point.hpp`:
- Around line 738-861: Move divide_decimal out of fixed_point.hpp into the
decimal_ops CUDA header used by decimal_ops.cu, and remove the unnecessary
CUDF_HOST_DEVICE annotation if it is device-only. Simplify the data-dependent
int32→int64→__int128 overflow chain into a clear, maintainable strategy,
explicitly handling Rep1 == __int128_t without claiming additional headroom or
allowing silent overflow. Rename local type aliases such as WiderRep and
WidestRep to the project’s lowercase convention.
- Around line 762-771: Handle positive rhs.scale() before computing scale_factor
in the fixed-point division path. Update the logic around detail::ipow in the
division operator to explicitly branch or validate when rhs.scale() is positive,
preventing a negative exponent from reaching ipow and preserving correct
division behavior for supported scales.
In `@cpp/src/binaryop/decimal_ops.cu`:
- Around line 172-184: Cache the scalar validity result once in
divide_decimal_scalar_impl and the scalar-column divide_decimal overload, then
reuse that local value in both the outer and nested conditionals. Apply the same
change symmetrically to rhs.is_valid(stream) and lhs.is_valid(stream),
preserving the existing null-result behavior.
- Around line 51-389: Restructure divide_decimal so all three public
overloads—column/column, column/scalar, and scalar/column—remain thin
CUDF_FUNC_RANGE() wrappers delegating to detail helpers. Add shared
type-dispatch helpers using cuda::constant_iterator for scalar operands,
eliminating duplicated type-id transform logic and moving scalar/column
implementation out of the public overload. Qualify numeric symbols explicitly
instead of using namespace numeric, and replace the scalar type-mismatch message
concatenation with std::format.
- Around line 83-85: Remove the unused lhs_dev and rhs_dev declarations created
by column_device_view::create in the decimal operation flow. Leave the existing
thrust::transform inputs and all other logic unchanged.
In `@cpp/tests/decimal/decimal_ops_test.cpp`:
- Around line 36-37: Expand DecimalTypes to include numeric::decimal128, and add
a DecimalOpsTest case using sufficiently large values to exceed the
scale-multiplication threshold and exercise numeric::divide_decimal’s
overflow-promotion fallback. Preserve the existing decimal32 and decimal64
coverage.
In `@python/cudf/cudf/core/column/decimal.py`:
- Around line 47-49: Replace the invalid `// existing` comment near the imports
with valid Python comment syntax or remove it, leaving the
`ColumnBinaryOperand`, `ColumnLike`, `Dtype`, `ScalarLike`, and `Buffer` imports
unchanged.
- Around line 316-320: Update divide_decimal’s handling of a NotImplemented
result from _normalize_binop_operand: because this public method must honor its
-> Self contract, raise an appropriate clear error instead of returning the bare
NotImplemented sentinel. Preserve normal division behavior for supported
operands.
In `@python/pylibcudf/pylibcudf/decimal_division.pyx`:
- Around line 47-84: Update the divide_decimal docstring’s Raises section to
document ValueError for an invalid rounding_mode instead of division by zero,
and revise the example to use the actual decimal-column construction API
available in pylibcudf rather than column_from_decimal_values.
- Around line 42-46: Guard the operands in all three decimal-division functions,
including divide_decimal, before any calls to .type(), .view(), or .get(). Add
explicit None validation for each typed Column or Scalar input and preserve the
existing behavior for valid operands, using the module’s established
invalid-argument error handling.
In `@python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd`:
- Around line 14-30: Import libcudf_exception_handler in the decimal bindings
and update all three divide_decimal overloads to use except
+libcudf_exception_handler instead of the generic exception specification,
preserving their existing signatures and arguments.
---
Nitpick comments:
In `@python/cudf/cudf/tests/test_decimal_division.py`:
- Around line 156-170: Implement test_different_decimal_types with Python-level
coverage for Decimal32, Decimal64, and Decimal128, replacing the pass stub and
validating decimal division results. Add focused tests in the same test module
for empty Series, fully-null Series, and single-element Series, using only
inline data and asserting the expected values or null behavior.
In `@python/pylibcudf/pylibcudf/decimal_division.pyx`:
- Around line 85-104: Create a shared helper for rounding-mode conversion and
decimal-family validation, then update all three overloads to call it instead of
duplicating the logic. Reuse the existing DecimalRoundingMode,
cpp_decimal_rounding_mode, and type_id checks, preserving the current ValueError
and TypeError behavior and messages.
In `@python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd`:
- Around line 14-30: Expose stream and memory-resource forwarding across the
decimal division API. In
python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd lines 14-30, add
Stream and mr parameters with the upstream C++ defaults to all three
divide_decimal declarations; in python/pylibcudf/pylibcudf/decimal_division.pyx
lines 42-115, 118-192, and 195-269, add optional Stream stream=None and
DeviceMemoryResource mr=None parameters to divide_decimal,
divide_decimal_column_scalar, and divide_decimal_scalar_column respectively,
forwarding both values to cpp_divide_decimal.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 0707e95b-e130-4b55-9fd4-fa94c9c84310
📒 Files selected for processing (12)
cpp/CMakeLists.txtcpp/include/cudf/decimal/decimal_ops.hppcpp/include/cudf/fixed_point/fixed_point.hppcpp/src/binaryop/decimal_ops.cucpp/tests/CMakeLists.txtcpp/tests/decimal/decimal_ops_test.cpppython/cudf/cudf/core/column/decimal.pypython/cudf/cudf/tests/test_decimal_division.pypython/pylibcudf/pylibcudf/CMakeLists.txtpython/pylibcudf/pylibcudf/__init__.pypython/pylibcudf/pylibcudf/decimal_division.pyxpython/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd
Comment on lines
+738
to
+861
| /** | ||
| * @brief Performs decimal division with scale preservation | ||
| * | ||
| * This function divides two fixed-point numbers while preserving the scale | ||
| * of the dividend (left-hand side). This behavior is similar to Java's | ||
| * BigDecimal.divide(divisor, roundingMode) which maintains the dividend's scale. | ||
| * | ||
| * @tparam Rep1 The representation type of the fixed-point numbers | ||
| * @tparam Rad1 The radix of the fixed-point numbers | ||
| * @param lhs The dividend (left-hand side of division) | ||
| * @param rhs The divisor (right-hand side of division) | ||
| * @param rounding_mode The rounding mode to use (default: HALF_UP) | ||
| * @return A fixed-point number with the same scale as the dividend | ||
| */ | ||
| template <typename Rep1, Radix Rad1> | ||
| CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> divide_decimal( | ||
| fixed_point<Rep1, Rad1> const& lhs, | ||
| fixed_point<Rep1, Rad1> const& rhs, | ||
| decimal_rounding_mode rounding_mode = decimal_rounding_mode::HALF_UP) | ||
| { | ||
| // Check for division by zero | ||
| // In CUDA device code, we cannot throw exceptions, so we assert | ||
| // In host code, this will cause undefined behavior (same as standard division) | ||
| #if defined(__CUDACC_DEBUG__) | ||
| assert(rhs.value() != 0 && "division by zero"); | ||
| assert(!detail::division_overflow<Rep1>(lhs.value(), rhs.value()) && "fixed_point overflow"); | ||
| #endif | ||
|
|
||
| // Scale up the dividend to maintain precision | ||
| // Result will have scale = lhs.scale() | ||
| // We need to compensate for the scale difference to preserve the dividend's scale | ||
| // Standard division would give us scale = lhs.scale() - rhs.scale() | ||
| // To preserve lhs.scale(), we need to scale up by 10^(-rhs.scale()) | ||
| auto const scale_factor = detail::ipow<Rep1, Rad1>(-static_cast<int>(rhs.scale())); | ||
|
|
||
| // Check for potential overflow when scaling | ||
| bool overflow = multiplication_overflow<Rep1>(lhs.value(), scale_factor); | ||
|
|
||
| if (!overflow) { | ||
| // Standard calculation without overflow | ||
| Rep1 scaled_dividend = lhs.value() * scale_factor; | ||
| Rep1 quotient = scaled_dividend / rhs.value(); | ||
| Rep1 remainder = scaled_dividend % rhs.value(); | ||
|
|
||
| // Apply rounding based on remainder | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| // Round half away from zero | ||
| // Avoid abs() ambiguity for __int128 by using conditional | ||
| auto abs_remainder = (remainder < 0) ? -remainder : remainder; | ||
| auto abs_divisor = (rhs.value() < 0) ? -rhs.value() : rhs.value(); | ||
| if (abs_remainder * 2 >= abs_divisor) { | ||
| // Round away from zero: if quotient is positive, add 1; if negative, subtract 1 | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| // Banker's rounding | ||
| // Avoid abs() ambiguity for __int128 by using conditional | ||
| auto abs_rem_2 = ((remainder < 0) ? -remainder : remainder) * 2; | ||
| auto abs_divisor = (rhs.value() < 0) ? -rhs.value() : rhs.value(); | ||
| if (abs_rem_2 > abs_divisor || (abs_rem_2 == abs_divisor && (quotient % 2) != 0)) { | ||
| // Round to nearest even: direction depends on quotient sign | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{quotient, lhs.scale()}}; | ||
| } | ||
|
|
||
| // Handle overflow cases with type promotion | ||
| if constexpr (cuda::std::is_same_v<Rep1, int32_t>) { | ||
| // Try int64_t first | ||
| using WiderRep = int64_t; | ||
| WiderRep wide_scale = static_cast<WiderRep>(scale_factor); | ||
| bool overflow_in_int64 = | ||
| multiplication_overflow<WiderRep>(static_cast<WiderRep>(lhs.value()), wide_scale); | ||
|
|
||
| if (!overflow_in_int64) { | ||
| WiderRep scaled_dividend = static_cast<WiderRep>(lhs.value()) * wide_scale; | ||
| WiderRep wide_divisor = static_cast<WiderRep>(rhs.value()); | ||
| WiderRep quotient = scaled_dividend / wide_divisor; | ||
| WiderRep remainder = scaled_dividend % wide_divisor; | ||
|
|
||
| // Apply rounding | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| auto abs_remainder = (remainder < 0) ? -remainder : remainder; | ||
| auto abs_divisor = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_remainder * 2 >= abs_divisor) { quotient += (quotient >= 0) ? 1 : -1; } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| auto abs_rem_2 = ((remainder < 0) ? -remainder : remainder) * 2; | ||
| auto abs_divisor = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_rem_2 > abs_divisor || (abs_rem_2 == abs_divisor && (quotient % 2) != 0)) { | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{ | ||
| scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; | ||
| } | ||
| } | ||
|
|
||
| // Fallback to __int128_t for severe overflow cases | ||
| using WidestRep = __int128_t; | ||
| WidestRep wide_scale = static_cast<WidestRep>(scale_factor); | ||
| WidestRep scaled_dividend = static_cast<WidestRep>(lhs.value()) * wide_scale; | ||
| WidestRep wide_divisor = static_cast<WidestRep>(rhs.value()); | ||
| WidestRep quotient = scaled_dividend / wide_divisor; | ||
| WidestRep remainder = scaled_dividend % wide_divisor; | ||
|
|
||
| // Apply rounding | ||
| if (rounding_mode == decimal_rounding_mode::HALF_UP) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| if (abs_rem * 2 >= abs_div) { quotient += (quotient >= 0) ? 1 : -1; } | ||
| } else if (rounding_mode == decimal_rounding_mode::HALF_EVEN) { | ||
| WidestRep abs_rem = (remainder < 0) ? -remainder : remainder; | ||
| WidestRep abs_div = (wide_divisor < 0) ? -wide_divisor : wide_divisor; | ||
| WidestRep abs_rem_2 = abs_rem * 2; | ||
| if (abs_rem_2 > abs_div || (abs_rem_2 == abs_div && (quotient % 2) != 0)) { | ||
| quotient += (quotient >= 0) ? 1 : -1; | ||
| } | ||
| } | ||
|
|
||
| return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{static_cast<Rep1>(quotient), lhs.scale()}}; | ||
| } |
There was a problem hiding this comment.
Placement, overflow-handling complexity, and __int128_t fallback gap already flagged.
This large divide_decimal block was previously flagged for: (1) belonging in a .cuh used only by decimal_ops.cu rather than the ubiquitously-included fixed_point.hpp; (2) CUDF_HOST_DEVICE likely unnecessary since only invoked from device code; (3) highly data-dependent overflow fallback chain (int32→int64→int128) that is hard to reason about/debug; (4) the final __int128_t fallback provides no additional headroom when Rep1 is already __int128_t (decimal128), so overflow can still silently occur there; (5) lowercase local type aliases per convention.
Based on past review comments: "Would it be possible to not put any of this into fixed_point.hpp? ... Perhaps decimal_ops.cuh", "The approach to handling overflow here ... would be difficult to debug and make changes to," "what if the existing rep is __int128_t?", and "For function-local aliases, we typically use lower-case."
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@cpp/include/cudf/fixed_point/fixed_point.hpp` around lines 738 - 861, Move
divide_decimal out of fixed_point.hpp into the decimal_ops CUDA header used by
decimal_ops.cu, and remove the unnecessary CUDF_HOST_DEVICE annotation if it is
device-only. Simplify the data-dependent int32→int64→__int128 overflow chain
into a clear, maintainable strategy, explicitly handling Rep1 == __int128_t
without claiming additional headroom or allowing silent overflow. Rename local
type aliases such as WiderRep and WidestRep to the project’s lowercase
convention.
Comment on lines
+762
to
+771
| assert(rhs.value() != 0 && "division by zero"); | ||
| assert(!detail::division_overflow<Rep1>(lhs.value(), rhs.value()) && "fixed_point overflow"); | ||
| #endif | ||
|
|
||
| // Scale up the dividend to maintain precision | ||
| // Result will have scale = lhs.scale() | ||
| // We need to compensate for the scale difference to preserve the dividend's scale | ||
| // Standard division would give us scale = lhs.scale() - rhs.scale() | ||
| // To preserve lhs.scale(), we need to scale up by 10^(-rhs.scale()) | ||
| auto const scale_factor = detail::ipow<Rep1, Rad1>(-static_cast<int>(rhs.scale())); |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
set -euo pipefail
git ls-files | rg 'cpp/.*/fixed_point|cpp/.*/decimal|cudf/fixed_point|fixed_point\.hpp|decimal.*hpp|type_id|scalar'Repository: rapidsai/cudf
Length of output: 2276
🏁 Script executed:
set -euo pipefail
# Locate relevant files
git ls-files | rg 'cpp/.*/fixed_point|cudf/fixed_point|decimal.*hpp|type_id|scalar'
# Quick text search for scale-related constructors/guards
rg -n "scale\(|fixed_point|decimal" cpp/include/cudf cpp/src cpp/tests -g '!**/build/**' | head -n 200Repository: rapidsai/cudf
Length of output: 26792
🏁 Script executed:
set -euo pipefail
# Read the relevant region in fixed_point.hpp
sed -n '700,820p' cpp/include/cudf/fixed_point/fixed_point.hpp | cat -n
# Search for any explicit restrictions on scale sign
rg -n "scale.*>=|scale.*<=|positive scale|negative scale|static_cast<int>\(.*scale\)|scale_type" cpp/include/cudf cpp/src cpp/tests | head -n 200Repository: rapidsai/cudf
Length of output: 29822
🏁 Script executed:
set -euo pipefail
# Inspect ipow implementation and any scale docs/constraints around fixed_point
sed -n '1,220p' cpp/include/cudf/fixed_point/fixed_point.hpp | cat -n
rg -n "ipow<|cudf_assert|negative exponent|exponent" cpp/include/cudf/fixed_point cpp/include/cudf/detail cpp/src cpp/tests | head -n 80Repository: rapidsai/cudf
Length of output: 18678
Handle positive rhs.scale() before calling ipow.
fixed_point scales are not limited to <= 0—public APIs and tests use positive scales—so detail::ipow<Rep1, Rad1>(-static_cast<int>(rhs.scale())) can receive a negative exponent here. In release builds ipow falls through and returns 10, which corrupts the division result instead of failing. Add an explicit check or branch for positive scales.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@cpp/include/cudf/fixed_point/fixed_point.hpp` around lines 762 - 771, Handle
positive rhs.scale() before computing scale_factor in the fixed-point division
path. Update the logic around detail::ipow in the division operator to
explicitly branch or validate when rhs.scale() is positive, preventing a
negative exponent from reaching ipow and preserving correct division behavior
for supported scales.
Comment on lines
+51
to
+389
| std::unique_ptr<column> divide_decimal_impl(column_view const& lhs, | ||
| column_view const& rhs, | ||
| numeric::decimal_rounding_mode rounding_mode, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| using namespace numeric; | ||
|
|
||
| auto const size = lhs.size(); | ||
| auto const lhs_type = lhs.type(); | ||
|
|
||
| // Create output column with same type as lhs (preserves scale) | ||
| // If there are nulls in inputs, create with null mask; otherwise create without | ||
| std::unique_ptr<column> result; | ||
| if (lhs.has_nulls() || rhs.has_nulls()) { | ||
| auto [null_mask, null_count] = cudf::detail::bitmask_and(table_view{{lhs, rhs}}, stream, mr); | ||
| result = | ||
| cudf::make_fixed_width_column(lhs_type, size, std::move(null_mask), null_count, stream, mr); | ||
| } else { | ||
| // Create non-nullable column when inputs have no nulls | ||
| // Use empty rmm::device_buffer{} directly to ensure column is non-nullable | ||
| result = | ||
| std::make_unique<column>(lhs_type, | ||
| size, | ||
| rmm::device_buffer{size * cudf::size_of(lhs_type), stream, mr}, | ||
| rmm::device_buffer{}, // Empty buffer = non-nullable | ||
| 0, // null_count = 0 | ||
| std::vector<std::unique_ptr<column>>{}); | ||
| } | ||
|
|
||
| auto result_view = result->mutable_view(); | ||
|
|
||
| // Get device views | ||
| auto const lhs_dev = column_device_view::create(lhs, stream); | ||
| auto const rhs_dev = column_device_view::create(rhs, stream); | ||
|
|
||
| // Perform element-wise divide_decimal | ||
| if (lhs_type.id() == type_id::DECIMAL32) { | ||
| using Type = int32_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (lhs_type.id() == type_id::DECIMAL64) { | ||
| using Type = int64_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (lhs_type.id() == type_id::DECIMAL128) { | ||
| using Type = __int128_t; | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
| using DecType = fixed_point<Type, numeric::Radix::BASE_10>; | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| rhs.begin<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rounding_mode] __device__(Type lhs_val, Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } | ||
|
|
||
| // Null mask already handled during column creation | ||
|
|
||
| return result; | ||
| } | ||
|
|
||
| template <typename DecimalType> | ||
| std::unique_ptr<column> divide_decimal_scalar_impl(column_view const& lhs, | ||
| scalar const& rhs, | ||
| numeric::decimal_rounding_mode rounding_mode, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| using namespace numeric; | ||
| using Type = typename DecimalType::rep; | ||
|
|
||
| auto const size = lhs.size(); | ||
| auto const lhs_type = lhs.type(); | ||
| auto const lhs_scale = lhs_type.scale(); | ||
| auto const rhs_scale = rhs.type().scale(); | ||
|
|
||
| // Get scalar value | ||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecimalType> const&>(rhs); | ||
| DecimalType rhs_fp = decimal_scalar.value(stream); | ||
| Type rhs_val = rhs_fp.value(); | ||
|
|
||
| // Create output column with same type as lhs | ||
| // If there are nulls in inputs, create with null mask; otherwise create without | ||
| std::unique_ptr<column> result; | ||
| if (lhs.has_nulls() || !rhs.is_valid(stream)) { | ||
| if (!rhs.is_valid(stream)) { | ||
| result = cudf::make_fixed_width_column( | ||
| lhs_type, | ||
| size, | ||
| cudf::detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr), | ||
| size, | ||
| stream, | ||
| mr); | ||
| } else { | ||
| result = cudf::make_fixed_width_column( | ||
| lhs_type, size, cudf::detail::copy_bitmask(lhs, stream, mr), lhs.null_count(), stream, mr); | ||
| } | ||
| } else { | ||
| // Create non-nullable column when inputs have no nulls | ||
| // Use empty rmm::device_buffer{} directly to ensure column is non-nullable | ||
| result = | ||
| std::make_unique<column>(lhs_type, | ||
| size, | ||
| rmm::device_buffer{size * cudf::size_of(lhs_type), stream, mr}, | ||
| rmm::device_buffer{}, // Empty buffer = non-nullable | ||
| 0, // null_count = 0 | ||
| std::vector<std::unique_ptr<column>>{}); | ||
| } | ||
|
|
||
| auto result_view = result->mutable_view(); | ||
|
|
||
| // Perform element-wise divide_decimal | ||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| lhs.begin<Type>(), | ||
| lhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, rhs_val, rounding_mode] __device__(Type lhs_val) { | ||
| DecimalType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecimalType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
|
|
||
| // Null mask already handled during column creation | ||
|
|
||
| return result; | ||
| } | ||
|
|
||
| } // namespace detail | ||
|
|
||
| std::unique_ptr<column> divide_decimal(column_view const& lhs, | ||
| column_view const& rhs, | ||
| numeric::decimal_rounding_mode rounding_mode, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| CUDF_FUNC_RANGE(); | ||
|
|
||
| CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes must match"); | ||
| // For decimal division, we only need the same base type (DECIMAL32/64/128) | ||
| // Different scales are allowed and expected | ||
| CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Column base types must match"); | ||
|
|
||
| CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL32 || lhs.type().id() == type_id::DECIMAL64 || | ||
| lhs.type().id() == type_id::DECIMAL128, | ||
| "Columns must be decimal type"); | ||
|
|
||
| // Note: Zero division check is handled in the kernel | ||
| // GPU kernels cannot throw exceptions, so they produce special values or assert in debug mode | ||
|
|
||
| if (lhs.is_empty()) { return make_empty_column(lhs.type()); } | ||
|
|
||
| return detail::divide_decimal_impl(lhs, rhs, rounding_mode, stream, mr); | ||
| } | ||
|
|
||
| std::unique_ptr<column> divide_decimal(column_view const& lhs, | ||
| scalar const& rhs, | ||
| numeric::decimal_rounding_mode rounding_mode, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| CUDF_FUNC_RANGE(); | ||
|
|
||
| CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL32 || lhs.type().id() == type_id::DECIMAL64 || | ||
| lhs.type().id() == type_id::DECIMAL128, | ||
| "Column must be decimal type"); | ||
| CUDF_EXPECTS(rhs.type() == lhs.type(), "Scalar type (" + std::to_string(int(rhs.type().id())) + ") must match column type (" + std::to_string(int(lhs.type().id())) + ")"); | ||
|
|
||
| if (lhs.is_empty()) { return make_empty_column(lhs.type()); } | ||
|
|
||
| using namespace numeric; | ||
|
|
||
| if (lhs.type().id() == type_id::DECIMAL32) { | ||
| using DecType = fixed_point<int32_t, Radix::BASE_10>; | ||
| return detail::divide_decimal_scalar_impl<DecType>(lhs, rhs, rounding_mode, stream, mr); | ||
| } else if (lhs.type().id() == type_id::DECIMAL64) { | ||
| using DecType = fixed_point<int64_t, Radix::BASE_10>; | ||
| return detail::divide_decimal_scalar_impl<DecType>(lhs, rhs, rounding_mode, stream, mr); | ||
| } else { | ||
| using DecType = fixed_point<__int128_t, Radix::BASE_10>; | ||
| return detail::divide_decimal_scalar_impl<DecType>(lhs, rhs, rounding_mode, stream, mr); | ||
| } | ||
| } | ||
|
|
||
| std::unique_ptr<column> divide_decimal(scalar const& lhs, | ||
| column_view const& rhs, | ||
| numeric::decimal_rounding_mode rounding_mode, | ||
| rmm::cuda_stream_view stream, | ||
| rmm::device_async_resource_ref mr) | ||
| { | ||
| CUDF_FUNC_RANGE(); | ||
|
|
||
| CUDF_EXPECTS(rhs.type().id() == type_id::DECIMAL32 || rhs.type().id() == type_id::DECIMAL64 || | ||
| rhs.type().id() == type_id::DECIMAL128, | ||
| "Column must be decimal type"); | ||
| CUDF_EXPECTS(lhs.type() == rhs.type(), "Scalar type must match column type"); | ||
|
|
||
| if (rhs.is_empty()) { return make_empty_column(rhs.type()); } | ||
|
|
||
| using namespace numeric; | ||
|
|
||
| // For scalar-column division, implement directly | ||
| auto const size = rhs.size(); | ||
| auto const rhs_type = rhs.type(); | ||
| auto const lhs_scale = lhs.type().scale(); | ||
| auto const rhs_scale = rhs_type.scale(); | ||
|
|
||
| // Create output column with same type as rhs | ||
| // If there are nulls in inputs, create with null mask; otherwise create without | ||
| std::unique_ptr<column> result; | ||
| if (!lhs.is_valid(stream) || rhs.has_nulls()) { | ||
| if (!lhs.is_valid(stream)) { | ||
| result = cudf::make_fixed_width_column( | ||
| rhs_type, | ||
| size, | ||
| cudf::detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr), | ||
| size, | ||
| stream, | ||
| mr); | ||
| } else { | ||
| result = cudf::make_fixed_width_column( | ||
| rhs_type, size, cudf::detail::copy_bitmask(rhs, stream, mr), rhs.null_count(), stream, mr); | ||
| } | ||
| } else { | ||
| // Create non-nullable column when inputs have no nulls | ||
| // Use empty rmm::device_buffer{} directly to ensure column is non-nullable | ||
| result = | ||
| std::make_unique<column>(rhs_type, | ||
| size, | ||
| rmm::device_buffer{size * cudf::size_of(rhs_type), stream, mr}, | ||
| rmm::device_buffer{}, // Empty buffer = non-nullable | ||
| 0, // null_count = 0 | ||
| std::vector<std::unique_ptr<column>>{}); | ||
| } | ||
|
|
||
| auto result_view = result->mutable_view(); | ||
|
|
||
| // Perform element-wise divide_decimal based on type | ||
| if (rhs_type.id() == type_id::DECIMAL32) { | ||
| using Type = int32_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else if (rhs_type.id() == type_id::DECIMAL64) { | ||
| using Type = int64_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } else { | ||
| using Type = __int128_t; | ||
| using DecType = fixed_point<Type, Radix::BASE_10>; | ||
|
|
||
| auto const& decimal_scalar = static_cast<fixed_point_scalar<DecType> const&>(lhs); | ||
| DecType lhs_fp = decimal_scalar.value(stream); | ||
| Type lhs_val = lhs_fp.value(); | ||
|
|
||
| thrust::transform( | ||
| rmm::exec_policy(stream), | ||
| rhs.begin<Type>(), | ||
| rhs.end<Type>(), | ||
| result_view.begin<Type>(), | ||
| [lhs_scale, rhs_scale, lhs_val, rounding_mode] __device__(Type rhs_val) { | ||
| DecType lhs_fp{numeric::scaled_integer<Type>{lhs_val, numeric::scale_type{lhs_scale}}}; | ||
| DecType rhs_fp{numeric::scaled_integer<Type>{rhs_val, numeric::scale_type{rhs_scale}}}; | ||
| auto result_fp = numeric::divide_decimal(lhs_fp, rhs_fp, rounding_mode); | ||
| return result_fp.value(); | ||
| }); | ||
| } | ||
|
|
||
| // Null mask already handled during column creation | ||
|
|
||
| return result; | ||
| } |
There was a problem hiding this comment.
Restructure into detail:: helpers with type-dispatch instead of per-overload giant type-id switches and inline scalar-column logic.
This mirrors previously-raised structural feedback: (1) the file should follow the standard pattern of thin public divide_decimal wrappers (CUDF_FUNC_RANGE() + delegate) over shared detail:: implementations across all three operand combinations (col/col, col/scalar, scalar/col), using cuda::constant_iterator for the scalar side instead of duplicating the type-dispatch/transform logic three times (lines 88-142, 327-383); (2) using namespace numeric; inside functions (lines 156, 259, 288) risks collisions/ambiguous resolution; (3) divide_decimal(scalar, column) (lines 273-389) inlines its whole implementation in the public function rather than delegating to a detail:: helper like the other two overloads do; (4) prefer std::format over +-concatenation for the type-mismatch message at line 255.
Based on past review comments recommending a detail::-based restructuring, flagging the giant non-exhaustive type-id switch, the using namespace risk, and the +-concatenation string building.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@cpp/src/binaryop/decimal_ops.cu` around lines 51 - 389, Restructure
divide_decimal so all three public overloads—column/column, column/scalar, and
scalar/column—remain thin CUDF_FUNC_RANGE() wrappers delegating to detail
helpers. Add shared type-dispatch helpers using cuda::constant_iterator for
scalar operands, eliminating duplicated type-id transform logic and moving
scalar/column implementation out of the public overload. Qualify numeric symbols
explicitly instead of using namespace numeric, and replace the scalar
type-mismatch message concatenation with std::format.
Comment on lines
+83
to
+85
| // Get device views | ||
| auto const lhs_dev = column_device_view::create(lhs, stream); | ||
| auto const rhs_dev = column_device_view::create(rhs, stream); |
There was a problem hiding this comment.
🚀 Performance & Scalability | 🟡 Minor | ⚡ Quick win
Unused device column views.
lhs_dev and rhs_dev are created via column_device_view::create(...) but never referenced — the subsequent thrust::transform uses lhs.begin<Type>()/rhs.begin<Type>() (host-side column_view iterators) directly. This creates two device view objects per call for no benefit.
🧹 Remove dead device views
- // Get device views
- auto const lhs_dev = column_device_view::create(lhs, stream);
- auto const rhs_dev = column_device_view::create(rhs, stream);
-
// Perform element-wise divide_decimal📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| // Get device views | |
| auto const lhs_dev = column_device_view::create(lhs, stream); | |
| auto const rhs_dev = column_device_view::create(rhs, stream); |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@cpp/src/binaryop/decimal_ops.cu` around lines 83 - 85, Remove the unused
lhs_dev and rhs_dev declarations created by column_device_view::create in the
decimal operation flow. Leave the existing thrust::transform inputs and all
other logic unchanged.
Comment on lines
+172
to
+184
| if (lhs.has_nulls() || !rhs.is_valid(stream)) { | ||
| if (!rhs.is_valid(stream)) { | ||
| result = cudf::make_fixed_width_column( | ||
| lhs_type, | ||
| size, | ||
| cudf::detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr), | ||
| size, | ||
| stream, | ||
| mr); | ||
| } else { | ||
| result = cudf::make_fixed_width_column( | ||
| lhs_type, size, cudf::detail::copy_bitmask(lhs, stream, mr), lhs.null_count(), stream, mr); | ||
| } |
There was a problem hiding this comment.
🚀 Performance & Scalability | 🟠 Major | ⚡ Quick win
Redundant is_valid() calls double the host-device synchronization per call.
In divide_decimal_scalar_impl (lines 172-184) and in the scalar-column divide_decimal overload (lines 299-311), rhs.is_valid(stream) / lhs.is_valid(stream) is each called twice in the same branch (once in the outer if, once again in the nested if). scalar::is_valid() performs a device-to-host read, so each call is a blocking sync point; calling it twice per invocation doubles avoidable synchronization on a path invoked for every scalar-column division.
⚡ Cache the validity check
- if (lhs.has_nulls() || !rhs.is_valid(stream)) {
- if (!rhs.is_valid(stream)) {
+ bool const rhs_valid = rhs.is_valid(stream);
+ if (lhs.has_nulls() || !rhs_valid) {
+ if (!rhs_valid) {and symmetrically for the lhs.is_valid(stream) checks at lines 299-300.
Also applies to: 299-311
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@cpp/src/binaryop/decimal_ops.cu` around lines 172 - 184, Cache the scalar
validity result once in divide_decimal_scalar_impl and the scalar-column
divide_decimal overload, then reuse that local value in both the outer and
nested conditionals. Apply the same change symmetrically to rhs.is_valid(stream)
and lhs.is_valid(stream), preserving the existing null-result behavior.
Source: Coding guidelines
Comment on lines
+47
to
+49
| // existing | ||
| from cudf._typing import ColumnBinaryOperand, ColumnLike, Dtype, ScalarLike | ||
| from cudf.core.buffer import Buffer |
There was a problem hiding this comment.
🎯 Functional Correctness | 🔴 Critical | ⚡ Quick win
Invalid syntax: // existing is not valid Python.
This line uses a JS/C-style comment (//) instead of Python's #, which is a syntax error that would break importing this module entirely. Confirmed independently by the Ruff hint at this exact line ("Expected a statement (invalid-syntax)").
🐛 Suggested fix
- // existing
from cudf._typing import ColumnBinaryOperand, ColumnLike, Dtype, ScalarLike
from cudf.core.buffer import Buffer📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| // existing | |
| from cudf._typing import ColumnBinaryOperand, ColumnLike, Dtype, ScalarLike | |
| from cudf.core.buffer import Buffer | |
| from cudf._typing import ColumnBinaryOperand, ColumnLike, Dtype, ScalarLike | |
| from cudf.core.buffer import Buffer |
🧰 Tools
🪛 Ruff (0.15.21)
[warning] 47-47: Expected a statement
(invalid-syntax)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@python/cudf/cudf/core/column/decimal.py` around lines 47 - 49, Replace the
invalid `// existing` comment near the imports with valid Python comment syntax
or remove it, leaving the `ColumnBinaryOperand`, `ColumnLike`, `Dtype`,
`ScalarLike`, and `Buffer` imports unchanged.
Source: Linters/SAST tools
Comment on lines
+316
to
+320
| # Type checking and normalization | ||
| reflect, _ = self._check_reflected_op("__div__") | ||
| other, other_cudf_dtype = self._normalize_binop_operand(other) # type: ignore[assignment] | ||
| if other is NotImplemented: | ||
| return NotImplemented |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
Returning NotImplemented from a non-dunder public method breaks the -> Self contract.
divide_decimal is called directly (not via Python's operator protocol), so returning the bare NotImplemented sentinel here just hands callers an unusable object instead of raising a clear error, and violates the declared -> Self return type.
🐛 Suggested fix
other, other_cudf_dtype = self._normalize_binop_operand(other) # type: ignore[assignment]
if other is NotImplemented:
- return NotImplemented
+ raise TypeError(
+ f"divide_decimal not supported between {self.dtype} and {type(other)}"
+ )📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| # Type checking and normalization | |
| reflect, _ = self._check_reflected_op("__div__") | |
| other, other_cudf_dtype = self._normalize_binop_operand(other) # type: ignore[assignment] | |
| if other is NotImplemented: | |
| return NotImplemented | |
| # Type checking and normalization | |
| reflect, _ = self._check_reflected_op("__div__") | |
| other, other_cudf_dtype = self._normalize_binop_operand(other) # type: ignore[assignment] | |
| if other is NotImplemented: | |
| raise TypeError( | |
| f"divide_decimal not supported between {self.dtype} and {type(other)}" | |
| ) |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@python/cudf/cudf/core/column/decimal.py` around lines 316 - 320, Update
divide_decimal’s handling of a NotImplemented result from
_normalize_binop_operand: because this public method must honor its -> Self
contract, raise an appropriate clear error instead of returning the bare
NotImplemented sentinel. Preserve normal division behavior for supported
operands.
Comment on lines
+42
to
+46
| cpdef Column divide_decimal( | ||
| Column lhs, | ||
| Column rhs, | ||
| rounding_mode=DecimalRoundingMode.HALF_UP | ||
| ): |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
# Check whether pylibcudf's Column/Scalar methods used here are cdef (crash-risk on None) or cpdef/def
ast-grep run --pattern 'cdef $RET view(self):' --lang cython python/pylibcudf/pylibcudf/column.pyx
ast-grep run --pattern 'cdef $RET type(self):' --lang cython python/pylibcudf/pylibcudf/column.pyx python/pylibcudf/pylibcudf/scalar.pyxRepository: rapidsai/cudf
Length of output: 377
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "=== decimal_division.pyx ==="
nl -ba python/pylibcudf/pylibcudf/decimal_division.pyx | sed -n '1,260p'
echo
echo "=== column.pyx: class and view/type methods ==="
rg -n "^(cdef|cpdef|def) .*class Column|^(cdef|cpdef|def) .*view|^(cdef|cpdef|def) .*type" python/pylibcudf/pylibcudf/column.pyx
echo
echo "=== scalar.pyx: class and type/get methods ==="
rg -n "^(cdef|cpdef|def) .*class Scalar|^(cdef|cpdef|def) .*type|^(cdef|cpdef|def) .*get" python/pylibcudf/pylibcudf/scalar.pyxRepository: rapidsai/cudf
Length of output: 220
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "=== decimal_division.pyx ==="
sed -n '1,260p' python/pylibcudf/pylibcudf/decimal_division.pyx
echo
echo "=== column.pyx: declarations around view/type ==="
rg -n -C 2 "view|type|get" python/pylibcudf/pylibcudf/column.pyx
echo
echo "=== scalar.pyx: declarations around type/get ==="
rg -n -C 2 "type|get" python/pylibcudf/pylibcudf/scalar.pyxRepository: rapidsai/cudf
Length of output: 50369
🌐 Web query:
Cython extension type function arguments None allowed not None semantics
💡 Result:
In Cython, when an argument is declared with an extension or Python builtin type (e.g., cdef func(MyExtType obj)), Cython defaults to allowing the value None in addition to instances of the declared type [1][2][3]. This is designed to behave like a C pointer, which can be NULL [1][4]. Because Cython prioritizes performance, it does not automatically check for None when accessing C-level attributes or methods of an extension type, which can lead to crashes or undefined behavior if the argument is None [1][5][3]. The not None clause and modern typing semantics are handled differently depending on the function type: 1. Python functions (def): - You can explicitly use the not None clause (e.g., def func(MyExtType x not None):) to forbid None values and raise a TypeError if one is passed [1][2]. - When using Python type annotations (e.g., x: MyExtType), Cython follows PEP-484 semantics [1][3]. Under this regime, the type is considered to not allow None unless typing.Optional (or tp | None / typing.Union[tp, None]) is used, or if a default value of None is provided (e.g., x: MyExtType = None) [1][3]. 2. C/Cython functions (cdef and cpdef): - The not None clause and typing.Optional annotations are generally not supported [1][2][3]. - If you need to ensure an argument is not None inside these functions, you must perform the check manually (e.g., if obj is None: raise TypeError(...)) [1][2][4]. For cpdef functions, which generate both a Python-accessible def interface and a fast C interface, these restrictions apply to the implementation; if you require explicit None handling for the Python interface, you may need to manage those checks or define separate logic to ensure safety across both call paths [6][7]. If performance allows, the compiler directive nonecheck can be enabled to force these checks automatically, though this incurs a performance penalty [5].
Citations:
- 1: http://cython.readthedocs.io/en/stable/src/userguide/extension_types.html
- 2: https://cython.readthedocs.io/en/3.0.x/src/userguide/extension_types.html
- 3: http://docs.cython.org/en/stable/src/userguide/extension%5Ftypes.html
- 4: http://docs.cython.org/en/latest/src/userguide/extension_types.html
- 5: https://cython.readthedocs.io/en/latest/src/tutorial/cdef_classes.html
- 6: https://cython-guidelines.readthedocs.io/en/latest/articles/cdef_cpdef_difference.html
- 7: https://cython.readthedocs.io/en/3.1.x/src/userguide/language_basics.html
🏁 Script executed:
#!/bin/bash
set -euo pipefail
python3 - <<'PY'
from pathlib import Path
path = Path("python/pylibcudf/pylibcudf/decimal_division.pyx")
lines = path.read_text().splitlines()
for start, end in [(1, 260)]:
print("=== decimal_division.pyx ===")
for i in range(start, min(end, len(lines)) + 1):
print(f"{i:4d}: {lines[i-1]}")
PYRepository: rapidsai/cudf
Length of output: 9704
Guard these inputs against None.
Typed Column/Scalar arguments here can still be None, so the first .type(), .view(), or .get() call can hit a null reference. Add explicit None checks in all three functions before using the operands.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@python/pylibcudf/pylibcudf/decimal_division.pyx` around lines 42 - 46, Guard
the operands in all three decimal-division functions, including divide_decimal,
before any calls to .type(), .view(), or .get(). Add explicit None validation
for each typed Column or Scalar input and preserve the existing behavior for
valid operands, using the module’s established invalid-argument error handling.
Comment on lines
+47
to
+84
| """ | ||
| Perform decimal division preserving the dividend's scale. | ||
|
|
||
| This function divides two decimal columns while maintaining the scale | ||
| of the dividend (left operand), similar to Java's BigDecimal.divide() | ||
| with a specified rounding mode. | ||
|
|
||
| Parameters | ||
| ---------- | ||
| lhs : Column | ||
| The dividend (left operand) - a decimal column | ||
| rhs : Column | ||
| The divisor (right operand) - a decimal column or scalar | ||
| rounding_mode : DecimalRoundingMode, optional | ||
| The rounding mode to use (default: HALF_UP) | ||
|
|
||
| Returns | ||
| ------- | ||
| Column | ||
| Result column with the same scale as the dividend | ||
|
|
||
| Raises | ||
| ------ | ||
| TypeError | ||
| If input columns are not decimal types | ||
| ValueError | ||
| If division by zero is attempted | ||
|
|
||
| Examples | ||
| -------- | ||
| >>> import pylibcudf | ||
| >>> # Create decimal columns | ||
| >>> lhs = pylibcudf.column_from_decimal_values([1.23, 4.56], scale=-2) | ||
| >>> rhs = pylibcudf.column_from_decimal_values([2.0, 3.0], scale=-1) | ||
| >>> # Divide preserving lhs scale | ||
| >>> result = divide_decimal(lhs, rhs, DecimalRoundingMode.HALF_UP) | ||
| >>> # Result has scale -2 (same as lhs) | ||
| """ |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
set -euo pipefail
printf '\n== File outline ==\n'
ast-grep outline python/pylibcudf/pylibcudf/decimal_division.pyx --view expanded || true
printf '\n== Search for column_from_decimal_values ==\n'
rg -n "column_from_decimal_values|divide_decimal|DecimalRoundingMode|division by zero|rounding_mode" python/pylibcudf -S
printf '\n== Search for decimal division C++/pyx implementation ==\n'
rg -n "divide_decimal|decimal division|division" cpp python/pylibcudf -S | head -n 200Repository: rapidsai/cudf
Length of output: 20276
🏁 Script executed:
set -euo pipefail
echo '== decimal_division.pyx =='
nl -ba python/pylibcudf/pylibcudf/decimal_division.pyx | sed -n '1,180p'
echo
echo '== search for column_from_decimal_values =='
rg -n "column_from_decimal_values" python -S || true
echo
echo '== search for divide_decimal implementation and error handling =='
rg -n "divide_decimal|rounding_mode|division by zero|invalid rounding_mode|ValueError" python/pylibcudf cpp -S || trueRepository: rapidsai/cudf
Length of output: 218
Fix the divide_decimal docstring
ValueError is raised for an invalid rounding_mode, not for division by zero, and the example uses pylibcudf.column_from_decimal_values(...), which doesn’t exist here.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@python/pylibcudf/pylibcudf/decimal_division.pyx` around lines 47 - 84, Update
the divide_decimal docstring’s Raises section to document ValueError for an
invalid rounding_mode instead of division by zero, and revise the example to use
the actual decimal-column construction API available in pylibcudf rather than
column_from_decimal_values.
Comment on lines
+14
to
+30
| cdef unique_ptr[column] divide_decimal( | ||
| const column_view& lhs, | ||
| const column_view& rhs, | ||
| decimal_rounding_mode rounding_mode | ||
| ) except + | ||
|
|
||
| cdef unique_ptr[column] divide_decimal( | ||
| const column_view& lhs, | ||
| const scalar& rhs, | ||
| decimal_rounding_mode rounding_mode | ||
| ) except + | ||
|
|
||
| cdef unique_ptr[column] divide_decimal( | ||
| const scalar& lhs, | ||
| const column_view& rhs, | ||
| decimal_rounding_mode rounding_mode | ||
| ) except + |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
# Confirm the conventional exception-handler import path used by other pylibcudf .pxd files
rg -n 'libcudf_exception_handler' python/pylibcudf/pylibcudf --type=cython -C2 2>/dev/null || \
rg -n 'libcudf_exception_handler' python/pylibcudf -C2Repository: rapidsai/cudf
Length of output: 50370
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "== decimal_ops.pxd =="
cat -n python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd
echo
echo "== related decimal binding files =="
for f in \
python/pylibcudf/pylibcudf/libcudf/round.pxd \
python/pylibcudf/pylibcudf/libcudf/scalar/scalar.pxd \
python/pylibcudf/pylibcudf/libcudf/decimal/*.pxd
do
[ -f "$f" ] && echo "--- $f" && sed -n '1,80p' "$f"
doneRepository: rapidsai/cudf
Length of output: 5869
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n -C 2 'divide_decimal|libcudf_exception_handler' python/pylibcudf/pylibcudf/libcudf/decimal python/pylibcudf/pylibcudf/libcudf/round.pxd python/pylibcudf/pylibcudf/libcudf/scalar/scalar.pxdRepository: rapidsai/cudf
Length of output: 6835
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "== exception_handler files =="
git ls-files python/pylibcudf/pylibcudf/exception_handler*
echo
echo "== exception_handler declaration =="
cat -n python/pylibcudf/pylibcudf/exception_handler.pxdRepository: rapidsai/cudf
Length of output: 3594
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "== error hierarchy =="
rg -n -C 3 'class logic_error|CUDF_EXPECTS|data_type_error|invalid_argument|std::logic_error' cpp include python/pylibcudf -g '!**/build/**'
echo
echo "== decimal divide implementation references =="
rg -n -C 3 'divide_decimal' cpp include python/pylibcudf -g '!**/build/**'Repository: rapidsai/cudf
Length of output: 50372
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n -C 2 'define CUDF_EXPECTS|CUDF_EXPECTS\(|class logic_error|struct logic_error|using logic_error|namespace cudf.*logic_error' cpp/include cpp/libcudf python/pylibcudf -g '!**/build/**' | head -n 300Repository: rapidsai/cudf
Length of output: 23698
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "== error.hpp definitions =="
rg -n -C 3 'class logic_error|class data_type_error|struct logic_error|struct data_type_error|using logic_error|using data_type_error' cpp/include/cudf/utilities/error.hpp
echo
echo "== decimal divide implementation =="
rg -n -C 4 'divide_decimal|CUDF_EXPECTS|throw' cpp/include cpp/src -g '*decimal*' -g '!**/build/**'Repository: rapidsai/cudf
Length of output: 22290
🏁 Script executed:
#!/bin/bash
set -euo pipefail
sed -n '1,140p' cpp/include/cudf/utilities/error.hppRepository: rapidsai/cudf
Length of output: 4278
Add libcudf_exception_handler to divide_decimal overloads
python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd should import libcudf_exception_handler and use except +libcudf_exception_handler on all three overloads so CUDF_EXPECTS failures map to the expected Python exception instead of falling back to RuntimeError.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@python/pylibcudf/pylibcudf/libcudf/decimal/decimal_ops.pxd` around lines 14 -
30, Import libcudf_exception_handler in the decimal bindings and update all
three divide_decimal overloads to use except +libcudf_exception_handler instead
of the generic exception specification, preserving their existing signatures and
arguments.
Source: Coding guidelines
Author
I haven't done any more with this yet. I just rebased it on master (trivial) and will work through the two reviews I've had so far, plus the quite-a-lot that CodeRabbit just said about it... |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
8 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Resurrection of #19861, rebased on
mainfor testing with Velox====
Implements divide_decimal for fixed-point decimal columns that preserves the dividend's scale, similar to Java BigDecimal.divide() with rounding mode.
Addresses issue #17448
(cherry picked from commit f8b1ff34ef9490b486ab76f2071b86c6d4d3423c)
Description
Checklist