Merge remote-tracking branch 'origin/main' into improve-jaccard

Author: davidwendt

cpp/include/cudf/fixed_point/detail/floating_conversion.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -17,46 +17,6 @@
 namespace numeric {
 namespace detail {
 
-/**
- * @brief Determine the number of significant bits in an integer
- *
- * @tparam T Type of input integer value. Must be either uint32_t, uint64_t, or __uint128_t
- * @param value The integer whose bits are being counted
- * @return The number of significant bits: the # of bits - # of leading zeroes
- */
-template <typename T,
-          CUDF_ENABLE_IF(cuda::std::is_same_v<T, uint32_t> || cuda::std::is_same_v<T, uint64_t> ||
-                         cuda::std::is_same_v<T, __uint128_t>)>
-CUDF_HOST_DEVICE inline constexpr int count_significant_bits(T value)
-{
-#ifdef __CUDA_ARCH__
-  if constexpr (cuda::std::is_same_v<T, uint64_t>) {
-    return 64 - __clzll(static_cast<int64_t>(value));
-  } else if constexpr (cuda::std::is_same_v<T, uint32_t>) {
-    return 32 - __clz(static_cast<int32_t>(value));
-  } else if constexpr (cuda::std::is_same_v<T, __uint128_t>) {
-    // 128 bit type, must break up into high and low components
-    auto const high_bits = static_cast<int64_t>(value >> 64);
-    auto const low_bits  = static_cast<int64_t>(value);
-    return 128 - (__clzll(high_bits) + static_cast<int>(high_bits == 0) * __clzll(low_bits));
-  }
-#else
-  if constexpr (cuda::std::is_same_v<T, uint64_t>) {
-    return 64 - cuda::std::countl_zero(value);
-  } else if constexpr (cuda::std::is_same_v<T, uint32_t>) {
-    return 32 - cuda::std::countl_zero(value);
-  } else if constexpr (cuda::std::is_same_v<T, __uint128_t>) {
-    // 128 bit type, must break up into high and low components
-    auto const high_bits = static_cast<uint64_t>(value >> 64);
-    if (high_bits == 0) {
-      return 64 - cuda::std::countl_zero(static_cast<uint64_t>(value));
-    } else {
-      return 128 - cuda::std::countl_zero(high_bits);
-    }
-  }
-#endif
-}
-
 /**
  * @brief Helper struct for getting and setting the components of a floating-point value
  *
@@ -189,7 +149,7 @@ struct floating_converter {
 
       // Line-up denormal to same (understood) bit as normal numbers
       // This is so bit-shifting starts at the same bit index
-      auto const lineup_shift = num_significand_bits - count_significant_bits(significand);
+      auto const lineup_shift = num_significand_bits - cuda::std::bit_width(significand);
       significand <<= lineup_shift;
       floating_pow2 -= lineup_shift;
     } else {
@@ -954,7 +914,7 @@ CUDF_HOST_DEVICE inline auto shift_to_binary_pospow(DecimalRep decimal_rep, int
   // except our first operation is a multiply, so not quite that far
   // We are bit-shifting down, so we need extra bits on the low-side, which this has.
   static constexpr int shift_up_to = sizeof(ShiftingRep) * 8 - Constants::max_bits_shift;
-  int const shift_from             = count_significant_bits(decimal_rep);
+  int const shift_from             = cuda::std::bit_width(decimal_rep);
   int const num_init_bit_shift     = shift_up_to - shift_from;
   int pow2                         = -num_init_bit_shift;
 
@@ -1017,7 +977,7 @@ CUDF_HOST_DEVICE inline auto shift_to_binary_negpow(DecimalRep decimal_rep, int
   // so that we lose minimal information we need on intermediary right-shifts.
   // Note that since we're shifting 2s up, we need num_2s_shift_buffer_bits space on the high side
   static constexpr int shift_up_to = sizeof(ShiftingRep) * 8 - Constants::num_2s_shift_buffer_bits;
-  int const shift_from             = count_significant_bits(decimal_rep);
+  int const shift_from             = cuda::std::bit_width(decimal_rep);
   int const num_init_bit_shift     = shift_up_to - shift_from;
   int pow2                         = -num_init_bit_shift;
 

cpp/include/cudf/strings/detail/scatter.cuh

@@ -9,11 +9,10 @@
 #include <cudf/strings/strings_column_view.hpp>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/memory_resource.hpp>
-#include <cudf/utilities/span.hpp>
 
-#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/functional>
 #include <cuda/iterator>
 #include <cuda/std/iterator>
@@ -59,7 +58,7 @@ std::unique_ptr<column> scatter(SourceIterator begin,
   if (target.is_empty()) return make_empty_column(type_id::STRING);
 
   // create vector of string_view's to scatter into
-  rmm::device_uvector<string_view> target_vector =
+  auto target_vector =
     create_string_vector_from_column(target, stream, cudf::get_current_device_resource_ref());
 
   // this ensures empty strings are not mapped to nulls in the make_strings_column function
@@ -77,8 +76,7 @@ std::unique_ptr<column> scatter(SourceIterator begin,
                   target_vector.begin());
 
   // build the output column
-  auto sv_span = cudf::device_span<string_view const>(target_vector);
-  return make_strings_column(sv_span, string_view{nullptr, 0}, stream, mr);
+  return make_strings_column(target_vector, string_view{nullptr, 0}, stream, mr);
 }
 
 }  // namespace detail

cpp/include/cudf/strings/detail/utilities.hpp

@@ -11,8 +11,7 @@
 #include <cudf/utilities/memory_resource.hpp>
 #include <cudf/utilities/span.hpp>
 
-#include <rmm/device_uvector.hpp>
-
+#include <cuda/buffer>
 #include <cuda/stream>
 
 namespace CUDF_EXPORT cudf {
@@ -37,14 +36,14 @@ std::unique_ptr<column> create_offsets_child_column(int64_t chars_bytes,
                                                     rmm::device_async_resource_ref mr);
 
 /**
- * @brief Creates a string_view vector from a strings column.
+ * @brief Creates a string_view buffer from a strings column.
  *
  * @param strings Strings column instance.
  * @param stream CUDA stream used for device memory operations and kernel launches.
- * @param mr Device memory resource used to allocate the returned vector's device memory.
- * @return Device vector of string_views
+ * @param mr Device memory resource used to allocate the returned buffer's device memory.
+ * @return Device buffer of string_views
  */
-rmm::device_uvector<string_view> create_string_vector_from_column(
+cuda::device_buffer<string_view> create_string_vector_from_column(
   cudf::strings_column_view const strings,
   cuda::stream_ref stream,
   rmm::device_async_resource_ref mr);

cpp/src/binaryop/compiled/binary_ops.cuh

@@ -100,8 +100,10 @@ struct ops_wrapper {
   column_device_view const& rhs;
   bool const& is_lhs_scalar;
   bool const& is_rhs_scalar;
+  // With bit_width, nvcc 13.3 can outline decimal128 conversions and add local memory setup
+  // before type dispatch, slowing float ADD. Keep this wrapper inline to avoid that setup.
   template <typename TypeCommon>
-  __device__ void operator()(size_type i)
+  __forceinline__ __device__ void operator()(size_type i)
   {
     if constexpr (cuda::std::is_invocable_v<BinaryOperator, TypeCommon, TypeCommon>) {
       TypeCommon x =

cpp/src/bitmask/null_mask.cu

@@ -238,14 +238,14 @@ void set_null_masks(cudf::host_span<bitmask_type*> bitmasks,
   average_nullmask_words +=
     cudf::util::div_rounding_up_safe<size_t>(cumulative_null_mask_words, num_bitmasks);
 
-  // Create device vectors from host spans
+  // Create device buffers from host spans
   auto const mr           = cudf::get_current_device_resource_ref();
-  auto destinations       = cudf::detail::make_device_uvector_async(bitmasks, stream, mr);
-  auto const d_begin_bits = cudf::detail::make_device_uvector_async(begin_bits, stream, mr);
-  auto const d_end_bits   = cudf::detail::make_device_uvector_async(end_bits, stream, mr);
-  auto const d_valids     = cudf::detail::make_device_uvector_async(valids, stream, mr);
+  auto destinations       = cudf::detail::make_device_buffer_async(bitmasks, stream, mr);
+  auto const d_begin_bits = cudf::detail::make_device_buffer_async(begin_bits, stream, mr);
+  auto const d_end_bits   = cudf::detail::make_device_buffer_async(end_bits, stream, mr);
+  auto const d_valids     = cudf::detail::make_device_buffer_async(valids, stream, mr);
   auto const number_of_mask_words =
-    cudf::detail::make_device_uvector_async(h_number_of_mask_words, stream, mr);
+    cudf::detail::make_device_buffer_async(h_number_of_mask_words, stream, mr);
 
   // Compute block size using heuristic and launch kernel
   constexpr size_t max_words_per_thread  = 64;
@@ -499,7 +499,7 @@ std::vector<size_type> batch_count_set_bits(host_span<bitmask_type const* const>
   if (!has_bitmask) { return output; }
 
   auto const tmp_mr     = cudf::get_current_device_resource_ref();
-  auto const d_bitmasks = cudf::detail::make_device_uvector_async(bitmasks, stream, tmp_mr);
+  auto const d_bitmasks = cudf::detail::make_device_buffer_async(bitmasks, stream, tmp_mr);
   auto const num_words  = num_bitmask_words(num_bits_to_count);
   auto d_non_zero_count =
     cudf::detail::make_zeroed_device_buffer_async<size_type>(num_bitmasks, stream, tmp_mr);

cpp/src/column/column_factories.cu

@@ -106,7 +106,6 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::stru
   cuda::stream_ref stream,
   rmm::device_async_resource_ref mr) const
 {
-  CUDF_EXPECTS(size != 0, "0-length struct column is unsupported.");
   auto& ss  = static_cast<scalar_type_t<cudf::struct_view> const&>(value);
   auto iter = cuda::make_constant_iterator(0);
 

cpp/src/copying/concatenate.cu

@@ -12,6 +12,7 @@
 #include <cudf/detail/get_value.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/utilities/buffer_factories.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/cuda.hpp>
 #include <cudf/detail/utilities/cuda_memcpy.hpp>
@@ -74,7 +75,7 @@ auto create_device_views(host_span<column_view const> views, cuda::stream_ref st
                  [](auto const& col) { return *col; });
 
   auto d_views =
-    make_device_uvector_async(device_views, stream, cudf::get_current_device_resource_ref());
+    make_device_buffer_async(device_views, stream, cudf::get_current_device_resource_ref());
 
   // Compute the partition offsets
   auto offsets = cudf::detail::make_pinned_vector_async<size_t>(views.size() + 1, stream);
@@ -86,7 +87,7 @@ auto create_device_views(host_span<column_view const> views, cuda::stream_ref st
     [](auto const& col) { return col.size(); },
     cuda::std::plus{});
   auto d_offsets =
-    make_device_uvector_async(offsets, stream, cudf::get_current_device_resource_ref());
+    make_device_buffer_async(offsets, stream, cudf::get_current_device_resource_ref());
   auto const output_size = offsets.back();
   cudf::detail::sync_stream(stream);
 

cpp/src/groupby/hash/hash_csr_kernels.cuh

@@ -16,6 +16,8 @@
 
 #include <cooperative_groups.h>
 #include <cuda/atomic>
+#include <cuda/cmath>
+#include <cuda/std/bit>
 #include <cuda/std/cstdint>
 #include <cuda/std/limits>
 #include <cuda/std/utility>
@@ -71,7 +73,7 @@ struct hash_set_ref {
   {
     representative = cudf::detail::CUDF_SIZE_TYPE_SENTINEL;
     // Multiply-high reduces the hash to any capacity without division or modulo.
-    auto slot = __umulhi(__brev(hash_value), capacity);
+    auto slot = cuda::mul_hi(cuda::std::bit_reverse(hash_value), capacity);
     for (cuda::std::uint32_t step = 0; step < max_probes; ++step) {
       auto slot_ref = cuda::atomic_ref<slot_type, cuda::thread_scope_device>{slots[slot]};
       auto current  = slot_ref.load(cuda::memory_order_relaxed);

cpp/src/groupby/sort/group_quantiles.cu

@@ -11,13 +11,12 @@
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/detail/device_scalar.hpp>
-#include <cudf/detail/utilities/vector_factories.hpp>
+#include <cudf/detail/utilities/buffer_factories.hpp>
 #include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 #include <cudf/utilities/span.hpp>
 
-#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <cuda/iterator>
@@ -156,7 +155,7 @@ std::unique_ptr<column> group_quantiles(column_view const& values,
                                         cuda::stream_ref stream,
                                         rmm::device_async_resource_ref mr)
 {
-  auto dv_quantiles = cudf::detail::make_device_uvector_async(
+  auto dv_quantiles = cudf::detail::make_device_buffer_async(
     quantiles, stream, cudf::get_current_device_resource_ref());
 
   auto values_type = cudf::is_dictionary(values.type())

cpp/src/io/comp/debrotli.cu

@@ -55,6 +55,7 @@ THE SOFTWARE.
 #include <cudf/utilities/memory_resource.hpp>
 
 #include <cuda/buffer>
+#include <cuda/std/bit>
 #include <cuda/stream>
 
 namespace cudf::io::detail {
@@ -78,11 +79,6 @@ inline __device__ uint32_t brotli_distance_alphabet_size(uint8_t npostfix,
   return brotli_num_distance_short_codes + ndirect + (maxnbits << (npostfix + 1));
 }
 
-inline __device__ uint32_t brev8(uint32_t x)
-{
-  return (__brev(x) >> 24u);  // kReverseBits[x]
-}
-
 #define CONSTANT static const __device__ __constant__
 #include "brotli_tables.hpp"
 
@@ -190,8 +186,6 @@ struct debrotli_state_s {
   __align__(8) char heap[local_heap_size];
 };
 
-inline __device__ uint32_t Log2Floor(uint32_t value) { return 32 - __clz(value); }
-
 /// @brief Safely read up to 4 bytes with bounds checking, handling end-of-buffer cases
 inline __device__ uint32_t safe_load_u32(uint8_t const* p, uint8_t const* end)
 {
@@ -615,7 +609,7 @@ static __device__ void BuildCodeLengthsHuffmanTable(huff_scratch_s* hs)
     for (int bits_count = hs->code_length_histo[bits]; bits_count != 0; --bits_count) {
       int end     = table_size;
       code        = huffcode(bits, hs->sorted[symbol++]);
-      uint16_t* p = &hs->lenvlctab[brev8(key)];
+      uint16_t* p = &hs->lenvlctab[cuda::std::bit_reverse(static_cast<uint8_t>(key))];
       do {
         end -= step;
         p[end] = code;
@@ -687,7 +681,7 @@ static __device__ uint32_t BuildHuffmanTable(uint16_t* root_lut,
     for (int bits_count = count[bits]; bits_count != 0; --bits_count) {
       symbol      = symbol_lists[symbol];
       code        = huffcode(bits, symbol);
-      uint16_t* p = &lut[brev8(key)];
+      uint16_t* p = &lut[cuda::std::bit_reverse(static_cast<uint8_t>(key))];
       int end     = table_size;
       do {
         end -= step;
@@ -717,15 +711,15 @@ static __device__ uint32_t BuildHuffmanTable(uint16_t* root_lut,
         table_bits = NextTableBitSize(count, len, root_bits);
         table_size = 1 << table_bits;
         total_size += table_size;
-        sub_key = brev8(key);
+        sub_key = cuda::std::bit_reverse(static_cast<uint8_t>(key));
         key += key_step;
         root_lut[sub_key] =
           huffcode(table_bits + root_bits, (((size_t)(lut - root_lut)) - sub_key));
         sub_key = 0;
       }
       symbol      = symbol_lists[symbol];
       code        = huffcode(len - root_bits, symbol);
-      uint16_t* p = &lut[brev8(sub_key)];
+      uint16_t* p = &lut[cuda::std::bit_reverse(static_cast<uint8_t>(sub_key))];
       int end     = table_size;
       do {
         end -= step;
@@ -904,7 +898,7 @@ static __device__ uint32_t DecodeHuffmanTree(debrotli_state_s* s,
   if (prefix_code_type == 1) {
     // Simple prefix Code
     uint32_t nsym          = getbits(s, 2);
-    uint32_t alphabet_bits = Log2Floor(alphabet_size - 1);
+    uint32_t alphabet_bits = cuda::std::bit_width(alphabet_size - 1);
     for (uint32_t n = 0; n <= nsym; n++) {
       uint32_t v = getbits(s, alphabet_bits);
       if (v >= max_symbol) {