[WIP] replace int_div_ceil with std::cuda::ceil

benchmarks/static_map/insert_or_apply_bench.cu

@@ -23,6 +23,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/std/cmath>
 #include <thrust/device_vector.h>
 #include <thrust/transform.h>
 
@@ -42,7 +43,9 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_insert_or_appl
   auto const occupancy    = state.get_float64("Occupancy");
   auto const multiplicity = state.get_int64("Multiplicity");
 
-  std::size_t const size = cuco::detail::int_div_ceil(num_keys, multiplicity) / occupancy;
+  std::size_t const size = static_cast<size_t>(cuda::std::ceil(static_cast<double>(num_keys) /
+                                                               static_cast<double>(multiplicity))) /
+                           occupancy;
 
   thrust::device_vector<Key> keys(num_keys);
 

include/cuco/detail/extent/extent.inl

@@ -18,12 +18,12 @@
 
 #include <cuco/detail/error.hpp>
 #include <cuco/detail/prime.hpp>  // TODO move to detail/extent/
-#include <cuco/detail/utility/math.cuh>
 #include <cuco/detail/utils.hpp>
 #include <cuco/probing_scheme.cuh>
 #include <cuco/storage.cuh>
 #include <cuco/utility/fast_int.cuh>
 
+#include <cuda/std/cmath>
 #include <cuda/std/type_traits>
 
 namespace cuco {
@@ -91,8 +91,9 @@ template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
     (static_cast<uint64_t>(std::numeric_limits<SizeType>::max()) < max_prime)
       ? std::numeric_limits<SizeType>::max()
       : static_cast<SizeType>(max_prime);
-  auto const size = cuco::detail::int_div_ceil(
-    std::max(static_cast<SizeType>(ext), static_cast<SizeType>(1)), CGSize * BucketSize);
+  auto const size = static_cast<SizeType>(cuda::std::ceil(
+    static_cast<double>(std::max(static_cast<SizeType>(ext), static_cast<SizeType>(1))) /
+    static_cast<double>(CGSize * BucketSize)));
   if (size > max_value) { CUCO_FAIL("Invalid input extent"); }
 
   if constexpr (N == dynamic_extent) {

include/cuco/detail/prime.hpp

@@ -16,7 +16,7 @@
 
 #pragma once
 
-#include <cuco/detail/utility/math.cuh>
+#include <cuda/std/cmath>
 
 #include <algorithm>
 #include <array>
@@ -20154,7 +20154,8 @@ constexpr T get_valid_capacity(T capacity) noexcept
     if constexpr (not uses_vector_load) { return cg_size; }
   }();
 
-  auto const c         = int_div_ceil(capacity, stride);
+  auto const c =
+    static_cast<T>(cuda::std::ceil(static_cast<double>(capacity) / static_cast<double>(stride)));
   auto const min_prime = std::lower_bound(primes.begin(), primes.end(), c);
   return *min_prime * stride;
 }

include/cuco/detail/utility/cuda.hpp

@@ -16,7 +16,8 @@
 #pragma once
 
 #include <cuco/detail/error.hpp>
-#include <cuco/detail/utility/math.cuh>
+
+#include <cuda/std/cmath>
 
 namespace cuco {
 namespace detail {
@@ -43,7 +44,8 @@ constexpr auto grid_size(index_type num,
                          int32_t stride     = default_stride(),
                          int32_t block_size = default_block_size()) noexcept
 {
-  return int_div_ceil(cg_size * num, stride * block_size);
+  return static_cast<int32_t>(
+    cuda::std::ceil(static_cast<double>(cg_size * num) / static_cast<double>(stride * block_size)));
 }
 
 /**