Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add support for cudf::binary_operation TRUE_DIV for decimal32 and decimal64 #7198

Merged
merged 21 commits into from
Jan 29, 2021
Merged

Add support for cudf::binary_operation TRUE_DIV for decimal32 and decimal64 #7198

Show file tree
Hide file tree
Changes from 18 commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
b2c4519
Simple unit test
codereport Jan 22, 2021
ae2ec1f
Start of implementation
codereport Jan 22, 2021
68a27f1
More unit tests
codereport Jan 23, 2021
b31652f
Column-scalar and unit tests
codereport Jan 23, 2021
23a57a8
Unit tests, enhancements + fix
codereport Jan 23, 2021
403da38
Add negative exponent assert to ipow
codereport Jan 25, 2021
4e274bd
Scalar + column & unit tests
codereport Jan 25, 2021
ee8f84d
Full scalar + column & unit tests
codereport Jan 25, 2021
9daeb0a
Slight error msg fix
codereport Jan 25, 2021
314d86c
Fix & unit tests simplified
codereport Jan 25, 2021
91715f3
Cleanup (unused_scale is explicit)
codereport Jan 26, 2021
9c4aca6
More unit tests
codereport Jan 26, 2021
f249025
Cleanup
codereport Jan 26, 2021
ab98c4e
Final fixes and unit tests
codereport Jan 26, 2021
cf92067
Update cpp/src/binaryop/binaryop.cpp
codereport Jan 27, 2021
5da8e50
Update cpp/src/binaryop/binaryop.cpp
codereport Jan 27, 2021
a0d5e35
Update cpp/src/binaryop/binaryop.cpp
codereport Jan 27, 2021
17a31c2
Move optional header
codereport Jan 27, 2021
f4c6dc8
Addressing PR comments
codereport Jan 27, 2021
cddd1f9
Use release_assert
codereport Jan 27, 2021
8cfbe11
Add CUDF_EXPPECTS
codereport Jan 28, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions cpp/include/cudf/fixed_point/fixed_point.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,6 +90,7 @@ template <typename Rep,
is_supported_representation_type<Rep>())>* = nullptr>
CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
{
assert(("integer exponentiation with negative exponent is not possible.", exponent >= 0));
codereport marked this conversation as resolved.
Show resolved Hide resolved
if (exponent == 0) return static_cast<Rep>(1);
auto extra = static_cast<Rep>(1);
auto square = static_cast<Rep>(Base);
Expand Down
147 changes: 118 additions & 29 deletions cpp/src/binaryop/binaryop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,8 @@

#include <string>

#include <thrust/optional.h>

namespace cudf {

namespace binops {
Expand Down Expand Up @@ -376,7 +378,8 @@ 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);
return is_basic_arithmetic_binop(op) or is_comparison_binop(op) or
op == binary_operator::TRUE_DIV;
}

/**
Expand All @@ -389,6 +392,8 @@ 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;
return std::min(left_scale, right_scale);
Expand Down Expand Up @@ -419,25 +424,50 @@ bool is_same_scale_necessary(binary_operator op)
std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
column_view const& rhs,
binary_operator op,
thrust::optional<cudf::data_type> output_type,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
using namespace numeric;

CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(),
"Both columns must be of the same fixed_point type");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch");
CUDF_EXPECTS(op != binary_operator::TRUE_DIV || output_type.has_value(),
"TRUE_DIV requires result_type.");

auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale());
auto const output_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, output_type, stream, mr);
auto const scale = op == binary_operator::TRUE_DIV
codereport marked this conversation as resolved.
Show resolved Hide resolved
? output_type.value().scale()
: compute_scale_for_binop(op, lhs.type().scale(), rhs.type().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 out = make_fixed_width_column_for_output(lhs, rhs, op, out_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)) {
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;
codereport marked this conversation as resolved.
Show resolved Hide resolved
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<int32_t, Radix::BASE_10>(std::abs(diff));
auto const val = static_cast<fixed_point_scalar<decimal32> const&>(lhs).value();
auto const scaled_value = diff < 0 ? val / factor : val * factor;
auto const scalar = make_fixed_point_scalar<decimal32>(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<int64_t, Radix::BASE_10>(std::abs(diff));
auto const val = static_cast<fixed_point_scalar<decimal64> const&>(lhs).value();
auto const scaled_value = diff < 0 ? val / factor : val * factor;
auto const scalar = make_fixed_point_scalar<decimal64>(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)) {
// 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();
Expand Down Expand Up @@ -495,25 +525,50 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
scalar const& rhs,
binary_operator op,
thrust::optional<cudf::data_type> output_type,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
using namespace numeric;

CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(),
"Both columns must be of the same fixed_point type");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch");
CUDF_EXPECTS(op != binary_operator::TRUE_DIV || output_type.has_value(),
"TRUE_DIV requires result_type.");

auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale());
auto const output_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, output_type, stream, mr);
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 out_type = output_type.value_or(
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);

if (lhs.is_empty()) return out;

auto out_view = out->mutable_view();

if (lhs.type().scale() != rhs.type().scale() && is_same_scale_necessary(op)) {
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;
codereport marked this conversation as resolved.
Show resolved Hide resolved
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<int32_t, Radix::BASE_10>(std::abs(diff));
auto const scalar = make_fixed_point_scalar<decimal32>(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<int64_t, Radix::BASE_10>(std::abs(diff));
auto const scalar = make_fixed_point_scalar<decimal64>(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)) {
// 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();
Expand Down Expand Up @@ -571,26 +626,51 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
column_view const& rhs,
binary_operator op,
thrust::optional<cudf::data_type> output_type,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
using namespace numeric;

CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(),
"Both columns must be of the same fixed_point type");
CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Data type mismatch");
CUDF_EXPECTS(op != binary_operator::TRUE_DIV || output_type.has_value(),
"TRUE_DIV requires result_type.");

auto const scale = compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale());
auto const output_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, output_type, stream, mr);
auto const scale = op == binary_operator::TRUE_DIV
? output_type.value().scale()
codereport marked this conversation as resolved.
Show resolved Hide resolved
: compute_scale_for_binop(op, lhs.type().scale(), rhs.type().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 out = make_fixed_width_column_for_output(lhs, rhs, op, out_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)) {
// Adjust columns so they have they same scale
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;
codereport marked this conversation as resolved.
Show resolved Hide resolved
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) {
codereport marked this conversation as resolved.
Show resolved Hide resolved
auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(std::abs(diff));
auto const scalar = make_fixed_point_scalar<decimal32>(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<int64_t, Radix::BASE_10>(std::abs(diff));
auto const scalar = make_fixed_point_scalar<decimal64>(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)) {
// 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 = [&] {
Expand Down Expand Up @@ -644,8 +724,11 @@ std::unique_ptr<column> 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()))
return fixed_point_binary_operation(lhs, rhs, op, 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<data_type>{thrust::nullopt};
return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr);
}

// Check for datatype
CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype");
Expand All @@ -671,8 +754,11 @@ std::unique_ptr<column> 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()))
return fixed_point_binary_operation(lhs, rhs, op, 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<data_type>{thrust::nullopt};
return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr);
}

// Check for datatype
CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype");
Expand Down Expand Up @@ -700,8 +786,11 @@ std::unique_ptr<column> 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()))
return fixed_point_binary_operation(lhs, rhs, op, 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<data_type>{thrust::nullopt};
return fixed_point_binary_operation(lhs, rhs, op, type, stream, mr);
}

// Check for datatype
CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype");
Expand Down
Loading