Work around issues with cccl main

cpp/src/groupby/sort/group_rank_scan.cu

@@ -133,10 +133,10 @@ std::unique_ptr<column> rank_generator(column_view const& grouped_values,
 
  auto [group_labels_begin, mutable_rank_begin] = [&]() {
  if constexpr (forward) {
- return thrust::pair{group_labels.begin(), mutable_ranks.begin<size_type>()};
+ return thrust::make_pair(group_labels.begin(), mutable_ranks.begin<size_type>());
  } else {
- return thrust::pair{thrust::reverse_iterator(group_labels.end()),
- thrust::reverse_iterator(mutable_ranks.end<size_type>())};
+ return thrust::make_pair(thrust::reverse_iterator(group_labels.end()),
+  thrust::reverse_iterator(mutable_ranks.end<size_type>()));
  }
  }();
  thrust::inclusive_scan_by_key(rmm::exec_policy(stream),

cpp/src/io/comp/nvcomp_adapter.cu

@@ -108,7 +108,7 @@ void skip_unsupported_inputs(device_span<size_t> input_sizes,
  input_sizes.begin(),
  status_size_it,
  [] __device__(auto const& status) {
- return thrust::pair{0, compression_result{0, compression_status::SKIPPED}};
+ return thrust::make_pair(0, compression_result{0, compression_status::SKIPPED});
  },
  [max_size = max_valid_input_size.value()] __device__(size_t input_size) {
  return input_size > max_size;

cpp/src/io/comp/statistics.cu

@@ -18,6 +18,7 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/functional>
 #include <thrust/transform_reduce.h>
 
 namespace cudf::io {
@@ -32,9 +33,9 @@ writer_compression_statistics collect_compression_statistics(
  rmm::exec_policy(stream),
  results.begin(),
  results.end(),
- [] __device__(auto& res) {
+ cuda::proclaim_return_type<size_t>([] __device__(const compression_result& res) {
  return res.status == compression_status::SUCCESS ? res.bytes_written : 0;
- },
+ }),
  0ul,
  thrust::plus<size_t>());
 
@@ -47,9 +48,9 @@ writer_compression_statistics collect_compression_statistics(
  rmm::exec_policy(stream),
  zipped_begin,
  zipped_end,
- [status] __device__(auto tup) {
+ cuda::proclaim_return_type<size_t>([status] __device__(auto tup) {
  return thrust::get<1>(tup).status == status ? thrust::get<0>(tup).size() : 0;
- },
+ }),
  0ul,
  thrust::plus<size_t>());
  };

cpp/src/io/parquet/page_string_decode.cu

@@ -1186,13 +1186,13 @@ void ComputePageStringSizes(cudf::detail::hostdevice_span<PageInfo> pages,
 
  // check for needed temp space for DELTA_BYTE_ARRAY
  auto const need_sizes = thrust::any_of(
- rmm::exec_policy(stream), pages.device_begin(), pages.device_end(), [] __device__(auto& page) {
+ rmm::exec_policy(stream), pages.device_begin(), pages.device_end(), cuda::proclaim_return_type<bool>([] __device__(auto& page) {
  return page.temp_string_size != 0;
- });
+ }));
 
  if (need_sizes) {
  // sum up all of the temp_string_sizes
- auto const page_sizes = [] __device__(PageInfo const& page) { return page.temp_string_size; };
+ auto const page_sizes = cuda::proclaim_return_type<int64_t>([] __device__(PageInfo const& page) { return page.temp_string_size; });
  auto const total_size = thrust::transform_reduce(rmm::exec_policy(stream),
  pages.device_begin(),
  pages.device_end(),

cpp/src/io/parquet/reader_impl_preprocess.cu

@@ -452,9 +452,9 @@ std::string encoding_to_string(Encoding encoding)
 [[nodiscard]] std::string list_unsupported_encodings(device_span<PageInfo const> pages,
  rmm::cuda_stream_view stream)
 {
- auto const to_mask = [] __device__(auto const& page) {
+ auto const to_mask = cuda::proclaim_return_type<uint32_t>([] __device__(auto const& page) {
  return is_supported_encoding(page.encoding) ? 0U : encoding_to_mask(page.encoding);
- };
+ });
  uint32_t const unsupported = thrust::transform_reduce(
  rmm::exec_policy(stream), pages.begin(), pages.end(), to_mask, 0U, thrust::bit_or<uint32_t>());
  return encoding_bitmask_to_str(unsupported);

cpp/src/io/utilities/data_casting.cu

@@ -33,6 +33,7 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/functional>
 #include <thrust/copy.h>
 #include <thrust/functional.h>
 #include <thrust/transform_reduce.h>
@@ -782,7 +783,7 @@ template <typename SymbolT>
 struct to_string_view_pair {
  SymbolT const* data;
  to_string_view_pair(SymbolT const* _data) : data(_data) {}
- __device__ auto operator()(thrust::tuple<size_type, size_type> ip)
+ __device__ thrust::pair<char const*, std::size_t> operator()(thrust::tuple<size_type, size_type> ip)
  {
  return thrust::pair<char const*, std::size_t>{data + thrust::get<0>(ip),
  static_cast<std::size_t>(thrust::get<1>(ip))};
@@ -804,7 +805,7 @@ static std::unique_ptr<column> parse_string(string_view_pair_it str_tuples,
  rmm::exec_policy(stream),
  str_tuples,
  str_tuples + col_size,
- [] __device__(auto t) { return t.second; },
+ cuda::proclaim_return_type<std::size_t>([] __device__(auto t) { return t.second; }),
  size_type{0},
  thrust::maximum<size_type>{});
 

cpp/src/lists/contains.cu

@@ -103,9 +103,9 @@ __device__ auto element_index_pair_iter(size_type const size)
  auto const end = thrust::make_counting_iterator(size);
 
  if constexpr (forward) {
- return thrust::pair{begin, end};
+ return thrust::make_pair(begin, end);
  } else {
- return thrust::pair{thrust::make_reverse_iterator(end), thrust::make_reverse_iterator(begin)};
+ return thrust::make_pair(thrust::make_reverse_iterator(end), thrust::make_reverse_iterator(begin));
  }
 }