Cleanup libcudf detail/aggregation.hpp/.cuh (#18642)

The detail/aggregation headers are cleaned up a bit to hopefully allow follow on work to address issue #18568 .
No functionality is changed -- mostly some device code is moved from `aggregation.cuh` to `aggregation.cu`.

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Yunsong Wang (https://github.com/PointKernel)

URL: #18642

Author: davidwendt

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

@@ -17,20 +17,14 @@
 #pragma once
 
 #include <cudf/aggregation.hpp>
-#include <cudf/column/column_device_view.cuh>
-#include <cudf/column/column_view.hpp>
 #include <cudf/detail/aggregation/aggregation.hpp>
-#include <cudf/detail/utilities/assert.cuh>
-#include <cudf/detail/utilities/device_atomics.cuh>
+#include <cudf/detail/utilities/device_operators.cuh>
 #include <cudf/table/table_view.hpp>
 #include <cudf/utilities/span.hpp>
-#include <cudf/utilities/traits.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/fill.h>
-
 #include <type_traits>
 
 namespace cudf {
@@ -121,107 +115,5 @@ constexpr bool has_corresponding_operator()
   return !std::is_same_v<typename corresponding_operator<k>::type, void>;
 }
 
-/**
- * @brief Dispatched functor to initialize a column with the identity of an
- * aggregation operation.
- *
- * Given a type `T` and `aggregation kind k`, determines and sets the value of
- * each element of the passed column to the appropriate initial value for the
- * aggregation.
- *
- * The initial values set as per aggregation are:
- * SUM: 0
- * COUNT_VALID: 0 and VALID
- * COUNT_ALL:   0 and VALID
- * MIN: Max value of type `T`
- * MAX: Min value of type `T`
- * ARGMAX: `ARGMAX_SENTINEL`
- * ARGMIN: `ARGMIN_SENTINEL`
- *
- * Only works on columns of fixed-width types.
- */
-struct identity_initializer {
- private:
-  template <typename T, aggregation::Kind k>
-  static constexpr bool is_supported()
-  {
-    return cudf::is_fixed_width<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 or
-            k == aggregation::SUM_OF_SQUARES or k == aggregation::STD or
-            k == aggregation::VARIANCE or
-            (k == aggregation::PRODUCT and is_product_supported<T>()));
-  }
-
-  template <typename T, aggregation::Kind k>
-  std::enable_if_t<not std::is_same_v<corresponding_operator_t<k>, void>, T>
-  identity_from_operator()
-  {
-    using DeviceType = device_storage_type_t<T>;
-    return corresponding_operator_t<k>::template identity<DeviceType>();
-  }
-
-  template <typename T, aggregation::Kind k>
-  std::enable_if_t<std::is_same_v<corresponding_operator_t<k>, void>, T> identity_from_operator()
-  {
-    CUDF_FAIL("Unable to get identity/sentinel from device operator");
-  }
-
-  template <typename T, aggregation::Kind k>
-  T get_identity()
-  {
-    if (k == aggregation::ARGMAX || k == aggregation::ARGMIN) {
-      if constexpr (cudf::is_timestamp<T>())
-        return k == aggregation::ARGMAX ? T{typename T::duration(ARGMAX_SENTINEL)}
-                                        : T{typename T::duration(ARGMIN_SENTINEL)};
-      else {
-        using DeviceType = device_storage_type_t<T>;
-        return k == aggregation::ARGMAX ? static_cast<DeviceType>(ARGMAX_SENTINEL)
-                                        : static_cast<DeviceType>(ARGMIN_SENTINEL);
-      }
-    }
-    return identity_from_operator<T, k>();
-  }
-
- public:
-  template <typename T, aggregation::Kind k>
-  std::enable_if_t<is_supported<T, k>(), void> operator()(mutable_column_view const& col,
-                                                          rmm::cuda_stream_view stream)
-  {
-    using DeviceType = device_storage_type_t<T>;
-    thrust::fill(rmm::exec_policy(stream),
-                 col.begin<DeviceType>(),
-                 col.end<DeviceType>(),
-                 get_identity<DeviceType, k>());
-  }
-
-  template <typename T, aggregation::Kind k>
-  std::enable_if_t<not is_supported<T, k>(), void> operator()(mutable_column_view const& col,
-                                                              rmm::cuda_stream_view stream)
-  {
-    CUDF_FAIL("Unsupported aggregation for initializing values");
-  }
-};
-
-/**
- * @brief Initializes each column in a table with a corresponding identity value
- * of an aggregation operation.
- *
- * The `i`th column will be initialized with the identity value of the `i`th
- * aggregation operation in `aggs`.
- *
- * @throw cudf::logic_error if column type and corresponding agg are incompatible
- * @throw cudf::logic_error if column type is not fixed-width
- *
- * @param table The table of columns to initialize.
- * @param aggs A span of aggregation operations corresponding to the table
- * columns. The aggregations determine the identity value for each column.
- * @param stream CUDA stream used for device memory operations and kernel launches.
- */
-void initialize_with_identity(mutable_table_view& table,
-                              host_span<cudf::aggregation::Kind const> aggs,
-                              rmm::cuda_stream_view stream);
-
 }  // namespace detail
 }  // namespace cudf

cpp/include/cudf/detail/aggregation/aggregation.hpp

@@ -1716,5 +1716,24 @@ constexpr inline bool is_valid_aggregation()
  */
 bool is_valid_aggregation(data_type source, aggregation::Kind k);
 
+/**
+ * @brief Initializes each column in a table with a corresponding identity value
+ * of an aggregation operation.
+ *
+ * The `i`th column will be initialized with the identity value of the `i`th
+ * aggregation operation in `aggs`.
+ *
+ * @throw cudf::logic_error if column type and corresponding agg are incompatible
+ * @throw cudf::logic_error if column type is not fixed-width
+ *
+ * @param table The table of columns to initialize.
+ * @param aggs A span of aggregation operations corresponding to the table
+ * columns. The aggregations determine the identity value for each column.
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+void initialize_with_identity(mutable_table_view& table,
+                              host_span<cudf::aggregation::Kind const> aggs,
+                              rmm::cuda_stream_view stream);
+
 }  // namespace detail
 }  // namespace CUDF_EXPORT cudf

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

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -23,7 +23,6 @@
 
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/fixed_point/temporary.hpp>
-#include <cudf/scalar/scalar.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
@@ -128,23 +127,17 @@ struct DeviceMin {
     return numeric::detail::min(lhs, rhs);
   }
 
-  template <typename T,
-            std::enable_if_t<!std::is_same_v<T, cudf::string_view> && !cudf::is_dictionary<T>() &&
-                             !cudf::is_fixed_point<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() && !cudf::is_fixed_point<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
-    // chrono types do not have std::numeric_limits specializations and should use T::max()
-    // https://eel.is/c++draft/numeric.limits.general#6
-    if constexpr (cudf::is_chrono<T>()) {
-      return T::max();
-    } else if constexpr (cuda::std::numeric_limits<T>::has_infinity) {
+    if constexpr (cuda::std::numeric_limits<T>::has_infinity) {
       return cuda::std::numeric_limits<T>::infinity();
     } else {
       return cuda::std::numeric_limits<T>::max();
     }
   }
 
-  template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
 #ifndef __CUDA_ARCH__
@@ -155,14 +148,17 @@ struct DeviceMin {
     return cuda::std::numeric_limits<T>::max();
   }
 
-  // @brief identity specialized for string_view
-  template <typename T, std::enable_if_t<std::is_same_v<T, cudf::string_view>>* = nullptr>
+  // identity specialized for string_view and chrono types
+  // chrono types do not have std::numeric_limits specializations and should use T::max()
+  // https://eel.is/c++draft/numeric.limits.general#6
+  template <typename T,
+            CUDF_ENABLE_IF(cuda::std::is_same_v<T, cudf::string_view> || cudf::is_chrono<T>())>
   CUDF_HOST_DEVICE inline static constexpr T identity()
   {
-    return string_view::max();
+    return T::max();
   }
 
-  template <typename T, std::enable_if_t<cudf::is_dictionary<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_dictionary<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
     return static_cast<T>(T::max_value());
@@ -180,23 +176,17 @@ struct DeviceMax {
     return numeric::detail::max(lhs, rhs);
   }
 
-  template <typename T,
-            std::enable_if_t<!std::is_same_v<T, cudf::string_view> && !cudf::is_dictionary<T>() &&
-                             !cudf::is_fixed_point<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_numeric<T>() && !cudf::is_fixed_point<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
-    // chrono types do not have std::numeric_limits specializations and should use T::min()
-    // https://eel.is/c++draft/numeric.limits.general#6
-    if constexpr (cudf::is_chrono<T>()) {
-      return T::min();
-    } else if constexpr (cuda::std::numeric_limits<T>::has_infinity) {
+    if constexpr (cuda::std::numeric_limits<T>::has_infinity) {
       return -cuda::std::numeric_limits<T>::infinity();
     } else {
       return cuda::std::numeric_limits<T>::lowest();
     }
   }
 
-  template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_fixed_point<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
 #ifndef __CUDA_ARCH__
@@ -207,13 +197,17 @@ struct DeviceMax {
     return cuda::std::numeric_limits<T>::lowest();
   }
 
-  template <typename T, std::enable_if_t<std::is_same_v<T, cudf::string_view>>* = nullptr>
+  // identity specialized for string_view and chrono types
+  // chrono types do not have std::numeric_limits specializations and should use T::min()
+  // https://eel.is/c++draft/numeric.limits.general#6
+  template <typename T,
+            CUDF_ENABLE_IF(cuda::std::is_same_v<T, cudf::string_view> || cudf::is_chrono<T>())>
   CUDF_HOST_DEVICE inline static constexpr T identity()
   {
-    return string_view::min();
+    return T::min();
   }
 
-  template <typename T, std::enable_if_t<cudf::is_dictionary<T>()>* = nullptr>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_dictionary<T>())>
   CUDF_HOST_DEVICE static constexpr T identity()
   {
     return static_cast<T>(T::lowest_value());

cpp/src/aggregation/aggregation.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,13 +16,99 @@
 
 #include <cudf/detail/aggregation/aggregation.cuh>
 #include <cudf/detail/aggregation/aggregation.hpp>
-#include <cudf/table/table_view.hpp>
-#include <cudf/utilities/span.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/fill.h>
 
 namespace cudf {
 namespace detail {
+namespace {
+
+/**
+ * @brief Dispatched functor to initialize a column with the identity of an
+ * aggregation operation.
+ *
+ * Given a type `T` and `aggregation kind k`, determines and sets the value of
+ * each element of the passed column to the appropriate initial value for the
+ * aggregation.
+ *
+ * The initial values set as per aggregation are:
+ * SUM: 0
+ * COUNT_VALID: 0 and VALID
+ * COUNT_ALL:   0 and VALID
+ * MIN: Max value of type `T`
+ * MAX: Min value of type `T`
+ * ARGMAX: `ARGMAX_SENTINEL`
+ * ARGMIN: `ARGMIN_SENTINEL`
+ *
+ * Only works on columns of fixed-width types.
+ */
+struct identity_initializer {
+ private:
+  template <typename T, aggregation::Kind k>
+  static constexpr bool is_supported()
+  {
+    return cudf::is_fixed_width<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 or
+            k == aggregation::SUM_OF_SQUARES or k == aggregation::STD or
+            k == aggregation::VARIANCE or
+            (k == aggregation::PRODUCT and is_product_supported<T>()));
+  }
+
+  template <typename T, aggregation::Kind k>
+  std::enable_if_t<not std::is_same_v<corresponding_operator_t<k>, void>, T>
+  identity_from_operator()
+  {
+    using DeviceType = device_storage_type_t<T>;
+    return corresponding_operator_t<k>::template identity<DeviceType>();
+  }
+
+  template <typename T, aggregation::Kind k>
+  std::enable_if_t<std::is_same_v<corresponding_operator_t<k>, void>, T> identity_from_operator()
+  {
+    CUDF_FAIL("Unable to get identity/sentinel from device operator");
+  }
+
+  template <typename T, aggregation::Kind k>
+  T get_identity()
+  {
+    if (k == aggregation::ARGMAX || k == aggregation::ARGMIN) {
+      if constexpr (cudf::is_timestamp<T>())
+        return k == aggregation::ARGMAX ? T{typename T::duration(ARGMAX_SENTINEL)}
+                                        : T{typename T::duration(ARGMIN_SENTINEL)};
+      else {
+        using DeviceType = device_storage_type_t<T>;
+        return k == aggregation::ARGMAX ? static_cast<DeviceType>(ARGMAX_SENTINEL)
+                                        : static_cast<DeviceType>(ARGMIN_SENTINEL);
+      }
+    }
+    return identity_from_operator<T, k>();
+  }
+
+ public:
+  template <typename T, aggregation::Kind k>
+  std::enable_if_t<is_supported<T, k>(), void> operator()(mutable_column_view const& col,
+                                                          rmm::cuda_stream_view stream)
+  {
+    using DeviceType = device_storage_type_t<T>;
+    thrust::fill(rmm::exec_policy(stream),
+                 col.begin<DeviceType>(),
+                 col.end<DeviceType>(),
+                 get_identity<DeviceType, k>());
+  }
+
+  template <typename T, aggregation::Kind k>
+  std::enable_if_t<not is_supported<T, k>(), void> operator()(mutable_column_view const& col,
+                                                              rmm::cuda_stream_view stream)
+  {
+    CUDF_FAIL("Unsupported aggregation for initializing values");
+  }
+};
+}  // namespace
+
 void initialize_with_identity(mutable_table_view& table,
                               host_span<cudf::aggregation::Kind const> aggs,
                               rmm::cuda_stream_view stream)

cpp/src/groupby/hash/create_sparse_results_table.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,7 +20,7 @@
 
 #include <cudf/aggregation.hpp>
 #include <cudf/column/column_factories.hpp>
-#include <cudf/detail/aggregation/aggregation.cuh>
+#include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/span.hpp>