Support for MOD, PMOD and PYMOD for decimal32/64/128 (#10179)

Resolves `libcudf` `MOD` and `PMOD` portion of #7132

To Do:
* [x] `MOD`
* [x] `PMOD`
* [x] `PYMOD`
* [x] Unit tests
* [x] Different scaled tests

Authors:
  - Conor Hoekstra (https://github.com/codereport)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Robert Maynard (https://github.com/robertmaynard)

URL: #10179

cpp/include/cudf/binaryop.hpp

@@ -45,7 +45,7 @@ enum class binary_operator : int32_t {
   PMOD,                  ///< positive modulo operator
                          ///< If remainder is negative, this returns (remainder + divisor) % divisor
                          ///< else, it returns (dividend % divisor)
-  PYMOD,                 ///< operator % but following python's sign rules for negatives
+  PYMOD,                 ///< operator % but following Python's sign rules for negatives
   POW,                   ///< lhs ^ rhs
   LOG_BASE,              ///< logarithm to the base
   ATAN2,                 ///< 2-argument arctangent

cpp/include/cudf/fixed_point/fixed_point.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -440,6 +440,21 @@ class fixed_point {
   CUDF_HOST_DEVICE inline friend fixed_point<Rep1, Rad1> operator/(
     fixed_point<Rep1, Rad1> const& lhs, fixed_point<Rep1, Rad1> const& rhs);
 
+  /**
+   * @brief operator % (for computing the modulo operation of two `fixed_point` numbers)
+   *
+   * If `_scale`s are equal, the modulus is computed directly.
+   * If `_scale`s are not equal, the number with larger `_scale` is shifted to the
+   * smaller `_scale`, and then the modulus is computed.
+   *
+   * @tparam Rep1 Representation type of number being modulo-ed to `this`
+   * @tparam Rad1 Radix (base) type of number being modulo-ed to `this`
+   * @return The resulting `fixed_point` number
+   */
+  template <typename Rep1, Radix Rad1>
+  CUDF_HOST_DEVICE inline friend fixed_point<Rep1, Rad1> operator%(
+    fixed_point<Rep1, Rad1> const& lhs, fixed_point<Rep1, Rad1> const& rhs);
+
   /**
    * @brief operator == (for comparing two `fixed_point` numbers)
    *
@@ -750,6 +765,16 @@ CUDF_HOST_DEVICE inline bool operator>(fixed_point<Rep1, Rad1> const& lhs,
   return lhs.rescaled(scale)._value > rhs.rescaled(scale)._value;
 }
 
+// MODULO OPERATION
+template <typename Rep1, Radix Rad1>
+CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> operator%(fixed_point<Rep1, Rad1> const& lhs,
+                                                          fixed_point<Rep1, Rad1> const& rhs)
+{
+  auto const scale     = std::min(lhs._scale, rhs._scale);
+  auto const remainder = lhs.rescaled(scale)._value % rhs.rescaled(scale)._value;
+  return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{remainder, scale}};
+}
+
 using decimal32  = fixed_point<int32_t, Radix::BASE_10>;
 using decimal64  = fixed_point<int64_t, Radix::BASE_10>;
 using decimal128 = fixed_point<__int128_t, Radix::BASE_10>;

cpp/src/binaryop/binaryop.cpp

@@ -88,7 +88,10 @@ bool is_basic_arithmetic_binop(binary_operator op)
          op == binary_operator::MUL or       // operator *
          op == binary_operator::DIV or       // operator / using common type of lhs and rhs
          op == binary_operator::NULL_MIN or  // 2 null = null, 1 null = value, else min
-         op == binary_operator::NULL_MAX;    // 2 null = null, 1 null = value, else max
+         op == binary_operator::NULL_MAX or  // 2 null = null, 1 null = value, else max
+         op == binary_operator::MOD or       // operator %
+         op == binary_operator::PMOD or      // positive modulo operator
+         op == binary_operator::PYMOD;  // operator % but following Python's negative sign rules
 }
 
 /**

cpp/src/binaryop/compiled/operation.cuh

@@ -162,12 +162,24 @@ struct PMod {
     if (rem < 0) rem = std::fmod(rem + yconv, yconv);
     return rem;
   }
+
+  template <typename TypeLhs,
+            typename TypeRhs,
+            std::enable_if_t<cudf::is_fixed_point<TypeLhs>() and
+                             std::is_same_v<TypeLhs, TypeRhs>>* = nullptr>
+  __device__ inline auto operator()(TypeLhs x, TypeRhs y)
+  {
+    auto const remainder = x % y;
+    return remainder.value() < 0 ? (remainder + y) % y : remainder;
+  }
 };
 
 struct PyMod {
   template <typename TypeLhs,
             typename TypeRhs,
-            std::enable_if_t<(std::is_integral_v<std::common_type_t<TypeLhs, TypeRhs>>)>* = nullptr>
+            std::enable_if_t<(std::is_integral_v<std::common_type_t<TypeLhs, TypeRhs>> or
+                              (cudf::is_fixed_point<TypeLhs>() and
+                               std::is_same_v<TypeLhs, TypeRhs>))>* = nullptr>
   __device__ inline auto operator()(TypeLhs x, TypeRhs y) -> decltype(((x % y) + y) % y)
   {
     return ((x % y) + y) % y;