From d07c35f7e41dabe6c15160b8b95af4e7312fab06 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Feb 2021 08:58:17 -0500 Subject: [PATCH 01/31] Initial changes for binary_v_v --- cpp/include/cudf/binaryop.hpp | 2 +- cpp/include/cudf/detail/binaryop.hpp | 2 +- cpp/src/binaryop/binaryop.cpp | 29 ++++++++++++------- cpp/tests/binaryop/binop-integration-test.cpp | 4 +++ 4 files changed, 25 insertions(+), 12 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 72abefef04f..71ac591b67e 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -146,7 +146,7 @@ std::unique_ptr binary_operation( column_view const& lhs, column_view const& rhs, binary_operator op, - data_type output_type, + thrust::optional output_type, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** diff --git a/cpp/include/cudf/detail/binaryop.hpp b/cpp/include/cudf/detail/binaryop.hpp index c12482967e1..147e5b71381 100644 --- a/cpp/include/cudf/detail/binaryop.hpp +++ b/cpp/include/cudf/detail/binaryop.hpp @@ -60,7 +60,7 @@ std::unique_ptr binary_operation( column_view const& lhs, column_view const& rhs, binary_operator op, - data_type output_type, + thrust::optional output_type, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index fc697267ca7..c6e8f506924 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -783,27 +783,36 @@ std::unique_ptr binary_operation(column_view const& lhs, std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, binary_operator op, - data_type output_type, + thrust::optional output_type, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes don't match"); - if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) - return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); - if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { - auto const type = - op == binary_operator::TRUE_DIV ? output_type : thrust::optional{thrust::nullopt}; - return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr); + if (op != binary_operator::TRUE_DIV) { + CUDF_EXPECTS( + not output_type.has_value(), + "Only TRUE_DIV supports specified output_type for fixed_point binary operations. For other " + "fixed_point binary operations, please pass {} or std::nullopt for output_type and the " + "cudf::data_type and numeric::scale_type will be automatically calculated."); + } + + return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); } + CUDF_EXPECTS(output_type.has_value(), "Must specify output_type of column."); + // Use output_type.value() for the rest of the function + + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) + return binops::compiled::binary_operation(lhs, rhs, op, output_type.value(), stream, mr); + // Check for datatype - CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); + CUDF_EXPECTS(is_fixed_width(type), "Invalid/Unsupported output datatype"); CUDF_EXPECTS(is_fixed_width(lhs.type()), "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(is_fixed_width(rhs.type()), "Invalid/Unsupported rhs datatype"); - auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type.value(), stream, mr); if (lhs.is_empty() or rhs.is_empty()) return out; @@ -868,7 +877,7 @@ std::unique_ptr binary_operation(column_view const& lhs, std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, binary_operator op, - data_type output_type, + thrust::optional output_type, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 2d17853a72b..b39f586b8ea 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -2123,6 +2124,9 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv2) auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, {}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); + auto output_type = cudf::data_type{type_to_id(), scale_type{1}}; + EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, output_type), + cudf::logic_error); } TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv3) From 0f528de86f4baf2ed9a00d3843c685406d92f89e Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 24 Feb 2021 10:41:33 -0500 Subject: [PATCH 02/31] Fix --- cpp/src/binaryop/binaryop.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index c6e8f506924..85fd38aec13 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -808,7 +808,7 @@ std::unique_ptr binary_operation(column_view const& lhs, return binops::compiled::binary_operation(lhs, rhs, op, output_type.value(), stream, mr); // Check for datatype - CUDF_EXPECTS(is_fixed_width(type), "Invalid/Unsupported output datatype"); + CUDF_EXPECTS(is_fixed_width(output_type.value()), "Invalid/Unsupported output datatype"); CUDF_EXPECTS(is_fixed_width(lhs.type()), "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(is_fixed_width(rhs.type()), "Invalid/Unsupported rhs datatype"); From 223fe40e22f0f97f4c5343856da756d403aa869d Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 13:54:15 -0500 Subject: [PATCH 03/31] Revert "Fix" This reverts commit 0f528de86f4baf2ed9a00d3843c685406d92f89e. --- cpp/src/binaryop/binaryop.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 85fd38aec13..c6e8f506924 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -808,7 +808,7 @@ std::unique_ptr binary_operation(column_view const& lhs, return binops::compiled::binary_operation(lhs, rhs, op, output_type.value(), stream, mr); // Check for datatype - CUDF_EXPECTS(is_fixed_width(output_type.value()), "Invalid/Unsupported output datatype"); + CUDF_EXPECTS(is_fixed_width(type), "Invalid/Unsupported output datatype"); CUDF_EXPECTS(is_fixed_width(lhs.type()), "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(is_fixed_width(rhs.type()), "Invalid/Unsupported rhs datatype"); From f8ec02b7bcfeb1b3f56ff2602e8e034ff5ab721b Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 13:54:26 -0500 Subject: [PATCH 04/31] Revert "Initial changes for binary_v_v" This reverts commit d07c35f7e41dabe6c15160b8b95af4e7312fab06. --- cpp/include/cudf/binaryop.hpp | 2 +- cpp/include/cudf/detail/binaryop.hpp | 2 +- cpp/src/binaryop/binaryop.cpp | 29 +++++++------------ cpp/tests/binaryop/binop-integration-test.cpp | 4 --- 4 files changed, 12 insertions(+), 25 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 71ac591b67e..72abefef04f 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -146,7 +146,7 @@ std::unique_ptr binary_operation( column_view const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + data_type output_type, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** diff --git a/cpp/include/cudf/detail/binaryop.hpp b/cpp/include/cudf/detail/binaryop.hpp index 147e5b71381..c12482967e1 100644 --- a/cpp/include/cudf/detail/binaryop.hpp +++ b/cpp/include/cudf/detail/binaryop.hpp @@ -60,7 +60,7 @@ std::unique_ptr binary_operation( column_view const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + data_type output_type, rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index c6e8f506924..fc697267ca7 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -783,36 +783,27 @@ std::unique_ptr binary_operation(column_view const& lhs, std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + data_type output_type, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(lhs.size() == rhs.size(), "Column sizes don't match"); - if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { - if (op != binary_operator::TRUE_DIV) { - CUDF_EXPECTS( - not output_type.has_value(), - "Only TRUE_DIV supports specified output_type for fixed_point binary operations. For other " - "fixed_point binary operations, please pass {} or std::nullopt for output_type and the " - "cudf::data_type and numeric::scale_type will be automatically calculated."); - } + if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) + return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); - return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { + auto const type = + op == binary_operator::TRUE_DIV ? output_type : thrust::optional{thrust::nullopt}; + return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr); } - CUDF_EXPECTS(output_type.has_value(), "Must specify output_type of column."); - // Use output_type.value() for the rest of the function - - if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) - return binops::compiled::binary_operation(lhs, rhs, op, output_type.value(), stream, mr); - // Check for datatype - CUDF_EXPECTS(is_fixed_width(type), "Invalid/Unsupported output datatype"); + CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); CUDF_EXPECTS(is_fixed_width(lhs.type()), "Invalid/Unsupported lhs datatype"); CUDF_EXPECTS(is_fixed_width(rhs.type()), "Invalid/Unsupported rhs datatype"); - auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type.value(), stream, mr); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty() or rhs.is_empty()) return out; @@ -877,7 +868,7 @@ std::unique_ptr binary_operation(column_view const& lhs, std::unique_ptr binary_operation(column_view const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + data_type output_type, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index b39f586b8ea..2d17853a72b 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -22,7 +22,6 @@ #include #include #include -#include #include #include @@ -2124,9 +2123,6 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv2) auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, {}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); - auto output_type = cudf::data_type{type_to_id(), scale_type{1}}; - EXPECT_THROW(cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, output_type), - cudf::logic_error); } TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv3) From e0bf2c6adcba1e53dc7526454e85e0e0d83b030a Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 15:23:13 -0500 Subject: [PATCH 05/31] Require TRUE_DIV scale to be == to lhs - rhs scales --- cpp/src/binaryop/binaryop.cpp | 130 +++++------------- cpp/tests/binaryop/binop-integration-test.cpp | 53 +++---- 2 files changed, 49 insertions(+), 134 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index fc697267ca7..8fb4b5ac26e 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -393,9 +393,9 @@ bool is_supported_fixed_point_binop(binary_operator op) int32_t compute_scale_for_binop(binary_operator op, int32_t left_scale, int32_t right_scale) { CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation."); - if (op == binary_operator::TRUE_DIV) CUDF_FAIL("TRUE_DIV scale cannot be computed."); if (op == binary_operator::MUL) return left_scale + right_scale; - if (op == binary_operator::DIV) return left_scale - right_scale; + if (op == binary_operator::DIV || op == binary_operator::TRUE_DIV) + return left_scale - right_scale; return std::min(left_scale, right_scale); } @@ -408,7 +408,8 @@ int32_t compute_scale_for_binop(binary_operator op, int32_t left_scale, int32_t */ bool is_same_scale_necessary(binary_operator op) { - return op != binary_operator::MUL && op != binary_operator::DIV; + return op != binary_operator::MUL && op != binary_operator::DIV && + op != binary_operator::TRUE_DIV; } /** @@ -424,7 +425,7 @@ bool is_same_scale_necessary(binary_operator op) std::unique_ptr fixed_point_binary_operation(scalar const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + cudf::data_type output_type, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { @@ -434,15 +435,14 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); - if (op == binary_operator::TRUE_DIV) - CUDF_EXPECTS(output_type.has_value(), "TRUE_DIV requires result_type."); - auto const scale = op == binary_operator::TRUE_DIV - ? output_type.value().scale() - : compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + if (op == binary_operator::TRUE_DIV) + CUDF_EXPECTS(scale == output_type.scale(), + "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - auto const out_type = output_type.value_or( - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}); + auto const out_type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -450,26 +450,7 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, auto out_view = out->mutable_view(); - if (op == binary_operator::TRUE_DIV) { - // Adjust scalar so lhs has the scale needed to get desired output data_type (scale) - auto const diff = (lhs.type().scale() - rhs.type().scale()) - scale; - auto const unused_scale = scale_type{0}; // scale of out_view already set - if (lhs.type().id() == type_id::DECIMAL32) { - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const val = static_cast const&>(lhs).value(); - auto const scaled_value = diff < 0 ? val / factor : val * factor; - auto const scalar = make_fixed_point_scalar(scaled_value, unused_scale); - binops::jit::binary_operation(out_view, *scalar, rhs, binary_operator::DIV, stream); - return out; - } else { - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const val = static_cast const&>(lhs).value(); - auto const scaled_value = diff < 0 ? val / factor : val * factor; - auto const scalar = make_fixed_point_scalar(scaled_value, unused_scale); - binops::jit::binary_operation(out_view, *scalar, rhs, binary_operator::DIV, stream); - return out; - } - } else if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale if (rhs.type().scale() < lhs.type().scale()) { auto const diff = lhs.type().scale() - rhs.type().scale(); @@ -527,7 +508,7 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, std::unique_ptr fixed_point_binary_operation(column_view const& lhs, scalar const& rhs, binary_operator op, - thrust::optional output_type, + cudf::data_type output_type, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { @@ -537,15 +518,14 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); - if (op == binary_operator::TRUE_DIV) - CUDF_EXPECTS(output_type.has_value(), "TRUE_DIV requires result_type."); - auto const scale = op == binary_operator::TRUE_DIV - ? output_type.value().scale() - : compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + if (op == binary_operator::TRUE_DIV) + CUDF_EXPECTS(scale == output_type.scale(), + "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - auto const out_type = output_type.value_or( - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}); + auto const out_type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -553,26 +533,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, auto out_view = out->mutable_view(); - if (op == binary_operator::TRUE_DIV) { - // Adjust columns so lhs has the scale needed to get desired output data_type (scale) - auto const diff = (lhs.type().scale() - rhs.type().scale()) - scale; - auto const interim_op = diff < 0 ? binary_operator::DIV : binary_operator::MUL; - auto const scalar_scale = scale_type{rhs.type().scale() + scale}; - auto const result = [&] { - if (lhs.type().id() == type_id::DECIMAL32) { - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const scalar = make_fixed_point_scalar(factor, scalar_scale); - return binary_operation(lhs, *scalar, interim_op, {}, stream, mr); - } else { - CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const scalar = make_fixed_point_scalar(factor, scalar_scale); - return binary_operation(lhs, *scalar, interim_op, {}, stream, mr); - } - }(); - binops::jit::binary_operation(out_view, result->view(), rhs, binary_operator::DIV, stream); - return out; - } else if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale if (rhs.type().scale() > lhs.type().scale()) { auto const diff = rhs.type().scale() - lhs.type().scale(); @@ -630,7 +591,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, std::unique_ptr fixed_point_binary_operation(column_view const& lhs, column_view const& rhs, binary_operator op, - thrust::optional output_type, + cudf::data_type output_type, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { @@ -640,15 +601,14 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); - if (op == binary_operator::TRUE_DIV) - CUDF_EXPECTS(output_type.has_value(), "TRUE_DIV requires result_type."); - auto const scale = op == binary_operator::TRUE_DIV - ? output_type.value().scale() - : compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); + if (op == binary_operator::TRUE_DIV) + CUDF_EXPECTS(scale == output_type.scale(), + "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - auto const out_type = output_type.value_or( - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}); + auto const out_type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -656,26 +616,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, auto out_view = out->mutable_view(); - if (op == binary_operator::TRUE_DIV) { - // Adjust columns so lhs has the scale needed to get desired output data_type (scale) - auto const diff = (lhs.type().scale() - rhs.type().scale()) - scale; - auto const interim_op = diff < 0 ? binary_operator::DIV : binary_operator::MUL; - auto const scalar_scale = scale_type{rhs.type().scale() + scale}; - auto const result = [&] { - if (lhs.type().id() == type_id::DECIMAL32) { - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const scalar = make_fixed_point_scalar(factor, scalar_scale); - return binary_operation(lhs, *scalar, interim_op, {}, stream, mr); - } else { - CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); - auto const factor = numeric::detail::ipow(std::abs(diff)); - auto const scalar = make_fixed_point_scalar(factor, scalar_scale); - return binary_operation(lhs, *scalar, interim_op, {}, stream, mr); - } - }(); - binops::jit::binary_operation(out_view, result->view(), rhs, binary_operator::DIV, stream); - return out; - } else if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { + if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust columns so they have they same scale TODO modify comment if (rhs.type().scale() < lhs.type().scale()) { auto const diff = lhs.type().scale() - rhs.type().scale(); @@ -731,9 +672,7 @@ std::unique_ptr binary_operation(scalar const& lhs, return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { - auto const type = - op == binary_operator::TRUE_DIV ? output_type : thrust::optional{thrust::nullopt}; - return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr); + return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); } // Check for datatype @@ -760,11 +699,8 @@ std::unique_ptr binary_operation(column_view const& lhs, if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); - if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { - auto const type = - op == binary_operator::TRUE_DIV ? output_type : thrust::optional{thrust::nullopt}; - return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr); - } + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) + return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); @@ -793,9 +729,7 @@ std::unique_ptr binary_operation(column_view const& lhs, return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { - auto const type = - op == binary_operator::TRUE_DIV ? output_type : thrust::optional{thrust::nullopt}; - return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr); + return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); } // Check for datatype diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 2d17853a72b..0d7a7a00e54 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2327,7 +2327,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{-1}}; + auto const lhs = fp_wrapper{{100, 300, 500, 700}, scale_type{-2}}; auto const rhs = fp_wrapper{{4, 4, 4, 4}, scale_type{0}}; auto const expected = fp_wrapper{{25, 75, 125, 175}, scale_type{-2}}; @@ -2343,7 +2343,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv2) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{1}}; + auto const lhs = fp_wrapper{{100000, 300000, 500000, 700000}, scale_type{-3}}; auto const rhs = fp_wrapper{{20, 20, 20, 20}, scale_type{-1}}; auto const expected = fp_wrapper{{5000, 15000, 25000, 35000}, scale_type{-2}}; @@ -2359,8 +2359,8 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{1}}; - auto const rhs = fp_wrapper{{300, 900, 300, 300}, scale_type{-2}}; + auto const lhs = fp_wrapper{{10000, 30000, 50000, 70000}, scale_type{-2}}; + auto const rhs = fp_wrapper{{3, 9, 3, 3}, scale_type{0}}; auto const expected = fp_wrapper{{3333, 3333, 16666, 23333}, scale_type{-2}}; auto const type = data_type{type_to_id(), -2}; @@ -2377,25 +2377,9 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv4) auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{1}}; auto const rhs = make_fixed_point_scalar(3, scale_type{0}); - auto const expected = fp_wrapper{{3333, 10000, 16666, 23333}, scale_type{-2}}; - - auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); + auto const expected = fp_wrapper{{3, 10, 16, 23}, scale_type{1}}; - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); -} - -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv5) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = device_storage_type_t; - - auto const lhs = fp_wrapper{{10, 30, 50, 70}, scale_type{1}}; - auto const rhs = make_fixed_point_scalar(30000, scale_type{-4}); - auto const expected = fp_wrapper{{3333, 10000, 16666, 23333}, scale_type{-2}}; - - auto const type = data_type{type_to_id(), -2}; + auto const type = data_type{type_to_id(), 1}; auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); @@ -2407,18 +2391,15 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6) using decimalXX = TypeParam; using RepType = device_storage_type_t; - for (auto const i : {0, 1, 2, 3, 4, 5, 6, 7}) { - auto const val = 3 * numeric::detail::ipow(i); - auto const lhs = make_fixed_point_scalar(val, scale_type{-i}); - auto const rhs = fp_wrapper{{10, 30, 50, 70}, scale_type{-1}}; + auto const lhs = make_fixed_point_scalar(3000, scale_type{-3}); + auto const rhs = fp_wrapper{{10, 30, 50, 70}, scale_type{-1}}; - auto const expected = fp_wrapper{{300, 100, 60, 42}, scale_type{-2}}; + auto const expected = fp_wrapper{{300, 100, 60, 42}, scale_type{-2}}; - auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const type = data_type{type_to_id(), -2}; + auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); - } + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7) @@ -2427,7 +2408,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = make_fixed_point_scalar(12000, scale_type{-1}); + auto const lhs = make_fixed_point_scalar(1200, scale_type{0}); auto const rhs = fp_wrapper{{100, 200, 300, 500, 600, 800, 1200, 1300}, scale_type{-2}}; auto const expected = fp_wrapper{{12, 6, 4, 2, 2, 1, 1, 0}, scale_type{2}}; @@ -2445,7 +2426,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8) using RepType = device_storage_type_t; auto const lhs = fp_wrapper{{4000, 6000, 80000}, scale_type{-1}}; - auto const rhs = make_fixed_point_scalar(500, scale_type{-2}); + auto const rhs = make_fixed_point_scalar(5000, scale_type{-3}); auto const expected = fp_wrapper{{0, 1, 16}, scale_type{2}}; auto const type = data_type{type_to_id(), 2}; @@ -2460,7 +2441,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv9) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{100000, 200000, 300000}, scale_type{-2}}; + auto const lhs = fp_wrapper{{10, 20, 30}, scale_type{2}}; auto const rhs = make_fixed_point_scalar(7, scale_type{1}); auto const expected = fp_wrapper{{1, 2, 4}, scale_type{1}}; @@ -2476,7 +2457,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv10) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{100000, 200000, 300000}, scale_type{-2}}; + auto const lhs = fp_wrapper{{100, 200, 300}, scale_type{1}}; auto const rhs = make_fixed_point_scalar(7, scale_type{0}); auto const expected = fp_wrapper{{14, 28, 42}, scale_type{1}}; @@ -2492,7 +2473,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv11) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const lhs = fp_wrapper{{1000000, 2000000, 3000000}, scale_type{-2}}; + auto const lhs = fp_wrapper{{1000, 2000, 3000}, scale_type{1}}; auto const rhs = fp_wrapper{{7, 7, 7}, scale_type{0}}; auto const expected = fp_wrapper{{142, 285, 428}, scale_type{1}}; From 2476eab4ff54ffc7d39f46fb88ea24ee91e69547 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 22:02:30 -0500 Subject: [PATCH 06/31] Use output_type --- cpp/include/cudf/binaryop.hpp | 24 ++++ cpp/src/binaryop/binaryop.cpp | 111 +++++++++-------- cpp/tests/binaryop/binop-integration-test.cpp | 114 +++++++++++++++--- 3 files changed, 184 insertions(+), 65 deletions(-) diff --git a/cpp/include/cudf/binaryop.hpp b/cpp/include/cudf/binaryop.hpp index 72abefef04f..7099c29b9df 100644 --- a/cpp/include/cudf/binaryop.hpp +++ b/cpp/include/cudf/binaryop.hpp @@ -178,5 +178,29 @@ std::unique_ptr binary_operation( data_type output_type, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Computes the `scale` for a `fixed_point` number based on given binary operator `op` + * + * @param op The binary_operator used for two `fixed_point` numbers + * @param left_scale Scale of left `fixed_point` number + * @param right_scale Scale of right `fixed_point` number + * @return The resulting `scale` of the computed `fixed_point` number + */ +int32_t binary_operation_fixed_point_scale(binary_operator op, + int32_t left_scale, + int32_t right_scale); + +/** + * @brief Computes the `data_type` for a `fixed_point` number based on given binary operator `op` + * + * @param op The binary_operator used for two `fixed_point` numbers + * @param lhs `cudf::data_type` of left `fixed_point` number + * @param rhs `cudf::data_type` of right `fixed_point` number + * @return The resulting `cudf::data_type` of the computed `fixed_point` number + */ +cudf::data_type binary_operation_fixed_point_output_type(binary_operator op, + cudf::data_type const& lhs, + cudf::data_type const& rhs); + /** @} */ // end of group } // namespace cudf diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 8fb4b5ac26e..105f2e80609 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -382,23 +382,6 @@ bool is_supported_fixed_point_binop(binary_operator op) op == binary_operator::TRUE_DIV; } -/** - * @brief Computes the scale for a `fixed_point` number based on given binary operator `op` - * - * @param op The binary_operator used for two `fixed_point` numbers - * @param left_scale Scale of left `fixed_point` number - * @param right_scale Scale of right `fixed_point` number - * @return int32_t The resulting `scale` of the computed `fixed_point` number - */ -int32_t compute_scale_for_binop(binary_operator op, int32_t left_scale, int32_t right_scale) -{ - CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation."); - if (op == binary_operator::MUL) return left_scale + right_scale; - if (op == binary_operator::DIV || op == binary_operator::TRUE_DIV) - return left_scale - right_scale; - return std::min(left_scale, right_scale); -} - /** * @brief Helper predicate function that identifies if `op` requires scales to be the same * @@ -412,6 +395,15 @@ bool is_same_scale_necessary(binary_operator op) op != binary_operator::TRUE_DIV; } +template +void fixed_point_binary_operation_validation(binary_operator op, Lhs lhs, Rhs rhs) +{ + CUDF_EXPECTS(is_fixed_point(lhs), "Input must have fixed_point data_type."); + CUDF_EXPECTS(is_fixed_point(rhs), "Input must have fixed_point data_type."); + CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); + CUDF_EXPECTS(lhs.id() == rhs.id(), "Data type mismatch"); +} + /** * @brief Function to compute binary operation of one `column_view` and one `scalar` * @@ -431,20 +423,19 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, { using namespace numeric; - CUDF_EXPECTS(is_fixed_point(lhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); - CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); - if (op == binary_operator::TRUE_DIV) + if (op == binary_operator::TRUE_DIV) { + auto const scale = + binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); CUDF_EXPECTS(scale == output_type.scale(), "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); + } - auto const out_type = - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; - - auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + auto out = [&] { + auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; + return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + }(); if (rhs.is_empty()) return out; @@ -514,20 +505,19 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, { using namespace numeric; - CUDF_EXPECTS(is_fixed_point(lhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); - CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); - if (op == binary_operator::TRUE_DIV) + if (op == binary_operator::TRUE_DIV) { + auto const scale = + binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); CUDF_EXPECTS(scale == output_type.scale(), "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); + } - auto const out_type = - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; - - auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + auto out = [&] { + auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; + return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + }(); if (lhs.is_empty()) return out; @@ -597,20 +587,19 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, { using namespace numeric; - CUDF_EXPECTS(is_fixed_point(lhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_fixed_point(rhs.type()), "Input must have fixed_point data_type."); - CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); - CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch"); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale()); - if (op == binary_operator::TRUE_DIV) + if (op == binary_operator::TRUE_DIV) { + auto const scale = + binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); CUDF_EXPECTS(scale == output_type.scale(), "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); + } - auto const out_type = - is_comparison_binop(op) ? data_type{type_id::BOOL8} : data_type{lhs.type().id(), scale}; - - auto out = make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + auto out = [&] { + auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; + return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); + }(); if (lhs.is_empty() or rhs.is_empty()) return out; @@ -671,9 +660,8 @@ std::unique_ptr binary_operation(scalar const& lhs, if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); - if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); - } // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); @@ -728,9 +716,8 @@ std::unique_ptr binary_operation(column_view const& lhs, if (lhs.type().id() == type_id::STRING and rhs.type().id() == type_id::STRING) return binops::compiled::binary_operation(lhs, rhs, op, output_type, stream, mr); - if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) { + if (is_fixed_point(lhs.type()) or is_fixed_point(rhs.type())) return fixed_point_binary_operation(lhs, rhs, op, output_type, stream, mr); - } // Check for datatype CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype"); @@ -779,6 +766,28 @@ std::unique_ptr binary_operation(column_view const& lhs, } // namespace detail +int32_t binary_operation_fixed_point_scale(binary_operator op, + int32_t left_scale, + int32_t right_scale) +{ + CUDF_EXPECTS(cudf::detail::is_supported_fixed_point_binop(op), + "Unsupported fixed_point binary operation."); + if (op == binary_operator::MUL) return left_scale + right_scale; + if (op == binary_operator::DIV || op == binary_operator::TRUE_DIV) + return left_scale - right_scale; + return std::min(left_scale, right_scale); +} + +cudf::data_type binary_operation_fixed_point_output_type(binary_operator op, + cudf::data_type const& lhs, + cudf::data_type const& rhs) +{ + cudf::detail::fixed_point_binary_operation_validation(op, lhs, rhs); + + auto const scale = binary_operation_fixed_point_scale(op, lhs.scale(), rhs.scale()); + return cudf::data_type{lhs.id(), scale}; +} + std::unique_ptr binary_operation(scalar const& lhs, column_view const& rhs, binary_operator op, diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 0d7a7a00e54..a98b468dc82 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2043,7 +2043,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd) auto const rhs = wrapper(vec2.begin(), vec2.end()); auto const expected_col = wrapper(expected.begin(), expected.end()); - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); } @@ -2072,7 +2076,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply) auto const rhs = wrapper(vec2.begin(), vec2.end()); auto const expected_col = wrapper(expected.begin(), expected.end()); - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_col, result->view()); } @@ -2090,7 +2098,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiply2) auto const rhs = fp_wrapper{{10, 10, 10, 10, 10}, scale_type{0}}; auto const expected = fp_wrapper{{110, 220, 330, 440, 550}, scale_type{-1}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::MUL, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::MUL, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2105,7 +2117,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv) auto const rhs = fp_wrapper{{4, 4, 4, 4}, scale_type{0}}; auto const expected = fp_wrapper{{2, 7, 12, 17}, scale_type{-1}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::DIV, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2120,7 +2136,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv2) auto const rhs = fp_wrapper{{4, 4, 4, 4}, scale_type{-2}}; auto const expected = fp_wrapper{{2, 7, 12, 17}, scale_type{1}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::DIV, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2135,7 +2155,9 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv3) auto const rhs = make_fixed_point_scalar(12, scale_type{-1}); auto const expected = fp_wrapper{{0, 2, 4, 5}, scale_type{0}}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, {}); + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::DIV, static_cast(lhs).type(), rhs->type()); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2153,7 +2175,9 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpDiv4) auto const rhs = make_fixed_point_scalar(12, scale_type{-1}); auto const expected = fp_wrapper(result_begin, result_begin + 1000, scale_type{0}); - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, {}); + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::DIV, static_cast(lhs).type(), rhs->type()); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2168,7 +2192,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd2) auto const rhs = fp_wrapper{{100, 200, 300, 400, 500}, scale_type{-2}}; auto const expected = fp_wrapper{{210, 420, 630, 840, 1050}, scale_type{-2}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2183,7 +2211,45 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd3) auto const rhs = fp_wrapper{{100, 200, 300, 400, 500}, scale_type{-2}}; auto const expected = fp_wrapper{{2100, 4200, 6300, 8400, 10500}, scale_type{-3}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, type); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd4) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const lhs = fp_wrapper{{11, 22, 33, 44, 55}, scale_type{-1}}; + auto const rhs = make_fixed_point_scalar(100, scale_type{-2}); + auto const expected = fp_wrapper{{210, 320, 430, 540, 650}, scale_type{-2}}; + + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::ADD, static_cast(lhs).type(), rhs->type()); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::ADD, type); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd5) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const lhs = make_fixed_point_scalar(100, scale_type{-2}); + auto const rhs = fp_wrapper{{11, 22, 33, 44, 55}, scale_type{-1}}; + auto const expected = fp_wrapper{{210, 320, 430, 540, 650}, scale_type{-2}}; + + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::ADD, lhs->type(), static_cast(rhs).type()); + auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::ADD, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2198,7 +2264,9 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiplyScalar) auto const rhs = make_fixed_point_scalar(100, scale_type{-1}); auto const expected = fp_wrapper{{1100, 2200, 3300, 4400, 5500}, scale_type{-2}}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::MUL, {}); + auto const type = cudf::binary_operation_fixed_point_output_type( + cudf::binary_operator::MUL, static_cast(lhs).type(), rhs->type()); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::MUL, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2213,7 +2281,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpSimplePlus) auto const rhs = fp_wrapper{{2250, 1005}, scale_type{-3}}; auto const expected = fp_wrapper{{3750, 3005}, scale_type{-3}}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD, + static_cast(lhs).type(), + static_cast(rhs).type()); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::ADD, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2229,6 +2301,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) auto const col2 = fp_wrapper{{100, 200, 300, 400}, scale_type{-2}}; auto const expected = wrapper(trues.begin(), trues.end()); + // TODO auto const result = cudf::binary_operation(col1, col2, binary_operator::EQUAL, {}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); @@ -2251,7 +2324,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) auto const iota_3 = fp_wrapper(vec1.begin(), vec1.end(), scale_type{-3}); auto const zeros_3 = fp_wrapper(vec2.begin(), vec2.end(), scale_type{-1}); - auto const iota_3_after_add = cudf::binary_operation(zeros_3, iota_3, binary_operator::ADD, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::ADD, + static_cast(iota_3).type(), + static_cast(zeros_3).type()); + auto const iota_3_after_add = cudf::binary_operation(zeros_3, iota_3, binary_operator::ADD, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(iota_3, iota_3_after_add->view()); @@ -2260,6 +2337,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) auto const trues = std::vector(sz, true); auto const true_col = wrapper(trues.begin(), trues.end()); + // TODO auto const equal_result = cudf::binary_operation(iota_3, iota_3_after_add->view(), binary_operator::EQUAL, {}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, equal_result->view()); @@ -2284,7 +2362,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpNullMaxSimple) auto const col2 = fp_wrapper{{10, 20, 30, 40, 0}, {1, 1, 1, 0, 0}, scale_type{-2}}; auto const expected = fp_wrapper{{40, 20, 30, 10, 0}, {1, 1, 1, 1, 0}, scale_type{-2}}; - auto const result = cudf::binary_operation(col1, col2, binary_operator::NULL_MAX, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::NULL_MAX, + static_cast(col1).type(), + static_cast(col2).type()); + auto const result = cudf::binary_operation(col1, col2, binary_operator::NULL_MAX, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2300,7 +2382,11 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpNullMinSimple) auto const col2 = fp_wrapper{{10, 20, 30, 40, 0}, {1, 0, 1, 1, 0}, scale_type{-1}}; auto const expected = fp_wrapper{{10, 30, 20, 40, 0}, {1, 1, 1, 1, 0}, scale_type{-1}}; - auto const result = cudf::binary_operation(col1, col2, binary_operator::NULL_MIN, {}); + auto const type = + cudf::binary_operation_fixed_point_output_type(cudf::binary_operator::NULL_MIN, + static_cast(col1).type(), + static_cast(col2).type()); + auto const result = cudf::binary_operation(col1, col2, binary_operator::NULL_MIN, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } From ba38c802bb35c3043995169ca63fd94cb99aa780 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 22:26:44 -0500 Subject: [PATCH 07/31] Remove TRUE_DIV --- cpp/src/binaryop/binaryop.cpp | 30 ++------------ cpp/tests/binaryop/binop-integration-test.cpp | 40 +++++++++---------- 2 files changed, 23 insertions(+), 47 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 105f2e80609..1857abb35ac 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -378,8 +378,7 @@ bool is_comparison_binop(binary_operator op) */ bool is_supported_fixed_point_binop(binary_operator op) { - return is_basic_arithmetic_binop(op) or is_comparison_binop(op) or - op == binary_operator::TRUE_DIV; + return is_basic_arithmetic_binop(op) or is_comparison_binop(op); } /** @@ -391,8 +390,7 @@ bool is_supported_fixed_point_binop(binary_operator op) */ bool is_same_scale_necessary(binary_operator op) { - return op != binary_operator::MUL && op != binary_operator::DIV && - op != binary_operator::TRUE_DIV; + return op != binary_operator::MUL && op != binary_operator::DIV; } template @@ -425,13 +423,6 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - if (op == binary_operator::TRUE_DIV) { - auto const scale = - binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - CUDF_EXPECTS(scale == output_type.scale(), - "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - } - auto out = [&] { auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -507,13 +498,6 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - if (op == binary_operator::TRUE_DIV) { - auto const scale = - binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - CUDF_EXPECTS(scale == output_type.scale(), - "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - } - auto out = [&] { auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -589,13 +573,6 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - if (op == binary_operator::TRUE_DIV) { - auto const scale = - binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - CUDF_EXPECTS(scale == output_type.scale(), - "TRUE_DIV requires output_type.scale() to match resulting lhs scale - rhs scale"); - } - auto out = [&] { auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); @@ -773,8 +750,7 @@ int32_t binary_operation_fixed_point_scale(binary_operator op, CUDF_EXPECTS(cudf::detail::is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation."); if (op == binary_operator::MUL) return left_scale + right_scale; - if (op == binary_operator::DIV || op == binary_operator::TRUE_DIV) - return left_scale - right_scale; + if (op == binary_operator::DIV) return left_scale - right_scale; return std::min(left_scale, right_scale); } diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index a98b468dc82..5a9502fd9b2 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2407,7 +2407,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpNullEqualsSimple) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div) { using namespace numeric; using decimalXX = TypeParam; @@ -2418,12 +2418,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv) auto const expected = fp_wrapper{{25, 75, 125, 175}, scale_type{-2}}; auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv2) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div2) { using namespace numeric; using decimalXX = TypeParam; @@ -2434,12 +2434,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv2) auto const expected = fp_wrapper{{5000, 15000, 25000, 35000}, scale_type{-2}}; auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div3) { using namespace numeric; using decimalXX = TypeParam; @@ -2450,12 +2450,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3) auto const expected = fp_wrapper{{3333, 3333, 16666, 23333}, scale_type{-2}}; auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv4) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div4) { using namespace numeric; using decimalXX = TypeParam; @@ -2466,12 +2466,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv4) auto const expected = fp_wrapper{{3, 10, 16, 23}, scale_type{1}}; auto const type = data_type{type_to_id(), 1}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div6) { using namespace numeric; using decimalXX = TypeParam; @@ -2483,12 +2483,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6) auto const expected = fp_wrapper{{300, 100, 60, 42}, scale_type{-2}}; auto const type = data_type{type_to_id(), -2}; - auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div7) { using namespace numeric; using decimalXX = TypeParam; @@ -2500,12 +2500,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7) auto const expected = fp_wrapper{{12, 6, 4, 2, 2, 1, 1, 0}, scale_type{2}}; auto const type = data_type{type_to_id(), 2}; - auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div8) { using namespace numeric; using decimalXX = TypeParam; @@ -2516,12 +2516,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8) auto const expected = fp_wrapper{{0, 1, 16}, scale_type{2}}; auto const type = data_type{type_to_id(), 2}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv9) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div9) { using namespace numeric; using decimalXX = TypeParam; @@ -2532,12 +2532,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv9) auto const expected = fp_wrapper{{1, 2, 4}, scale_type{1}}; auto const type = data_type{type_to_id(), 1}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv10) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div10) { using namespace numeric; using decimalXX = TypeParam; @@ -2548,12 +2548,12 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv10) auto const expected = fp_wrapper{{14, 28, 42}, scale_type{1}}; auto const type = data_type{type_to_id(), 1}; - auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } -TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv11) +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div11) { using namespace numeric; using decimalXX = TypeParam; @@ -2564,7 +2564,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv11) auto const expected = fp_wrapper{{142, 285, 428}, scale_type{1}}; auto const type = data_type{type_to_id(), 1}; - auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type); + auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } From f1a62b334a33dce7c9587b00e561c7869a196c59 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 25 Feb 2021 22:39:27 -0500 Subject: [PATCH 08/31] Don't hardcode BOOL --- cpp/src/binaryop/binaryop.cpp | 24 +++---------------- cpp/tests/binaryop/binop-integration-test.cpp | 15 ++++++------ 2 files changed, 11 insertions(+), 28 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 1857abb35ac..6c8fdf37e37 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -422,14 +422,8 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - - auto out = [&] { - auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; - return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); - }(); - + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (rhs.is_empty()) return out; - auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { @@ -497,14 +491,8 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - - auto out = [&] { - auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; - return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); - }(); - + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty()) return out; - auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { @@ -572,14 +560,8 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); - - auto out = [&] { - auto const out_type = is_comparison_binop(op) ? data_type{type_id::BOOL8} : output_type; - return make_fixed_width_column_for_output(lhs, rhs, op, out_type, stream, mr); - }(); - + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty() or rhs.is_empty()) return out; - auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 5a9502fd9b2..4e31c90f94c 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2301,8 +2301,8 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) auto const col2 = fp_wrapper{{100, 200, 300, 400}, scale_type{-2}}; auto const expected = wrapper(trues.begin(), trues.end()); - // TODO - auto const result = cudf::binary_operation(col1, col2, binary_operator::EQUAL, {}); + auto const result = + cudf::binary_operation(col1, col2, binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } @@ -2337,17 +2337,17 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) auto const trues = std::vector(sz, true); auto const true_col = wrapper(trues.begin(), trues.end()); - // TODO + auto const btype = cudf::data_type{type_id::BOOL8}; auto const equal_result = - cudf::binary_operation(iota_3, iota_3_after_add->view(), binary_operator::EQUAL, {}); + cudf::binary_operation(iota_3, iota_3_after_add->view(), binary_operator::EQUAL, btype); CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, equal_result->view()); auto const less_result = - cudf::binary_operation(zeros_3, iota_3_after_add->view(), binary_operator::LESS, {}); + cudf::binary_operation(zeros_3, iota_3_after_add->view(), binary_operator::LESS, btype); CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, less_result->view()); auto const greater_result = - cudf::binary_operation(iota_3_after_add->view(), zeros_3, binary_operator::GREATER, {}); + cudf::binary_operation(iota_3_after_add->view(), zeros_3, binary_operator::GREATER, btype); CUDF_TEST_EXPECT_COLUMNS_EQUAL(true_col, greater_result->view()); } @@ -2402,7 +2402,8 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpNullEqualsSimple) auto const col2 = fp_wrapper{{40, 200, 20, 400}, {1, 0, 1, 0}, scale_type{-1}}; auto const expected = wrapper{{1, 0, 0, 1}, {1, 1, 1, 1}}; - auto const result = cudf::binary_operation(col1, col2, binary_operator::NULL_EQUALS, {}); + auto const result = cudf::binary_operation( + col1, col2, binary_operator::NULL_EQUALS, cudf::data_type{type_id::BOOL8}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } From c0769f11550db723a7be795e6bafde71a390946e Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 1 Mar 2021 09:43:38 -0500 Subject: [PATCH 09/31] Add unit tests + CUDF_EXPECTS --- cpp/src/binaryop/binaryop.cpp | 19 ++++++++++++++---- cpp/tests/binaryop/binop-integration-test.cpp | 20 +++++++++++++++++++ 2 files changed, 35 insertions(+), 4 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 6c8fdf37e37..17695ee0a61 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -394,12 +394,23 @@ bool is_same_scale_necessary(binary_operator op) } template -void fixed_point_binary_operation_validation(binary_operator op, Lhs lhs, Rhs rhs) +void fixed_point_binary_operation_validation(binary_operator op, + Lhs lhs, + Rhs rhs, + thrust::optional output_type = {}) { CUDF_EXPECTS(is_fixed_point(lhs), "Input must have fixed_point data_type."); CUDF_EXPECTS(is_fixed_point(rhs), "Input must have fixed_point data_type."); CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation"); CUDF_EXPECTS(lhs.id() == rhs.id(), "Data type mismatch"); + if (output_type.has_value()) { + if (is_comparison_binop(op)) + CUDF_EXPECTS(output_type == cudf::data_type{type_id::BOOL8}, + "Comparison operations require boolean output type."); + else + CUDF_EXPECTS(is_fixed_point(output_type.value()), + "fixed_point binary operations require fixed_point output type."); + } } /** @@ -421,7 +432,7 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, { using namespace numeric; - fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (rhs.is_empty()) return out; auto out_view = out->mutable_view(); @@ -490,7 +501,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, { using namespace numeric; - fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty()) return out; auto out_view = out->mutable_view(); @@ -559,7 +570,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, { using namespace numeric; - fixed_point_binary_operation_validation(op, lhs.type(), rhs.type()); + fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty() or rhs.is_empty()) return out; auto out_view = out->mutable_view(); diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 4e31c90f94c..2ee62594cdd 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -30,6 +30,7 @@ #include #include +#include "cudf/utilities/error.hpp" namespace cudf { namespace test { @@ -2570,6 +2571,25 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOp_Div11) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpThrows) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const col = fp_wrapper{{100, 300, 500, 700}, scale_type{-2}}; + // auto const expected = fp_wrapper{{25, 75, 125, 175}, scale_type{-2}}; + + auto const non_bool_type = data_type{type_to_id(), -2}; + auto const float_type = data_type{type_id::FLOAT32}; + EXPECT_THROW(cudf::binary_operation(col, col, cudf::binary_operator::LESS, non_bool_type), + cudf::logic_error); + EXPECT_THROW(cudf::binary_operation(col, col, cudf::binary_operator::MUL, float_type), + cudf::logic_error); + + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + } // namespace binop } // namespace test } // namespace cudf From 65d071f64ab6c30133b6cb905fc04b6996e0eaf9 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 1 Mar 2021 09:48:14 -0500 Subject: [PATCH 10/31] Remove dead code --- cpp/tests/binaryop/binop-integration-test.cpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 2ee62594cdd..67e9eddb6b5 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2577,17 +2577,13 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpThrows) using decimalXX = TypeParam; using RepType = device_storage_type_t; - auto const col = fp_wrapper{{100, 300, 500, 700}, scale_type{-2}}; - // auto const expected = fp_wrapper{{25, 75, 125, 175}, scale_type{-2}}; - + auto const col = fp_wrapper{{100, 300, 500, 700}, scale_type{-2}}; auto const non_bool_type = data_type{type_to_id(), -2}; auto const float_type = data_type{type_id::FLOAT32}; EXPECT_THROW(cudf::binary_operation(col, col, cudf::binary_operator::LESS, non_bool_type), cudf::logic_error); EXPECT_THROW(cudf::binary_operation(col, col, cudf::binary_operator::MUL, float_type), cudf::logic_error); - - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } } // namespace binop From 83070e8f1fe56559cc2bc44b785ecb67a2497d32 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Mar 2021 08:57:39 -0500 Subject: [PATCH 11/31] Fix tests --- cpp/tests/fixed_point/fixed_point_tests.cu | 18 ++++++++++++------ 1 file changed, 12 insertions(+), 6 deletions(-) diff --git a/cpp/tests/fixed_point/fixed_point_tests.cu b/cpp/tests/fixed_point/fixed_point_tests.cu index 5f969098b48..5f74e459bb1 100644 --- a/cpp/tests/fixed_point/fixed_point_tests.cu +++ b/cpp/tests/fixed_point/fixed_point_tests.cu @@ -596,8 +596,9 @@ TEST_F(FixedPointTest, PositiveScaleWithValuesOutsideUnderlyingType32) auto const expected1 = fp_wrapper{{150000000}, scale_type{6}}; auto const expected2 = fp_wrapper{{50000000}, scale_type{6}}; - auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, {}); - auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, {}); + auto const type = cudf::data_type{cudf::type_id::DECIMAL32, 6}; + auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, type); + auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected1, result1->view()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected2, result2->view()); @@ -618,8 +619,9 @@ TEST_F(FixedPointTest, PositiveScaleWithValuesOutsideUnderlyingType64) auto const expected1 = fp_wrapper{{150000000}, scale_type{100}}; auto const expected2 = fp_wrapper{{50000000}, scale_type{100}}; - auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, {}); - auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, {}); + auto const type = cudf::data_type{cudf::type_id::DECIMAL64, 100}; + auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, type); + auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, type); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected1, result1->view()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected2, result2->view()); @@ -630,6 +632,7 @@ TYPED_TEST(FixedPointTestBothReps, ExtremelyLargeNegativeScale) // This is testing fixed_point values with an extremely large negative scale. The fixed_point // implementation should be able to handle any scale representable by an int32_t + using decimalXX = fixed_point; using fp_wrapper = cudf::test::fixed_point_column_wrapper; auto const a = fp_wrapper{{10}, scale_type{-201}}; @@ -639,8 +642,11 @@ TYPED_TEST(FixedPointTestBothReps, ExtremelyLargeNegativeScale) auto const expected1 = fp_wrapper{{150}, scale_type{-202}}; auto const expected2 = fp_wrapper{{5}, scale_type{-201}}; - auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, {}); - auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, {}); + auto const type1 = cudf::data_type{cudf::type_to_id(), -202}; + auto const result1 = cudf::binary_operation(a, b, cudf::binary_operator::ADD, type1); + + auto const type2 = cudf::data_type{cudf::type_to_id(), -201}; + auto const result2 = cudf::binary_operation(a, c, cudf::binary_operator::DIV, type2); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected1, result1->view()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected2, result2->view()); From 4860f2205d5474c34c940df7f2060dd76a921780 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Mar 2021 11:44:29 -0500 Subject: [PATCH 12/31] Fix rescale --- cpp/src/binaryop/binaryop.cpp | 3 +++ cpp/src/unary/cast_ops.cu | 6 ++++-- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 17695ee0a61..250ca959645 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -433,6 +433,7 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (rhs.is_empty()) return out; auto out_view = out->mutable_view(); @@ -502,6 +503,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty()) return out; auto out_view = out->mutable_view(); @@ -571,6 +573,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, using namespace numeric; fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); + auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); if (lhs.is_empty() or rhs.is_empty()) return out; auto out_view = out->mutable_view(); diff --git a/cpp/src/unary/cast_ops.cu b/cpp/src/unary/cast_ops.cu index 15db8e6a3dd..7e3a4050b4f 100644 --- a/cpp/src/unary/cast_ops.cu +++ b/cpp/src/unary/cast_ops.cu @@ -178,11 +178,13 @@ std::unique_ptr rescale(column_view input, if (input.type().scale() > scale) { auto const scalar = make_fixed_point_scalar(0, scale_type{scale}); - return detail::binary_operation(input, *scalar, binary_operator::ADD, {}, stream, mr); + auto const type = cudf::data_type{cudf::type_to_id(), scale}; + return detail::binary_operation(input, *scalar, binary_operator::ADD, type, stream, mr); } else { auto const diff = input.type().scale() - scale; auto const scalar = make_fixed_point_scalar(std::pow(10, -diff), scale_type{diff}); - return detail::binary_operation(input, *scalar, binary_operator::DIV, {}, stream, mr); + auto const type = cudf::data_type{cudf::type_to_id(), scale}; + return detail::binary_operation(input, *scalar, binary_operator::DIV, type, stream, mr); } }; From aa2b5fe0a7e744c2a680ff530b7e60938772fa56 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 3 Mar 2021 23:42:13 -0500 Subject: [PATCH 13/31] Python changes (incomplete) --- python/cudf/cudf/_lib/binaryop.pyx | 11 ++--------- python/cudf/cudf/_lib/column.pyx | 1 + python/cudf/cudf/_lib/types.pxd | 6 +++++- python/cudf/cudf/_lib/types.pyx | 19 ++++++++++++++++++- python/cudf/cudf/core/column/decimal.py | 12 +++++++++++- 5 files changed, 37 insertions(+), 12 deletions(-) diff --git a/python/cudf/cudf/_lib/binaryop.pyx b/python/cudf/cudf/_lib/binaryop.pyx index 5228c34af67..59a6b876961 100644 --- a/python/cudf/cudf/_lib/binaryop.pyx +++ b/python/cudf/cudf/_lib/binaryop.pyx @@ -13,7 +13,7 @@ from cudf._lib.replace import replace_nulls from cudf._lib.scalar import as_device_scalar from cudf._lib.scalar cimport DeviceScalar from cudf._lib.types import np_to_cudf_types -from cudf._lib.types cimport underlying_type_t_type_id +from cudf._lib.types cimport underlying_type_t_type_id, dtype_to_data_type from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.scalar.scalar cimport scalar @@ -174,15 +174,8 @@ def binaryop(lhs, rhs, op, dtype): cdef binary_operator c_op = ( op ) - cdef type_id tid = ( - ( - ( - np_to_cudf_types[np.dtype(dtype)] - ) - ) - ) - cdef data_type c_dtype = data_type(tid) + cdef data_type c_dtype = dtype_to_data_type(dtype) if is_scalar(lhs) or lhs is None: is_string_col = is_string_dtype(rhs.dtype) diff --git a/python/cudf/cudf/_lib/column.pyx b/python/cudf/cudf/_lib/column.pyx index f9b5c859ff2..3f0923b1ea6 100644 --- a/python/cudf/cudf/_lib/column.pyx +++ b/python/cudf/cudf/_lib/column.pyx @@ -384,6 +384,7 @@ cdef class Column: data_dtype = col.dtype cdef libcudf_types.type_id tid + # TODO refactor this? both self.dtype and data_dtype used if is_list_dtype(self.dtype): tid = libcudf_types.type_id.LIST elif is_struct_dtype(self.dtype): diff --git a/python/cudf/cudf/_lib/types.pxd b/python/cudf/cudf/_lib/types.pxd index 9b35ca2e80c..3b74b4182d3 100644 --- a/python/cudf/cudf/_lib/types.pxd +++ b/python/cudf/cudf/_lib/types.pxd @@ -4,7 +4,8 @@ from libc.stdint cimport int32_t from libcpp cimport bool from cudf._lib.cpp.column.column_view cimport column_view from cudf._lib.cpp.lists.lists_column_view cimport lists_column_view - +cimport cudf._lib.cpp.types as libcudf_types +# from cudf._typing import Dtype ctypedef bool underlying_type_t_order ctypedef bool underlying_type_t_null_order @@ -14,3 +15,6 @@ ctypedef int32_t underlying_type_t_type_id ctypedef bool underlying_type_t_null_policy cdef dtype_from_column_view(column_view cv) + +# TODO not sure if i need the import above ^ +cdef libcudf_types.data_type dtype_to_data_type(dtype) diff --git a/python/cudf/cudf/_lib/types.pyx b/python/cudf/cudf/_lib/types.pyx index 370d083d7ac..25710ffd1d7 100644 --- a/python/cudf/cudf/_lib/types.pyx +++ b/python/cudf/cudf/_lib/types.pyx @@ -15,10 +15,10 @@ from cudf._lib.types cimport ( from cudf._lib.cpp.column.column_view cimport column_view from cudf._lib.cpp.lists.lists_column_view cimport lists_column_view from cudf.core.dtypes import ListDtype, StructDtype, Decimal64Dtype +from cudf.utils.dtypes import is_decimal_dtype, is_list_dtype, is_struct_dtype cimport cudf._lib.cpp.types as libcudf_types - class TypeId(IntEnum): EMPTY = libcudf_types.type_id.EMPTY INT8 = libcudf_types.type_id.INT8 @@ -208,3 +208,20 @@ cdef dtype_from_column_view(column_view cv): "Use decimal64 instead") else: return cudf_to_np_types[(tid)] + +cdef libcudf_types.data_type dtype_to_data_type(dtype): + if is_list_dtype(dtype): + tid = libcudf_types.type_id.LIST + elif is_struct_dtype(dtype): + tid = libcudf_types.type_id.STRUCT + elif is_decimal_dtype(dtype): + tid = libcudf_types.type_id.DECIMAL64 + else: + tid = ( + ( + np_to_cudf_types[np.dtype(dtype)])) + + if tid == libcudf_types.type_id.DECIMAL64: + return libcudf_types.data_type(tid, -dtype.scale) + else: + return libcudf_types.data_type(tid) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index 58156c3826c..fce5af8ef8e 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -55,10 +55,20 @@ def to_arrow(self): def binary_operator(self, op, other, reflect=False): if reflect: self, other = other, self - result = libcudf.binaryop.binaryop(self, other, op, "int32") + scale = _binop_scale(self.dtype, other.dtype, op) + output_type = Decimal64Dtype(scale=scale, precision=18) # precision will be ignored + result = libcudf.binaryop.binaryop(self, other, op, output_type) result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) return result +def _binop_scale(l_dtype, r_dtype, op): + s1, s2 = l_dtype.scale, r_dtype.scale + if op in ("add", "sub"): + return max(s1, s2) + elif op == "mul": + return s1 + s2 + else: + raise NotImplementedError() def _binop_precision(l_dtype, r_dtype, op): """ From e76979b52b5572820b06771a255d79a8d16536d9 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 10:41:54 -0500 Subject: [PATCH 14/31] Python fix --- python/cudf/cudf/utils/dtypes.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/utils/dtypes.py b/python/cudf/cudf/utils/dtypes.py index 20c86b2a4b7..19d6b3af6d1 100644 --- a/python/cudf/cudf/utils/dtypes.py +++ b/python/cudf/cudf/utils/dtypes.py @@ -13,7 +13,7 @@ from pandas.core.dtypes.dtypes import CategoricalDtype, CategoricalDtypeType import cudf -from cudf._lib.scalar import DeviceScalar, _is_null_host_scalar +from cudf._lib.scalar import DeviceScalar from cudf.core._compat import PANDAS_GE_120 _NA_REP = "" @@ -331,6 +331,7 @@ def to_cudf_compatible_scalar(val, dtype=None): If `val` is None, returns None. """ + from cudf._lib.scalar import DeviceScalar, _is_null_host_scalar if _is_null_host_scalar(val) or isinstance(val, cudf.Scalar): return val From 0f9a043e7995add61ba9d209d745846f227b4aa7 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 10:51:12 -0500 Subject: [PATCH 15/31] Black python formatting --- python/cudf/cudf/core/column/decimal.py | 6 +++++- python/cudf/cudf/utils/dtypes.py | 1 + 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index fce5af8ef8e..38601d384d1 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -56,11 +56,14 @@ def binary_operator(self, op, other, reflect=False): if reflect: self, other = other, self scale = _binop_scale(self.dtype, other.dtype, op) - output_type = Decimal64Dtype(scale=scale, precision=18) # precision will be ignored + output_type = Decimal64Dtype( + scale=scale, precision=18 + ) # precision will be ignored result = libcudf.binaryop.binaryop(self, other, op, output_type) result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) return result + def _binop_scale(l_dtype, r_dtype, op): s1, s2 = l_dtype.scale, r_dtype.scale if op in ("add", "sub"): @@ -70,6 +73,7 @@ def _binop_scale(l_dtype, r_dtype, op): else: raise NotImplementedError() + def _binop_precision(l_dtype, r_dtype, op): """ Returns the result precision when performing the diff --git a/python/cudf/cudf/utils/dtypes.py b/python/cudf/cudf/utils/dtypes.py index 19d6b3af6d1..0854d96bff5 100644 --- a/python/cudf/cudf/utils/dtypes.py +++ b/python/cudf/cudf/utils/dtypes.py @@ -332,6 +332,7 @@ def to_cudf_compatible_scalar(val, dtype=None): If `val` is None, returns None. """ from cudf._lib.scalar import DeviceScalar, _is_null_host_scalar + if _is_null_host_scalar(val) or isinstance(val, cudf.Scalar): return val From 938874824598d9a5440322b710b1226813447817 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 11:00:16 -0500 Subject: [PATCH 16/31] Flake8 python --- python/cudf/cudf/utils/dtypes.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/cudf/cudf/utils/dtypes.py b/python/cudf/cudf/utils/dtypes.py index 0854d96bff5..2d86d822afa 100644 --- a/python/cudf/cudf/utils/dtypes.py +++ b/python/cudf/cudf/utils/dtypes.py @@ -331,7 +331,7 @@ def to_cudf_compatible_scalar(val, dtype=None): If `val` is None, returns None. """ - from cudf._lib.scalar import DeviceScalar, _is_null_host_scalar + from cudf._lib.scalar import _is_null_host_scalar if _is_null_host_scalar(val) or isinstance(val, cudf.Scalar): return val From 4f85dd4452f4e7b41fc613eadb6aa2cc6168def8 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 11:06:36 -0500 Subject: [PATCH 17/31] flake8 --- python/cudf/cudf/_lib/types.pyx | 1 + 1 file changed, 1 insertion(+) diff --git a/python/cudf/cudf/_lib/types.pyx b/python/cudf/cudf/_lib/types.pyx index 25710ffd1d7..46521b3137c 100644 --- a/python/cudf/cudf/_lib/types.pyx +++ b/python/cudf/cudf/_lib/types.pyx @@ -19,6 +19,7 @@ from cudf.utils.dtypes import is_decimal_dtype, is_list_dtype, is_struct_dtype cimport cudf._lib.cpp.types as libcudf_types + class TypeId(IntEnum): EMPTY = libcudf_types.type_id.EMPTY INT8 = libcudf_types.type_id.INT8 From fada90b82e2b4149d9ba4565782cbb324a8e5a06 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 15:25:46 -0500 Subject: [PATCH 18/31] Python cleanup/refactoring --- python/cudf/cudf/_lib/column.pyx | 26 ++++--------------------- python/cudf/cudf/_lib/types.pxd | 2 -- python/cudf/cudf/core/column/decimal.py | 4 ++-- 3 files changed, 6 insertions(+), 26 deletions(-) diff --git a/python/cudf/cudf/_lib/column.pyx b/python/cudf/cudf/_lib/column.pyx index 3f0923b1ea6..e6e4b479fde 100644 --- a/python/cudf/cudf/_lib/column.pyx +++ b/python/cudf/cudf/_lib/column.pyx @@ -32,7 +32,8 @@ from rmm._lib.device_buffer cimport DeviceBuffer from cudf._lib.types import np_to_cudf_types, cudf_to_np_types from cudf._lib.types cimport ( underlying_type_t_type_id, - dtype_from_column_view + dtype_from_column_view, + dtype_to_data_type ) from cudf._lib.null_mask import bitmask_allocation_size_bytes @@ -381,27 +382,8 @@ cdef class Column: else: col = self - data_dtype = col.dtype - cdef libcudf_types.type_id tid - - # TODO refactor this? both self.dtype and data_dtype used - if is_list_dtype(self.dtype): - tid = libcudf_types.type_id.LIST - elif is_struct_dtype(self.dtype): - tid = libcudf_types.type_id.STRUCT - elif is_decimal_dtype(self.dtype): - tid = libcudf_types.type_id.DECIMAL64 - else: - tid = ( - ( - np_to_cudf_types[np.dtype(data_dtype)] - ) - ) - cdef libcudf_types.data_type dtype = ( - libcudf_types.data_type(tid, -self.dtype.scale) - if tid == libcudf_types.type_id.DECIMAL64 - else libcudf_types.data_type(tid) - ) + data_dtype = self.dtype if not is_categorical_dtype(self.dtype) else col.dtype + cdef libcudf_types.data_type dtype = dtype_to_data_type(data_dtype) cdef libcudf_types.size_type offset = self.offset cdef vector[column_view] children cdef void* data diff --git a/python/cudf/cudf/_lib/types.pxd b/python/cudf/cudf/_lib/types.pxd index 3b74b4182d3..4e2202cbe6e 100644 --- a/python/cudf/cudf/_lib/types.pxd +++ b/python/cudf/cudf/_lib/types.pxd @@ -5,7 +5,6 @@ from libcpp cimport bool from cudf._lib.cpp.column.column_view cimport column_view from cudf._lib.cpp.lists.lists_column_view cimport lists_column_view cimport cudf._lib.cpp.types as libcudf_types -# from cudf._typing import Dtype ctypedef bool underlying_type_t_order ctypedef bool underlying_type_t_null_order @@ -16,5 +15,4 @@ ctypedef bool underlying_type_t_null_policy cdef dtype_from_column_view(column_view cv) -# TODO not sure if i need the import above ^ cdef libcudf_types.data_type dtype_to_data_type(dtype) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index 38601d384d1..c8dc52bdb1c 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -58,12 +58,12 @@ def binary_operator(self, op, other, reflect=False): scale = _binop_scale(self.dtype, other.dtype, op) output_type = Decimal64Dtype( scale=scale, precision=18 - ) # precision will be ignored + ) # precision will be ignored, libcudf has no notion of precision result = libcudf.binaryop.binaryop(self, other, op, output_type) result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) return result - +# This should at some point be hooked up to libcudf's binary_operation_fixed_point_scale def _binop_scale(l_dtype, r_dtype, op): s1, s2 = l_dtype.scale, r_dtype.scale if op in ("add", "sub"): From 7517f7376e6c42ed0ce4f20a5edaa750d5baf2dd Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 15:39:00 -0500 Subject: [PATCH 19/31] flake8 fix --- python/cudf/cudf/_lib/column.pyx | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/_lib/column.pyx b/python/cudf/cudf/_lib/column.pyx index e6e4b479fde..9f6413a2bab 100644 --- a/python/cudf/cudf/_lib/column.pyx +++ b/python/cudf/cudf/_lib/column.pyx @@ -379,10 +379,11 @@ cdef class Column: cdef column_view _view(self, libcudf_types.size_type null_count) except *: if is_categorical_dtype(self.dtype): col = self.base_children[0] + data_dtype = col.dtype else: col = self + data_dtype = self.dtype - data_dtype = self.dtype if not is_categorical_dtype(self.dtype) else col.dtype cdef libcudf_types.data_type dtype = dtype_to_data_type(data_dtype) cdef libcudf_types.size_type offset = self.offset cdef vector[column_view] children From 64c3f407fc9d2a05141e3a97cec30a2e00fe5d02 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 15:46:28 -0500 Subject: [PATCH 20/31] flake8 fix --- python/cudf/cudf/_lib/column.pyx | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/cudf/cudf/_lib/column.pyx b/python/cudf/cudf/_lib/column.pyx index 9f6413a2bab..ffc3fdfd70a 100644 --- a/python/cudf/cudf/_lib/column.pyx +++ b/python/cudf/cudf/_lib/column.pyx @@ -379,10 +379,10 @@ cdef class Column: cdef column_view _view(self, libcudf_types.size_type null_count) except *: if is_categorical_dtype(self.dtype): col = self.base_children[0] - data_dtype = col.dtype + data_dtype = col.dtype else: col = self - data_dtype = self.dtype + data_dtype = self.dtype cdef libcudf_types.data_type dtype = dtype_to_data_type(data_dtype) cdef libcudf_types.size_type offset = self.offset From 3d325f80a8a900530e9457dedc9d4f3ad660311c Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 15:55:40 -0500 Subject: [PATCH 21/31] flake8 fix --- python/cudf/cudf/core/column/decimal.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index c8dc52bdb1c..661ef907715 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -63,8 +63,9 @@ def binary_operator(self, op, other, reflect=False): result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) return result -# This should at some point be hooked up to libcudf's binary_operation_fixed_point_scale def _binop_scale(l_dtype, r_dtype, op): + # This should at some point be hooked up to libcudf's + # binary_operation_fixed_point_scale s1, s2 = l_dtype.scale, r_dtype.scale if op in ("add", "sub"): return max(s1, s2) From 5063738de7040be03ea623b1bf3473d1e9e68a20 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 16:03:43 -0500 Subject: [PATCH 22/31] flake8 fix --- python/cudf/cudf/core/column/decimal.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index 661ef907715..0fdd592bc26 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -63,6 +63,7 @@ def binary_operator(self, op, other, reflect=False): result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) return result + def _binop_scale(l_dtype, r_dtype, op): # This should at some point be hooked up to libcudf's # binary_operation_fixed_point_scale From 055bf1af979b414cbc58a6c5e51547e5698a5bfc Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 17:25:02 -0500 Subject: [PATCH 23/31] Adding cudf::cast for APIs --- cpp/src/binaryop/binaryop.cpp | 73 +++++++++---------- cpp/tests/binaryop/binop-integration-test.cpp | 34 +++++++++ 2 files changed, 70 insertions(+), 37 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 250ca959645..54ac1fb8dfa 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -42,6 +42,7 @@ #include #include #include +#include #include #include @@ -434,9 +435,13 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); - auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); - if (rhs.is_empty()) return out; - auto out_view = out->mutable_view(); + if (rhs.is_empty()) + return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); + + auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); + auto const type = cudf::data_type{rhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale @@ -448,7 +453,6 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, auto const scale = scale_type{rhs.type().scale()}; auto const scalar = make_fixed_point_scalar(val * factor, scale); binops::jit::binary_operation(out_view, *scalar, rhs, op, stream); - return out; } else { CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); @@ -456,31 +460,28 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, auto const scale = scale_type{rhs.type().scale()}; auto const scalar = make_fixed_point_scalar(val * factor, scale); binops::jit::binary_operation(out_view, *scalar, rhs, op, stream); - return out; } } else { auto const diff = rhs.type().scale() - lhs.type().scale(); auto const result = [&] { if (lhs.type().id() == type_id::DECIMAL32) { auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type(), stream, mr); } else { CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type(), stream, mr); } }(); binops::jit::binary_operation(out_view, lhs, result->view(), op, stream); - return out; } } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); - return out; } + if (output_type.scale() != scale) return cudf::cast(out_view, output_type); + return out; } /** @@ -504,9 +505,13 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); - auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); - if (lhs.is_empty()) return out; - auto out_view = out->mutable_view(); + if (lhs.is_empty()) + return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); + + auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); + auto const type = cudf::data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale @@ -518,7 +523,6 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, auto const scale = scale_type{lhs.type().scale()}; auto const scalar = make_fixed_point_scalar(val * factor, scale); binops::jit::binary_operation(out_view, lhs, *scalar, op, stream); - return out; } else { CUDF_EXPECTS(rhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); @@ -526,31 +530,28 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, auto const scale = scale_type{rhs.type().scale()}; auto const scalar = make_fixed_point_scalar(val * factor, scale); binops::jit::binary_operation(out_view, lhs, *scalar, op, stream); - return out; } } else { auto const diff = lhs.type().scale() - rhs.type().scale(); auto const result = [&] { if (rhs.type().id() == type_id::DECIMAL32) { auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type(), stream, mr); } else { CUDF_EXPECTS(rhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type(), stream, mr); } }(); binops::jit::binary_operation(out_view, result->view(), rhs, op, stream); - return out; } } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); - return out; } + if (output_type.scale() != scale) return cudf::cast(out_view, output_type); + return out; } /** @@ -574,53 +575,51 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, fixed_point_binary_operation_validation(op, lhs.type(), rhs.type(), output_type); - auto out = make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); - if (lhs.is_empty() or rhs.is_empty()) return out; - auto out_view = out->mutable_view(); + if (lhs.is_empty() or rhs.is_empty()) + return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); + + auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); + auto const type = cudf::data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { - // Adjust columns so they have they same scale TODO modify comment if (rhs.type().scale() < lhs.type().scale()) { auto const diff = lhs.type().scale() - rhs.type().scale(); auto const result = [&] { if (lhs.type().id() == type_id::DECIMAL32) { auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type(), stream, mr); } else { CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{lhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, lhs, binary_operator::MUL, rhs.type(), stream, mr); } }(); binops::jit::binary_operation(out_view, result->view(), rhs, op, stream); - return out; } else { auto const diff = rhs.type().scale() - lhs.type().scale(); auto const result = [&] { if (lhs.type().id() == type_id::DECIMAL32) { auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type(), stream, mr); } else { CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE"); auto const factor = numeric::detail::ipow(diff); - auto const scalar = - make_fixed_point_scalar(factor, scale_type{rhs.type().scale()}); + auto const scalar = make_fixed_point_scalar(factor, scale_type{-diff}); return binary_operation(*scalar, rhs, binary_operator::MUL, lhs.type(), stream, mr); } }(); binops::jit::binary_operation(out_view, lhs, result->view(), op, stream); - return out; } } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); - return out; } + if (output_type.scale() != scale) return cudf::cast(out_view, output_type); + return out; } std::unique_ptr binary_operation(scalar const& lhs, diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 67e9eddb6b5..4f4b01331d2 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -2255,6 +2256,39 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd5) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpAdd6) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const col = fp_wrapper{{3, 4, 5, 6, 7, 8}, scale_type{0}}; + + auto const expected1 = fp_wrapper{{6, 8, 10, 12, 14, 16}, scale_type{0}}; + auto const expected2 = fp_wrapper{{0, 0, 1, 1, 1, 1}, scale_type{1}}; + auto const type1 = cudf::data_type{cudf::type_to_id(), 0}; + auto const type2 = cudf::data_type{cudf::type_to_id(), 1}; + auto const result1 = cudf::binary_operation(col, col, cudf::binary_operator::ADD, type1); + auto const result2 = cudf::binary_operation(col, col, cudf::binary_operator::ADD, type2); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected1, result1->view()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected2, result2->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointCast) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const col = fp_wrapper{{6, 8, 10, 12, 14, 16}, scale_type{0}}; + auto const expected = fp_wrapper{{0, 0, 1, 1, 1, 1}, scale_type{1}}; + auto const type = cudf::data_type{cudf::type_to_id(), 1}; + auto const result = cudf::cast(col, type); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpMultiplyScalar) { using namespace numeric; From ff1926f22c78fc313b5cbb80a97fdf46ddd43c5d Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Mar 2021 17:36:20 -0500 Subject: [PATCH 24/31] Add except * --- python/cudf/cudf/_lib/types.pxd | 2 +- python/cudf/cudf/_lib/types.pyx | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/python/cudf/cudf/_lib/types.pxd b/python/cudf/cudf/_lib/types.pxd index 4e2202cbe6e..383b3665bd9 100644 --- a/python/cudf/cudf/_lib/types.pxd +++ b/python/cudf/cudf/_lib/types.pxd @@ -15,4 +15,4 @@ ctypedef bool underlying_type_t_null_policy cdef dtype_from_column_view(column_view cv) -cdef libcudf_types.data_type dtype_to_data_type(dtype) +cdef libcudf_types.data_type dtype_to_data_type(dtype) except * diff --git a/python/cudf/cudf/_lib/types.pyx b/python/cudf/cudf/_lib/types.pyx index 46521b3137c..d49623d6c67 100644 --- a/python/cudf/cudf/_lib/types.pyx +++ b/python/cudf/cudf/_lib/types.pyx @@ -210,7 +210,7 @@ cdef dtype_from_column_view(column_view cv): else: return cudf_to_np_types[(tid)] -cdef libcudf_types.data_type dtype_to_data_type(dtype): +cdef libcudf_types.data_type dtype_to_data_type(dtype) except *: if is_list_dtype(dtype): tid = libcudf_types.type_id.LIST elif is_struct_dtype(dtype): From f86f72095b06ad380853f64f9502db48de06ad4a Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 13:01:43 -0500 Subject: [PATCH 25/31] Fix + unit test for same scale comparison op --- cpp/src/binaryop/binaryop.cpp | 21 +++++++++++-------- cpp/tests/binaryop/binop-integration-test.cpp | 16 ++++++++++++++ 2 files changed, 28 insertions(+), 9 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 54ac1fb8dfa..05dd5ad6863 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -439,9 +439,10 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - auto const type = cudf::data_type{rhs.type().id(), scale}; - auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); - auto out_view = out->mutable_view(); + auto const type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : cudf::data_type{rhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale @@ -509,9 +510,10 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - auto const type = cudf::data_type{lhs.type().id(), scale}; - auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); - auto out_view = out->mutable_view(); + auto const type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : cudf::data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { // Adjust scalar/column so they have they same scale @@ -579,9 +581,10 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, return make_fixed_width_column_for_output(lhs, rhs, op, output_type, stream, mr); auto const scale = binary_operation_fixed_point_scale(op, lhs.type().scale(), rhs.type().scale()); - auto const type = cudf::data_type{lhs.type().id(), scale}; - auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); - auto out_view = out->mutable_view(); + auto const type = + is_comparison_binop(op) ? data_type{type_id::BOOL8} : cudf::data_type{lhs.type().id(), scale}; + auto out = make_fixed_width_column_for_output(lhs, rhs, op, type, stream, mr); + auto out_view = out->mutable_view(); if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) { if (rhs.type().scale() < lhs.type().scale()) { diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 4f4b01331d2..2f074364ca2 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2342,6 +2342,22 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimple) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimpleScale0) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const trues = std::vector(4, true); + auto const col = fp_wrapper{{1, 2, 3, 4}, scale_type{0}}; + auto const expected = wrapper(trues.begin(), trues.end()); + + auto const result = + cudf::binary_operation(col, col, binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) { using namespace numeric; From ff540c399becb5a78bae4496647d6bf7b8da373c Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 16:30:04 -0500 Subject: [PATCH 26/31] Use MAX_PRECISION --- python/cudf/cudf/_lib/types.pyx | 3 +-- python/cudf/cudf/core/dtypes.py | 6 +++--- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/python/cudf/cudf/_lib/types.pyx b/python/cudf/cudf/_lib/types.pyx index d49623d6c67..e9ed4f21ddd 100644 --- a/python/cudf/cudf/_lib/types.pyx +++ b/python/cudf/cudf/_lib/types.pyx @@ -193,8 +193,7 @@ cdef dtype_from_structs_column_view(column_view cv): cdef dtype_from_decimal_column_view(column_view cv): scale = -cv.type().scale() - precision = 18 # max of 64 bit integer - return Decimal64Dtype(precision=precision, scale=scale) + return Decimal64Dtype(precision=Decimal64Dtype.MAX_PRECISION, scale=scale) cdef dtype_from_column_view(column_view cv): cdef libcudf_types.type_id tid = cv.type().id() diff --git a/python/cudf/cudf/core/dtypes.py b/python/cudf/cudf/core/dtypes.py index 8b7d54b6715..a18aad3872b 100644 --- a/python/cudf/cudf/core/dtypes.py +++ b/python/cudf/cudf/core/dtypes.py @@ -235,7 +235,7 @@ class Decimal64Dtype(ExtensionDtype): name = "decimal" _metadata = ("precision", "scale") - _MAX_PRECISION = np.floor(np.log10(np.iinfo("int64").max)) + MAX_PRECISION = np.floor(np.log10(np.iinfo("int64").max)) def __init__(self, precision, scale=0): """ @@ -303,10 +303,10 @@ def __hash__(self): @classmethod def _validate(cls, precision, scale=0): - if precision > Decimal64Dtype._MAX_PRECISION: + if precision > Decimal64Dtype.MAX_PRECISION: raise ValueError( f"Cannot construct a {cls.__name__}" - f" with precision > {cls._MAX_PRECISION}" + f" with precision > {cls.MAX_PRECISION}" ) if abs(scale) > precision: raise ValueError(f"scale={scale} exceeds precision={precision}") From 98f4411f5458a30d551788444454ba5427769ccf Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 18:36:32 -0500 Subject: [PATCH 27/31] Use more declarative ternary operator --- cpp/src/binaryop/binaryop.cpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 05dd5ad6863..6b5afa69300 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -481,8 +481,7 @@ std::unique_ptr fixed_point_binary_operation(scalar const& lhs, } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); } - if (output_type.scale() != scale) return cudf::cast(out_view, output_type); - return out; + return output_type.scale() != scale ? cudf::cast(out_view, output_type) : std::move(out); } /** @@ -552,8 +551,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); } - if (output_type.scale() != scale) return cudf::cast(out_view, output_type); - return out; + return output_type.scale() != scale ? cudf::cast(out_view, output_type) : std::move(out); } /** @@ -621,8 +619,7 @@ std::unique_ptr fixed_point_binary_operation(column_view const& lhs, } else { binops::jit::binary_operation(out_view, lhs, rhs, op, stream); } - if (output_type.scale() != scale) return cudf::cast(out_view, output_type); - return out; + return output_type.scale() != scale ? cudf::cast(out_view, output_type) : std::move(out); } std::unique_ptr binary_operation(scalar const& lhs, From 81783200da665bd211cabb8abe9691e64c270baf Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 19:43:16 -0500 Subject: [PATCH 28/31] Use absolute path and remove local import --- python/cudf/cudf/utils/dtypes.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/python/cudf/cudf/utils/dtypes.py b/python/cudf/cudf/utils/dtypes.py index 2d86d822afa..e3d2ba2ea4a 100644 --- a/python/cudf/cudf/utils/dtypes.py +++ b/python/cudf/cudf/utils/dtypes.py @@ -331,9 +331,8 @@ def to_cudf_compatible_scalar(val, dtype=None): If `val` is None, returns None. """ - from cudf._lib.scalar import _is_null_host_scalar - if _is_null_host_scalar(val) or isinstance(val, cudf.Scalar): + if cudf._lib.scalar._is_null_host_scalar(val) or isinstance(val, cudf.Scalar): return val if not is_scalar(val): From c63870772f0672263856063a17f5a335eccfd975 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 19:50:58 -0500 Subject: [PATCH 29/31] black / flake8 fix --- python/cudf/cudf/utils/dtypes.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/utils/dtypes.py b/python/cudf/cudf/utils/dtypes.py index e3d2ba2ea4a..1438421bb12 100644 --- a/python/cudf/cudf/utils/dtypes.py +++ b/python/cudf/cudf/utils/dtypes.py @@ -332,7 +332,9 @@ def to_cudf_compatible_scalar(val, dtype=None): If `val` is None, returns None. """ - if cudf._lib.scalar._is_null_host_scalar(val) or isinstance(val, cudf.Scalar): + if cudf._lib.scalar._is_null_host_scalar(val) or isinstance( + val, cudf.Scalar + ): return val if not is_scalar(val): From efe92a3ae85c0806c6f2594df3d298961b48db5c Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 23:08:00 -0500 Subject: [PATCH 30/31] Use MAX_PRECISION --- python/cudf/cudf/core/column/decimal.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index ec956fea2ef..99af884eace 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -65,7 +65,7 @@ def binary_operator(self, op, other, reflect=False): self, other = other, self scale = _binop_scale(self.dtype, other.dtype, op) output_type = Decimal64Dtype( - scale=scale, precision=18 + scale=scale, precision=Decimal64Dtype.MAX_PRECISION ) # precision will be ignored, libcudf has no notion of precision result = libcudf.binaryop.binaryop(self, other, op, output_type) result.dtype.precision = _binop_precision(self.dtype, other.dtype, op) From a1e38878403db4070aa2baccabd9cf60449d252d Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Mon, 8 Mar 2021 23:30:36 -0500 Subject: [PATCH 31/31] Unit tests --- cpp/tests/binaryop/binop-integration-test.cpp | 32 +++++++++++++++++++ 1 file changed, 32 insertions(+) diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp index 2f074364ca2..019e72d3d3f 100644 --- a/cpp/tests/binaryop/binop-integration-test.cpp +++ b/cpp/tests/binaryop/binop-integration-test.cpp @@ -2358,6 +2358,38 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimpleScale0) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); } +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimpleScale0Null) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const col1 = fp_wrapper{{1, 2, 3, 4}, {1, 1, 1, 1}, scale_type{0}}; + auto const col2 = fp_wrapper{{1, 2, 3, 4}, {0, 0, 0, 0}, scale_type{0}}; + auto const expected = wrapper{{0, 1, 0, 1}, {0, 0, 0, 0}}; + + auto const result = + cudf::binary_operation(col1, col2, binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + +TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualSimpleScale2Null) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = device_storage_type_t; + + auto const col1 = fp_wrapper{{1, 2, 3, 4}, {1, 1, 1, 1}, scale_type{-2}}; + auto const col2 = fp_wrapper{{1, 2, 3, 4}, {0, 0, 0, 0}, scale_type{0}}; + auto const expected = wrapper{{0, 1, 0, 1}, {0, 0, 0, 0}}; + + auto const result = + cudf::binary_operation(col1, col2, binary_operator::EQUAL, cudf::data_type{type_id::BOOL8}); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view()); +} + TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpEqualLessGreater) { using namespace numeric;