Dropped duplicated math function from Thrust

cub/cub/device/dispatch/dispatch_merge_sort.cuh

@@ -48,6 +48,7 @@
 #include <thrust/detail/integer_math.h>
 
 #include <cuda/__cmath/ceil_div.h>
+#include <cuda/__cmath/ilog.h>
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__type_traits/is_same.h>
 #include <cuda/std/cstdint>
@@ -282,7 +283,7 @@ struct DispatchMergeSort
         break;
       }
 
-      const int num_passes = static_cast<int>(THRUST_NS_QUALIFIER::detail::log2_ri(num_tiles));
+      const int num_passes = static_cast<int>(::cuda::ceil_ilog2(num_tiles));
 
       /*
        * The algorithm consists of stages. At each stage, there are input and output arrays. There are two pairs of

thrust/thrust/detail/execute_with_allocator.h

@@ -33,6 +33,8 @@
 #include <thrust/detail/type_traits/pointer_traits.h>
 #include <thrust/pair.h>
 
+#include <cuda/__cmath/ceil_div.h>
+
 THRUST_NAMESPACE_BEGIN
 
 namespace detail
@@ -50,7 +52,7 @@ get_temporary_buffer(thrust::detail::execute_with_allocator<Allocator, BaseSyste
 
   // How many elements of type value_type do we need to accommodate n elements
   // of type T?
-  size_type num_elements = divide_ri(sizeof(T) * n, sizeof(value_type));
+  size_type num_elements = ::cuda::ceil_div(sizeof(T) * n, sizeof(value_type));
 
   void_pointer ptr = alloc_traits::allocate(system.get_allocator(), num_elements);
 
@@ -69,7 +71,7 @@ _CCCL_HOST void return_temporary_buffer(
   using value_type      = typename alloc_traits::value_type;
   using T               = typename thrust::detail::pointer_traits<Pointer>::element_type;
 
-  size_type num_elements = divide_ri(sizeof(T) * n, sizeof(value_type));
+  size_type num_elements = ::cuda::ceil_div(sizeof(T) * n, sizeof(value_type));
 
   pointer to_ptr = thrust::reinterpret_pointer_cast<pointer>(p);
   alloc_traits::deallocate(system.get_allocator(), to_ptr, num_elements);

thrust/thrust/detail/integer_math.h

@@ -27,6 +27,9 @@
 #endif // no system header
 #include <thrust/detail/type_deduction.h>
 
+#include <cuda/__cmath/ceil_div.h>
+#include <cuda/__cmath/round_down.h>
+#include <cuda/__cmath/round_up.h>
 #include <cuda/std/__bit/countl.h>
 #include <cuda/std/__type_traits/make_unsigned.h>
 #include <cuda/std/limits>
@@ -38,12 +41,6 @@ THRUST_NAMESPACE_BEGIN
 namespace detail
 {
 
-template <typename Integer>
-_CCCL_HOST_DEVICE _CCCL_FORCEINLINE bool is_power_of_2(Integer x)
-{
-  return 0 == (x & (x - 1));
-}
-
 template <typename T>
 _CCCL_HOST_DEVICE _CCCL_FORCEINLINE typename std::enable_if<std::is_signed<T>::value, bool>::type is_negative(T x)
 {
@@ -61,53 +58,6 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE bool is_odd(Integer x)
 {
   return 1 & x;
 }
-
-template <typename Integer>
-_CCCL_HOST_DEVICE _CCCL_FORCEINLINE Integer log2(Integer x)
-{
-  Integer num_bits           = 8 * sizeof(Integer);
-  Integer num_bits_minus_one = num_bits - 1;
-
-  return num_bits_minus_one - ::cuda::std::countl_zero(::cuda::std::__to_unsigned_like(x));
-}
-
-template <typename Integer>
-_CCCL_HOST_DEVICE _CCCL_FORCEINLINE Integer log2_ri(Integer x)
-{
-  Integer result = log2(x);
-
-  // This is where we round up to the nearest log.
-  if (!is_power_of_2(x))
-  {
-    ++result;
-  }
-
-  return result;
-}
-
-// x/y rounding towards +infinity for integers
-// Used to determine # of blocks/warps etc.
-template <typename Integer0, typename Integer1>
-_CCCL_HOST_DEVICE _CCCL_FORCEINLINE
-  // FIXME: Should use common_type.
-  auto
-  divide_ri(Integer0 const x, Integer1 const y) THRUST_DECLTYPE_RETURNS((x + (y - 1)) / y)
-
-  // x/y rounding towards zero for integers.
-  // Used to determine # of blocks/warps etc.
-  template <typename Integer0, typename Integer1>
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto divide_rz(Integer0 const x, Integer1 const y) THRUST_DECLTYPE_RETURNS(x / y)
-
-  // Round x towards infinity to the next multiple of y.
-  template <typename Integer0, typename Integer1>
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto round_i(Integer0 const x, Integer1 const y)
-    THRUST_DECLTYPE_RETURNS(y* divide_ri(x, y))
-
-  // Round x towards 0 to the next multiple of y.
-  template <typename Integer0, typename Integer1>
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE auto round_z(Integer0 const x, Integer1 const y)
-    THRUST_DECLTYPE_RETURNS(y* divide_rz(x, y))
-
 } // namespace detail
 
 THRUST_NAMESPACE_END

thrust/thrust/mr/disjoint_pool.h

@@ -41,6 +41,8 @@
 #include <thrust/mr/memory_resource.h>
 #include <thrust/mr/pool_options.h>
 
+#include <cuda/__cmath/ilog.h>
+#include <cuda/__cmath/pow2.h>
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__algorithm/min.h>
 #include <cuda/std/__cccl/algorithm_wrapper.h>
@@ -123,7 +125,7 @@ class disjoint_unsynchronized_pool_resource final
       : m_upstream(upstream)
       , m_bookkeeper(bookkeeper)
       , m_options(options)
-      , m_smallest_block_log2(detail::log2_ri(m_options.smallest_block_size))
+      , m_smallest_block_log2(::cuda::ceil_ilog2(m_options.smallest_block_size))
       , m_pools(m_bookkeeper)
       , m_allocated(m_bookkeeper)
       , m_cached_oversized(m_bookkeeper)
@@ -133,7 +135,7 @@ class disjoint_unsynchronized_pool_resource final
 
     pointer_vector free(m_bookkeeper);
     pool p(free);
-    m_pools.resize(detail::log2_ri(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
+    m_pools.resize(::cuda::ceil_ilog2(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
   }
 
   // TODO: C++11: use delegating constructors
@@ -146,7 +148,7 @@ class disjoint_unsynchronized_pool_resource final
       : m_upstream(get_global_resource<Upstream>())
       , m_bookkeeper(get_global_resource<Bookkeeper>())
       , m_options(options)
-      , m_smallest_block_log2(detail::log2_ri(m_options.smallest_block_size))
+      , m_smallest_block_log2(::cuda::ceil_ilog2(m_options.smallest_block_size))
       , m_pools(m_bookkeeper)
       , m_allocated(m_bookkeeper)
       , m_cached_oversized(m_bookkeeper)
@@ -156,7 +158,7 @@ class disjoint_unsynchronized_pool_resource final
 
     pointer_vector free(m_bookkeeper);
     pool p(free);
-    m_pools.resize(detail::log2_ri(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
+    m_pools.resize(::cuda::ceil_ilog2(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
   }
 
   /*! Destructor. Releases all held memory to upstream.
@@ -370,7 +372,7 @@ class disjoint_unsynchronized_pool_resource final
   [[nodiscard]] void_ptr do_allocate_impl(std::size_t bytes, std::size_t alignment)
   {
     bytes = (std::max) (bytes, m_options.smallest_block_size);
-    assert(detail::is_power_of_2(alignment));
+    assert(::cuda::ceil_ilog2(alignment));
 
     // an oversized and/or overaligned allocation requested; needs to be allocated separately
     if (bytes > m_options.largest_block_size || alignment > m_options.alignment)
@@ -430,7 +432,7 @@ class disjoint_unsynchronized_pool_resource final
 
     // the request is NOT for oversized and/or overaligned memory
     // allocate a block from an appropriate bucket
-    std::size_t bytes_log2 = thrust::detail::log2_ri(bytes);
+    std::size_t bytes_log2 = ::cuda::ceil_ilog2(bytes);
     std::size_t pool_idx   = bytes_log2 - m_smallest_block_log2;
     pool& bucket           = m_pools[pool_idx];
 
@@ -482,7 +484,7 @@ class disjoint_unsynchronized_pool_resource final
   virtual void do_deallocate(void_ptr p, std::size_t n, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
   {
     n = (std::max) (n, m_options.smallest_block_size);
-    assert(detail::is_power_of_2(alignment));
+    assert(::cuda::is_power_of_two(alignment));
 
     // verify that the pointer is at least as aligned as claimed
     assert(reinterpret_cast<::cuda::std::intmax_t>(detail::pointer_traits<void_ptr>::get(p)) % alignment == 0);
@@ -511,7 +513,7 @@ class disjoint_unsynchronized_pool_resource final
     }
 
     // push the block to the front of the appropriate bucket's free list
-    std::size_t n_log2   = thrust::detail::log2_ri(n);
+    std::size_t n_log2   = ::cuda::ceil_ilog2(n);
     std::size_t pool_idx = n_log2 - m_smallest_block_log2;
     pool& bucket         = m_pools[pool_idx];
 

thrust/thrust/mr/pool.h

@@ -37,6 +37,8 @@
 #include <thrust/mr/memory_resource.h>
 #include <thrust/mr/pool_options.h>
 
+#include <cuda/__cmath/ilog.h>
+#include <cuda/__cmath/pow2.h>
 #include <cuda/std/__cccl/algorithm_wrapper.h>
 #include <cuda/std/cassert>
 #include <cuda/std/cstdint>
@@ -112,7 +114,7 @@ class unsynchronized_pool_resource final
   unsynchronized_pool_resource(Upstream* upstream, pool_options options = get_default_options())
       : m_upstream(upstream)
       , m_options(options)
-      , m_smallest_block_log2(detail::log2_ri(m_options.smallest_block_size))
+      , m_smallest_block_log2(::cuda::ceil_ilog2(m_options.smallest_block_size))
       , m_pools(upstream)
       , m_allocated()
       , m_oversized()
@@ -121,7 +123,7 @@ class unsynchronized_pool_resource final
     assert(m_options.validate());
 
     pool p = {block_descriptor_ptr(), 0};
-    m_pools.resize(detail::log2_ri(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
+    m_pools.resize(::cuda::ceil_ilog2(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
   }
 
   // TODO: C++11: use delegating constructors
@@ -133,7 +135,7 @@ class unsynchronized_pool_resource final
   unsynchronized_pool_resource(pool_options options = get_default_options())
       : m_upstream(get_global_resource<Upstream>())
       , m_options(options)
-      , m_smallest_block_log2(detail::log2_ri(m_options.smallest_block_size))
+      , m_smallest_block_log2(::cuda::ceil_ilog2(m_options.smallest_block_size))
       , m_pools(get_global_resource<Upstream>())
       , m_allocated()
       , m_oversized()
@@ -142,7 +144,7 @@ class unsynchronized_pool_resource final
     assert(m_options.validate());
 
     pool p = {block_descriptor_ptr(), 0};
-    m_pools.resize(detail::log2_ri(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
+    m_pools.resize(::cuda::ceil_ilog2(m_options.largest_block_size) - m_smallest_block_log2 + 1, p);
   }
 
   /*! Destructor. Releases all held memory to upstream.
@@ -260,7 +262,7 @@ class unsynchronized_pool_resource final
   do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
   {
     bytes = (std::max) (bytes, m_options.smallest_block_size);
-    assert(detail::is_power_of_2(alignment));
+    assert(::cuda::is_power_of_two(alignment));
 
     // an oversized and/or overaligned allocation requested; needs to be allocated separately
     if (bytes > m_options.largest_block_size || alignment > m_options.alignment)
@@ -371,7 +373,7 @@ class unsynchronized_pool_resource final
 
     // the request is NOT for oversized and/or overaligned memory
     // allocate a block from an appropriate bucket
-    std::size_t bytes_log2 = thrust::detail::log2_ri(bytes);
+    std::size_t bytes_log2 = ::cuda::ceil_ilog2(bytes);
     std::size_t bucket_idx = bytes_log2 - m_smallest_block_log2;
     pool& bucket           = thrust::raw_reference_cast(m_pools[bucket_idx]);
 
@@ -439,7 +441,7 @@ class unsynchronized_pool_resource final
   virtual void do_deallocate(void_ptr p, std::size_t n, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
   {
     n = (std::max) (n, m_options.smallest_block_size);
-    assert(detail::is_power_of_2(alignment));
+    assert(::cuda::is_power_of_two(alignment));
 
     // verify that the pointer is at least as aligned as claimed
     assert(reinterpret_cast<::cuda::std::intmax_t>(void_ptr_traits::get(p)) % alignment == 0);
@@ -504,7 +506,7 @@ class unsynchronized_pool_resource final
     }
 
     // push the block to the front of the appropriate bucket's free list
-    std::size_t n_log2     = thrust::detail::log2_ri(n);
+    std::size_t n_log2     = ::cuda::ceil_ilog2(n);
     std::size_t bucket_idx = n_log2 - m_smallest_block_log2;
     pool& bucket           = thrust::raw_reference_cast(m_pools[bucket_idx]);
 

thrust/thrust/mr/pool_options.h

@@ -34,6 +34,7 @@
 #include <thrust/detail/config/memory_resource.h>
 #include <thrust/detail/integer_math.h>
 
+#include <cuda/__cmath/pow2.h>
 #include <cuda/std/cstddef>
 
 THRUST_NAMESPACE_BEGIN
@@ -104,15 +105,15 @@ struct pool_options
    */
   bool validate() const
   {
-    if (!detail::is_power_of_2(smallest_block_size))
+    if (!::cuda::is_power_of_two(smallest_block_size))
     {
       return false;
     }
-    if (!detail::is_power_of_2(largest_block_size))
+    if (!::cuda::is_power_of_two(largest_block_size))
     {
       return false;
     }
-    if (!detail::is_power_of_2(alignment))
+    if (!::cuda::is_power_of_two(alignment))
     {
       return false;
     }