From b2c451991e6be1d728bb1949073b1d7b4a093bed Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Fri, 22 Jan 2021 17:54:24 -0500
Subject: [PATCH 01/21] Simple unit test

---
 cpp/tests/binaryop/binop-integration-test.cpp | 16 ++++++++++++++++
 1 file changed, 16 insertions(+)

diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index f2f68c7601b..50fe0b2801e 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2303,6 +2303,22 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpNullEqualsSimple)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+  auto const rhs      = fp_wrapper<RepType>{{4, 4, 4, 4}, scale_type{0}};
+  auto const expected = fp_wrapper<RepType>{{25, 75, 125, 175}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From ae2ec1f4b1731c04ffa43ce3fc2373741839c4c4 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Fri, 22 Jan 2021 17:54:34 -0500
Subject: [PATCH 02/21] Start of implementation

---
 cpp/src/binaryop/binaryop.cpp | 54 ++++++++++++++++++++++++++++-------
 1 file changed, 43 insertions(+), 11 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 2b1bc348d45..f2dbf8e1a73 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -20,6 +20,7 @@
 #include "compiled/binary_ops.hpp"
 #include "jit/code/code.h"
 #include "jit/util.hpp"
+#include "thrust/optional.h"
 
 #include <jit/launcher.h>
 #include <jit/parser.h>
@@ -376,7 +377,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;
 }
 
 /**
@@ -389,6 +391,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);
@@ -571,26 +575,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(), "Both columns must be of the same type");
+  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);
+  // TODO cleanup
+  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([&] {
+    return 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) {
+    auto const diff   = lhs.type().scale() - output_type.value().scale();
+    auto const result = [&] {
+      if (lhs.type().id() == type_id::DECIMAL32) {
+        auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
+        auto const scalar = make_fixed_point_scalar<decimal32>(factor, scale_type{scale});
+        return binary_operation(*scalar, lhs, binary_operator::MUL, lhs.type(), stream, mr);
+      } else {
+        CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE");
+        auto const factor = numeric::detail::ipow<int64_t, Radix::BASE_10>(diff);
+        auto const scalar = make_fixed_point_scalar<decimal64>(factor, scale_type{scale});
+        return binary_operation(*scalar, lhs, binary_operator::MUL, lhs.type(), 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 = [&] {
@@ -700,8 +729,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 optional_type =
+      op == binary_operator::TRUE_DIV ? output_type : thrust::optional<data_type>{thrust::nullopt};
+    return fixed_point_binary_operation(lhs, rhs, op, optional_type, stream, mr);
+  }
 
   // Check for datatype
   CUDF_EXPECTS(is_fixed_width(output_type), "Invalid/Unsupported output datatype");

From 68a27f104704d1f3d9f952d57c008201f054a6b8 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Sat, 23 Jan 2021 00:54:28 -0500
Subject: [PATCH 03/21] More unit tests

---
 cpp/src/binaryop/binaryop.cpp                 | 12 +++----
 cpp/tests/binaryop/binop-integration-test.cpp | 32 +++++++++++++++++++
 2 files changed, 37 insertions(+), 7 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index f2dbf8e1a73..831c38039cf 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -591,10 +591,8 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
                        ? output_type.value().scale()
                        : compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale());
 
-  auto const out_type = output_type.value_or([&] {
-    return is_comparison_binop(op) ? data_type{type_id::BOOL8}  //
-                                   : data_type{lhs.type().id(), 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);
 
@@ -603,17 +601,17 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
   auto out_view = out->mutable_view();
 
   if (op == binary_operator::TRUE_DIV) {
-    auto const diff   = lhs.type().scale() - output_type.value().scale();
+    auto const diff   = lhs.type().scale() - scale;
     auto const result = [&] {
       if (lhs.type().id() == type_id::DECIMAL32) {
         auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
         auto const scalar = make_fixed_point_scalar<decimal32>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, lhs.type(), stream, mr);
+        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
       } else {
         CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE");
         auto const factor = numeric::detail::ipow<int64_t, Radix::BASE_10>(diff);
         auto const scalar = make_fixed_point_scalar<decimal64>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, lhs.type(), stream, mr);
+        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
       }
     }();
     binops::jit::binary_operation(out_view, result->view(), rhs, binary_operator::DIV, stream);
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index 50fe0b2801e..898284a9946 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2319,6 +2319,38 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv2)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
+  auto const rhs      = fp_wrapper<RepType>{{2, 2, 2, 2}, scale_type{0}};
+  auto const expected = fp_wrapper<RepType>{{5000, 15000, 25000, 35000}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
+  auto const rhs      = fp_wrapper<RepType>{{3, 9, 3, 3}, scale_type{0}};
+  auto const expected = fp_wrapper<RepType>{{3333, 3333, 16666, 23333}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From b31652f74b8eb4eba44ac24dedd7fb219d3e5b4f Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Sat, 23 Jan 2021 01:34:20 -0500
Subject: [PATCH 04/21] Column-scalar and unit tests

---
 cpp/src/binaryop/binaryop.cpp                 | 50 ++++++++++++++-----
 cpp/tests/binaryop/binop-integration-test.cpp | 16 ++++++
 2 files changed, 54 insertions(+), 12 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 831c38039cf..e5e2890f08e 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -499,25 +499,48 @@ 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(), "Both columns must be of the same type");
+  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() - scale;
+    auto const result = [&] {
+      if (lhs.type().id() == type_id::DECIMAL32) {
+        auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
+        auto const scalar = make_fixed_point_scalar<decimal32>(factor, scale_type{scale});
+        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
+      } else {
+        CUDF_EXPECTS(lhs.type().id() == type_id::DECIMAL64, "Unexpected DTYPE");
+        auto const factor = numeric::detail::ipow<int64_t, Radix::BASE_10>(diff);
+        auto const scalar = make_fixed_point_scalar<decimal64>(factor, scale_type{scale});
+        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, 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();
@@ -586,7 +609,6 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
   CUDF_EXPECTS(op != binary_operator::TRUE_DIV || output_type.has_value(),
                "TRUE_DIV requires result_type.");
 
-  // TODO cleanup
   auto const scale = op == binary_operator::TRUE_DIV
                        ? output_type.value().scale()
                        : compute_scale_for_binop(op, lhs.type().scale(), rhs.type().scale());
@@ -601,6 +623,7 @@ std::unique_ptr<column> 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() - scale;
     auto const result = [&] {
       if (lhs.type().id() == type_id::DECIMAL32) {
@@ -698,8 +721,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");
@@ -728,9 +754,9 @@ std::unique_ptr<column> 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 optional_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, optional_type, stream, mr);
+    return fixed_point_binary_operation(lhs, rhs, op, 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 898284a9946..c4d3c919340 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2351,6 +2351,22 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv4)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
+  auto const rhs      = make_fixed_point_scalar<decimalXX>(3, scale_type{0});
+  auto const expected = fp_wrapper<RepType>{{3333, 10000, 16666, 23333}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From 23a57a8b5b2307289e3632298a7b33f33fbee8bf Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Sat, 23 Jan 2021 01:43:20 -0500
Subject: [PATCH 05/21] Unit tests, enhancements + fix

---
 cpp/src/binaryop/binaryop.cpp                 |  4 ++--
 cpp/tests/binaryop/binop-integration-test.cpp | 20 +++++++++++++++++--
 2 files changed, 20 insertions(+), 4 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index e5e2890f08e..9ba66aab57a 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -525,7 +525,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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() - scale;
+    auto const diff   = lhs.type().scale() - 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>(diff);
@@ -624,7 +624,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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() - scale;
+    auto const diff   = lhs.type().scale() - 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>(diff);
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index c4d3c919340..8fca5a30f4e 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2326,7 +2326,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv2)
   using RepType   = device_storage_type_t<decimalXX>;
 
   auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
-  auto const rhs      = fp_wrapper<RepType>{{2, 2, 2, 2}, scale_type{0}};
+  auto const rhs      = fp_wrapper<RepType>{{20, 20, 20, 20}, scale_type{-1}};
   auto const expected = fp_wrapper<RepType>{{5000, 15000, 25000, 35000}, scale_type{-2}};
 
   auto const type   = data_type{type_to_id<decimalXX>(), -2};
@@ -2342,7 +2342,7 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv3)
   using RepType   = device_storage_type_t<decimalXX>;
 
   auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
-  auto const rhs      = fp_wrapper<RepType>{{3, 9, 3, 3}, scale_type{0}};
+  auto const rhs      = fp_wrapper<RepType>{{300, 900, 300, 300}, scale_type{-2}};
   auto const expected = fp_wrapper<RepType>{{3333, 3333, 16666, 23333}, scale_type{-2}};
 
   auto const type   = data_type{type_to_id<decimalXX>(), -2};
@@ -2367,6 +2367,22 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv4)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv5)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{1}};
+  auto const rhs      = make_fixed_point_scalar<decimalXX>(30000, scale_type{-4});
+  auto const expected = fp_wrapper<RepType>{{3333, 10000, 16666, 23333}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From 403da38e6b091215517a94908c6f593243f42e12 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 25 Jan 2021 09:52:21 -0500
Subject: [PATCH 06/21] Add negative exponent assert to  ipow

---
 cpp/include/cudf/fixed_point/fixed_point.hpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp
index 0524b8e336a..627b0188965 100644
--- a/cpp/include/cudf/fixed_point/fixed_point.hpp
+++ b/cpp/include/cudf/fixed_point/fixed_point.hpp
@@ -27,6 +27,7 @@
 #include <cassert>
 #include <cmath>
 #include <string>
+#include "cudf/utilities/error.hpp"
 
 //! `fixed_point` and supporting types
 namespace numeric {
@@ -90,6 +91,7 @@ template <typename Rep,
                                            is_supported_representation_type<Rep>())>* = nullptr>
 CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
 {
+  assert(("integer power on negative exponent is not possible.", exponent >= 0));
   if (exponent == 0) return static_cast<Rep>(1);
   auto extra  = static_cast<Rep>(1);
   auto square = static_cast<Rep>(Base);

From 4e274bd794f85b80c3eb7066e982ac25b91eac83 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 25 Jan 2021 09:53:41 -0500
Subject: [PATCH 07/21] Scalar + column & unit tests

---
 cpp/src/binaryop/binaryop.cpp                 | 49 ++++++++++---
 cpp/tests/binaryop/binop-integration-test.cpp | 68 +++++++++++++++++++
 2 files changed, 108 insertions(+), 9 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 9ba66aab57a..22f8694df83 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -423,25 +423,53 @@ 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(), "Both columns must be of the same type");
+  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 (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;
+    if (lhs.type().id() == type_id::DECIMAL32) {
+      auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
+      auto const val    = static_cast<fixed_point_scalar<decimal32> const&>(lhs).value();
+      auto const scalar = lhs.type().scale() < rhs.type().scale()
+                            ? make_fixed_point_scalar<decimal32>(val / factor, scale_type{scale})
+                            : make_fixed_point_scalar<decimal32>(val * factor, scale_type{scale});
+      printf("%i %i\n", factor, val);
+      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>(diff);
+      auto const val    = static_cast<fixed_point_scalar<decimal64> const&>(lhs).value();
+      auto const scalar = lhs.type().scale() < rhs.type().scale()
+                            ? make_fixed_point_scalar<decimal64>(val / factor, scale_type{scale})
+                            : make_fixed_point_scalar<decimal64>(val * factor, scale_type{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();
@@ -694,8 +722,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");
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index 8fca5a30f4e..ea769ad48ef 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2383,6 +2383,74 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv5)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs = make_fixed_point_scalar<decimalXX>(30000, scale_type{-4});
+  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+
+  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
+// {
+//   using namespace numeric;
+//   using decimalXX = TypeParam;
+//   using RepType   = device_storage_type_t<decimalXX>;
+
+//   auto const lhs = make_fixed_point_scalar<decimalXX>(300, scale_type{-2});
+//   auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+
+//   auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+
+//   auto const type   = data_type{type_to_id<decimalXX>(), -2};
+//   auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+//   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+// }
+
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs = make_fixed_point_scalar<decimalXX>(30, scale_type{-1});
+  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+
+  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv9)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs = make_fixed_point_scalar<decimalXX>(3, scale_type{0});
+  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+
+  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), -2};
+  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From ee8f84dac37f7e8fed67cdf431b2510953d06f37 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 25 Jan 2021 14:59:27 -0500
Subject: [PATCH 08/21] Full scalar + column & unit tests

---
 cpp/src/binaryop/binaryop.cpp                 | 14 +++++------
 cpp/tests/binaryop/binop-integration-test.cpp | 24 +++++++++----------
 2 files changed, 18 insertions(+), 20 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 22f8694df83..fea49a9c910 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -449,23 +449,21 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
 
   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 diff         = lhs.type().scale() - rhs.type().scale() - scale;
+    auto const scalar_scale = scale_type{rhs.type().scale() + scale};
     if (lhs.type().id() == type_id::DECIMAL32) {
       auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
       auto const val    = static_cast<fixed_point_scalar<decimal32> const&>(lhs).value();
-      auto const scalar = lhs.type().scale() < rhs.type().scale()
-                            ? make_fixed_point_scalar<decimal32>(val / factor, scale_type{scale})
-                            : make_fixed_point_scalar<decimal32>(val * factor, scale_type{scale});
-      printf("%i %i\n", factor, val);
+      auto const scalar = diff < 0 ? make_fixed_point_scalar<decimal32>(val / factor, scalar_scale)
+                                   : make_fixed_point_scalar<decimal32>(val * factor, scalar_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>(diff);
       auto const val    = static_cast<fixed_point_scalar<decimal64> const&>(lhs).value();
-      auto const scalar = lhs.type().scale() < rhs.type().scale()
-                            ? make_fixed_point_scalar<decimal64>(val / factor, scale_type{scale})
-                            : make_fixed_point_scalar<decimal64>(val * factor, scale_type{scale});
+      auto const scalar = diff < 0 ? make_fixed_point_scalar<decimal64>(val / factor, scalar_scale)
+                                   : make_fixed_point_scalar<decimal64>(val * factor, scalar_scale);
       binops::jit::binary_operation(out_view, *scalar, rhs, binary_operator::DIV, stream);
       return out;
     }
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index ea769ad48ef..1bddf82ec06 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2400,22 +2400,22 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
-// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
-// {
-//   using namespace numeric;
-//   using decimalXX = TypeParam;
-//   using RepType   = device_storage_type_t<decimalXX>;
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
 
-//   auto const lhs = make_fixed_point_scalar<decimalXX>(300, scale_type{-2});
-//   auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+  auto const lhs = make_fixed_point_scalar<decimalXX>(300, scale_type{-2});
+  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
 
-//   auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
 
-//   auto const type   = data_type{type_to_id<decimalXX>(), -2};
-//   auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+  auto const type   = data_type{type_to_id<decimalXX>(), -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, FixedPointBinaryOpTrueDiv8)
 {

From 9daeb0ac6f2085349e3fd75dce5ff0b5a8ced693 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 25 Jan 2021 15:01:45 -0500
Subject: [PATCH 09/21] Slight error msg fix

---
 cpp/include/cudf/fixed_point/fixed_point.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp
index 627b0188965..b79b9694b89 100644
--- a/cpp/include/cudf/fixed_point/fixed_point.hpp
+++ b/cpp/include/cudf/fixed_point/fixed_point.hpp
@@ -91,7 +91,7 @@ template <typename Rep,
                                            is_supported_representation_type<Rep>())>* = nullptr>
 CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
 {
-  assert(("integer power on negative exponent is not possible.", exponent >= 0));
+  assert(("integer power with negative exponent is not possible.", exponent >= 0));
   if (exponent == 0) return static_cast<Rep>(1);
   auto extra  = static_cast<Rep>(1);
   auto square = static_cast<Rep>(Base);

From 314d86c38cbdf8eecc4fc1ca7fa333a143bae50e Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 25 Jan 2021 15:52:00 -0500
Subject: [PATCH 10/21] Fix & unit tests simplified

---
 cpp/src/binaryop/binaryop.cpp                 |  4 +-
 cpp/tests/binaryop/binop-integration-test.cpp | 66 +++----------------
 2 files changed, 11 insertions(+), 59 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index fea49a9c910..f9c1a090b78 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -452,7 +452,7 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
     auto const diff         = lhs.type().scale() - rhs.type().scale() - scale;
     auto const scalar_scale = scale_type{rhs.type().scale() + scale};
     if (lhs.type().id() == type_id::DECIMAL32) {
-      auto const factor = numeric::detail::ipow<int32_t, Radix::BASE_10>(diff);
+      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 scalar = diff < 0 ? make_fixed_point_scalar<decimal32>(val / factor, scalar_scale)
                                    : make_fixed_point_scalar<decimal32>(val * factor, scalar_scale);
@@ -460,7 +460,7 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
       return out;
 
     } else {
-      auto const factor = numeric::detail::ipow<int64_t, Radix::BASE_10>(diff);
+      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 scalar = diff < 0 ? make_fixed_point_scalar<decimal64>(val / factor, scalar_scale)
                                    : make_fixed_point_scalar<decimal64>(val * factor, scalar_scale);
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index 1bddf82ec06..0d50e56b1fb 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2389,66 +2389,18 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6)
   using decimalXX = TypeParam;
   using RepType   = device_storage_type_t<decimalXX>;
 
-  auto const lhs = make_fixed_point_scalar<decimalXX>(30000, scale_type{-4});
-  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
+  for (auto const i : {0, 1, 2, 3, 4, 5, 6, 7}) {
+    auto const val = 3 * numeric::detail::ipow<int32_t, Radix::BASE_10>(i);
+    auto const lhs = make_fixed_point_scalar<decimalXX>(val, scale_type{-i});
+    auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
 
-  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
+    auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
 
-  auto const type   = data_type{type_to_id<decimalXX>(), -2};
-  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
-
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
-}
-
-TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
-{
-  using namespace numeric;
-  using decimalXX = TypeParam;
-  using RepType   = device_storage_type_t<decimalXX>;
-
-  auto const lhs = make_fixed_point_scalar<decimalXX>(300, scale_type{-2});
-  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
-
-  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
-
-  auto const type   = data_type{type_to_id<decimalXX>(), -2};
-  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
-
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
-}
-
-TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8)
-{
-  using namespace numeric;
-  using decimalXX = TypeParam;
-  using RepType   = device_storage_type_t<decimalXX>;
-
-  auto const lhs = make_fixed_point_scalar<decimalXX>(30, scale_type{-1});
-  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
-
-  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
-
-  auto const type   = data_type{type_to_id<decimalXX>(), -2};
-  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+    auto const type   = data_type{type_to_id<decimalXX>(), -2};
+    auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
 
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
-}
-
-TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv9)
-{
-  using namespace numeric;
-  using decimalXX = TypeParam;
-  using RepType   = device_storage_type_t<decimalXX>;
-
-  auto const lhs = make_fixed_point_scalar<decimalXX>(3, scale_type{0});
-  auto const rhs = fp_wrapper<RepType>{{10, 30, 50, 70}, scale_type{-1}};
-
-  auto const expected = fp_wrapper<RepType>{{300, 100, 60, 42}, scale_type{-2}};
-
-  auto const type   = data_type{type_to_id<decimalXX>(), -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());
+  }
 }
 
 }  // namespace binop

From 91715f3fd1fa2afe8c05177c7e11b3f750dbde25 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Tue, 26 Jan 2021 10:58:41 -0500
Subject: [PATCH 11/21] Cleanup (unused_scale is explicit)

---
 cpp/src/binaryop/binaryop.cpp | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index f9c1a090b78..f70f70ad1a2 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -450,20 +450,20 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
   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 scalar_scale = scale_type{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<int32_t, Radix::BASE_10>(std::abs(diff));
-      auto const val    = static_cast<fixed_point_scalar<decimal32> const&>(lhs).value();
-      auto const scalar = diff < 0 ? make_fixed_point_scalar<decimal32>(val / factor, scalar_scale)
-                                   : make_fixed_point_scalar<decimal32>(val * factor, scalar_scale);
+      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 scalar = diff < 0 ? make_fixed_point_scalar<decimal64>(val / factor, scalar_scale)
-                                   : make_fixed_point_scalar<decimal64>(val * factor, scalar_scale);
+      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;
     }

From 9c4aca6acb0ff52798b7e5af806eac97e8e0d259 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Tue, 26 Jan 2021 11:06:20 -0500
Subject: [PATCH 12/21] More unit tests

---
 cpp/src/binaryop/binaryop.cpp                 |  1 -
 cpp/tests/binaryop/binop-integration-test.cpp | 33 +++++++++++++++++++
 2 files changed, 33 insertions(+), 1 deletion(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index f70f70ad1a2..7315dc80602 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -458,7 +458,6 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
       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();
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index 0d50e56b1fb..a76a613c162 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2403,6 +2403,39 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv6)
   }
 }
 
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs = make_fixed_point_scalar<decimalXX>(12000, scale_type{-1});
+  auto const rhs = fp_wrapper<RepType>{{100, 200, 300, 500, 600, 800, 1200, 1300}, scale_type{-2}};
+
+  auto const expected = fp_wrapper<RepType>{{12, 6, 4, 2, 2, 1, 1, 0}, scale_type{2}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), 2};
+  auto const result = cudf::binary_operation(*lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8)
+// {
+//   using namespace numeric;
+//   using decimalXX = TypeParam;
+//   using RepType   = device_storage_type_t<decimalXX>;
+
+//   auto const lhs      = fp_wrapper<RepType>{{4000, 6000, 80000}, scale_type{-1}};
+//   auto const rhs      = make_fixed_point_scalar<decimalXX>(500, scale_type{-2});
+//   auto const expected = fp_wrapper<RepType>{{0, 1, 40}, scale_type{2}};
+
+//   auto const type   = data_type{type_to_id<decimalXX>(), 2};
+//   auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+
+//   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+// }
+
 }  // namespace binop
 }  // namespace test
 }  // namespace cudf

From f2490254489dff773b07c1622b0ffe64a5ae177e Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Tue, 26 Jan 2021 11:07:26 -0500
Subject: [PATCH 13/21] Cleanup

---
 cpp/include/cudf/fixed_point/fixed_point.hpp | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp
index b79b9694b89..b3fa9b7be0a 100644
--- a/cpp/include/cudf/fixed_point/fixed_point.hpp
+++ b/cpp/include/cudf/fixed_point/fixed_point.hpp
@@ -27,7 +27,6 @@
 #include <cassert>
 #include <cmath>
 #include <string>
-#include "cudf/utilities/error.hpp"
 
 //! `fixed_point` and supporting types
 namespace numeric {
@@ -91,7 +90,7 @@ template <typename Rep,
                                            is_supported_representation_type<Rep>())>* = nullptr>
 CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
 {
-  assert(("integer power with negative exponent is not possible.", exponent >= 0));
+  assert(("integer exponentiation with negative exponent is not possible.", exponent >= 0));
   if (exponent == 0) return static_cast<Rep>(1);
   auto extra  = static_cast<Rep>(1);
   auto square = static_cast<Rep>(Base);

From ab98c4ee857b718393526e76d112cdc9361e4fd5 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Tue, 26 Jan 2021 13:09:12 -0500
Subject: [PATCH 14/21] Final fixes and unit tests

---
 cpp/src/binaryop/binaryop.cpp                 | 36 +++++-----
 cpp/tests/binaryop/binop-integration-test.cpp | 72 +++++++++++++++----
 2 files changed, 80 insertions(+), 28 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 7315dc80602..32a7f4fa149 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -550,17 +550,19 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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 result = [&] {
+    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<int32_t, Radix::BASE_10>(diff);
-        auto const scalar = make_fixed_point_scalar<decimal32>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
+        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>(diff);
-        auto const scalar = make_fixed_point_scalar<decimal64>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
+        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);
@@ -649,17 +651,19 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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 result = [&] {
+    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<int32_t, Radix::BASE_10>(diff);
-        auto const scalar = make_fixed_point_scalar<decimal32>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
+        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>(diff);
-        auto const scalar = make_fixed_point_scalar<decimal64>(factor, scale_type{scale});
-        return binary_operation(*scalar, lhs, binary_operator::MUL, {}, stream, mr);
+        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);
diff --git a/cpp/tests/binaryop/binop-integration-test.cpp b/cpp/tests/binaryop/binop-integration-test.cpp
index a76a613c162..087a65f6c66 100644
--- a/cpp/tests/binaryop/binop-integration-test.cpp
+++ b/cpp/tests/binaryop/binop-integration-test.cpp
@@ -2420,21 +2420,69 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv7)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
 }
 
-// TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8)
-// {
-//   using namespace numeric;
-//   using decimalXX = TypeParam;
-//   using RepType   = device_storage_type_t<decimalXX>;
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv8)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
 
-//   auto const lhs      = fp_wrapper<RepType>{{4000, 6000, 80000}, scale_type{-1}};
-//   auto const rhs      = make_fixed_point_scalar<decimalXX>(500, scale_type{-2});
-//   auto const expected = fp_wrapper<RepType>{{0, 1, 40}, scale_type{2}};
+  auto const lhs      = fp_wrapper<RepType>{{4000, 6000, 80000}, scale_type{-1}};
+  auto const rhs      = make_fixed_point_scalar<decimalXX>(500, scale_type{-2});
+  auto const expected = fp_wrapper<RepType>{{0, 1, 16}, scale_type{2}};
 
-//   auto const type   = data_type{type_to_id<decimalXX>(), 2};
-//   auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+  auto const type   = data_type{type_to_id<decimalXX>(), 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, FixedPointBinaryOpTrueDiv9)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{100000, 200000, 300000}, scale_type{-2}};
+  auto const rhs      = make_fixed_point_scalar<decimalXX>(7, scale_type{1});
+  auto const expected = fp_wrapper<RepType>{{1, 2, 4}, scale_type{1}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), 1};
+  auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv10)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{100000, 200000, 300000}, scale_type{-2}};
+  auto const rhs      = make_fixed_point_scalar<decimalXX>(7, scale_type{0});
+  auto const expected = fp_wrapper<RepType>{{14, 28, 42}, scale_type{1}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), 1};
+  auto const result = cudf::binary_operation(lhs, *rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
+TYPED_TEST(FixedPointTestBothReps, FixedPointBinaryOpTrueDiv11)
+{
+  using namespace numeric;
+  using decimalXX = TypeParam;
+  using RepType   = device_storage_type_t<decimalXX>;
+
+  auto const lhs      = fp_wrapper<RepType>{{1000000, 2000000, 3000000}, scale_type{-2}};
+  auto const rhs      = fp_wrapper<RepType>{{7, 7, 7}, scale_type{0}};
+  auto const expected = fp_wrapper<RepType>{{142, 285, 428}, scale_type{1}};
+
+  auto const type   = data_type{type_to_id<decimalXX>(), 1};
+  auto const result = cudf::binary_operation(lhs, rhs, cudf::binary_operator::TRUE_DIV, type);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
 
 }  // namespace binop
 }  // namespace test

From cf920677e0c6bd89024215b913192771a1bb1938 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <36027403+codereport@users.noreply.github.com>
Date: Wed, 27 Jan 2021 09:30:17 -0500
Subject: [PATCH 15/21] Update cpp/src/binaryop/binaryop.cpp

Co-authored-by: Mark Harris <mharris@nvidia.com>
---
 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 32a7f4fa149..a75e463b17e 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -632,7 +632,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
   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 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.");
 

From 5da8e50aa46b71c15fc7330fd72013c4d1baf742 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <36027403+codereport@users.noreply.github.com>
Date: Wed, 27 Jan 2021 09:30:28 -0500
Subject: [PATCH 16/21] Update cpp/src/binaryop/binaryop.cpp

Co-authored-by: Mark Harris <mharris@nvidia.com>
---
 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 a75e463b17e..2397497cbf3 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -430,7 +430,7 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
   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 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.");
 

From a0d5e351052746ecb3c66ca1f5fa339a40c63d5f Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <36027403+codereport@users.noreply.github.com>
Date: Wed, 27 Jan 2021 09:30:35 -0500
Subject: [PATCH 17/21] Update cpp/src/binaryop/binaryop.cpp

Co-authored-by: Mark Harris <mharris@nvidia.com>
---
 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 2397497cbf3..690ada10421 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -531,7 +531,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
   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 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.");
 

From 17a31c26203c1268d296efe0739d1c3106255b88 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Wed, 27 Jan 2021 09:38:32 -0500
Subject: [PATCH 18/21] Move optional header

---
 cpp/src/binaryop/binaryop.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 690ada10421..a14eea5ccc4 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -20,7 +20,6 @@
 #include "compiled/binary_ops.hpp"
 #include "jit/code/code.h"
 #include "jit/util.hpp"
-#include "thrust/optional.h"
 
 #include <jit/launcher.h>
 #include <jit/parser.h>
@@ -50,6 +49,8 @@
 
 #include <string>
 
+#include <thrust/optional.h>
+
 namespace cudf {
 
 namespace binops {

From f4c6dc89966bda6c272aab495f4095146c67ea48 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Wed, 27 Jan 2021 15:11:41 -0500
Subject: [PATCH 19/21] Addressing PR comments

---
 cpp/src/binaryop/binaryop.cpp | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index a14eea5ccc4..3d70946f098 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -432,8 +432,8 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
 
   CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
   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.");
+  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()
@@ -450,7 +450,7 @@ std::unique_ptr<column> fixed_point_binary_operation(scalar const& lhs,
 
   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 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<int32_t, Radix::BASE_10>(std::abs(diff));
@@ -533,8 +533,8 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
   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.");
+  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()
@@ -551,7 +551,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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 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       = [&] {
@@ -634,8 +634,8 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   CUDF_EXPECTS(is_supported_fixed_point_binop(op), "Unsupported fixed_point binary operation");
   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.");
+  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()
@@ -652,7 +652,7 @@ std::unique_ptr<column> fixed_point_binary_operation(column_view const& lhs,
 
   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 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       = [&] {

From cddd1f952d28f831fcf3fdcc898dcd53991fe50e Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Wed, 27 Jan 2021 16:08:25 -0500
Subject: [PATCH 20/21] Use release_assert

---
 cpp/include/cudf/fixed_point/fixed_point.hpp | 21 ++++++++++----------
 1 file changed, 11 insertions(+), 10 deletions(-)

diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp
index b3fa9b7be0a..c8e2ac01f41 100644
--- a/cpp/include/cudf/fixed_point/fixed_point.hpp
+++ b/cpp/include/cudf/fixed_point/fixed_point.hpp
@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cudf/detail/utilities/release_assert.cuh>
 #include <cudf/types.hpp>
 
 // Note: The <cuda/std/*> versions are used in order for Jitify to work with our fixed_point type.
@@ -24,7 +25,6 @@
 #include <cuda/std/type_traits>  // add cuda namespace
 
 #include <algorithm>
-#include <cassert>
 #include <cmath>
 #include <string>
 
@@ -90,7 +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));
+  release_assert(exponent >= 0 && "integer exponentiation with negative exponent is not possible.");
   if (exponent == 0) return static_cast<Rep>(1);
   auto extra  = static_cast<Rep>(1);
   auto square = static_cast<Rep>(Base);
@@ -641,8 +641,8 @@ CUDA_HOST_DEVICE_CALLABLE fixed_point<Rep1, Rad1> operator+(fixed_point<Rep1, Ra
 
 #if defined(__CUDACC_DEBUG__)
 
-  assert(("fixed_point overflow of underlying representation type " + print_rep<Rep1>(),
-          !addition_overflow<Rep1>(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value)));
+  release_assert(!addition_overflow<Rep1>(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value) &&
+                 "fixed_point overflow of underlying representation type " + print_rep<Rep1>());
 
 #endif
 
@@ -659,8 +659,9 @@ CUDA_HOST_DEVICE_CALLABLE fixed_point<Rep1, Rad1> operator-(fixed_point<Rep1, Ra
 
 #if defined(__CUDACC_DEBUG__)
 
-  assert(("fixed_point overflow of underlying representation type " + print_rep<Rep1>(),
-          !subtraction_overflow<Rep1>(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value)));
+  release_assert(
+    !subtraction_overflow<Rep1>(lhs.rescaled(scale)._value, rhs.rescaled(scale)._value) &&
+    "fixed_point overflow of underlying representation type " + print_rep<Rep1>());
 
 #endif
 
@@ -674,8 +675,8 @@ CUDA_HOST_DEVICE_CALLABLE fixed_point<Rep1, Rad1> operator*(fixed_point<Rep1, Ra
 {
 #if defined(__CUDACC_DEBUG__)
 
-  assert(("fixed_point overflow of underlying representation type " + print_rep<Rep1>(),
-          !multiplication_overflow<Rep1>(lhs._value, rhs._value)));
+  release_assert(!multiplication_overflow<Rep1>(lhs._value, rhs._value) &&
+                 "fixed_point overflow of underlying representation type " + print_rep<Rep1>());
 
 #endif
 
@@ -690,8 +691,8 @@ CUDA_HOST_DEVICE_CALLABLE fixed_point<Rep1, Rad1> operator/(fixed_point<Rep1, Ra
 {
 #if defined(__CUDACC_DEBUG__)
 
-  assert(("fixed_point overflow of underlying representation type " + print_rep<Rep1>(),
-          !division_overflow<Rep1>(lhs._value, rhs._value)));
+  release_assert(!division_overflow<Rep1>(lhs._value, rhs._value) &&
+                 "fixed_point overflow of underlying representation type " + print_rep<Rep1>());
 
 #endif
 

From 8cfbe1114c2885fb0229052a8c69e3415a02a2e6 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Thu, 28 Jan 2021 13:27:18 -0500
Subject: [PATCH 21/21] Add CUDF_EXPPECTS

---
 cpp/src/binaryop/binaryop.cpp | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp
index 3d70946f098..aa7e300d715 100644
--- a/cpp/src/binaryop/binaryop.cpp
+++ b/cpp/src/binaryop/binaryop.cpp
@@ -430,6 +430,8 @@ std::unique_ptr<column> 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");
   if (op == binary_operator::TRUE_DIV)
@@ -531,6 +533,8 @@ std::unique_ptr<column> 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");
   if (op == binary_operator::TRUE_DIV)
@@ -632,6 +636,8 @@ std::unique_ptr<column> 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");
   if (op == binary_operator::TRUE_DIV)