Move HLL strong types to HLL header

Author: PointKernel

include/cuco/detail/hyperloglog/hyperloglog_impl.cuh

@@ -19,9 +19,9 @@
 #include <cuco/detail/error.hpp>
 #include <cuco/detail/hyperloglog/finalizer.cuh>
 #include <cuco/detail/hyperloglog/kernels.cuh>
+#include <cuco/detail/utility/strong_type.cuh>
 #include <cuco/detail/utils.hpp>
 #include <cuco/hash_functions.cuh>
-#include <cuco/types.cuh>
 #include <cuco/utility/cuda_thread_scope.cuh>
 #include <cuco/utility/traits.hpp>
 
@@ -42,6 +42,9 @@
 #include <vector>
 
 namespace cuco::detail {
+CUCO_DEFINE_STRONG_TYPE(sketch_size_kb, double);
+CUCO_DEFINE_STRONG_TYPE(standard_deviation, double);
+CUCO_DEFINE_STRONG_TYPE(precision, int32_t);
 
 /**
  * @brief A GPU-accelerated utility for approximating the number of distinct items in a multiset.
@@ -84,9 +87,9 @@ class hyperloglog_impl {
   __host__ __device__ constexpr hyperloglog_impl(cuda::std::span<cuda::std::byte> sketch_span,
                                                  Hash const& hash)
     : hash_{hash},
-      precision_{cuda::std::countr_zero(
-        sketch_bytes(cuco::sketch_size_kb(static_cast<double>(sketch_span.size() / 1024.0))) /
-        sizeof(register_type))},
+      precision_{cuda::std::countr_zero(sketch_bytes(cuco::detail::sketch_size_kb(
+                                          static_cast<double>(sketch_span.size() / 1024.0))) /
+                                        sizeof(register_type))},
       sketch_{reinterpret_cast<register_type*>(sketch_span.data()),
               this->sketch_bytes() / sizeof(register_type)}
   {
@@ -520,7 +523,7 @@ class hyperloglog_impl {
    * @return The number of bytes required for the sketch
    */
   [[nodiscard]] __host__ __device__ static constexpr cuda::std::size_t sketch_bytes(
-    cuco::sketch_size_kb sketch_size_kb) noexcept
+    cuco::detail::sketch_size_kb sketch_size_kb) noexcept
   {
     // minimum precision is 4 or 64 bytes
     return cuda::std::max(
@@ -536,7 +539,7 @@ class hyperloglog_impl {
    * @return The number of bytes required for the sketch
    */
   [[nodiscard]] __host__ __device__ static constexpr cuda::std::size_t sketch_bytes(
-    cuco::standard_deviation standard_deviation) noexcept
+    cuco::detail::standard_deviation standard_deviation) noexcept
   {
     // implementation taken from
     // https://github.com/apache/spark/blob/6a27789ad7d59cd133653a49be0bb49729542abe/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/util/HyperLogLogPlusPlusHelper.scala#L43
@@ -561,7 +564,7 @@ class hyperloglog_impl {
    * @return The number of bytes required for the sketch
    */
   [[nodiscard]] __host__ __device__ static constexpr cuda::std::size_t sketch_bytes(
-    cuco::precision precision) noexcept
+    cuco::detail::precision precision) noexcept
   {
     // minimum precision is 4 or 64 bytes
     auto const clamped_precision =

include/cuco/hyperloglog.cuh

@@ -18,7 +18,6 @@
 #include <cuco/detail/storage/storage_base.cuh>
 #include <cuco/hash_functions.cuh>
 #include <cuco/hyperloglog_ref.cuh>
-#include <cuco/types.cuh>
 #include <cuco/utility/allocator.hpp>
 #include <cuco/utility/cuda_thread_scope.cuh>
 

include/cuco/hyperloglog_ref.cuh

@@ -17,14 +17,50 @@
 
 #include <cuco/detail/hyperloglog/hyperloglog_impl.cuh>
 #include <cuco/hash_functions.cuh>
-#include <cuco/types.cuh>
 #include <cuco/utility/cuda_thread_scope.cuh>
 
 #include <cuda/std/cstddef>
 #include <cuda/stream_ref>
 
 #include <cooperative_groups.h>
 
+namespace cuco {
+/**
+ * @brief A strong type wrapper for specifying the upper-bound sketch size of
+ * `cuco::hyperloglog(_ref)` in KB.
+ *
+ * @note Underlying type is `double`. Values can also be specified as literals, e.g., 64.3_KB.
+ */
+using sketch_size_kb = detail::sketch_size_kb;
+
+/**
+ * @brief A strong type wrapper for specifying the desired standard deviation for the cardinality
+ * estimate of `cuco::hyperloglog(_ref)`.
+ *
+ * @note Underlying type is `double`.
+ */
+using standard_deviation = detail::standard_deviation;
+
+/**
+ * @brief A strong type wrapper for specifying the HyperLogLog precision parameter of
+ * `cuco::hyperloglog(_ref)`.
+ *
+ * @note Underlying type is `int32_t`. Precision `p` determines the number of registers as `2^p`.
+ * Valid range is typically [4, 18].
+ */
+using precision = detail::precision;
+}  // namespace cuco
+
+__host__ __device__ constexpr cuco::sketch_size_kb operator""_KB(long double value)
+{
+  return cuco::sketch_size_kb{static_cast<double>(value)};
+}
+
+__host__ __device__ constexpr cuco::sketch_size_kb operator""_KB(unsigned long long int value)
+{
+  return cuco::sketch_size_kb{static_cast<double>(value)};
+}
+
 namespace cuco {
 /**
  * @brief A GPU-accelerated utility for approximating the number of distinct items in a multiset.

include/cuco/types.cuh

@@ -44,38 +44,4 @@ CUCO_DEFINE_TEMPLATE_STRONG_TYPE(empty_value);
  */
 CUCO_DEFINE_TEMPLATE_STRONG_TYPE(erased_key);
 
-/**
- * @brief A strong type wrapper `cuco::sketch_size_kb` for specifying the upper-bound sketch size of
- * `cuco::hyperloglog(_ref)` in KB.
- *
- * @note Values can also be specified as literals, e.g., 64.3_KB.
- */
-CUCO_DEFINE_STRONG_TYPE(sketch_size_kb, double);
-
-/**
- * @brief A strong type wrapper `cuco::standard_deviation` for specifying the desired standard
- * deviation for the cardinality estimate of `cuco::hyperloglog(_ref)`.
- */
-CUCO_DEFINE_STRONG_TYPE(standard_deviation, double);
-
-/**
- * @brief A strong type wrapper `cuco::precision` for specifying the HyperLogLog precision
- * parameter of `cuco::hyperloglog(_ref)`.
- *
- * @note Precision `p` determines the number of registers as `2^p`. Valid range is typically [4,
- * 18].
- */
-CUCO_DEFINE_STRONG_TYPE(precision, int32_t);
-
 }  // namespace cuco
-
-// User-defined literal operators for `cuco::sketch_size_KB`
-__host__ __device__ constexpr cuco::sketch_size_kb operator""_KB(long double value)
-{
-  return cuco::sketch_size_kb{static_cast<double>(value)};
-}
-
-__host__ __device__ constexpr cuco::sketch_size_kb operator""_KB(unsigned long long int value)
-{
-  return cuco::sketch_size_kb{static_cast<double>(value)};
-}