Dropped duplicated math function from Thrust (#6188)

Co-authored-by: Michael Schellenberger Costa <[email protected]>
Co-authored-by: Federico Busato <[email protected]>

Author: 3 people

cub/cub/device/dispatch/dispatch_merge_sort.cuh

@@ -45,9 +45,8 @@
 #include <cub/util_namespace.cuh>
 #include <cub/util_vsmem.cuh>
 
-#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 +281,7 @@ struct DispatchMergeSort
         break;
       }
 
-      const int num_passes = static_cast<int>(THRUST_NS_QUALIFIER::detail::log2_ri(num_tiles));
+      const int num_passes = ::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

@@ -28,11 +28,12 @@
 
 #include <thrust/detail/allocator/allocator_traits.h>
 #include <thrust/detail/execute_with_allocator_fwd.h>
-#include <thrust/detail/integer_math.h>
 #include <thrust/detail/raw_pointer_cast.h>
 #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 +51,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));
+  const size_type num_elements = static_cast<size_type>(::cuda::ceil_div(sizeof(T) * n, sizeof(value_type)));
 
   void_pointer ptr = alloc_traits::allocate(system.get_allocator(), num_elements);
 
@@ -69,7 +70,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

@@ -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::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)
@@ -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/disjoint_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.h

@@ -32,8 +32,8 @@
 #endif // no system header
 
 #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 +104,15 @@ struct pool_options
    */
   bool validate() const
   {
-    if (!detail::is_power_of_2(smallest_block_size))
+    if (smallest_block_size != 0 && !::cuda::is_power_of_two(smallest_block_size))
     {
       return false;
     }
-    if (!detail::is_power_of_2(largest_block_size))
+    if (largest_block_size != 0 && !::cuda::is_power_of_two(largest_block_size))
     {
       return false;
     }
-    if (!detail::is_power_of_2(alignment))
+    if (alignment != 0 && !::cuda::is_power_of_two(alignment))
     {
       return false;
     }

thrust/thrust/mr/pool_options.h

@@ -26,16 +26,39 @@
 #  pragma system_header
 #endif // no system header
 
-#include <thrust/detail/integer_math.h>
 #include <thrust/detail/preprocessor.h>
 
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__type_traits/is_arithmetic.h>
+#include <cuda/std/__type_traits/is_unsigned.h>
 #include <cuda/std/cstdint>
 #include <cuda/std/detail/libcxx/include/stdexcept>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
 
 #include <string>
 
+THRUST_NAMESPACE_BEGIN
+namespace detail
+{
+
+_CCCL_TEMPLATE(typename T)
+_CCCL_REQUIRES(::cuda::std::is_arithmetic_v<T>)
+[[nodiscard]] _CCCL_API constexpr bool is_negative([[maybe_unused]] T x) noexcept
+{
+  if constexpr (::cuda::std::is_unsigned_v<T>)
+  {
+    return false;
+  }
+  else
+  {
+    return x < 0;
+  }
+}
+
+} // namespace detail
+THRUST_NAMESPACE_END
+
 #if defined(THRUST_FORCE_32_BIT_OFFSET_TYPE) && defined(THRUST_FORCE_64_BIT_OFFSET_TYPE)
 #  error "Only THRUST_FORCE_32_BIT_OFFSET_TYPE or THRUST_FORCE_64_BIT_OFFSET_TYPE may be defined!"
 #endif // THRUST_FORCE_32_BIT_OFFSET_TYPE && THRUST_FORCE_64_BIT_OFFSET_TYPE