Add decimal support to VARIANT casting - #23858
Open
vuule wants to merge 14 commits into
Open
Conversation
cast_variant and extract_variant_field now accept DECIMAL32/64/128 targets. The VARIANT encoding scales every value individually while a cuDF column carries a single scale, so each value is rescaled to the requested scale, truncating toward zero, and a value that no longer fits the target representation is nulled with the OVERFLOW status.
…hmark decimals Adds a DECIMAL16 test at the int128 limits, which the previous cases left the high half of the payload zeroed for, and a sliced 512-row case so the decimal kernel's grid-stride loop and slice offset are covered. Factors the incoming-status and null-bit preamble the cast paths share into should_decode_row, so the protocol lives in one place instead of three, and extends the variant nvbench with decimal32 and decimal128 cases.
vuule
added
feature request
Adds a DECIMAL64 arm to the overflow test, the only place the int64_t range check is reachable, and a decimal64 case to the cast benchmark's type axis.
# Conflicts: # cpp/tests/io/experimental/variant_extract_test.cpp
The cast target scale and the expected column scale must agree for these tests to mean anything, so route both through one named constant instead of repeating the literal.
Derive the overflow bounds from the target type's limits instead of literals a reviewer has to count digits in, fold the empty-input loops together, and make the interchangeable-widths case a typed test over the three fixed-point types.
vuule
commented
Sep 1, 2026
| d_output[row] = T{}; | ||
| continue; | ||
| } | ||
| if (!should_decode_row(row, d_null_mask, d_status)) { |
Contributor
Author
There was a problem hiding this comment.
No behavioral changes, just using the new helper
The per-digit loop paid a full 128-bit software division for every digit of rescale distance. Computing the divisor with ipow and dividing once, narrowed to 64 bits when both operands fit, cuts a two-digit decimal32 rescale from 168 to 107 us on 2M rows, against a 102 us baseline for a cast that needs no rescale. Also tightens a few comments in the shared row helper.
Give width 16 its own case so an unexpected width yields zero instead of reading 16 bytes, and fix a comment indent in the cast matrix test.
Contributor
Author
|
/ok to test e818b78 |
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (4)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesThe PR adds DECIMAL32, DECIMAL64, and DECIMAL128 decoding for Parquet VARIANT extraction and casting. It handles source-scale conversion, truncation, overflow, malformed payloads, operation statuses, and expanded benchmark and test coverage. VARIANT decimal casting
Estimated code review effort: 4 (Complex) | ~45 minutes Merge Risk: ⚪ Minimal · up to This change adds decimal VARIANT casting with documented overflow and malformed-input handling; no actionable merge-blocking risk remains after normal checks and review. Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
Full details: Linked Issues checkExplanation The changes address issue
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
mhaseeb123
approved these changes
Sep 1, 2026
mhaseeb123
left a comment
mhaseeb123
left a comment
Contributor
There was a problem hiding this comment.
Couldn't find anything worth requesting changes except some of the constexprs / magic numbers used here could use a one liner comment.
LGTM with some optional comments that you can take liberty in addressing
Comment on lines
+118
to
+119
| * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128` (see `cast_variant` | ||
| * for decimal rescaling) |
Contributor
There was a problem hiding this comment.
very optional nit: just reads better
Suggested change
| * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128` (see `cast_variant` | |
| * for decimal rescaling) | |
| * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128` | |
| * (see `cast_variant` for decimal rescaling) |
Comment on lines
+809
to
+815
| constexpr __int128_t max_over_10 = cuda::std::numeric_limits<__int128_t>::max() / 10; | ||
| constexpr __int128_t min_over_10 = cuda::std::numeric_limits<__int128_t>::min() / 10; | ||
| for (int i = 0; i < exp && value != 0; ++i) { | ||
| if (value > max_over_10 || value < min_over_10) { return cuda::std::nullopt; } | ||
| value *= 10; | ||
| } | ||
| return value; |
Contributor
There was a problem hiding this comment.
We can use cuda::mul_overflow here Something like the following but it requires __int128_t to satisfy the integer concept in concepts.cuh. Please check if this works and ignore if it errors out.
__device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t value, int exp)
{
for (int i = 0; i < exp && value != 0; ++i) {
auto r = ops::mul_overflow<__int128_t>(value, __int128_t{10});
if (!r) { return cuda::std::nullopt; }
value = *r;
}
return value;
}| return value / ipow<__int128_t, numeric::Radix::BASE_10>(exp); | ||
| } | ||
|
|
||
| __device__ int variant_decimal_unscaled_width(primitive_type ptype) |
Contributor
There was a problem hiding this comment.
Suggested change
| __device__ int variant_decimal_unscaled_width(primitive_type ptype) | |
| __device__ int constexpr variant_decimal_unscaled_width(primitive_type ptype) |
| } | ||
|
|
||
| // Divide `value` by 10^exp, truncating toward zero. | ||
| __device__ __int128_t divide_pow10(__int128_t value, int exp) |
Contributor
There was a problem hiding this comment.
Suggested change
| __device__ __int128_t divide_pow10(__int128_t value, int exp) | |
| __device__ __int128_t constexpr divide_pow10(__int128_t value, int exp) |
| constexpr int variant_decimal_max_scale = 38; | ||
|
|
||
| // Multiply `value` by 10^exp, or return nullopt if the result does not fit in `__int128_t`. | ||
| __device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int exp) |
Contributor
There was a problem hiding this comment.
Suggested change
| __device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int exp) | |
| __device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t value, int exp) |
Comment on lines
+979
to
+1011
| /** | ||
| * @brief Per-row kernel: decode each VARIANT decimal value blob into a fixed-point representation | ||
| * of type `Rep`, rescaled to `desired_scale`. Same null and status protocol as | ||
| * `cast_variant_primitive_kernel`. | ||
| */ | ||
| template <typename Rep> | ||
| CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_decimal_kernel( | ||
| cudf::lists_column_device_view values, | ||
| device_span<Rep> d_output, | ||
| int desired_scale, | ||
| bitmask_type* d_null_mask, | ||
| op_status* d_status) // nullptr when no status was requested | ||
| { | ||
| auto const num_rows = static_cast<size_type>(d_output.size()); | ||
| auto const tid = cudf::detail::grid_1d::global_thread_id<block_size>(); | ||
| auto const stride = cudf::detail::grid_1d::grid_stride<block_size>(); | ||
|
|
||
| for (auto row = tid; row < num_rows; row += stride) { | ||
| if (!should_decode_row(row, d_null_mask, d_status)) { | ||
| d_output[row] = Rep{}; | ||
| continue; | ||
| } | ||
|
|
||
| auto const [value, status] = decode_decimal<Rep>(list_row_span(values, row), desired_scale); | ||
| if (status == op_status::SUCCESS) { | ||
| d_output[row] = value; | ||
| } else { | ||
| d_output[row] = Rep{}; | ||
| cudf::clear_bit(d_null_mask, row); | ||
| } | ||
| if (d_status != nullptr) { d_status[row] = status; } | ||
| } | ||
| } |
Contributor
There was a problem hiding this comment.
Optional: Since this kernel processes one element per thread, we could make this a functor and launch via thrust::transform
| { | ||
| constexpr __int128_t max_over_10 = cuda::std::numeric_limits<__int128_t>::max() / 10; | ||
| constexpr __int128_t min_over_10 = cuda::std::numeric_limits<__int128_t>::min() / 10; | ||
| for (int i = 0; i < exp && value != 0; ++i) { |
Contributor
There was a problem hiding this comment.
Exit early once exp exceeds the max spec scale similar to divide_pow10, so a large user-supplied desired_scale can't force many redundant iterations after overflow is inevitable.
Suggested change
| for (int i = 0; i < exp && value != 0; ++i) { | |
| if (exp > variant_decimal_max_scale) { return cuda::std::nullopt; } | |
| for (int i = 0; i < exp && value != 0; ++i) { |
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
3 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.
Description
closes #23817
cast_variantandextract_variant_fieldnow acceptDECIMAL32/DECIMAL64/DECIMAL128target types, decoding the DECIMAL4/8/16 VARIANT primitives.The encoding stores a scale per value while a cuDF column carries a single scale, so each value is rescaled to
desired_type.scale(). A value that does not fit the target after rescaling is nulled and reported asvariant_operation_status::OVERFLOW. An out-of-range scale byte or a truncated payload reportsMALFORMED_VARIANT, and a non-decimal encoding reportsTYPE_MISMATCH.Also added test and benchmark coverage.
Checklist