From 64be3f0a23a3af845f4ea9fe1d13578836ad0847 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 01:49:03 +0000 Subject: [PATCH 01/19] Add decimal support to VARIANT casting 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. --- cpp/include/cudf/io/experimental/variant.hpp | 13 +- .../parquet/experimental/variant_extract.cu | 201 ++++++++++++- .../io/experimental/variant_extract_test.cpp | 265 +++++++++++++++++- 3 files changed, 466 insertions(+), 13 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 08313ae21a71..db420f48f695 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -77,9 +77,14 @@ namespace io::parquet::experimental { * A null value is produced when the input row is null or the encoded type does not match * `desired_type`. * + * For a decimal `desired_type`, any of the encoded decimal widths decodes into the requested type + * and each value is rescaled from its own encoded scale to `desired_type.scale()`, truncating + * toward zero. A value that does not fit the target after rescaling produces a null row with + * `variant_operation_status::OVERFLOW`. + * * @param values `list` column of VARIANT-encoded value bytes * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, - * `FLOAT32`/`FLOAT64`, or `BOOL8`) + * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128`) * @param status Optional in-out parameter, `variant_operation_status` values, one per row. Must be * non-nullable, `UINT8`, and have the same row count as `values`. On input, its existing * values are treated as status from a prior `get_variant_field` call: rows already marked @@ -92,7 +97,8 @@ namespace io::parquet::experimental { * * @throws std::invalid_argument if `values` is not a `list` column; if `desired_type` * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, - * `FLOAT32`/`FLOAT64`, or `BOOL8`); or if `status` is provided but is nullable, not + * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128`); or if `status` + * is provided but is nullable, not * `UINT8`, or has a different row count than `values` */ [[nodiscard]] std::unique_ptr cast_variant( @@ -111,7 +117,8 @@ namespace io::parquet::experimental { * @param variant_column Struct column (VARIANT materialization) * @param path JSONPath-like path string (see `get_variant_field` for syntax) * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, - * `FLOAT32`/`FLOAT64`, or `BOOL8` + * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128` (see `cast_variant` + * for decimal rescaling) * @param status Optional. When provided, filled with `variant_operation_status` values, one per * row. Must be non-nullable, `UINT8`, and have the same row count as * `variant_column` diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index b5ef8d9ac654..0cd5a6ef1267 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -507,11 +507,16 @@ constexpr bool is_variant_int = template constexpr bool is_variant_numerical = is_variant_int || cudf::is_floating_point(); -// The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, bool, -// and strings. +// The fixed-point types a VARIANT decimal value can be decoded into: DECIMAL32/64/128. template -constexpr bool is_variant_castable = is_variant_numerical || cuda::std::is_same_v || - cuda::std::is_same_v; +constexpr bool is_variant_decimal = cudf::is_fixed_point(); + +// The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, +// decimals, bool, and strings. +template +constexpr bool is_variant_castable = + is_variant_numerical || is_variant_decimal || cuda::std::is_same_v || + cuda::std::is_same_v; // Maps a fixed-width output type to the VARIANT primitive type header id that encodes it. template @@ -775,6 +780,116 @@ __device__ op_status cast_status_for_primitive(device_span val) : op_status::MALFORMED_VARIANT; } +// Largest number of fractional digits a VARIANT decimal may declare (the spec caps the scale at the +// maximum precision of the widest decimal, DECIMAL16). +constexpr int variant_decimal_max_scale = 38; + +// Multiply `value` by 10^exp, or return nullopt if the result does not fit in `__int128_t`. The +// step-by-step check keeps the bound exact for any `exp`, including values far beyond the range of +// a decimal scale. +__device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int exp) +{ + 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; +} + +// Divide `value` by 10^exp, truncating toward zero. Iterating instead of dividing by a single power +// of ten keeps the result exact for an `exp` whose power of ten would itself overflow, where every +// value truncates to zero. +__device__ __int128_t divide_pow10(__int128_t value, int exp) +{ + for (int i = 0; i < exp && value != 0; ++i) { + value /= 10; + } + return value; +} + +// Byte width of a VARIANT decimal's unscaled integer payload, or 0 if `ptype` is not a decimal. +__device__ int variant_decimal_unscaled_width(primitive_type ptype) +{ + switch (ptype) { + case primitive_type::DECIMAL4: return 4; + case primitive_type::DECIMAL8: return 8; + case primitive_type::DECIMAL16: return 16; + default: return 0; + } +} + +/** + * @brief Decode a single VARIANT decimal value blob into the representation of a cuDF fixed-point + * type, rescaled to `desired_scale`. + * + * A decimal value is encoded as the value metadata byte, a one-byte unsigned scale (the number of + * fractional digits), and the little-endian two's-complement unscaled integer, 4, 8, or 16 bytes + * wide depending on the primitive type id. Any of the three widths decodes into any `Rep`, provided + * the rescaled value fits. + * + * cuDF carries a single scale for the whole column while the encoding scales each value + * individually, so every value is rescaled to `desired_scale` (a base-10 exponent, hence the + * negation of the encoded fractional digit count). Digits below the target scale are truncated + * toward zero. + * + * @return The rescaled representation, valid only when the returned status is `SUCCESS` + */ +template +__device__ cuda::std::pair decode_decimal(device_span enc, + int desired_scale) +{ + auto const fail = [](op_status status) { return cuda::std::pair{Rep{}, status}; }; + + if (enc.empty()) { return fail(op_status::MALFORMED_VARIANT); } + if (is_variant_null(enc)) { return fail(op_status::VARIANT_NULL); } + if (decode_basic_type(enc[0]) != basic_type::PRIMITIVE) { return fail(op_status::TYPE_MISMATCH); } + + auto const ptype = static_cast(variant_value_header(enc[0])); + auto const width = variant_decimal_unscaled_width(ptype); + if (width == 0) { + return fail(is_recognized_primitive_type(ptype) ? op_status::TYPE_MISMATCH + : op_status::MALFORMED_VARIANT); + } + + // Header byte, scale byte, then the unscaled integer. + constexpr size_type scale_bytes = 1; + if (cuda::std::cmp_less(enc.size(), variant_header_bytes + scale_bytes + width)) { + return fail(op_status::MALFORMED_VARIANT); + } + int const encoded_scale = enc[variant_header_bytes]; + if (encoded_scale > variant_decimal_max_scale) { return fail(op_status::MALFORMED_VARIANT); } + + auto const* unscaled_data = enc.data() + variant_header_bytes + scale_bytes; + auto const unscaled = [&]() -> __int128_t { + switch (width) { + case 4: return cudf::io::unaligned_load(unscaled_data); + case 8: return cudf::io::unaligned_load(unscaled_data); + default: return cudf::io::unaligned_load<__int128_t>(unscaled_data); + } + }(); + + // The encoded value is `unscaled * 10^-encoded_scale` and the output is `rep * 10^desired_scale`. + auto const shift = -encoded_scale - desired_scale; + __int128_t rescaled{}; + if (shift >= 0) { + auto const scaled = multiply_pow10(unscaled, shift); + if (!scaled.has_value()) { return fail(op_status::OVERFLOW); } + rescaled = scaled.value(); + } else { + rescaled = divide_pow10(unscaled, -shift); + } + + if constexpr (!cuda::std::is_same_v) { + if (rescaled < static_cast<__int128_t>(cuda::std::numeric_limits::min()) || + rescaled > static_cast<__int128_t>(cuda::std::numeric_limits::max())) { + return fail(op_status::OVERFLOW); + } + } + return {static_cast(rescaled), op_status::SUCCESS}; +} + /** * @brief Per-row kernel: decode each VARIANT value blob into a fixed-width primitive of type `T`. * @@ -833,6 +948,57 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_kernel( } } +/** + * @brief Per-row kernel: decode each VARIANT decimal value blob into a fixed-point representation + * of type `Rep`, rescaled to `desired_scale`. + * + * Follows the same null and status protocol as `cast_variant_primitive_kernel`; see + * `decode_decimal` for the accepted encodings. + */ +template +CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_decimal_kernel( + cudf::lists_column_device_view values, + device_span 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(d_output.size()); + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + for (auto row = tid; row < num_rows; row += stride) { + if (d_status != nullptr) { + // Status column is always non-nullable; row_null replaces the null bit. + auto const s = d_status[row]; + if (s != op_status::SUCCESS) { + d_output[row] = Rep{}; + if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } + continue; + } + if (!cudf::bit_is_set(d_null_mask, row)) { + d_output[row] = Rep{}; + d_status[row] = op_status::ROW_NULL; + continue; + } + } else { + if (!cudf::bit_is_set(d_null_mask, row)) { + d_output[row] = Rep{}; + continue; + } + } + + auto const [value, status] = decode_decimal(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; } + } +} + __device__ op_status cast_status_for_bool(device_span val) { if (val.empty()) { return op_status::MALFORMED_VARIANT; } @@ -974,6 +1140,28 @@ struct cast_variant_fn { null_count); } + template + std::unique_ptr operator()() + requires(is_variant_decimal) + { + using Rep = typename T::rep; + rmm::device_buffer data{num_rows * sizeof(Rep), stream, mr}; + auto const grid = cudf::detail::grid_1d{num_rows, block_size}; + auto const d_out = + device_span{static_cast(data.data()), static_cast(num_rows)}; + cast_variant_decimal_kernel<<>>( + values, d_out, desired_type.scale(), d_null_mask, d_status); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const null_count = + num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); + return std::make_unique(desired_type, + num_rows, + std::move(data), + null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, + null_count); + } + template std::unique_ptr operator()() requires(cuda::std::is_same_v) @@ -1239,7 +1427,10 @@ std::unique_ptr cast_variant(column_view const& values, case type_id::FLOAT32: case type_id::FLOAT64: case type_id::BOOL8: - case type_id::STRING: break; + case type_id::STRING: + case type_id::DECIMAL32: + case type_id::DECIMAL64: + case type_id::DECIMAL128: break; default: CUDF_FAIL("unsupported type for variant cast", std::invalid_argument); } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 0d5f12e63938..b311def675ba 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -491,6 +491,31 @@ inline std::vector enc_float64(double v) return out; } +// Decimal primitive blobs: header + 1-byte scale (the number of fractional digits) + the +// little-endian two's-complement unscaled integer, 4, 8, or 16 bytes wide. +inline std::vector enc_decimal4(int32_t unscaled, uint8_t scale) +{ + std::vector out{make_variant_primitive(variant_primitive_type::DECIMAL4), scale}; + append_le(out, static_cast(unscaled), 4); + return out; +} + +inline std::vector enc_decimal8(int64_t unscaled, uint8_t scale) +{ + std::vector out{make_variant_primitive(variant_primitive_type::DECIMAL8), scale}; + append_le(out, static_cast(unscaled), 8); + return out; +} + +inline std::vector enc_decimal16(__int128_t unscaled, uint8_t scale) +{ + std::vector out{make_variant_primitive(variant_primitive_type::DECIMAL16), scale}; + auto const bits = static_cast<__uint128_t>(unscaled); + append_le(out, static_cast(bits), 8); + append_le(out, static_cast(bits >> 64), 8); + return out; +} + // Long-string primitive blob: header + 4-byte LE length + payload. inline std::vector enc_long_string(std::string_view s) { @@ -655,6 +680,32 @@ TEST_F(ExtractVariantFieldTest, NestedPathMultiRowMixedNulls) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } +TEST_F(ExtractVariantFieldTest, NestedDecimalField) +{ + // Row 0: { price: DECIMAL8(123456, scale 3) } -> 123.456 -> 123.45 at scale -2 (truncated) + auto const m0 = build_metadata({"price"}); + auto const v0 = build_single_field_object(/*fid=price*/ 0, enc_decimal8(123456, 3)); + // Row 1: { price: DECIMAL4(-5, scale 1) } -> -0.5 + auto const m1 = build_metadata({"price"}); + auto const v1 = build_single_field_object(/*fid=price*/ 0, enc_decimal4(-5, 1)); + // Row 2: { other: DECIMAL4(1, scale 0) } -> key "price" missing from dict -> null + auto const m2 = build_metadata({"other"}); + auto const v2 = build_single_field_object(/*fid=other*/ 0, enc_decimal4(1, 0)); + + auto col = wrap_multi_row_variant({m0, m1, m2}, {v0, v1, v2}); + + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, + "price", + cudf::data_type{cudf::type_id::DECIMAL64, -2}, + std::nullopt, + cudf::test::get_default_stream()); + + cudf::test::fixed_point_column_wrapper expected{ + {12345, -50, 0}, {true, true, false}, numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + TEST_F(ExtractVariantFieldTest, EmptyPathRejected) { auto col = wrap_single_variant(build_metadata({}), enc_int32(1)); @@ -1326,6 +1377,172 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) } } +TEST_F(CastVariantTest, ApachePrimitiveDecimals) +{ + // The Apache fixtures all encode a scale of 2, so they decode exactly into a scale -2 column. + auto const stream = cudf::test::get_default_stream(); + auto const cast = [&](auto const& fixture, cudf::type_id id) { + auto col = make_apache_variant(fixture); + auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); + return cudf::io::parquet::experimental::cast_variant( + value, cudf::data_type{id, -2}, std::nullopt, stream); + }; + + { + auto got = cast(avf::primitive_decimal4, cudf::type_id::DECIMAL32); + cudf::test::fixed_point_column_wrapper expected{{1234}, numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + { + auto got = cast(avf::primitive_decimal8, cudf::type_id::DECIMAL64); + cudf::test::fixed_point_column_wrapper expected{{1234567890}, numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + { + auto got = cast(avf::primitive_decimal16, cudf::type_id::DECIMAL128); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234567891234567890}, + numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } +} + +TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) +{ + // Unlike the integer targets, a decimal target accepts every encoded decimal width: writers pick + // the narrowest width that holds each value, so one column can mix all three. + auto const stream = cudf::test::get_default_stream(); + std::vector> const val_rows{ + enc_decimal4(1234, 2), enc_decimal8(1234, 2), enc_decimal16(1234, 2)}; + auto col = + wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + { + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, + numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + { + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, + numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + { + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL128, -2}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234, 1234, 1234}, + numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } +} + +TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) +{ + // A column carries one scale while the encoding scales each value individually, so every value is + // rescaled. Digits below the requested scale are truncated toward zero, never rounded. + auto const stream = cudf::test::get_default_stream(); + std::vector> const val_rows{ + enc_decimal4(125, 2), // 1.25 -> 1.2 at scale -1 (truncated, not rounded to 1.3) + enc_decimal4(-125, 2), // -1.25 -> -1.2, truncation is toward zero for negatives too + enc_decimal4(199, 2), // 1.99 -> 1.9 + enc_decimal4(-199, 2), // -1.99 -> -1.9 + enc_decimal4(7, 0), // 7 -> 7.0 + enc_decimal8(30, 3), // 0.030 -> 0.0 + enc_decimal16(1, 38)}; // 1e-38 -> 0.0 + auto col = + wrap_multi_row_variant(std::vector>(7, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL32, -1}, std::nullopt, stream); + + cudf::test::fixed_point_column_wrapper expected{{12, -12, 19, -19, 70, 0, 0}, + numeric::scale_type{-1}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantTest, DecimalRescaledUp) +{ + // A requested scale below the encoded one multiplies the unscaled value up. + auto const stream = cudf::test::get_default_stream(); + auto col = wrap_multi_row_variant({build_metadata({})}, {enc_decimal4(1234, 2)}); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL64, -4}, std::nullopt, stream); + + cudf::test::fixed_point_column_wrapper expected{{123400}, numeric::scale_type{-4}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantTest, DecimalOverflowYieldsNull) +{ + // A value that does not fit the target representation after rescaling is dropped. + auto const stream = cudf::test::get_default_stream(); + std::vector> const val_rows{ + enc_decimal8(1234567890123, 2), // too large for an int32 representation + enc_decimal4(3000000, 0), // fits int32 as encoded, but rescaling up by 10^3 does not + enc_decimal4(1234, 2)}; // fits, to show the overflow is per row + auto col = + wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL32, -3}, std::nullopt, stream); + + cudf::test::fixed_point_column_wrapper expected{ + {0, 0, 12340}, {false, false, true}, numeric::scale_type{-3}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantTest, DecimalMalformedYieldsNull) +{ + // A scale beyond the spec maximum and a truncated unscaled payload are both malformed. + auto const stream = cudf::test::get_default_stream(); + auto out_of_range_scale = enc_decimal4(1234, 39); + auto truncated = enc_decimal8(1234, 2); + truncated.pop_back(); + + std::vector> const val_rows{out_of_range_scale, truncated}; + auto col = + wrap_multi_row_variant(std::vector>(2, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); + + EXPECT_EQ(got->size(), 2); + EXPECT_EQ(got->null_count(), 2); +} + +TEST_F(CastVariantTest, NonDecimalSourceYieldsNullForDecimalTarget) +{ + // Only the decimal primitives decode into a decimal target; there is no conversion from the + // integer, float, string, bool, or object encodings. + auto const stream = cudf::test::get_default_stream(); + std::vector> const val_rows{ + enc_int32(1234), + enc_int64(1234), + enc_float64(12.34), + enc_bool(true), + enc_short_string("12.34"), + enc_null(), + build_single_field_object(0, enc_decimal4(1, 0))}; + auto col = + wrap_multi_row_variant(std::vector>(7, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); + + EXPECT_EQ(got->size(), static_cast(val_rows.size())); + EXPECT_EQ(got->null_count(), static_cast(val_rows.size())); +} + TEST_F(CastVariantTest, ApacheShortString) { auto col = make_apache_variant(avf::short_string); @@ -1388,6 +1605,16 @@ TEST_F(CastVariantTest, EmptyInput) EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); } + + for (auto const id : + {cudf::type_id::DECIMAL32, cudf::type_id::DECIMAL64, cudf::type_id::DECIMAL128}) { + auto got = cudf::io::parquet::experimental::cast_variant( + *values, cudf::data_type{id, -2}, std::nullopt, stream); + EXPECT_EQ(got->type().id(), id); + EXPECT_EQ(got->type().scale(), -2); + EXPECT_EQ(got->size(), 0); + EXPECT_EQ(got->null_count(), 0); + } } TEST_F(CastVariantTest, UnsupportedTypeThrows) @@ -1402,10 +1629,7 @@ TEST_F(CastVariantTest, UnsupportedTypeThrows) cudf::type_id::TIMESTAMP_DAYS, cudf::type_id::TIMESTAMP_SECONDS, cudf::type_id::TIMESTAMP_MICROSECONDS, - cudf::type_id::DURATION_SECONDS, - cudf::type_id::DECIMAL32, - cudf::type_id::DECIMAL64, - cudf::type_id::DECIMAL128}; + cudf::type_id::DURATION_SECONDS}; // Empty input: the early-return path must still validate the type. auto const empty_values = @@ -1434,7 +1658,8 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) // are INT8/16/32/64 and STRING. Expected behaviour: // - integer targets: only a source whose physical type has the *exact* same width decodes; // every - // other source (including narrower/wider ints) yields null — cast_variant does not widen. + // other source (including narrower/wider ints and the decimals) yields null — cast_variant + // does not widen or convert between logical types. // - STRING target: short_string and long_string sources decode; every other source yields null. auto const stream = cudf::test::get_default_stream(); @@ -1453,6 +1678,9 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) {"float64", enc_float64(2.5)}, {"short_string", enc_short_string("hi")}, {"long_string", enc_long_string(std::string(70, 'a'))}, + {"decimal4", enc_decimal4(1234, 2)}, + {"decimal8", enc_decimal8(1234, 2)}, + {"decimal16", enc_decimal16(1234, 2)}, }; auto values_of = [](std::vector const& b) { @@ -2162,6 +2390,7 @@ constexpr uint8_t ST_MISSING = static_cast(op_status::MISSING_PATH); constexpr uint8_t ST_VNULL = static_cast(op_status::VARIANT_NULL); constexpr uint8_t ST_MISMATCH = static_cast(op_status::TYPE_MISMATCH); constexpr uint8_t ST_MALFORMED = static_cast(op_status::MALFORMED_VARIANT); +constexpr uint8_t ST_OVERFLOW = static_cast(op_status::OVERFLOW); // --------------------------------------------------------------------------- // GetVariantField status tests @@ -2507,6 +2736,32 @@ TEST_F(CastVariantStatusTest, BoolStatusTracking) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } +TEST_F(CastVariantStatusTest, DecimalStatusTracking) +{ + // Status for a decimal target: a decode, a variant null, a non-decimal primitive, a value that + // overflows the target representation, and a scale beyond the spec maximum. + auto stream = cudf::test::get_default_stream(); + + std::vector> const val_rows{ + enc_decimal4(1234, 2), // success + enc_null(), // variant_null + enc_int32(1234), // type_mismatch (recognized non-decimal primitive) + enc_decimal8(1234567890123, 2), // overflow (does not fit an int32 representation) + enc_decimal4(1234, 39)}; // malformed (scale above the spec maximum) + auto col = + wrap_multi_row_variant(std::vector>(5, build_metadata({})), val_rows); + auto values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto status = make_status_buffer(values.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_OVERFLOW, ST_MALFORMED}); + cudf::test::fixed_point_column_wrapper expected{ + {1234, 0, 0, 0, 0}, {true, false, false, false, false}, numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + TEST_F(CastVariantStatusTest, StringStatusTracking) { // Status for string target: short_string, variant_null, type_mismatch, malformed long_string, From 13700718b1f11cd07af277956bc13755462e78b5 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 02:03:08 +0000 Subject: [PATCH 02/19] Address review: int128 and sliced coverage, share cast preamble, benchmark 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. --- .../parquet/experimental/variant/extract.cpp | 39 ++++++++- .../parquet/experimental/variant_extract.cu | 81 ++++++++--------- .../io/experimental/variant_extract_test.cpp | 87 +++++++++++++++++++ 3 files changed, 165 insertions(+), 42 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 00c224c24651..c0076cf3ce3c 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -35,7 +35,15 @@ using cudf::io::parquet::experimental::variant_basic_type; using cudf::io::parquet::experimental::variant_primitive_type; // The leaf value type exercised by the benchmark (nvbench "type" string axis). -enum class bench_variant_type : uint8_t { INT32, FLOAT, BOOL, STRING, ARRAY }; +enum class bench_variant_type : uint8_t { + INT32, + FLOAT, + BOOL, + STRING, + ARRAY, + DECIMAL32, + DECIMAL128 +}; bench_variant_type parse_bench_variant_type(std::string const& type_str) { @@ -44,9 +52,15 @@ bench_variant_type parse_bench_variant_type(std::string const& type_str) if (type_str == "bool") { return bench_variant_type::BOOL; } if (type_str == "string") { return bench_variant_type::STRING; } if (type_str == "array") { return bench_variant_type::ARRAY; } + if (type_str == "decimal32") { return bench_variant_type::DECIMAL32; } + if (type_str == "decimal128") { return bench_variant_type::DECIMAL128; } CUDF_FAIL("Unrecognized benchmark type: " + type_str); } +// Number of fractional digits encoded by the decimal leaf values, and the matching cuDF column +// scale. The two agree so the cast measures decoding rather than rescaling. +constexpr uint8_t bench_decimal_scale = 2; + // Compose a value-metadata header byte from a basic type and its 6-bit value_header. // See cpp/tests/io/experimental/variant_extract_test.cpp for the header byte layout. constexpr uint8_t make_variant_header(variant_basic_type basic, uint8_t value_header) @@ -168,6 +182,23 @@ std::vector build_leaf_value(bench_variant_type type) out.insert(out.end(), s.begin(), s.end()); return out; } + case bench_variant_type::DECIMAL32: { + // Narrowest decimal encoding: 1-byte scale + 4-byte little-endian unscaled value. + std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL4), + bench_decimal_scale}; + append_le(out, 1234u, 4); + return out; + } + case bench_variant_type::DECIMAL128: { + // Widest decimal encoding: 1-byte scale + 16-byte little-endian unscaled value. The value + // needs more than 64 bits so the decode is not measured on an all-zero high half. + std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL16), + bench_decimal_scale}; + auto const unscaled = (static_cast<__uint128_t>(1234u) << 64) | 5678u; + append_le(out, static_cast(unscaled), 8); + append_le(out, static_cast(unscaled >> 64), 8); + return out; + } case bench_variant_type::ARRAY: { // VARIANT array of two INT32 values [42, 99]; element [1] is accessed in the benchmark. // 2 elements, offsets [0, 5, 10], then INT32(42) and INT32(99) (5 bytes each). @@ -368,6 +399,10 @@ cudf::data_type get_target_type(bench_variant_type type) case bench_variant_type::FLOAT: return cudf::data_type{cudf::type_id::FLOAT32}; case bench_variant_type::BOOL: return cudf::data_type{cudf::type_id::BOOL8}; case bench_variant_type::STRING: return cudf::data_type{cudf::type_id::STRING}; + case bench_variant_type::DECIMAL32: + return cudf::data_type{cudf::type_id::DECIMAL32, -bench_decimal_scale}; + case bench_variant_type::DECIMAL128: + return cudf::data_type{cudf::type_id::DECIMAL128, -bench_decimal_scale}; // "array": element access yields INT32. case bench_variant_type::INT32: case bench_variant_type::ARRAY: return cudf::data_type{cudf::type_id::INT32}; @@ -435,7 +470,7 @@ static void bench_variant_cast(nvbench::state& state) NVBENCH_BENCH(bench_variant_cast) .set_name("bench_variant_cast") .add_int64_axis("num_rows", {32768, 262144, 2097152}) - .add_string_axis("type", {"string", "float", "bool", "int32_t"}) + .add_string_axis("type", {"string", "float", "bool", "int32_t", "decimal32", "decimal128"}) .add_int64_axis("hit_rate", {20, 80}); // Benchmarks get_variant_field with varying path depth (nesting >= 1). Casting is exercised diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 0cd5a6ef1267..02b4aa4e7c16 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -890,6 +890,35 @@ __device__ cuda::std::pair decode_decimal(device_span(rescaled), op_status::SUCCESS}; } +/** + * @brief Shared per-row preamble for the cast paths: decides whether row `row` should be decoded. + * + * When a row must be skipped, the null bit is cleared and the status recorded before returning: + * a row carrying non-success status from a prior `get_variant_field` call keeps that status, and a + * SQL-null row becomes `ROW_NULL`. Callers are left to zero the output element themselves, since + * only they know its type. + * + * `d_status`, when present, is the in-out status buffer described on `cast_variant`; it is nullptr + * when no status was requested, in which case the null bit alone decides. + * + * @return True when the row's value blob should be decoded + */ +__device__ bool should_decode_row(size_type row, bitmask_type* d_null_mask, op_status* d_status) +{ + if (d_status == nullptr) { return cudf::bit_is_set(d_null_mask, row); } + + // Status column is always non-nullable; ROW_NULL replaces the null bit. + if (d_status[row] != op_status::SUCCESS) { + if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } + return false; + } + if (!cudf::bit_is_set(d_null_mask, row)) { + d_status[row] = op_status::ROW_NULL; + return false; + } + return true; +} + /** * @brief Per-row kernel: decode each VARIANT value blob into a fixed-width primitive of type `T`. * @@ -915,24 +944,9 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_kernel( auto const stride = cudf::detail::grid_1d::grid_stride(); for (auto row = tid; row < num_rows; row += stride) { - if (d_status != nullptr) { - // Status column is always non-nullable; row_null replaces the null bit. - auto const s = d_status[row]; - if (s != op_status::SUCCESS) { - d_output[row] = T{}; - if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } - continue; - } - if (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = T{}; - d_status[row] = op_status::ROW_NULL; - continue; - } - } else { - if (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = T{}; - continue; - } + if (!should_decode_row(row, d_null_mask, d_status)) { + d_output[row] = T{}; + continue; } auto const val = list_row_span(values, row); @@ -968,24 +982,9 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_decimal_kernel( auto const stride = cudf::detail::grid_1d::grid_stride(); for (auto row = tid; row < num_rows; row += stride) { - if (d_status != nullptr) { - // Status column is always non-nullable; row_null replaces the null bit. - auto const s = d_status[row]; - if (s != op_status::SUCCESS) { - d_output[row] = Rep{}; - if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } - continue; - } - if (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = Rep{}; - d_status[row] = op_status::ROW_NULL; - continue; - } - } else { - if (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = Rep{}; - continue; - } + if (!should_decode_row(row, d_null_mask, d_status)) { + d_output[row] = Rep{}; + continue; } auto const [value, status] = decode_decimal(list_row_span(values, row), desired_scale); @@ -1177,14 +1176,16 @@ struct cast_variant_fn { d_out = static_cast(data.data()), dnm = this->d_null_mask, dp_s] __device__(size_type row) { + if (!should_decode_row(row, dnm, dp_s)) { + d_out[row] = false; + return; + } + // The row is known live here, so a decode failure always clears its bit. auto const fail = [&](op_status s) { d_out[row] = false; - if (cudf::bit_is_set(dnm, row)) { cudf::clear_bit(dnm, row); } + cudf::clear_bit(dnm, row); if (dp_s) { dp_s[row] = s; } }; - if (dp_s and dp_s[row] != op_status::SUCCESS) { return fail(dp_s[row]); } - // Status column is always non-nullable; ROW_NULL replaces the null bit. - if (!cudf::bit_is_set(dnm, row)) { return fail(op_status::ROW_NULL); } auto const val = list_row_span(vals, row); auto const decoded = decode_bool(val); if (!decoded) { return fail(cast_status_for_bool(val)); } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index b311def675ba..458895c6ea8e 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1440,6 +1440,44 @@ TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) } } +TEST_F(CastVariantTest, Decimal16FullRange) +{ + // Exercises the high half of a 16-byte unscaled payload, which the values above and the Apache + // fixture all leave zeroed: every one of them fits in an int64_t. + auto const stream = cudf::test::get_default_stream(); + constexpr __int128_t int128_max = + static_cast<__int128_t>((~static_cast<__uint128_t>(0)) >> 1); // 2^127 - 1 + constexpr __int128_t int128_min = -int128_max - 1; + + std::vector> const val_rows{enc_decimal16(int128_max, 0), + enc_decimal16(int128_min, 0), + enc_decimal16(int128_max, 38), + enc_decimal16(int128_min, 38)}; + auto col = + wrap_multi_row_variant(std::vector>(4, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + // Scale 0 target: the scale-0 rows pass through untouched, while the scale-38 rows lose all their + // fractional digits (|int128_max| is just over 1.7e38, so it truncates to 1). + { + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL128, 0}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{int128_max, int128_min, 1, -1}, + numeric::scale_type{0}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + + // Scale -38 target: the scale-0 rows would need 10^38 more digits than the representation holds, + // while the scale-38 rows are already at that scale and pass through. + { + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL128, -38}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{ + {0, 0, int128_max, int128_min}, {false, false, true, true}, numeric::scale_type{-38}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } +} + TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) { // A column carries one scale while the encoding scales each value individually, so every value is @@ -1499,6 +1537,55 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } +TEST_F(CastVariantTest, DecimalSlicedMultiBlock) +{ + // The sliced window (slice_end - slice_beg = 512) spans more than one + // cast_variant_decimal_kernel block (block_size = 256), so this covers the grid-stride loop and + // a non-zero slice offset. Each row encodes a distinct value, so a row-indexing mistake cannot + // pass by coincidence, and the rows vary in encoded width to keep the value offsets irregular. + auto const stream = cudf::test::get_default_stream(); + constexpr int num_rows = 516; + constexpr int slice_beg = 3; + constexpr int slice_end = 515; + + std::vector> val_rows(num_rows); + std::vector exp_reps(num_rows); + std::vector exp_valid(num_rows); + for (int i = 0; i < num_rows; ++i) { + switch (i % 3) { + case 0: + val_rows[i] = enc_decimal4(i, 2); + exp_reps[i] = i; + exp_valid[i] = true; + break; + case 1: + val_rows[i] = enc_decimal8(i, 2); + exp_reps[i] = i; + exp_valid[i] = true; + break; + default: + val_rows[i] = enc_int32(i); // not a decimal encoding, so the row drops out + exp_reps[i] = 0; + exp_valid[i] = false; + break; + } + } + + auto col = wrap_multi_row_variant(std::vector>(num_rows, build_metadata({})), + val_rows); + auto const sliced = cudf::slice(col, {slice_beg, slice_end}).front(); + auto const values = cudf::structs_column_view{sliced}.get_sliced_child(1, stream); + + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, std::nullopt, stream); + + cudf::test::fixed_point_column_wrapper expected(exp_reps.begin() + slice_beg, + exp_reps.begin() + slice_end, + exp_valid.begin() + slice_beg, + numeric::scale_type{-2}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + TEST_F(CastVariantTest, DecimalMalformedYieldsNull) { // A scale beyond the spec maximum and a truncated unscaled payload are both malformed. From b93a33f6b1d776915034f8bbd85238f750a36f70 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 02:22:28 +0000 Subject: [PATCH 03/19] Trim verbose comments in variant decimal cast --- .../parquet/experimental/variant/extract.cpp | 7 +-- cpp/include/cudf/io/experimental/variant.hpp | 10 ++--- .../parquet/experimental/variant_extract.cu | 44 ++++++------------- .../io/experimental/variant_extract_test.cpp | 34 +++++--------- 4 files changed, 31 insertions(+), 64 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index c0076cf3ce3c..72f8dd329f7b 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -57,8 +57,7 @@ bench_variant_type parse_bench_variant_type(std::string const& type_str) CUDF_FAIL("Unrecognized benchmark type: " + type_str); } -// Number of fractional digits encoded by the decimal leaf values, and the matching cuDF column -// scale. The two agree so the cast measures decoding rather than rescaling. +// Shared by the encoded leaf values and the target column, so the cast measures decoding only. constexpr uint8_t bench_decimal_scale = 2; // Compose a value-metadata header byte from a basic type and its 6-bit value_header. @@ -183,15 +182,13 @@ std::vector build_leaf_value(bench_variant_type type) return out; } case bench_variant_type::DECIMAL32: { - // Narrowest decimal encoding: 1-byte scale + 4-byte little-endian unscaled value. std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL4), bench_decimal_scale}; append_le(out, 1234u, 4); return out; } case bench_variant_type::DECIMAL128: { - // Widest decimal encoding: 1-byte scale + 16-byte little-endian unscaled value. The value - // needs more than 64 bits so the decode is not measured on an all-zero high half. + // The value needs more than 64 bits, so the decode is not measured on an all-zero high half. std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL16), bench_decimal_scale}; auto const unscaled = (static_cast<__uint128_t>(1234u) << 64) | 5678u; diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index db420f48f695..507421341985 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -77,10 +77,9 @@ namespace io::parquet::experimental { * A null value is produced when the input row is null or the encoded type does not match * `desired_type`. * - * For a decimal `desired_type`, any of the encoded decimal widths decodes into the requested type - * and each value is rescaled from its own encoded scale to `desired_type.scale()`, truncating - * toward zero. A value that does not fit the target after rescaling produces a null row with - * `variant_operation_status::OVERFLOW`. + * For a decimal `desired_type`, every encoded width is accepted and each value is rescaled from its + * own encoded scale to `desired_type.scale()`, truncating toward zero; a value that no longer fits + * produces a null row with `variant_operation_status::OVERFLOW`. * * @param values `list` column of VARIANT-encoded value bytes * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, @@ -98,8 +97,7 @@ namespace io::parquet::experimental { * @throws std::invalid_argument if `values` is not a `list` column; if `desired_type` * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, * `FLOAT32`/`FLOAT64`, `BOOL8`, or `DECIMAL32`/`DECIMAL64`/`DECIMAL128`); or if `status` - * is provided but is nullable, not - * `UINT8`, or has a different row count than `values` + * is provided but is nullable, not `UINT8`, or has a different row count than `values` */ [[nodiscard]] std::unique_ptr cast_variant( column_view const& values, diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 02b4aa4e7c16..5a050e8fa466 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -780,13 +780,10 @@ __device__ op_status cast_status_for_primitive(device_span val) : op_status::MALFORMED_VARIANT; } -// Largest number of fractional digits a VARIANT decimal may declare (the spec caps the scale at the -// maximum precision of the widest decimal, DECIMAL16). +// The spec allows a scale in [0, 38] for every decimal width. constexpr int variant_decimal_max_scale = 38; -// Multiply `value` by 10^exp, or return nullopt if the result does not fit in `__int128_t`. The -// step-by-step check keeps the bound exact for any `exp`, including values far beyond the range of -// a decimal scale. +// 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) { constexpr __int128_t max_over_10 = cuda::std::numeric_limits<__int128_t>::max() / 10; @@ -798,9 +795,8 @@ __device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int return value; } -// Divide `value` by 10^exp, truncating toward zero. Iterating instead of dividing by a single power -// of ten keeps the result exact for an `exp` whose power of ten would itself overflow, where every -// value truncates to zero. +// Divide `value` by 10^exp, truncating toward zero. Iterating keeps an `exp` whose power of ten +// would itself overflow exact, since every value truncates to zero there. __device__ __int128_t divide_pow10(__int128_t value, int exp) { for (int i = 0; i < exp && value != 0; ++i) { @@ -809,7 +805,7 @@ __device__ __int128_t divide_pow10(__int128_t value, int exp) return value; } -// Byte width of a VARIANT decimal's unscaled integer payload, or 0 if `ptype` is not a decimal. +// Returns 0 for a `ptype` that is not a decimal. __device__ int variant_decimal_unscaled_width(primitive_type ptype) { switch (ptype) { @@ -824,15 +820,9 @@ __device__ int variant_decimal_unscaled_width(primitive_type ptype) * @brief Decode a single VARIANT decimal value blob into the representation of a cuDF fixed-point * type, rescaled to `desired_scale`. * - * A decimal value is encoded as the value metadata byte, a one-byte unsigned scale (the number of - * fractional digits), and the little-endian two's-complement unscaled integer, 4, 8, or 16 bytes - * wide depending on the primitive type id. Any of the three widths decodes into any `Rep`, provided - * the rescaled value fits. - * - * cuDF carries a single scale for the whole column while the encoding scales each value - * individually, so every value is rescaled to `desired_scale` (a base-10 exponent, hence the - * negation of the encoded fractional digit count). Digits below the target scale are truncated - * toward zero. + * Encoded as the value metadata byte, a one-byte scale (fractional digit count), then the + * little-endian two's-complement unscaled integer. Any of the three widths decodes into any `Rep` + * whose range fits the rescaled value; digits below `desired_scale` truncate toward zero. * * @return The rescaled representation, valid only when the returned status is `SUCCESS` */ @@ -853,7 +843,6 @@ __device__ cuda::std::pair decode_decimal(device_span decode_decimal(device_span= 0) { @@ -893,13 +882,8 @@ __device__ cuda::std::pair decode_decimal(device_span CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_decimal_kernel( diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 458895c6ea8e..8a47af323e87 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -491,8 +491,7 @@ inline std::vector enc_float64(double v) return out; } -// Decimal primitive blobs: header + 1-byte scale (the number of fractional digits) + the -// little-endian two's-complement unscaled integer, 4, 8, or 16 bytes wide. +// Decimal primitive blobs: header + 1-byte scale + little-endian two's-complement unscaled integer. inline std::vector enc_decimal4(int32_t unscaled, uint8_t scale) { std::vector out{make_variant_primitive(variant_primitive_type::DECIMAL4), scale}; @@ -1408,8 +1407,7 @@ TEST_F(CastVariantTest, ApachePrimitiveDecimals) TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) { - // Unlike the integer targets, a decimal target accepts every encoded decimal width: writers pick - // the narrowest width that holds each value, so one column can mix all three. + // Writers pick the narrowest width per value, so one column can mix all three. auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(1234, 2), enc_decimal8(1234, 2), enc_decimal16(1234, 2)}; @@ -1442,8 +1440,8 @@ TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) TEST_F(CastVariantTest, Decimal16FullRange) { - // Exercises the high half of a 16-byte unscaled payload, which the values above and the Apache - // fixture all leave zeroed: every one of them fits in an int64_t. + // Exercises the high half of a 16-byte payload, which every other decimal case here leaves + // zeroed. auto const stream = cudf::test::get_default_stream(); constexpr __int128_t int128_max = static_cast<__int128_t>((~static_cast<__uint128_t>(0)) >> 1); // 2^127 - 1 @@ -1457,8 +1455,7 @@ TEST_F(CastVariantTest, Decimal16FullRange) wrap_multi_row_variant(std::vector>(4, build_metadata({})), val_rows); auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); - // Scale 0 target: the scale-0 rows pass through untouched, while the scale-38 rows lose all their - // fractional digits (|int128_max| is just over 1.7e38, so it truncates to 1). + // The scale-38 rows lose every fractional digit; |int128_max| is just over 1.7e38. { auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL128, 0}, std::nullopt, stream); @@ -1467,8 +1464,7 @@ TEST_F(CastVariantTest, Decimal16FullRange) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } - // Scale -38 target: the scale-0 rows would need 10^38 more digits than the representation holds, - // while the scale-38 rows are already at that scale and pass through. + // The scale-0 rows would need 10^38 more digits than the representation holds. { auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL128, -38}, std::nullopt, stream); @@ -1480,8 +1476,7 @@ TEST_F(CastVariantTest, Decimal16FullRange) TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) { - // A column carries one scale while the encoding scales each value individually, so every value is - // rescaled. Digits below the requested scale are truncated toward zero, never rounded. + // Digits below the requested scale are truncated toward zero, never rounded. auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(125, 2), // 1.25 -> 1.2 at scale -1 (truncated, not rounded to 1.3) @@ -1539,10 +1534,8 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) TEST_F(CastVariantTest, DecimalSlicedMultiBlock) { - // The sliced window (slice_end - slice_beg = 512) spans more than one - // cast_variant_decimal_kernel block (block_size = 256), so this covers the grid-stride loop and - // a non-zero slice offset. Each row encodes a distinct value, so a row-indexing mistake cannot - // pass by coincidence, and the rows vary in encoded width to keep the value offsets irregular. + // The 512-row sliced window spans several kernel blocks (block_size = 256) at a non-zero offset. + // Values and encoded widths differ per row, so a row-indexing mistake cannot pass by coincidence. auto const stream = cudf::test::get_default_stream(); constexpr int num_rows = 516; constexpr int slice_beg = 3; @@ -1608,8 +1601,7 @@ TEST_F(CastVariantTest, DecimalMalformedYieldsNull) TEST_F(CastVariantTest, NonDecimalSourceYieldsNullForDecimalTarget) { - // Only the decimal primitives decode into a decimal target; there is no conversion from the - // integer, float, string, bool, or object encodings. + // Only the decimal primitives decode into a decimal target; nothing is converted. auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_int32(1234), @@ -1745,8 +1737,7 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) // are INT8/16/32/64 and STRING. Expected behaviour: // - integer targets: only a source whose physical type has the *exact* same width decodes; // every - // other source (including narrower/wider ints and the decimals) yields null — cast_variant - // does not widen or convert between logical types. + // other source (including narrower/wider ints and the decimals) yields null. // - STRING target: short_string and long_string sources decode; every other source yields null. auto const stream = cudf::test::get_default_stream(); @@ -2825,8 +2816,7 @@ TEST_F(CastVariantStatusTest, BoolStatusTracking) TEST_F(CastVariantStatusTest, DecimalStatusTracking) { - // Status for a decimal target: a decode, a variant null, a non-decimal primitive, a value that - // overflows the target representation, and a scale beyond the spec maximum. + // A decode, a variant null, a non-decimal primitive, an overflow, and an out-of-spec scale. auto stream = cudf::test::get_default_stream(); std::vector> const val_rows{ From 44c87493e61b55e3f755cec0818785c28fcf24a5 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 02:50:06 +0000 Subject: [PATCH 04/19] Remove redundant decimal cast tests, extend cast matrix with decimal target --- .../io/experimental/variant_extract_test.cpp | 94 ++++++------------- 1 file changed, 30 insertions(+), 64 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 8a47af323e87..95e942134484 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1498,20 +1498,6 @@ TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } -TEST_F(CastVariantTest, DecimalRescaledUp) -{ - // A requested scale below the encoded one multiplies the unscaled value up. - auto const stream = cudf::test::get_default_stream(); - auto col = wrap_multi_row_variant({build_metadata({})}, {enc_decimal4(1234, 2)}); - auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); - - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL64, -4}, std::nullopt, stream); - - cudf::test::fixed_point_column_wrapper expected{{123400}, numeric::scale_type{-4}}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); -} - TEST_F(CastVariantTest, DecimalOverflowYieldsNull) { // A value that does not fit the target representation after rescaling is dropped. @@ -1579,49 +1565,6 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } -TEST_F(CastVariantTest, DecimalMalformedYieldsNull) -{ - // A scale beyond the spec maximum and a truncated unscaled payload are both malformed. - auto const stream = cudf::test::get_default_stream(); - auto out_of_range_scale = enc_decimal4(1234, 39); - auto truncated = enc_decimal8(1234, 2); - truncated.pop_back(); - - std::vector> const val_rows{out_of_range_scale, truncated}; - auto col = - wrap_multi_row_variant(std::vector>(2, build_metadata({})), val_rows); - auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); - - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); - - EXPECT_EQ(got->size(), 2); - EXPECT_EQ(got->null_count(), 2); -} - -TEST_F(CastVariantTest, NonDecimalSourceYieldsNullForDecimalTarget) -{ - // Only the decimal primitives decode into a decimal target; nothing is converted. - auto const stream = cudf::test::get_default_stream(); - std::vector> const val_rows{ - enc_int32(1234), - enc_int64(1234), - enc_float64(12.34), - enc_bool(true), - enc_short_string("12.34"), - enc_null(), - build_single_field_object(0, enc_decimal4(1, 0))}; - auto col = - wrap_multi_row_variant(std::vector>(7, build_metadata({})), val_rows); - auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); - - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); - - EXPECT_EQ(got->size(), static_cast(val_rows.size())); - EXPECT_EQ(got->null_count(), static_cast(val_rows.size())); -} - TEST_F(CastVariantTest, ApacheShortString) { auto col = make_apache_variant(avf::short_string); @@ -1733,12 +1676,12 @@ TEST_F(CastVariantTest, UnsupportedTypeThrows) TEST_F(CastVariantTest, CastSourceTargetMatrix) { - // Exhaustively covers (source physical type) x (supported target) casts. The supported targets - // are INT8/16/32/64 and STRING. Expected behaviour: + // Exhaustively covers (source physical type) x (supported target) casts. Expected behaviour: // - integer targets: only a source whose physical type has the *exact* same width decodes; // every // other source (including narrower/wider ints and the decimals) yields null. // - STRING target: short_string and long_string sources decode; every other source yields null. + // - decimal target: every decimal width decodes; every other source yields null. auto const stream = cudf::test::get_default_stream(); struct source_blob { @@ -1759,6 +1702,7 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) {"decimal4", enc_decimal4(1234, 2)}, {"decimal8", enc_decimal8(1234, 2)}, {"decimal16", enc_decimal16(1234, 2)}, + {"object", build_single_field_object(0, enc_int32(1234))}, }; auto values_of = [](std::vector const& b) { @@ -1804,6 +1748,23 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) EXPECT_EQ(got->null_count(), 1); } } + + // Decimal target: all three encoded widths decode, since the sources share the encoded scale. + auto const decimal_type = cudf::data_type{cudf::type_id::DECIMAL32, -2}; + for (auto const& src : sources) { + SCOPED_TRACE(std::string{"decimal target, source "} + src.label); + auto values = values_of(src.bytes); + auto got = + cudf::io::parquet::experimental::cast_variant(values, decimal_type, std::nullopt, stream); + if (src.label.starts_with("decimal")) { + cudf::test::fixed_point_column_wrapper const expected{{1234}, + numeric::scale_type{-2}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } else { + ASSERT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 1); + } + } } TEST_F(CastVariantTest, ShortStringLengthZero) @@ -2816,26 +2777,31 @@ TEST_F(CastVariantStatusTest, BoolStatusTracking) TEST_F(CastVariantStatusTest, DecimalStatusTracking) { - // A decode, a variant null, a non-decimal primitive, an overflow, and an out-of-spec scale. + // A decode, a variant null, a non-decimal primitive, an overflow, and the two malformed forms. auto stream = cudf::test::get_default_stream(); + auto truncated_payload = enc_decimal8(1234, 2); + truncated_payload.pop_back(); + std::vector> const val_rows{ enc_decimal4(1234, 2), // success enc_null(), // variant_null enc_int32(1234), // type_mismatch (recognized non-decimal primitive) enc_decimal8(1234567890123, 2), // overflow (does not fit an int32 representation) - enc_decimal4(1234, 39)}; // malformed (scale above the spec maximum) + enc_decimal4(1234, 39), // malformed (scale above the spec maximum) + truncated_payload}; // malformed (unscaled payload shorter than the width implies) auto col = - wrap_multi_row_variant(std::vector>(5, build_metadata({})), val_rows); + wrap_multi_row_variant(std::vector>(6, build_metadata({})), val_rows); auto values = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto status = make_status_buffer(values.size()); auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, status->mutable_view(), stream, cmr()); - expect_status_values(*status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_OVERFLOW, ST_MALFORMED}); + expect_status_values( + *status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_OVERFLOW, ST_MALFORMED, ST_MALFORMED}); cudf::test::fixed_point_column_wrapper expected{ - {1234, 0, 0, 0, 0}, {true, false, false, false, false}, numeric::scale_type{-2}}; + {1234, 0, 0, 0, 0, 0}, {true, false, false, false, false, false}, numeric::scale_type{-2}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } From df1bcbcd4c36f73849258e46fc88a967106de0ea Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 23:19:43 +0000 Subject: [PATCH 05/19] Cover the int64 decimal representation in tests and benchmarks 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. --- .../parquet/experimental/variant/extract.cpp | 13 ++++++- .../io/experimental/variant_extract_test.cpp | 38 +++++++++++++------ 2 files changed, 38 insertions(+), 13 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 72f8dd329f7b..053eba7822e6 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -42,6 +42,7 @@ enum class bench_variant_type : uint8_t { STRING, ARRAY, DECIMAL32, + DECIMAL64, DECIMAL128 }; @@ -53,6 +54,7 @@ bench_variant_type parse_bench_variant_type(std::string const& type_str) if (type_str == "string") { return bench_variant_type::STRING; } if (type_str == "array") { return bench_variant_type::ARRAY; } if (type_str == "decimal32") { return bench_variant_type::DECIMAL32; } + if (type_str == "decimal64") { return bench_variant_type::DECIMAL64; } if (type_str == "decimal128") { return bench_variant_type::DECIMAL128; } CUDF_FAIL("Unrecognized benchmark type: " + type_str); } @@ -187,6 +189,12 @@ std::vector build_leaf_value(bench_variant_type type) append_le(out, 1234u, 4); return out; } + case bench_variant_type::DECIMAL64: { + std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL8), + bench_decimal_scale}; + append_le(out, (static_cast(1234u) << 32) | 5678u, 8); + return out; + } case bench_variant_type::DECIMAL128: { // The value needs more than 64 bits, so the decode is not measured on an all-zero high half. std::vector out{make_variant_primitive_header(variant_primitive_type::DECIMAL16), @@ -398,6 +406,8 @@ cudf::data_type get_target_type(bench_variant_type type) case bench_variant_type::STRING: return cudf::data_type{cudf::type_id::STRING}; case bench_variant_type::DECIMAL32: return cudf::data_type{cudf::type_id::DECIMAL32, -bench_decimal_scale}; + case bench_variant_type::DECIMAL64: + return cudf::data_type{cudf::type_id::DECIMAL64, -bench_decimal_scale}; case bench_variant_type::DECIMAL128: return cudf::data_type{cudf::type_id::DECIMAL128, -bench_decimal_scale}; // "array": element access yields INT32. @@ -467,7 +477,8 @@ static void bench_variant_cast(nvbench::state& state) NVBENCH_BENCH(bench_variant_cast) .set_name("bench_variant_cast") .add_int64_axis("num_rows", {32768, 262144, 2097152}) - .add_string_axis("type", {"string", "float", "bool", "int32_t", "decimal32", "decimal128"}) + .add_string_axis("type", + {"string", "float", "bool", "int32_t", "decimal32", "decimal64", "decimal128"}) .add_int64_axis("hit_rate", {20, 80}); // Benchmarks get_variant_field with varying path depth (nesting >= 1). Casting is exercised diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 95e942134484..6e43478b6582 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1502,20 +1502,34 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) { // A value that does not fit the target representation after rescaling is dropped. auto const stream = cudf::test::get_default_stream(); - std::vector> const val_rows{ - enc_decimal8(1234567890123, 2), // too large for an int32 representation - enc_decimal4(3000000, 0), // fits int32 as encoded, but rescaling up by 10^3 does not - enc_decimal4(1234, 2)}; // fits, to show the overflow is per row - auto col = - wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); - auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + auto const cast = [&](std::vector> const& rows, cudf::data_type target) { + auto col = wrap_multi_row_variant( + std::vector>(rows.size(), build_metadata({})), rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + return cudf::io::parquet::experimental::cast_variant(values, target, std::nullopt, stream); + }; - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, -3}, std::nullopt, stream); + { + auto got = cast({enc_decimal8(1234567890123, 2), // too large for an int32 representation + enc_decimal4(3000000, 0), // fits int32 as encoded, rescaling up by 10^3 does + // not + enc_decimal4(1234, 2)}, // fits, to show the overflow is per row + cudf::data_type{cudf::type_id::DECIMAL32, -3}); + cudf::test::fixed_point_column_wrapper expected{ + {0, 0, 12340}, {false, false, true}, numeric::scale_type{-3}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } - cudf::test::fixed_point_column_wrapper expected{ - {0, 0, 12340}, {false, false, true}, numeric::scale_type{-3}}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + // The int64 representation has its own bound, reachable only from a 16-byte encoded value. + { + constexpr __int128_t past_int64_max = static_cast<__int128_t>(10000000000000000000ULL); + auto got = cast( + {enc_decimal16(past_int64_max, 0), enc_decimal16(-past_int64_max, 0), enc_decimal8(1234, 2)}, + cudf::data_type{cudf::type_id::DECIMAL64, 0}); + cudf::test::fixed_point_column_wrapper expected{ + {0, 0, 12}, {false, false, true}, numeric::scale_type{0}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } } TEST_F(CastVariantTest, DecimalSlicedMultiBlock) From 305e6d747be92d97702df1f4347e2644b663eff5 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 27 Aug 2026 23:36:25 +0000 Subject: [PATCH 06/19] Use cudf::is_fixed_point directly instead of an alias that narrows nothing --- cpp/src/io/parquet/experimental/variant_extract.cu | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 5a050e8fa466..4c7ca8a3b429 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -507,15 +507,11 @@ constexpr bool is_variant_int = template constexpr bool is_variant_numerical = is_variant_int || cudf::is_floating_point(); -// The fixed-point types a VARIANT decimal value can be decoded into: DECIMAL32/64/128. -template -constexpr bool is_variant_decimal = cudf::is_fixed_point(); - // The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, // decimals, bool, and strings. template constexpr bool is_variant_castable = - is_variant_numerical || is_variant_decimal || cuda::std::is_same_v || + is_variant_numerical || cudf::is_fixed_point() || cuda::std::is_same_v || cuda::std::is_same_v; // Maps a fixed-width output type to the VARIANT primitive type header id that encodes it. @@ -1123,7 +1119,7 @@ struct cast_variant_fn { template std::unique_ptr operator()() - requires(is_variant_decimal) + requires(cudf::is_fixed_point()) { using Rep = typename T::rep; rmm::device_buffer data{num_rows * sizeof(Rep), stream, mr}; From 8327c7c2439c5782bbb7bdea04a17ac0723f79f8 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 17:52:45 +0000 Subject: [PATCH 07/19] Name the expected scale in the decimal cast tests 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. --- .../io/experimental/variant_extract_test.cpp | 106 +++++++++++------- 1 file changed, 63 insertions(+), 43 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 51305459f467..8989079786f3 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -722,6 +722,7 @@ TEST_F(ExtractVariantFieldTest, NestedPathMultiRowMixedNulls) TEST_F(ExtractVariantFieldTest, NestedDecimalField) { + constexpr int32_t expected_scale = -2; // Row 0: { price: DECIMAL8(123456, scale 3) } -> 123.456 -> 123.45 at scale -2 (truncated) auto const m0 = build_metadata({"price"}); auto const v0 = build_single_field_object(/*fid=price*/ 0, enc_decimal8(123456, 3)); @@ -737,12 +738,12 @@ TEST_F(ExtractVariantFieldTest, NestedDecimalField) auto got = cudf::io::parquet::experimental::extract_variant_field( col, "price", - cudf::data_type{cudf::type_id::DECIMAL64, -2}, + cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}, std::nullopt, cudf::test::get_default_stream()); cudf::test::fixed_point_column_wrapper expected{ - {12345, -50, 0}, {true, true, false}, numeric::scale_type{-2}}; + {12345, -50, 0}, {true, true, false}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1452,29 +1453,32 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) TEST_F(CastVariantTest, ApachePrimitiveDecimals) { - // The Apache fixtures all encode a scale of 2, so they decode exactly into a scale -2 column. - auto const stream = cudf::test::get_default_stream(); - auto const cast = [&](auto const& fixture, cudf::type_id id) { + // The Apache fixtures all encode a scale of 2, so they decode exactly into this column scale. + constexpr int32_t expected_scale = -2; + auto const stream = cudf::test::get_default_stream(); + auto const cast = [&](auto const& fixture, cudf::type_id id) { auto col = make_apache_variant(fixture); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); return cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{id, -2}, std::nullopt, stream); + value, cudf::data_type{id, expected_scale}, std::nullopt, stream); }; { auto got = cast(avf::primitive_decimal4, cudf::type_id::DECIMAL32); - cudf::test::fixed_point_column_wrapper expected{{1234}, numeric::scale_type{-2}}; + cudf::test::fixed_point_column_wrapper expected{{1234}, + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cast(avf::primitive_decimal8, cudf::type_id::DECIMAL64); - cudf::test::fixed_point_column_wrapper expected{{1234567890}, numeric::scale_type{-2}}; + cudf::test::fixed_point_column_wrapper expected{{1234567890}, + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cast(avf::primitive_decimal16, cudf::type_id::DECIMAL128); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234567891234567890}, - numeric::scale_type{-2}}; + cudf::test::fixed_point_column_wrapper<__int128_t> expected{ + {1234567891234567890}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1482,7 +1486,8 @@ TEST_F(CastVariantTest, ApachePrimitiveDecimals) TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) { // Writers pick the narrowest width per value, so one column can mix all three. - auto const stream = cudf::test::get_default_stream(); + constexpr int32_t expected_scale = -2; + auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(1234, 2), enc_decimal8(1234, 2), enc_decimal16(1234, 2)}; auto col = @@ -1491,23 +1496,23 @@ TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) { auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, std::nullopt, stream); + values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, - numeric::scale_type{-2}}; + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL64, -2}, std::nullopt, stream); + values, cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, - numeric::scale_type{-2}}; + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL128, -2}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234, 1234, 1234}, - numeric::scale_type{-2}}; + values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{ + {1234, 1234, 1234}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1531,19 +1536,23 @@ TEST_F(CastVariantTest, Decimal16FullRange) // The scale-38 rows lose every fractional digit; |int128_max| is just over 1.7e38. { - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL128, 0}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{{int128_max, int128_min, 1, -1}, - numeric::scale_type{0}}; + constexpr int32_t expected_scale = 0; + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); + cudf::test::fixed_point_column_wrapper<__int128_t> expected{ + {int128_max, int128_min, 1, -1}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } // The scale-0 rows would need 10^38 more digits than the representation holds. { - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL128, -38}, std::nullopt, stream); + constexpr int32_t expected_scale = -38; + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper<__int128_t> expected{ - {0, 0, int128_max, int128_min}, {false, false, true, true}, numeric::scale_type{-38}}; + {0, 0, int128_max, int128_min}, + {false, false, true, true}, + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1551,7 +1560,8 @@ TEST_F(CastVariantTest, Decimal16FullRange) TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) { // Digits below the requested scale are truncated toward zero, never rounded. - auto const stream = cudf::test::get_default_stream(); + constexpr int32_t expected_scale = -1; + auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(125, 2), // 1.25 -> 1.2 at scale -1 (truncated, not rounded to 1.3) enc_decimal4(-125, 2), // -1.25 -> -1.2, truncation is toward zero for negatives too @@ -1565,10 +1575,10 @@ TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, -1}, std::nullopt, stream); + values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper expected{{12, -12, 19, -19, 70, 0, 0}, - numeric::scale_type{-1}}; + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1584,24 +1594,26 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) }; { + constexpr int32_t expected_scale = -3; auto got = cast({enc_decimal8(1234567890123, 2), // too large for an int32 representation enc_decimal4(3000000, 0), // fits int32 as encoded, rescaling up by 10^3 does // not enc_decimal4(1234, 2)}, // fits, to show the overflow is per row - cudf::data_type{cudf::type_id::DECIMAL32, -3}); + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ - {0, 0, 12340}, {false, false, true}, numeric::scale_type{-3}}; + {0, 0, 12340}, {false, false, true}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } // The int64 representation has its own bound, reachable only from a 16-byte encoded value. { + constexpr int32_t expected_scale = 0; constexpr __int128_t past_int64_max = static_cast<__int128_t>(10000000000000000000ULL); auto got = cast( {enc_decimal16(past_int64_max, 0), enc_decimal16(-past_int64_max, 0), enc_decimal8(1234, 2)}, - cudf::data_type{cudf::type_id::DECIMAL64, 0}); + cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ - {0, 0, 12}, {false, false, true}, numeric::scale_type{0}}; + {0, 0, 12}, {false, false, true}, numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1610,10 +1622,11 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) { // The 512-row sliced window spans several kernel blocks (block_size = 256) at a non-zero offset. // Values and encoded widths differ per row, so a row-indexing mistake cannot pass by coincidence. - auto const stream = cudf::test::get_default_stream(); - constexpr int num_rows = 516; - constexpr int slice_beg = 3; - constexpr int slice_end = 515; + constexpr int32_t expected_scale = -2; + auto const stream = cudf::test::get_default_stream(); + constexpr int num_rows = 516; + constexpr int slice_beg = 3; + constexpr int slice_end = 515; std::vector> val_rows(num_rows); std::vector exp_reps(num_rows); @@ -1644,12 +1657,12 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) auto const values = cudf::structs_column_view{sliced}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, std::nullopt, stream); + values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper expected(exp_reps.begin() + slice_beg, exp_reps.begin() + slice_end, exp_valid.begin() + slice_beg, - numeric::scale_type{-2}); + numeric::scale_type{expected_scale}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1845,8 +1858,8 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) auto got = cudf::io::parquet::experimental::cast_variant(values, decimal_type, std::nullopt, stream); if (src.label.starts_with("decimal")) { - cudf::test::fixed_point_column_wrapper const expected{{1234}, - numeric::scale_type{-2}}; + cudf::test::fixed_point_column_wrapper const expected{ + {1234}, numeric::scale_type{decimal_type.scale()}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } else { ASSERT_EQ(got->size(), 1); @@ -2825,7 +2838,8 @@ TEST_F(CastVariantStatusTest, BoolStatusTracking) TEST_F(CastVariantStatusTest, DecimalStatusTracking) { // A decode, a variant null, a non-decimal primitive, an overflow, and the two malformed forms. - auto stream = cudf::test::get_default_stream(); + constexpr int32_t expected_scale = -2; + auto stream = cudf::test::get_default_stream(); auto truncated_payload = enc_decimal8(1234, 2); truncated_payload.pop_back(); @@ -2843,12 +2857,18 @@ TEST_F(CastVariantStatusTest, DecimalStatusTracking) auto status = make_status_buffer(values.size()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, -2}, status->mutable_view(), stream, cmr()); + values, + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, + status->mutable_view(), + stream, + cmr()); expect_status_values( *status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_OVERFLOW, ST_MALFORMED, ST_MALFORMED}); cudf::test::fixed_point_column_wrapper expected{ - {1234, 0, 0, 0, 0, 0}, {true, false, false, false, false, false}, numeric::scale_type{-2}}; + {1234, 0, 0, 0, 0, 0}, + {true, false, false, false, false, false}, + numeric::scale_type{expected_scale}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } From 405e29df8be13c55d634a2c580d5faa9912e4c40 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 17:59:25 +0000 Subject: [PATCH 08/19] Give expected_scale the scale_type so expectations pass it through --- .../io/experimental/variant_extract_test.cpp | 83 +++++++++---------- 1 file changed, 37 insertions(+), 46 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 8989079786f3..7da32a1ace8f 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -722,7 +722,7 @@ TEST_F(ExtractVariantFieldTest, NestedPathMultiRowMixedNulls) TEST_F(ExtractVariantFieldTest, NestedDecimalField) { - constexpr int32_t expected_scale = -2; + constexpr auto expected_scale = numeric::scale_type{-2}; // Row 0: { price: DECIMAL8(123456, scale 3) } -> 123.456 -> 123.45 at scale -2 (truncated) auto const m0 = build_metadata({"price"}); auto const v0 = build_single_field_object(/*fid=price*/ 0, enc_decimal8(123456, 3)); @@ -743,7 +743,7 @@ TEST_F(ExtractVariantFieldTest, NestedDecimalField) cudf::test::get_default_stream()); cudf::test::fixed_point_column_wrapper expected{ - {12345, -50, 0}, {true, true, false}, numeric::scale_type{expected_scale}}; + {12345, -50, 0}, {true, true, false}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1454,9 +1454,9 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) TEST_F(CastVariantTest, ApachePrimitiveDecimals) { // The Apache fixtures all encode a scale of 2, so they decode exactly into this column scale. - constexpr int32_t expected_scale = -2; - auto const stream = cudf::test::get_default_stream(); - auto const cast = [&](auto const& fixture, cudf::type_id id) { + constexpr auto expected_scale = numeric::scale_type{-2}; + auto const stream = cudf::test::get_default_stream(); + auto const cast = [&](auto const& fixture, cudf::type_id id) { auto col = make_apache_variant(fixture); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); return cudf::io::parquet::experimental::cast_variant( @@ -1465,20 +1465,18 @@ TEST_F(CastVariantTest, ApachePrimitiveDecimals) { auto got = cast(avf::primitive_decimal4, cudf::type_id::DECIMAL32); - cudf::test::fixed_point_column_wrapper expected{{1234}, - numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper expected{{1234}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cast(avf::primitive_decimal8, cudf::type_id::DECIMAL64); - cudf::test::fixed_point_column_wrapper expected{{1234567890}, - numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper expected{{1234567890}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cast(avf::primitive_decimal16, cudf::type_id::DECIMAL128); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{ - {1234567891234567890}, numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234567891234567890}, + expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1486,8 +1484,8 @@ TEST_F(CastVariantTest, ApachePrimitiveDecimals) TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) { // Writers pick the narrowest width per value, so one column can mix all three. - constexpr int32_t expected_scale = -2; - auto const stream = cudf::test::get_default_stream(); + constexpr auto expected_scale = numeric::scale_type{-2}; + auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(1234, 2), enc_decimal8(1234, 2), enc_decimal16(1234, 2)}; auto col = @@ -1497,22 +1495,19 @@ TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) { auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, - numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, - numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{ - {1234, 1234, 1234}, numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234, 1234, 1234}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1536,23 +1531,21 @@ TEST_F(CastVariantTest, Decimal16FullRange) // The scale-38 rows lose every fractional digit; |int128_max| is just over 1.7e38. { - constexpr int32_t expected_scale = 0; - auto got = cudf::io::parquet::experimental::cast_variant( + constexpr auto expected_scale = numeric::scale_type{0}; + auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{ - {int128_max, int128_min, 1, -1}, numeric::scale_type{expected_scale}}; + cudf::test::fixed_point_column_wrapper<__int128_t> expected{{int128_max, int128_min, 1, -1}, + expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } // The scale-0 rows would need 10^38 more digits than the representation holds. { - constexpr int32_t expected_scale = -38; - auto got = cudf::io::parquet::experimental::cast_variant( + constexpr auto expected_scale = numeric::scale_type{-38}; + auto got = cudf::io::parquet::experimental::cast_variant( values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper<__int128_t> expected{ - {0, 0, int128_max, int128_min}, - {false, false, true, true}, - numeric::scale_type{expected_scale}}; + {0, 0, int128_max, int128_min}, {false, false, true, true}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1560,8 +1553,8 @@ TEST_F(CastVariantTest, Decimal16FullRange) TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) { // Digits below the requested scale are truncated toward zero, never rounded. - constexpr int32_t expected_scale = -1; - auto const stream = cudf::test::get_default_stream(); + constexpr auto expected_scale = numeric::scale_type{-1}; + auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ enc_decimal4(125, 2), // 1.25 -> 1.2 at scale -1 (truncated, not rounded to 1.3) enc_decimal4(-125, 2), // -1.25 -> -1.2, truncation is toward zero for negatives too @@ -1578,7 +1571,7 @@ TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); cudf::test::fixed_point_column_wrapper expected{{12, -12, 19, -19, 70, 0, 0}, - numeric::scale_type{expected_scale}}; + expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1594,26 +1587,26 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) }; { - constexpr int32_t expected_scale = -3; + constexpr auto expected_scale = numeric::scale_type{-3}; auto got = cast({enc_decimal8(1234567890123, 2), // too large for an int32 representation enc_decimal4(3000000, 0), // fits int32 as encoded, rescaling up by 10^3 does // not enc_decimal4(1234, 2)}, // fits, to show the overflow is per row cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ - {0, 0, 12340}, {false, false, true}, numeric::scale_type{expected_scale}}; + {0, 0, 12340}, {false, false, true}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } // The int64 representation has its own bound, reachable only from a 16-byte encoded value. { - constexpr int32_t expected_scale = 0; + constexpr auto expected_scale = numeric::scale_type{0}; constexpr __int128_t past_int64_max = static_cast<__int128_t>(10000000000000000000ULL); auto got = cast( {enc_decimal16(past_int64_max, 0), enc_decimal16(-past_int64_max, 0), enc_decimal8(1234, 2)}, cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ - {0, 0, 12}, {false, false, true}, numeric::scale_type{expected_scale}}; + {0, 0, 12}, {false, false, true}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } @@ -1622,11 +1615,11 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) { // The 512-row sliced window spans several kernel blocks (block_size = 256) at a non-zero offset. // Values and encoded widths differ per row, so a row-indexing mistake cannot pass by coincidence. - constexpr int32_t expected_scale = -2; - auto const stream = cudf::test::get_default_stream(); - constexpr int num_rows = 516; - constexpr int slice_beg = 3; - constexpr int slice_end = 515; + constexpr auto expected_scale = numeric::scale_type{-2}; + auto const stream = cudf::test::get_default_stream(); + constexpr int num_rows = 516; + constexpr int slice_beg = 3; + constexpr int slice_end = 515; std::vector> val_rows(num_rows); std::vector exp_reps(num_rows); @@ -1662,7 +1655,7 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) cudf::test::fixed_point_column_wrapper expected(exp_reps.begin() + slice_beg, exp_reps.begin() + slice_end, exp_valid.begin() + slice_beg, - numeric::scale_type{expected_scale}); + expected_scale); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -2838,8 +2831,8 @@ TEST_F(CastVariantStatusTest, BoolStatusTracking) TEST_F(CastVariantStatusTest, DecimalStatusTracking) { // A decode, a variant null, a non-decimal primitive, an overflow, and the two malformed forms. - constexpr int32_t expected_scale = -2; - auto stream = cudf::test::get_default_stream(); + constexpr auto expected_scale = numeric::scale_type{-2}; + auto stream = cudf::test::get_default_stream(); auto truncated_payload = enc_decimal8(1234, 2); truncated_payload.pop_back(); @@ -2866,9 +2859,7 @@ TEST_F(CastVariantStatusTest, DecimalStatusTracking) expect_status_values( *status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_OVERFLOW, ST_MALFORMED, ST_MALFORMED}); cudf::test::fixed_point_column_wrapper expected{ - {1234, 0, 0, 0, 0, 0}, - {true, false, false, false, false, false}, - numeric::scale_type{expected_scale}}; + {1234, 0, 0, 0, 0, 0}, {true, false, false, false, false, false}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } From c8c166047c4a46834f02a2022e7ab7e698674032 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 18:48:40 +0000 Subject: [PATCH 09/19] Tidy up the decimal cast tests 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. --- .../io/experimental/variant_extract_test.cpp | 103 ++++++++---------- 1 file changed, 48 insertions(+), 55 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 7da32a1ace8f..97a4da1d5b01 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -8,6 +8,7 @@ #include #include #include +#include #include #include @@ -23,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -1400,9 +1402,9 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } - // Sliced multi-row column exercises grid-stride paths with a non-zero slice offset. - // num_rows / slice range is chosen so the sliced window (slice_end - slice_beg = 512) - // spans more than one cast_variant_bool_kernel block (block_size = 256). + // A 512-row sliced window at a non-zero offset, covering the row-to-blob mapping the bool path + // shares with the cast kernels. The bool path itself is a thrust::for_each, so it has no + // grid-stride loop of its own. { constexpr int num_rows = 516; constexpr int slice_beg = 3; @@ -1481,9 +1483,15 @@ TEST_F(CastVariantTest, ApachePrimitiveDecimals) } } -TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) +template +struct CastVariantDecimalTest : public cudf::test::BaseFixture {}; +TYPED_TEST_SUITE(CastVariantDecimalTest, cudf::test::FixedPointTypes); + +// Writers pick the narrowest width per value, so one column can mix all three, and every encoded +// width decodes into whichever decimal target was asked for. +TYPED_TEST(CastVariantDecimalTest, WidthsAreInterchangeable) { - // Writers pick the narrowest width per value, so one column can mix all three. + using Rep = typename TypeParam::rep; constexpr auto expected_scale = numeric::scale_type{-2}; auto const stream = cudf::test::get_default_stream(); std::vector> const val_rows{ @@ -1492,24 +1500,11 @@ TEST_F(CastVariantTest, DecimalWidthsAreInterchangeable) wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); - { - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, expected_scale}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); - } - { - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, expected_scale}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); - } - { - auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::DECIMAL128, expected_scale}, std::nullopt, stream); - cudf::test::fixed_point_column_wrapper<__int128_t> expected{{1234, 1234, 1234}, expected_scale}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); - } + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_to_id(), expected_scale}, std::nullopt, stream); + + cudf::test::fixed_point_column_wrapper expected{{1234, 1234, 1234}, expected_scale}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(CastVariantTest, Decimal16FullRange) @@ -1517,7 +1512,7 @@ TEST_F(CastVariantTest, Decimal16FullRange) // Exercises the high half of a 16-byte payload, which every other decimal case here leaves // zeroed. auto const stream = cudf::test::get_default_stream(); - constexpr __int128_t int128_max = + constexpr auto int128_max = static_cast<__int128_t>((~static_cast<__uint128_t>(0)) >> 1); // 2^127 - 1 constexpr __int128_t int128_min = -int128_max - 1; @@ -1577,7 +1572,7 @@ TEST_F(CastVariantTest, DecimalRescaledToRequestedScale) TEST_F(CastVariantTest, DecimalOverflowYieldsNull) { - // A value that does not fit the target representation after rescaling is dropped. + // A value that does not fit the target representation after rescaling becomes a null row. auto const stream = cudf::test::get_default_stream(); auto const cast = [&](std::vector> const& rows, cudf::data_type target) { auto col = wrap_multi_row_variant( @@ -1588,11 +1583,13 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) { constexpr auto expected_scale = numeric::scale_type{-3}; - auto got = cast({enc_decimal8(1234567890123, 2), // too large for an int32 representation - enc_decimal4(3000000, 0), // fits int32 as encoded, rescaling up by 10^3 does - // not - enc_decimal4(1234, 2)}, // fits, to show the overflow is per row - cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}); + constexpr auto int32_max = std::numeric_limits::max(); + auto got = + cast({enc_decimal8(int64_t{int32_max} + 1, 2), // past the int32 representation as encoded + enc_decimal4(int32_max / 1000 + 1, 0), // fits as encoded, but not after the 10^3 + // scale-up + enc_decimal4(1234, 2)}, // fits, to show the overflow is per row + cudf::data_type{cudf::type_id::DECIMAL32, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ {0, 0, 12340}, {false, false, true}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -1600,11 +1597,13 @@ TEST_F(CastVariantTest, DecimalOverflowYieldsNull) // The int64 representation has its own bound, reachable only from a 16-byte encoded value. { - constexpr auto expected_scale = numeric::scale_type{0}; - constexpr __int128_t past_int64_max = static_cast<__int128_t>(10000000000000000000ULL); - auto got = cast( - {enc_decimal16(past_int64_max, 0), enc_decimal16(-past_int64_max, 0), enc_decimal8(1234, 2)}, - cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}); + constexpr auto expected_scale = numeric::scale_type{0}; + // 2^64 rather than 2^63, so that the negated row is out of range too: -2^63 is int64_min. + constexpr auto out_of_int64_range = static_cast<__int128_t>(1) << 64; + auto got = cast({enc_decimal16(out_of_int64_range, 0), + enc_decimal16(-out_of_int64_range, 0), + enc_decimal8(1234, 2)}, + cudf::data_type{cudf::type_id::DECIMAL64, expected_scale}); cudf::test::fixed_point_column_wrapper expected{ {0, 0, 12}, {false, false, true}, expected_scale}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -1637,7 +1636,7 @@ TEST_F(CastVariantTest, DecimalSlicedMultiBlock) exp_valid[i] = true; break; default: - val_rows[i] = enc_int32(i); // not a decimal encoding, so the row drops out + val_rows[i] = enc_int32(i); // not a decimal encoding, so the row casts to null exp_reps[i] = 0; exp_valid[i] = false; break; @@ -1710,24 +1709,18 @@ TEST_F(CastVariantTest, EmptyInput) auto const values = cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); - for (auto const id : {cudf::type_id::INT32, - cudf::type_id::STRING, - cudf::type_id::FLOAT32, - cudf::type_id::FLOAT64, - cudf::type_id::BOOL8}) { - auto got = cudf::io::parquet::experimental::cast_variant( - *values, cudf::data_type{id}, std::nullopt, stream); - EXPECT_EQ(got->type().id(), id); - EXPECT_EQ(got->size(), 0); - EXPECT_EQ(got->null_count(), 0); - } - - for (auto const id : - {cudf::type_id::DECIMAL32, cudf::type_id::DECIMAL64, cudf::type_id::DECIMAL128}) { - auto got = cudf::io::parquet::experimental::cast_variant( - *values, cudf::data_type{id, -2}, std::nullopt, stream); - EXPECT_EQ(got->type().id(), id); - EXPECT_EQ(got->type().scale(), -2); + for (auto const target : {cudf::data_type{cudf::type_id::INT32}, + cudf::data_type{cudf::type_id::STRING}, + cudf::data_type{cudf::type_id::FLOAT32}, + cudf::data_type{cudf::type_id::FLOAT64}, + cudf::data_type{cudf::type_id::BOOL8}, + cudf::data_type{cudf::type_id::DECIMAL32, -2}, + cudf::data_type{cudf::type_id::DECIMAL64, -2}, + cudf::data_type{cudf::type_id::DECIMAL128, -2}}) { + SCOPED_TRACE(static_cast(target.id())); + auto got = cudf::io::parquet::experimental::cast_variant(*values, target, std::nullopt, stream); + EXPECT_EQ(got->type().id(), target.id()); + EXPECT_EQ(got->type().scale(), target.scale()); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); } @@ -1843,13 +1836,13 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) } } - // Decimal target: all three encoded widths decode, since the sources share the encoded scale. auto const decimal_type = cudf::data_type{cudf::type_id::DECIMAL32, -2}; for (auto const& src : sources) { SCOPED_TRACE(std::string{"decimal target, source "} + src.label); auto values = values_of(src.bytes); auto got = cudf::io::parquet::experimental::cast_variant(values, decimal_type, std::nullopt, stream); + // Decimal target: all three encoded widths decode, since the sources share the encoded scale. if (src.label.starts_with("decimal")) { cudf::test::fixed_point_column_wrapper const expected{ {1234}, numeric::scale_type{decimal_type.scale()}}; From 3fde3951117e3e197680d1ab6b155241e9b3a0d7 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 20:20:25 +0000 Subject: [PATCH 10/19] Divide once when rescaling a decimal, in 64 bits where possible 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. --- .../parquet/experimental/variant_extract.cu | 38 ++++++++++++------- 1 file changed, 24 insertions(+), 14 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 04bcc11101e3..1e8f8d0e6335 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -814,17 +815,24 @@ __device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int return value; } -// Divide `value` by 10^exp, truncating toward zero. Iterating keeps an `exp` whose power of ten -// would itself overflow exact, since every value truncates to zero there. +// Divide `value` by 10^exp, truncating toward zero. __device__ __int128_t divide_pow10(__int128_t value, int exp) { - for (int i = 0; i < exp && value != 0; ++i) { - value /= 10; + using numeric::detail::ipow; + + // Any `__int128_t` is smaller than 10^39, so a larger divisor truncates it away + if (exp > variant_decimal_max_scale) { return 0; } + + // 128-bit division is a slow software sequence; use the 64-bit one when both operands fit. + constexpr __int128_t i64_max = cuda::std::numeric_limits::max(); + constexpr __int128_t i64_min = cuda::std::numeric_limits::min(); + constexpr int max_int64_pow10 = 18; + if (exp <= max_int64_pow10 && value <= i64_max && value >= i64_min) { + return static_cast(value) / ipow(exp); } - return value; + return value / ipow<__int128_t, numeric::Radix::BASE_10>(exp); } -// Returns 0 for a `ptype` that is not a decimal. __device__ int variant_decimal_unscaled_width(primitive_type ptype) { switch (ptype) { @@ -899,23 +907,25 @@ __device__ cuda::std::pair decode_decimal(device_span Date: Tue, 1 Sep 2026 20:54:35 +0000 Subject: [PATCH 11/19] Load the unscaled payload only for a known decimal width 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. --- cpp/src/io/parquet/experimental/variant_extract.cu | 3 ++- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 1e8f8d0e6335..456219f6b7ff 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -882,7 +882,8 @@ __device__ cuda::std::pair decode_decimal(device_span(unscaled_data); case 8: return cudf::io::unaligned_load(unscaled_data); - default: return cudf::io::unaligned_load<__int128_t>(unscaled_data); + case 16: return cudf::io::unaligned_load<__int128_t>(unscaled_data); + default: return 0; // should be unreachable } }(); diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 97a4da1d5b01..2d7459243faf 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1842,7 +1842,7 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) auto values = values_of(src.bytes); auto got = cudf::io::parquet::experimental::cast_variant(values, decimal_type, std::nullopt, stream); - // Decimal target: all three encoded widths decode, since the sources share the encoded scale. + // Decimal target: all three encoded widths decode, since the sources share the encoded scale. if (src.label.starts_with("decimal")) { cudf::test::fixed_point_column_wrapper const expected{ {1234}, numeric::scale_type{decimal_type.scale()}}; From b15b724d14c8a56ab696fb09988027942f4be8b6 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 23:01:06 -0700 Subject: [PATCH 12/19] comment update Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/include/cudf/io/experimental/variant.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 507421341985..5677e22ea7ae 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -115,8 +115,8 @@ namespace io::parquet::experimental { * @param variant_column Struct column (VARIANT materialization) * @param path JSONPath-like path string (see `get_variant_field` for syntax) * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, - * `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) * @param status Optional. When provided, filled with `variant_operation_status` values, one per * row. Must be non-nullable, `UINT8`, and have the same row count as * `variant_column` From 370bd42b382ee1f3fa5247bd458b3f4c149dc087 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 23:17:45 -0700 Subject: [PATCH 13/19] constexpr1 Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/src/io/parquet/experimental/variant_extract.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 456219f6b7ff..8a5746e9bd57 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -833,7 +833,7 @@ __device__ __int128_t divide_pow10(__int128_t value, int exp) return value / ipow<__int128_t, numeric::Radix::BASE_10>(exp); } -__device__ int variant_decimal_unscaled_width(primitive_type ptype) +__device__ int constexpr variant_decimal_unscaled_width(primitive_type ptype) { switch (ptype) { case primitive_type::DECIMAL4: return 4; From 3e5e445316d2b64998849d5859606b0f6beb9ad7 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 23:18:11 -0700 Subject: [PATCH 14/19] constexpr2 Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/src/io/parquet/experimental/variant_extract.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 8a5746e9bd57..f91c4dae450e 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -816,7 +816,7 @@ __device__ cuda::std::optional<__int128_t> multiply_pow10(__int128_t value, int } // Divide `value` by 10^exp, truncating toward zero. -__device__ __int128_t divide_pow10(__int128_t value, int exp) +__device__ __int128_t constexpr divide_pow10(__int128_t value, int exp) { using numeric::detail::ipow; From d53483a8ef64a3b4ef7159a3f2063d3d132ddb5d Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 1 Sep 2026 23:18:26 -0700 Subject: [PATCH 15/19] constexpr3 Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/src/io/parquet/experimental/variant_extract.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index f91c4dae450e..0d4e51362178 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -804,7 +804,7 @@ __device__ op_status cast_status_for_primitive(device_span val) 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) +__device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t value, int exp) { 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; From 6f348af7e912f639357b30faf56baaa35163104b Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 2 Sep 2026 06:34:20 +0000 Subject: [PATCH 16/19] Answer the two decided cases of a decimal scale-up up front Zero fits at any scale and a nonzero value cannot survive a scale-up past 10^38, so neither needs the loop. Handling both up front also bounds the requested scale, which the public API does not validate. --- cpp/src/io/parquet/experimental/variant_extract.cu | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 456219f6b7ff..d3d6e4d9243d 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -806,9 +806,13 @@ 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) { + // Zero is representable at every scale, while any other value overflows past 10^38. + if (value == 0) { return 0; } + if (exp > variant_decimal_max_scale) { return cuda::std::nullopt; } + 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) { + for (int i = 0; i < exp; ++i) { if (value > max_over_10 || value < min_over_10) { return cuda::std::nullopt; } value *= 10; } From ba7a2a22205c51e6dcdce2a7ed9cd88d1b62fd03 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 2 Sep 2026 06:38:48 +0000 Subject: [PATCH 17/19] Fix clang-format spacing after the constexpr suggestions --- cpp/src/io/parquet/experimental/variant_extract.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 1690df080ffe..0dc431d8c998 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -837,7 +837,7 @@ __device__ __int128_t constexpr divide_pow10(__int128_t value, int exp) return value / ipow<__int128_t, numeric::Radix::BASE_10>(exp); } -__device__ int constexpr variant_decimal_unscaled_width(primitive_type ptype) +__device__ int constexpr variant_decimal_unscaled_width(primitive_type ptype) { switch (ptype) { case primitive_type::DECIMAL4: return 4; From 3fbdfe4c868d535b7209d645a3b8f70f6e0b631c Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 2 Sep 2026 06:55:54 +0000 Subject: [PATCH 18/19] paranoid overflow protection --- .../io/parquet/experimental/variant_extract.cu | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 0dc431d8c998..a6e6802da1d3 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -804,7 +804,7 @@ __device__ op_status cast_status_for_primitive(device_span val) 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> constexpr multiply_pow10(__int128_t value, int exp) +__device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t value, int64_t exp) { // Zero is representable at every scale, while any other value overflows past 10^38. if (value == 0) { return 0; } @@ -812,7 +812,7 @@ __device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t v 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; ++i) { + for (int64_t i = 0; i < exp; ++i) { if (value > max_over_10 || value < min_over_10) { return cuda::std::nullopt; } value *= 10; } @@ -820,21 +820,22 @@ __device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t v } // Divide `value` by 10^exp, truncating toward zero. -__device__ __int128_t constexpr divide_pow10(__int128_t value, int exp) +__device__ __int128_t constexpr divide_pow10(__int128_t value, int64_t exp) { using numeric::detail::ipow; // Any `__int128_t` is smaller than 10^39, so a larger divisor truncates it away if (exp > variant_decimal_max_scale) { return 0; } + auto const exponent = static_cast(exp); // 128-bit division is a slow software sequence; use the 64-bit one when both operands fit. constexpr __int128_t i64_max = cuda::std::numeric_limits::max(); constexpr __int128_t i64_min = cuda::std::numeric_limits::min(); constexpr int max_int64_pow10 = 18; - if (exp <= max_int64_pow10 && value <= i64_max && value >= i64_min) { - return static_cast(value) / ipow(exp); + if (exponent <= max_int64_pow10 && value <= i64_max && value >= i64_min) { + return static_cast(value) / ipow(exponent); } - return value / ipow<__int128_t, numeric::Radix::BASE_10>(exp); + return value / ipow<__int128_t, numeric::Radix::BASE_10>(exponent); } __device__ int constexpr variant_decimal_unscaled_width(primitive_type ptype) @@ -892,7 +893,7 @@ __device__ cuda::std::pair decode_decimal(device_span(encoded_scale) - desired_scale; __int128_t rescaled{}; if (shift >= 0) { auto const scaled = multiply_pow10(unscaled, shift); From e5b6dc69275488256320559fdbe3d890f8f10373 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 2 Sep 2026 22:45:54 +0000 Subject: [PATCH 19/19] Derive the decimal overflow bounds from the representation types The exponent guards and the 64-bit fast path depend on how many decimal digits the representation holds, so take those from digits10 instead of literals. The spec's cap on the encoded scale stays a literal, since no type provides it. --- cpp/src/io/parquet/experimental/variant_extract.cu | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index a6e6802da1d3..f68610b74924 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -803,12 +803,16 @@ __device__ op_status cast_status_for_primitive(device_span val) // The spec allows a scale in [0, 38] for every decimal width. constexpr int variant_decimal_max_scale = 38; +// The largest power of ten each representation holds; no value fits past it. +constexpr int max_int128_pow10 = cuda::std::numeric_limits<__int128_t>::digits10; +constexpr int max_int64_pow10 = cuda::std::numeric_limits::digits10; + // Multiply `value` by 10^exp, or return nullopt if the result does not fit in `__int128_t`. __device__ cuda::std::optional<__int128_t> constexpr multiply_pow10(__int128_t value, int64_t exp) { // Zero is representable at every scale, while any other value overflows past 10^38. if (value == 0) { return 0; } - if (exp > variant_decimal_max_scale) { return cuda::std::nullopt; } + if (exp > max_int128_pow10) { return cuda::std::nullopt; } 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; @@ -825,13 +829,12 @@ __device__ __int128_t constexpr divide_pow10(__int128_t value, int64_t exp) using numeric::detail::ipow; // Any `__int128_t` is smaller than 10^39, so a larger divisor truncates it away - if (exp > variant_decimal_max_scale) { return 0; } + if (exp > max_int128_pow10) { return 0; } auto const exponent = static_cast(exp); // 128-bit division is a slow software sequence; use the 64-bit one when both operands fit. - constexpr __int128_t i64_max = cuda::std::numeric_limits::max(); - constexpr __int128_t i64_min = cuda::std::numeric_limits::min(); - constexpr int max_int64_pow10 = 18; + constexpr __int128_t i64_max = cuda::std::numeric_limits::max(); + constexpr __int128_t i64_min = cuda::std::numeric_limits::min(); if (exponent <= max_int64_pow10 && value <= i64_max && value >= i64_min) { return static_cast(value) / ipow(exponent); }