Merge branch 'branch-25.02' into enh-nosync_write_json

.github/copy-pr-bot.yaml

@@ -2,3 +2,4 @@
 # https://docs.gha-runners.nvidia.com/apps/copy-pr-bot/
 
 enabled: true
+auto_sync_draft: false

.github/workflows/trigger-breaking-change-alert.yaml

@@ -0,0 +1,26 @@
+name: Trigger Breaking Change Notifications
+
+on:
+  pull_request_target:
+    types:
+      - closed
+      - reopened
+      - labeled
+      - unlabeled
+
+jobs:
+  trigger-notifier:
+    if: contains(github.event.pull_request.labels.*.name, 'breaking')
+    secrets: inherit
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
+    with:
+      sender_login: ${{ github.event.sender.login }}
+      sender_avatar: ${{ github.event.sender.avatar_url }}
+      repo: ${{ github.repository }}
+      pr_number: ${{ github.event.pull_request.number }}
+      pr_title: "${{ github.event.pull_request.title }}"
+      pr_body: "${{ github.event.pull_request.body || '_Empty PR description_' }}"
+      pr_base_ref: ${{ github.event.pull_request.base.ref }}
+      pr_author: ${{ github.event.pull_request.user.login }}
+      event_action: ${{ github.event.action }}
+      pr_merged: ${{ github.event.pull_request.merged }}

cpp/include/cudf/detail/utilities/cuda.cuh

@@ -16,7 +16,6 @@
 
 #pragma once
 
-#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/default_stream.hpp>
@@ -25,8 +24,7 @@
 #include <rmm/cuda_stream_view.hpp>
 
 #include <cub/cub.cuh>
-
-#include <type_traits>
+#include <cuda/std/type_traits>
 
 namespace cudf {
 namespace detail {
@@ -164,7 +162,7 @@ template <int32_t block_size, int32_t leader_lane = 0, typename T>
 __device__ T single_lane_block_sum_reduce(T lane_value)
 {
   static_assert(block_size <= 1024, "Invalid block size.");
-  static_assert(std::is_arithmetic_v<T>, "Invalid non-arithmetic type.");
+  static_assert(cuda::std::is_arithmetic_v<T>, "Invalid non-arithmetic type.");
   constexpr auto warps_per_block{block_size / warp_size};
   auto const lane_id{threadIdx.x % warp_size};
   auto const warp_id{threadIdx.x / warp_size};

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

@@ -85,15 +85,15 @@ CUDF_KERNEL void gather_chars_fn_string_parallel(StringIterator strings_begin,
   constexpr size_t out_datatype_size = sizeof(uint4);
   constexpr size_t in_datatype_size  = sizeof(uint);
 
-  int global_thread_id = blockIdx.x * blockDim.x + threadIdx.x;
-  int global_warp_id   = global_thread_id / cudf::detail::warp_size;
-  int warp_lane        = global_thread_id % cudf::detail::warp_size;
-  int nwarps           = gridDim.x * blockDim.x / cudf::detail::warp_size;
+  auto const global_thread_id = cudf::detail::grid_1d::global_thread_id();
+  auto const global_warp_id   = global_thread_id / cudf::detail::warp_size;
+  auto const warp_lane        = global_thread_id % cudf::detail::warp_size;
+  auto const nwarps           = cudf::detail::grid_1d::grid_stride() / cudf::detail::warp_size;
 
   auto const alignment_offset = reinterpret_cast<std::uintptr_t>(out_chars) % out_datatype_size;
   uint4* out_chars_aligned    = reinterpret_cast<uint4*>(out_chars - alignment_offset);
 
-  for (size_type istring = global_warp_id; istring < total_out_strings; istring += nwarps) {
+  for (auto istring = global_warp_id; istring < total_out_strings; istring += nwarps) {
     auto const out_start = out_offsets[istring];
     auto const out_end   = out_offsets[istring + 1];
 

cpp/src/binaryop/jit/kernel.cu

@@ -51,15 +51,10 @@ CUDF_KERNEL void kernel_v_v(cudf::size_type size,
                             TypeLhs* lhs_data,
                             TypeRhs* rhs_data)
 {
-  int tid    = threadIdx.x;
-  int blkid  = blockIdx.x;
-  int blksz  = blockDim.x;
-  int gridsz = gridDim.x;
+  auto const start = threadIdx.x + static_cast<cudf::thread_index_type>(blockIdx.x) * blockDim.x;
+  auto const step  = static_cast<cudf::thread_index_type>(blockDim.x) * gridDim.x;
 
-  int start = tid + blkid * blksz;
-  int step  = blksz * gridsz;
-
-  for (cudf::size_type i = start; i < size; i += step) {
+  for (auto i = start; i < size; i += step) {
     out_data[i] = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(lhs_data[i], rhs_data[i]);
   }
 }
@@ -75,15 +70,10 @@ CUDF_KERNEL void kernel_v_v_with_validity(cudf::size_type size,
                                           cudf::bitmask_type const* rhs_mask,
                                           cudf::size_type rhs_offset)
 {
-  int tid    = threadIdx.x;
-  int blkid  = blockIdx.x;
-  int blksz  = blockDim.x;
-  int gridsz = gridDim.x;
-
-  int start = tid + blkid * blksz;
-  int step  = blksz * gridsz;
+  auto const start = threadIdx.x + static_cast<cudf::thread_index_type>(blockIdx.x) * blockDim.x;
+  auto const step  = static_cast<cudf::thread_index_type>(blockDim.x) * gridDim.x;
 
-  for (cudf::size_type i = start; i < size; i += step) {
+  for (auto i = start; i < size; i += step) {
     bool output_valid = false;
     out_data[i]       = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(
       lhs_data[i],

cpp/src/copying/contiguous_split.cu

@@ -21,6 +21,7 @@
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/lists/lists_column_view.hpp>

cpp/src/io/json/write_json.cu

@@ -244,6 +244,7 @@ struct validity_fn {
  *
  * @param strings_columns Table of strings columns
  * @param column_names Column of names for each column in the table
+ * @param num_rows Number of rows in the table
  * @param row_prefix Prepend this string to each row
  * @param row_suffix  Append this string to each row
  * @param value_separator Separator between values
@@ -255,6 +256,7 @@ struct validity_fn {
  */
 std::unique_ptr<column> struct_to_strings(table_view const& strings_columns,
                                           column_view const& column_names,
+                                          size_type const num_rows,
                                           string_view const row_prefix,
                                           string_view const row_suffix,
                                           string_view const value_separator,
@@ -268,40 +270,54 @@ std::unique_ptr<column> struct_to_strings(table_view const& strings_columns,
   auto const num_columns = strings_columns.num_columns();
   CUDF_EXPECTS(num_columns == column_names.size(),
                "Number of column names should be equal to number of columns in the table");
-  auto const strings_count = strings_columns.num_rows();
-  if (strings_count == 0)  // empty begets empty
+  if (num_rows == 0)  // empty begets empty
     return make_empty_column(type_id::STRING);
   // check all columns are of type string
   CUDF_EXPECTS(std::all_of(strings_columns.begin(),
                            strings_columns.end(),
                            [](auto const& c) { return c.type().id() == type_id::STRING; }),
                "All columns must be of type string");
   auto constexpr strviews_per_column = 3;  // (for each "column_name:", "value",  "separator")
-  auto const num_strviews_per_row    = strings_columns.num_columns() * strviews_per_column + 1;
+  auto const num_strviews_per_row    = strings_columns.num_columns() == 0
+                                         ? 2
+                                         : (1 + strings_columns.num_columns() * strviews_per_column);
   // e.g. {col1: value, col2: value, col3: value} = 1 + 3 + 3 + (3-1) + 1 = 10
 
   auto tbl_device_view = cudf::table_device_view::create(strings_columns, stream);
   auto d_column_names  = column_device_view::create(column_names, stream);
 
   // Note for future: chunk it but maximize parallelism, if memory usage is high.
-  auto const total_strings = num_strviews_per_row * strings_columns.num_rows();
-  auto const total_rows    = strings_columns.num_rows() * strings_columns.num_columns();
+  auto const total_strings = num_strviews_per_row * num_rows;
+  auto const total_rows    = num_rows * strings_columns.num_columns();
   rmm::device_uvector<string_view> d_strviews(total_strings, stream);
-  struct_scatter_strings_fn scatter_fn{*tbl_device_view,
-                                       *d_column_names,
-                                       strviews_per_column,
-                                       num_strviews_per_row,
-                                       row_prefix,
-                                       row_suffix,
-                                       value_separator,
-                                       narep.value(stream),
-                                       include_nulls,
-                                       d_strviews.begin()};
-  // scatter row_prefix, row_suffix, column_name:, value, value_separator as string_views
-  thrust::for_each(rmm::exec_policy(stream),
-                   thrust::make_counting_iterator<size_type>(0),
-                   thrust::make_counting_iterator<size_type>(total_rows),
-                   scatter_fn);
+  if (strings_columns.num_columns() > 0) {
+    struct_scatter_strings_fn scatter_fn{*tbl_device_view,
+                                         *d_column_names,
+                                         strviews_per_column,
+                                         num_strviews_per_row,
+                                         row_prefix,
+                                         row_suffix,
+                                         value_separator,
+                                         narep.value(stream),
+                                         include_nulls,
+                                         d_strviews.begin()};
+    // scatter row_prefix, row_suffix, column_name:, value, value_separator as string_views
+    thrust::for_each(rmm::exec_policy_nosync(stream),
+                     thrust::make_counting_iterator<size_type>(0),
+                     thrust::make_counting_iterator<size_type>(total_rows),
+                     scatter_fn);
+  } else {
+    thrust::for_each(
+      rmm::exec_policy_nosync(stream),
+      thrust::make_counting_iterator<size_type>(0),
+      thrust::make_counting_iterator<size_type>(num_rows),
+      [d_strviews = d_strviews.begin(), row_prefix, row_suffix, num_strviews_per_row] __device__(
+        auto idx) {
+        auto const this_index                             = idx * num_strviews_per_row;
+        d_strviews[this_index]                            = row_prefix;
+        d_strviews[this_index + num_strviews_per_row - 1] = row_suffix;
+      });
+  }
   if (!include_nulls) {
     // if previous column was null, then we skip the value separator
     rmm::device_uvector<bool> d_str_separator(total_rows, stream);
@@ -341,7 +357,7 @@ std::unique_ptr<column> struct_to_strings(table_view const& strings_columns,
 
   // gather from offset and create a new string column
   auto old_offsets = strings_column_view(joined_col->view()).offsets();
-  rmm::device_uvector<size_type> row_string_offsets(strings_columns.num_rows() + 1, stream, mr);
+  rmm::device_uvector<size_type> row_string_offsets(num_rows + 1, stream, mr);
   auto const d_strview_offsets = cudf::detail::make_counting_transform_iterator(
     0, cuda::proclaim_return_type<size_type>([num_strviews_per_row] __device__(size_type const i) {
       return i * num_strviews_per_row;
@@ -353,7 +369,7 @@ std::unique_ptr<column> struct_to_strings(table_view const& strings_columns,
                  row_string_offsets.begin());
   auto chars_data = joined_col->release().data;
   return make_strings_column(
-    strings_columns.num_rows(),
+    num_rows,
     std::make_unique<cudf::column>(std::move(row_string_offsets), rmm::device_buffer{}, 0),
     std::move(chars_data.release()[0]),
     0,
@@ -677,6 +693,7 @@ struct column_to_strings_fn {
     auto col_string = operator()(child_it,
                                  child_it + column.num_children(),
                                  children_names,
+                                 column.size(),
                                  struct_row_end_wrap.value(stream_));
     col_string->set_null_mask(cudf::detail::copy_bitmask(column, stream_, mr_),
                               column.null_count());
@@ -688,6 +705,7 @@ struct column_to_strings_fn {
   std::unique_ptr<column> operator()(column_iterator column_begin,
                                      column_iterator column_end,
                                      host_span<column_name_info const> children_names,
+                                     size_type num_rows,
                                      cudf::string_view const row_end_wrap_value) const
   {
     auto const num_columns = std::distance(column_begin, column_end);
@@ -733,6 +751,7 @@ struct column_to_strings_fn {
     //
     return struct_to_strings(str_table_view,
                              column_names_view,
+                             num_rows,
                              struct_row_begin_wrap.value(stream_),
                              row_end_wrap_value,
                              struct_value_separator.value(stream_),
@@ -908,6 +927,7 @@ void write_json_uncompressed(data_sink* out_sink,
       auto str_concat_col = converter(sub_view.begin(),
                                       sub_view.end(),
                                       user_column_names,
+                                      sub_view.num_rows(),
                                       d_line_terminator_with_row_end.value(stream));
 
       // Needs line_terminator at the end, to separate from next chunk

cpp/src/join/conditional_join.cu

@@ -22,6 +22,7 @@
 #include <cudf/ast/detail/expression_parser.hpp>
 #include <cudf/ast/expressions.hpp>
 #include <cudf/detail/device_scalar.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/join.hpp>
 #include <cudf/table/table.hpp>

cpp/src/join/hash_join.cu

@@ -18,6 +18,7 @@
 #include <cudf/copying.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/join.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/structs/utilities.hpp>
 #include <cudf/hashing/detail/helper_functions.cuh>
 #include <cudf/join.hpp>

cpp/src/join/join.cu

@@ -16,6 +16,7 @@
 #include "join_common_utils.hpp"
 
 #include <cudf/detail/gather.cuh>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/dictionary/detail/update_keys.hpp>
 #include <cudf/join.hpp>
 #include <cudf/table/table.hpp>

cpp/src/join/join_common_utils.hpp

@@ -50,11 +50,6 @@ using mixed_multimap_type =
                         cudf::detail::cuco_allocator<char>,
                         cuco::legacy::double_hashing<1, hash_type, hash_type>>;
 
-using row_hash_legacy =
-  cudf::row_hasher<cudf::hashing::detail::default_hash, cudf::nullate::DYNAMIC>;
-
-using row_equality_legacy = cudf::row_equality_comparator<cudf::nullate::DYNAMIC>;
-
 bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type);
 }  // namespace detail
 }  // namespace cudf

cpp/src/join/mixed_join.cu

@@ -21,6 +21,7 @@
 
 #include <cudf/ast/detail/expression_parser.hpp>
 #include <cudf/ast/expressions.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/hashing/detail/helper_functions.cuh>
 #include <cudf/join.hpp>

cpp/src/join/mixed_join_semi.cu

@@ -22,6 +22,7 @@
 #include <cudf/ast/expressions.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/hashing/detail/helper_functions.cuh>
 #include <cudf/join.hpp>

cpp/src/json/json_path.cu

@@ -22,6 +22,7 @@
 #include <cudf/detail/device_scalar.hpp>
 #include <cudf/detail/get_value.cuh>
 #include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/offsets_iterator_factory.cuh>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>

cpp/src/labeling/label_bins.cu

@@ -19,6 +19,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/label_bins.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/labeling/label_bins.hpp>
 #include <cudf/types.hpp>

cpp/src/lists/contains.cu

@@ -17,6 +17,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/lists/contains.hpp>
 #include <cudf/lists/detail/contains.hpp>

cpp/src/lists/copying/segmented_gather.cu

@@ -19,6 +19,7 @@
 #include <cudf/detail/indexalator.cuh>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/lists/detail/gather.cuh>
 #include <cudf/lists/gather.hpp>
 #include <cudf/utilities/default_stream.hpp>

cpp/src/quantiles/tdigest/tdigest.cu

@@ -18,6 +18,7 @@
 
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/tdigest/tdigest.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/valid_if.cuh>

cpp/src/reductions/minmax.cu

@@ -19,6 +19,7 @@
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/device_scalar.hpp>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/device_operators.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>

cpp/src/rolling/detail/rolling_fixed_window.cu

@@ -19,6 +19,7 @@
 #include <cudf_test/column_utilities.hpp>
 
 #include <cudf/detail/aggregation/aggregation.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 

cpp/src/rolling/detail/rolling_variable_window.cu

@@ -17,6 +17,7 @@
 #include "rolling.cuh"
 
 #include <cudf/detail/aggregation/aggregation.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 

cpp/src/rolling/grouped_rolling.cu

@@ -21,6 +21,7 @@
 #include "detail/rolling_jit.hpp"
 
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/rolling.hpp>
 #include <cudf/detail/utilities/assert.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>