From 3964950ba2fecf7f962917276058a6381d396246 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Tue, 3 Oct 2023 15:11:15 -0500 Subject: [PATCH] Fix inaccurate ceil/floor and inaccurate rescaling casts of fixed-point values. (#14242) This is a follow-up PR to #14233. This PR fixes a bug where floating-point values were used as intermediates in ceil/floor unary operations and cast operations that require rescaling for fixed-point types, giving inaccurate results. See also: - https://github.com/rapidsai/cudf/pull/14233#discussion_r1340786769 - https://github.com/rapidsai/cudf/issues/14243 Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Mike Wilson (https://github.com/hyperbolic2346) - Vukasin Milovanovic (https://github.com/vuule) --- cpp/src/unary/cast_ops.cu | 8 +++++- cpp/src/unary/math_ops.cu | 8 ++++-- cpp/tests/unary/cast_tests.cpp | 40 ++++++++++++++++++++++++++++++ cpp/tests/unary/unary_ops_test.cpp | 33 ++++++++++++++++++++++++ 4 files changed, 86 insertions(+), 3 deletions(-) diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index f40ace8d10b..1c81f266200 100644 --- a/cpp/src/unary/cast_ops.cu +++ b/cpp/src/unary/cast_ops.cu @@ -199,7 +199,13 @@ std::unique_ptr rescale(column_view input, } return output_column; } - auto const scalar = make_fixed_point_scalar(std::pow(10, -diff), scale_type{diff}, stream); + + RepType scalar_value = 10; + for (int i = 1; i < -diff; ++i) { + scalar_value *= 10; + } + + auto const scalar = make_fixed_point_scalar(scalar_value, scale_type{diff}, stream); return detail::binary_operation(input, *scalar, binary_operator::DIV, type, stream, mr); } }; diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu index 961f3a9e720..d0cae81a9c8 100644 --- a/cpp/src/unary/math_ops.cu +++ b/cpp/src/unary/math_ops.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -295,7 +295,11 @@ std::unique_ptr unary_op_with(column_view const& input, input.type(), input.size(), copy_bitmask(input, stream, mr), input.null_count(), stream, mr); auto out_view = result->mutable_view(); - Type const n = std::pow(10, -input.type().scale()); + + Type n = 10; + for (int i = 1; i < -input.type().scale(); ++i) { + n *= 10; + } thrust::transform(rmm::exec_policy(stream), input.begin(), diff --git a/cpp/tests/unary/cast_tests.cpp b/cpp/tests/unary/cast_tests.cpp index 9506e1918c0..d565359a4ea 100644 --- a/cpp/tests/unary/cast_tests.cpp +++ b/cpp/tests/unary/cast_tests.cpp @@ -30,6 +30,8 @@ #include #include +#include + #include #include @@ -967,6 +969,44 @@ TYPED_TEST(FixedPointTests, Decimal128ToDecimalXXWithLargerScale) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointTests, ValidateCastRescalePrecision) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + // This test is designed to protect against floating point conversion + // introducing errors in fixed-point arithmetic. The rescaling that occurs + // during casting to different scales should only use fixed-precision math. + // Realistically, we are only able to show precision failures due to floating + // conversion in a few very specific circumstances where dividing by specific + // powers of 10 works against us. Some examples: 10^23, 10^25, 10^26, 10^27, + // 10^30, 10^32, 10^36. See https://godbolt.org/z/cP1MddP8P for a derivation. + // For completeness and to ensure that we are not missing any other cases, we + // test casting to/from all scales in the range of each decimal type. Values + // that are powers of ten show this error more readily than non-powers of 10 + // because the rescaling factor is a power of 10, meaning that errors in + // division are more visible. + constexpr auto min_scale = -cuda::std::numeric_limits::digits10; + for (int input_scale = 0; input_scale >= min_scale; --input_scale) { + for (int result_scale = 0; result_scale >= min_scale; --result_scale) { + RepType input_value = 1; + for (int k = 0; k > input_scale; --k) { + input_value *= 10; + } + RepType result_value = 1; + for (int k = 0; k > result_scale; --k) { + result_value *= 10; + } + auto const input = fp_wrapper{{input_value}, scale_type{input_scale}}; + auto const expected = fp_wrapper{{result_value}, scale_type{result_scale}}; + auto const result = cudf::cast(input, make_fixed_point_data_type(result_scale)); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); + } + } +} + TYPED_TEST(FixedPointTests, Decimal32ToDecimalXXWithLargerScaleAndNullMask) { using namespace numeric; diff --git a/cpp/tests/unary/unary_ops_test.cpp b/cpp/tests/unary/unary_ops_test.cpp index 49764f22373..76d1f769856 100644 --- a/cpp/tests/unary/unary_ops_test.cpp +++ b/cpp/tests/unary/unary_ops_test.cpp @@ -24,6 +24,8 @@ #include +#include + template cudf::test::fixed_width_column_wrapper create_fixed_columns(cudf::size_type start, cudf::size_type size, @@ -372,4 +374,35 @@ TYPED_TEST(FixedPointUnaryTests, FixedPointUnaryFloorLarge) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointUnaryTests, ValidateCeilFloorPrecision) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + // This test is designed to protect against floating point conversion + // introducing errors in fixed-point arithmetic. The rounding that occurs + // during ceil/floor should only use fixed-precision math. Realistically, + // we are only able to show precision failures due to floating conversion in + // a few very specific circumstances where dividing by specific powers of 10 + // works against us. Some examples: 10^23, 10^25, 10^26, 10^27, 10^30, + // 10^32, 10^36. See https://godbolt.org/z/cP1MddP8P for a derivation. For + // completeness and to ensure that we are not missing any other cases, we + // test all scales representable in the range of each decimal type. + constexpr auto min_scale = -cuda::std::numeric_limits::digits10; + for (int input_scale = 0; input_scale >= min_scale; --input_scale) { + RepType input_value = 1; + for (int k = 0; k > input_scale; --k) { + input_value *= 10; + } + auto const input = fp_wrapper{{input_value}, scale_type{input_scale}}; + auto const expected = fp_wrapper{{input_value}, scale_type{input_scale}}; + auto const ceil_result = cudf::unary_operation(input, cudf::unary_operator::CEIL); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, ceil_result->view()); + auto const floor_result = cudf::unary_operation(input, cudf::unary_operator::FLOOR); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, floor_result->view()); + } +} + CUDF_TEST_PROGRAM_MAIN()