[REVIEW] Initial fixed_point Column Support

CHANGELOG.md

@@ -50,6 +50,7 @@
 - PR #5658 Add `filter_tokens` nvtext API
 - PR #5666 Add `filter_characters_of_type` strings API
 - PR #5673 Always build and test with per-thread default stream enabled in the GPU CI build
+- PR #5704 Initial `fixed_point` Column Support
 - PR #5716 Add `double_type_dispatcher` to libcudf
 - PR #5739 Add `nvtext::detokenize` API
 - PR #5645 Enforce pd.NA and Pandas nullable dtype parity

cpp/CMakeLists.txt

@@ -623,6 +623,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                           ${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit
                           ${CMAKE_BINARY_DIR}/include/bit.hpp.jit
                           ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
+                          ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit
                           ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
@@ -632,6 +633,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit
+                          ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit
+                          ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cfloat.jit
@@ -640,10 +643,13 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit
                           ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit
+                          ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit
+                          ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit
                    MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.h
                                    ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.hpp
                                    ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/utilities/bit.hpp
                                    ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/timestamps.hpp
+                                   ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/fixed_point/fixed_point.hpp
                                    ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/durations.hpp
                                    ${LIBCUDACXX_INCLUDE_DIR}/details/__config
                                    ${LIBCUDACXX_INCLUDE_DIR}/simt/limits
@@ -653,6 +659,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                                    ${LIBCUDACXX_INCLUDE_DIR}/simt/ratio
                                    ${LIBCUDACXX_INCLUDE_DIR}/simt/type_traits
                                    ${LIBCUDACXX_INCLUDE_DIR}/simt/version
+                                   ${LIBCUDACXX_INCLUDE_DIR}/simt/cmath
+                                   ${LIBCUDACXX_INCLUDE_DIR}/simt/cassert
                                    ${LIBCXX_INCLUDE_DIR}/__config
                                    ${LIBCXX_INCLUDE_DIR}/__undef_macros
                                    ${LIBCXX_INCLUDE_DIR}/cfloat
@@ -661,13 +669,16 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                                    ${LIBCXX_INCLUDE_DIR}/limits
                                    ${LIBCXX_INCLUDE_DIR}/ratio
                                    ${LIBCXX_INCLUDE_DIR}/type_traits
+                                   ${LIBCXX_INCLUDE_DIR}/cmath
+                                   ${LIBCXX_INCLUDE_DIR}/cassert
 
                    # stringified headers are placed underneath the bin include jit directory and end in ".jit"
                    COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/types.h > ${CMAKE_BINARY_DIR}/include/jit/types.h.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/types.hpp > ${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/utilities/bit.hpp > ${CMAKE_BINARY_DIR}/include/bit.hpp.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ../src/rolling/rolling_jit_detail.hpp > ${CMAKE_BINARY_DIR}/include/rolling_jit_detail.hpp.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/timestamps.hpp > ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/fixed_point/fixed_point.hpp > ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/durations.hpp > ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/details/__config libcudacxx_details_config > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/limits libcudacxx_simt_limits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
@@ -676,6 +687,8 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/ctime libcudacxx_simt_ctime > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ctime.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/ratio libcudacxx_simt_ratio > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/type_traits libcudacxx_simt_type_traits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cmath libcudacxx_simt_cmath > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cassert libcudacxx_simt_cassert > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/version libcudacxx_simt_version > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/__config libcxx_config > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/__undef_macros libcxx_undef_macros > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit
@@ -685,13 +698,16 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/limits libcxx_limits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/ratio libcxx_ratio > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit
                    COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/type_traits libcxx_type_traits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/cmath libcxx_cmath > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCXX_INCLUDE_DIR}/cassert libcxx_cassert > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit
                    )
 
 add_custom_target(stringify_run DEPENDS
                   ${CMAKE_BINARY_DIR}/include/jit/types.h.jit
                   ${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit
                   ${CMAKE_BINARY_DIR}/include/bit.hpp.jit
                   ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
+                  ${CMAKE_BINARY_DIR}/include/jit/fixed_point.hpp.jit
                   ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
@@ -700,6 +716,8 @@ add_custom_target(stringify_run DEPENDS
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ctime.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/ratio.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/type_traits.jit
+                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cmath.jit
+                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cassert.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/version.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__config.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/__undef_macros.jit
@@ -708,7 +726,10 @@ add_custom_target(stringify_run DEPENDS
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ctime.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/limits.jit
                   ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/ratio.jit
-                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit)
+                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/type_traits.jit
+                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cmath.jit
+                  ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/libcxx/include/cassert.jit
+                  )
 
 add_dependencies(cudf stringify_run)
 

cpp/include/cudf/column/column_factories.hpp

@@ -86,7 +86,7 @@ std::unique_ptr<column> make_numeric_column(
   cudaStream_t stream                 = 0,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource())
 {
-  CUDF_EXPECTS(is_numeric(type), "Invalid, non-numeric type.");
+  CUDF_EXPECTS(is_numeric(type) || is_fixed_point(type), "Invalid, non-numeric type.");
   return std::make_unique<column>(type,
                                   size,
                                   rmm::device_buffer{size * cudf::size_of(type), stream, mr},

cpp/include/cudf/detail/aggregation/aggregation.cuh

@@ -96,11 +96,12 @@ struct update_target_element {
 };
 
 template <typename Source, bool target_has_nulls, bool source_has_nulls>
-struct update_target_element<Source,
-                             aggregation::MIN,
-                             target_has_nulls,
-                             source_has_nulls,
-                             std::enable_if_t<is_fixed_width<Source>()>> {
+struct update_target_element<
+  Source,
+  aggregation::MIN,
+  target_has_nulls,
+  source_has_nulls,
+  std::enable_if_t<is_fixed_width<Source>() && !is_fixed_point<Source>()>> {
   __device__ void operator()(mutable_column_device_view target,
                              size_type target_index,
                              column_device_view source,
@@ -117,11 +118,12 @@ struct update_target_element<Source,
 };
 
 template <typename Source, bool target_has_nulls, bool source_has_nulls>
-struct update_target_element<Source,
-                             aggregation::MAX,
-                             target_has_nulls,
-                             source_has_nulls,
-                             std::enable_if_t<is_fixed_width<Source>()>> {
+struct update_target_element<
+  Source,
+  aggregation::MAX,
+  target_has_nulls,
+  source_has_nulls,
+  std::enable_if_t<is_fixed_width<Source>() && !is_fixed_point<Source>()>> {
   __device__ void operator()(mutable_column_device_view target,
                              size_type target_index,
                              column_device_view source,
@@ -138,11 +140,12 @@ struct update_target_element<Source,
 };
 
 template <typename Source, bool target_has_nulls, bool source_has_nulls>
-struct update_target_element<Source,
-                             aggregation::SUM,
-                             target_has_nulls,
-                             source_has_nulls,
-                             std::enable_if_t<is_fixed_width<Source>()>> {
+struct update_target_element<
+  Source,
+  aggregation::SUM,
+  target_has_nulls,
+  source_has_nulls,
+  std::enable_if_t<is_fixed_width<Source>() && !is_fixed_point<Source>()>> {
   __device__ void operator()(mutable_column_device_view target,
                              size_type target_index,
                              column_device_view source,
@@ -359,7 +362,7 @@ struct identity_initializer {
   template <typename T, aggregation::Kind k>
   static constexpr bool is_supported()
   {
-    return cudf::is_fixed_width<T>() and
+    return cudf::is_fixed_width<T>() && !is_fixed_point<T>() and
            (k == aggregation::SUM or k == aggregation::MIN or k == aggregation::MAX or
             k == aggregation::COUNT_VALID or k == aggregation::COUNT_ALL or
             k == aggregation::ARGMAX or k == aggregation::ARGMIN);

cpp/include/cudf/detail/reduction.cuh

@@ -79,9 +79,8 @@ std::unique_ptr<scalar> reduce(InputIterator d_in,
                             identity,
                             stream);
 
-  using ScalarType = cudf::scalar_type_t<OutputType>;
-  auto s           = new ScalarType(
-    std::move(dev_result), true, stream, mr);  // only for string_view, data is copied
+  // only for string_view, data is copied
+  auto s = new cudf::scalar_type_t<OutputType>(std::move(dev_result), true, stream, mr);
   return std::unique_ptr<scalar>(s);
 }
 

cpp/include/cudf/detail/utilities/device_operators.cuh

@@ -27,6 +27,10 @@
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/traits.hpp>
+
+// will fail to compile if grouped with the includes above
+#include <cudf/fixed_point/fixed_point.hpp>
+
 #include <type_traits>
 
 namespace cudf {
@@ -40,11 +44,18 @@ struct DeviceSum {
     return lhs + rhs;
   }
 
-  template <typename T>
+  template <typename T, typename std::enable_if_t<!cudf::is_fixed_point<T>()>* = nullptr>
   static constexpr T identity()
   {
     return T{0};
   }
+
+  template <typename T, typename std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
+  static constexpr T identity()
+  {
+    CUDF_FAIL("fixed_point does not yet support device operator identity");
+    return T{};
+  }
 };
 
 /* @brief `count` operator - used in rolling windows */
@@ -88,9 +99,17 @@ struct DeviceMin {
   }
 
   template <typename T,
-            typename std::enable_if_t<!std::is_same<T, cudf::string_view>::value>* = nullptr>
+            typename std::enable_if_t<!std::is_same<T, cudf::string_view>::value &&
+                                      !cudf::is_fixed_point<T>()>* = nullptr>
+  static constexpr T identity()
+  {
+    return std::numeric_limits<T>::max();
+  }
+
+  template <typename T, typename std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
   static constexpr T identity()
   {
+    CUDF_FAIL("fixed_point does not yet support DeviceMin identity");
     return std::numeric_limits<T>::max();
   }
 
@@ -118,11 +137,20 @@ struct DeviceMax {
   }
 
   template <typename T,
-            typename std::enable_if_t<!std::is_same<T, cudf::string_view>::value>* = nullptr>
+            typename std::enable_if_t<!std::is_same<T, cudf::string_view>::value &&
+                                      !cudf::is_fixed_point<T>()>* = nullptr>
+  static constexpr T identity()
+  {
+    return std::numeric_limits<T>::lowest();
+  }
+
+  template <typename T, typename std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
   static constexpr T identity()
   {
+    CUDF_FAIL("fixed_point does not yet support DeviceMax identity");
     return std::numeric_limits<T>::lowest();
   }
+
   template <typename T,
             typename std::enable_if_t<std::is_same<T, cudf::string_view>::value>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE static constexpr T identity()
@@ -145,11 +173,18 @@ struct DeviceProduct {
     return lhs * rhs;
   }
 
-  template <typename T>
+  template <typename T, typename std::enable_if_t<!cudf::is_fixed_point<T>()>* = nullptr>
   static constexpr T identity()
   {
     return T{1};
   }
+
+  template <typename T, typename std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
+  static constexpr T identity()
+  {
+    CUDF_FAIL("fixed_point does not yet support DeviceProduct identity");
+    return T{1, numeric::scale_type{0}};
+  }
 };
 
 /* @brief binary `and` operator */