Merge remote-tracking branch 'origin/branch-23.10' into bot-submodule…

…-sync-branch-23.10

src/main/cpp/src/cast_decimal_to_string.cu

@@ -52,7 +52,7 @@ namespace {
 template <typename DecimalType>
 struct decimal_to_non_ansi_string_fn {
   column_device_view d_decimals;
-  offset_type* d_offsets{};
+  size_type* d_offsets{};
   char* d_chars{};
 
   /**

src/main/cpp/src/cast_string.cu

@@ -159,7 +159,7 @@ template <typename T>
 void __global__ string_to_integer_kernel(T* out,
                                          bitmask_type* validity,
                                          const char* const chars,
-                                         offset_type const* offsets,
+                                         size_type const* offsets,
                                          bitmask_type const* incoming_null_mask,
                                          size_type num_rows,
                                          bool ansi_mode,
@@ -392,7 +392,7 @@ template <typename T>
 __global__ void string_to_decimal_kernel(T* out,
                                          bitmask_type* validity,
                                          const char* const chars,
-                                         offset_type const* offsets,
+                                         size_type const* offsets,
                                          bitmask_type const* incoming_null_mask,
                                          size_type num_rows,
                                          int32_t scale,
@@ -612,10 +612,10 @@ void validate_ansi_column(column_view const& col,
                                              thrust::make_counting_iterator(col.size()),
                                              row_valid_fn{col.null_mask(), source_col.null_mask()});
 
-    offset_type string_bounds[2];
+    size_type string_bounds[2];
     cudaMemcpyAsync(&string_bounds,
-                    &source_col.offsets().data<offset_type>()[*first_error],
-                    sizeof(offset_type) * 2,
+                    &source_col.offsets().data<size_type>()[*first_error],
+                    sizeof(size_type) * 2,
                     cudaMemcpyDeviceToHost,
                     stream.value());
     stream.synchronize();
@@ -668,7 +668,7 @@ struct string_to_integer_impl {
       data.data(),
       null_mask.data(),
       string_col.chars().data<char const>(),
-      string_col.offsets().data<offset_type>(),
+      string_col.offsets().data<size_type>(),
       string_col.null_mask(),
       string_col.size(),
       ansi_mode,
@@ -737,7 +737,7 @@ struct string_to_decimal_impl {
       data.data(),
       null_mask.data(),
       string_col.chars().data<char const>(),
-      string_col.offsets().data<offset_type>(),
+      string_col.offsets().data<size_type>(),
       string_col.null_mask(),
       string_col.size(),
       dtype.scale(),

src/main/cpp/src/cast_string_to_float.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-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.
@@ -59,7 +59,7 @@ class string_to_float {
                              int32_t* ansi_except,
                              size_type* valid_count,
                              const char* const chars,
-                             offset_type const* offsets,
+                             size_type const* offsets,
                              uint64_t const* const ipow,
                              bitmask_type const* incoming_null_mask,
                              size_type const num_rows)
@@ -595,7 +595,7 @@ __global__ void string_to_float_kernel(T* out,
                                        int32_t* ansi_except,
                                        size_type* valid_count,
                                        const char* const chars,
-                                       offset_type const* offsets,
+                                       size_type const* offsets,
                                        bitmask_type const* incoming_null_mask,
                                        size_type const num_rows)
 {
@@ -676,7 +676,7 @@ std::unique_ptr<column> string_to_float(data_type dtype,
         ansi_mode ? static_cast<ScalarType*>(ansi_count.get())->data() : nullptr,
         static_cast<ScalarType*>(valid_count.get())->data(),
         string_col.chars().begin<char>(),
-        string_col.offsets().begin<offset_type>(),
+        string_col.offsets().begin<size_type>(),
         string_col.null_mask(),
         num_rows);
   } else {
@@ -687,7 +687,7 @@ std::unique_ptr<column> string_to_float(data_type dtype,
         ansi_mode ? static_cast<ScalarType*>(ansi_count.get())->data() : nullptr,
         static_cast<ScalarType*>(valid_count.get())->data(),
         string_col.chars().begin<char>(),
-        string_col.offsets().begin<offset_type>(),
+        string_col.offsets().begin<size_type>(),
         string_col.null_mask(),
         num_rows);
   }
@@ -698,10 +698,10 @@ std::unique_ptr<column> string_to_float(data_type dtype,
     auto const val = static_cast<ScalarType*>(ansi_count.get())->value(stream);
     if (val >= 0) {
       auto const error_row = num_rows - val;
-      offset_type string_bounds[2];
+      size_type string_bounds[2];
       cudaMemcpyAsync(&string_bounds,
-                      &string_col.offsets().data<offset_type>()[error_row],
-                      sizeof(offset_type) * 2,
+                      &string_col.offsets().data<size_type>()[error_row],
+                      sizeof(size_type) * 2,
                       cudaMemcpyDeviceToHost,
                       stream.value());
       stream.synchronize();

src/main/cpp/src/map_utils.cu

@@ -500,7 +500,7 @@ struct substring_fn {
   cudf::device_span<char const> const d_string;
   cudf::device_span<thrust::pair<SymbolOffsetT, SymbolOffsetT> const> const d_ranges;
 
-  cudf::offset_type* d_offsets{};
+  cudf::size_type* d_offsets{};
   char* d_chars{};
 
   __device__ void operator()(cudf::size_type const idx)
@@ -557,7 +557,7 @@ std::unique_ptr<cudf::column> extract_keys_or_values(
 }
 
 // Compute the offsets for the final lists of Struct<String,String>.
-rmm::device_uvector<cudf::offset_type> compute_list_offsets(
+rmm::device_uvector<cudf::size_type> compute_list_offsets(
   cudf::size_type n_lists,
   rmm::device_uvector<NodeIndexT> const& parent_node_ids,
   rmm::device_uvector<int8_t> const& key_or_value,
@@ -599,7 +599,7 @@ rmm::device_uvector<cudf::offset_type> compute_list_offsets(
   print_debug(node_child_counts, "Nodes' child keys counts", ", ", stream);
 #endif
 
-  auto list_offsets   = rmm::device_uvector<cudf::offset_type>(n_lists + 1, stream, mr);
+  auto list_offsets   = rmm::device_uvector<cudf::size_type>(n_lists + 1, stream, mr);
   auto const copy_end = cudf::detail::copy_if_safe(
     node_child_counts.begin(),
     node_child_counts.end(),

src/main/cpp/src/map_utils_debug.cuh

@@ -108,7 +108,7 @@ void print_pair_debug(rmm::device_uvector<T> const& input,
 }
 
 // Print the final output map data (Spark's MapType, i.e., List<Struct<String,String>>).
-void print_output_spark_map(rmm::device_uvector<cudf::offset_type> const& list_offsets,
+void print_output_spark_map(rmm::device_uvector<cudf::size_type> const& list_offsets,
                             std::unique_ptr<cudf::column> const& extracted_keys,
                             std::unique_ptr<cudf::column> const& extracted_values,
                             rmm::cuda_stream_view stream)
@@ -138,7 +138,7 @@ void print_output_spark_map(rmm::device_uvector<cudf::offset_type> const& list_o
     stream);
 
   auto const h_list_offsets = cudf::detail::make_host_vector_sync(
-    cudf::device_span<cudf::offset_type const>{list_offsets.data(), list_offsets.size()}, stream);
+    cudf::device_span<cudf::size_type const>{list_offsets.data(), list_offsets.size()}, stream);
   CUDF_EXPECTS(h_list_offsets.back() == extracted_keys->size(),
                "Invalid list offsets computation.");
 

src/main/cpp/src/row_conversion.cu

@@ -183,7 +183,7 @@ struct tile_info {
 struct row_batch {
   size_type num_bytes;                      // number of bytes in this batch
   size_type row_count;                      // number of rows in the batch
-  device_uvector<offset_type> row_offsets;  // offsets column of output cudf column
+  device_uvector<size_type> row_offsets;  // offsets column of output cudf column
 };
 
 /**

src/main/cpp/tests/bloom_filter.cu

@@ -118,7 +118,7 @@ TEST_F(BloomFilterTest, BuildAndProbeWithNulls)
 struct bloom_filter_stride_transform {
   int const stride;
 
-  cudf::offset_type __device__ operator()(cudf::size_type i) { return i * stride; }
+  cudf::size_type __device__ operator()(cudf::size_type i) { return i * stride; }
 };
 
 TEST_F(BloomFilterTest, ProbeMerged)
@@ -154,7 +154,7 @@ TEST_F(BloomFilterTest, ProbeMerged)
   thrust::transform(rmm::exec_policy(cudf::get_default_stream()),
                     thrust::make_counting_iterator(0),
                     thrust::make_counting_iterator(0) + 4,
-                    premerge_offsets->mutable_view().begin<cudf::offset_type>(),
+                    premerge_offsets->mutable_view().begin<cudf::size_type>(),
                     bloom_filter_stride_transform{bloom_filter_a->view().size()});
   auto premerged = cudf::make_lists_column(
     3, std::move(premerge_offsets), std::move(premerge_children), 0, rmm::device_buffer{});