diff --git a/benchmarks/static_map/insert_or_apply_bench.cu b/benchmarks/static_map/insert_or_apply_bench.cu
index 4633a8b0f..ad668b6ac 100644
--- a/benchmarks/static_map/insert_or_apply_bench.cu
+++ b/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);
 
diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl
index bb5145c70..bdbb8f56a 100644
--- a/include/cuco/detail/extent/extent.inl
+++ b/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) {
diff --git a/include/cuco/detail/prime.hpp b/include/cuco/detail/prime.hpp
index 4ef0a35e3..26e876ab3 100644
--- a/include/cuco/detail/prime.hpp
+++ b/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;
 }
diff --git a/include/cuco/detail/utility/cuda.hpp b/include/cuco/detail/utility/cuda.hpp
index 49445e02f..7cde2d572 100644
--- a/include/cuco/detail/utility/cuda.hpp
+++ b/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)));
 }
 
 /**
diff --git a/include/cuco/detail/utility/math.cuh b/include/cuco/detail/utility/math.cuh
deleted file mode 100644
index c2715f6fa..000000000
--- a/include/cuco/detail/utility/math.cuh
+++ /dev/null
@@ -1,46 +0,0 @@
-/*
- * Copyright (c) 2023, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- */
-
-#pragma once
-
-#include <cuda/std/type_traits>
-
-namespace cuco {
-namespace detail {
-
-/**
- * @brief Ceiling of an integer division
- *
- * @tparam T Type of dividend
- * @tparam U Type of divisor
- *
- * @throw If `T` is not an integral type
- * @throw If `U` is not an integral type
- *
- * @param dividend Numerator
- * @param divisor Denominator
- *
- * @return Ceiling of the integer division
- */
-template <typename T, typename U>
-__host__ __device__ constexpr T int_div_ceil(T dividend, U divisor) noexcept
-{
-  static_assert(cuda::std::is_integral_v<T>);
-  static_assert(cuda::std::is_integral_v<U>);
-  return (dividend + divisor - 1) / divisor;
-}
-
-}  // namespace detail
-}  // namespace cuco