Merge pull request #11758 from rapidsai/branch-22.10

[gpuCI] Forward-merge branch-22.10 to branch-22.12 [skip gpuci]

cpp/tests/column/column_test.cu

@@ -31,6 +31,8 @@
 #include <cudf_test/type_list_utilities.hpp>
 #include <cudf_test/type_lists.hpp>
 
+#include <rmm/exec_policy.hpp>
+
 #include <thrust/execution_policy.h>
 #include <thrust/sequence.h>
 
@@ -46,8 +48,10 @@ struct TypedColumnTest : public cudf::test::BaseFixture {
   {
     auto typed_data = static_cast<char*>(data.data());
     auto typed_mask = static_cast<char*>(mask.data());
-    thrust::sequence(thrust::device, typed_data, typed_data + data.size());
-    thrust::sequence(thrust::device, typed_mask, typed_mask + mask.size());
+    thrust::sequence(
+      rmm::exec_policy(cudf::default_stream_value), typed_data, typed_data + data.size());
+    thrust::sequence(
+      rmm::exec_policy(cudf::default_stream_value), typed_mask, typed_mask + mask.size());
   }
 
   cudf::size_type num_elements() { return _num_elements; }
@@ -349,7 +353,7 @@ TYPED_TEST(TypedColumnTest, DeviceUvectorConstructorNoMask)
 {
   rmm::device_uvector<TypeParam> original{static_cast<std::size_t>(this->num_elements()),
                                           cudf::default_stream_value};
-  thrust::copy(thrust::device,
+  thrust::copy(rmm::exec_policy(cudf::default_stream_value),
                static_cast<TypeParam*>(this->data.data()),
                static_cast<TypeParam*>(this->data.data()) + this->num_elements(),
                original.begin());
@@ -366,7 +370,7 @@ TYPED_TEST(TypedColumnTest, DeviceUvectorConstructorWithMask)
 {
   rmm::device_uvector<TypeParam> original{static_cast<std::size_t>(this->num_elements()),
                                           cudf::default_stream_value};
-  thrust::copy(thrust::device,
+  thrust::copy(rmm::exec_policy(cudf::default_stream_value),
                static_cast<TypeParam*>(this->data.data()),
                static_cast<TypeParam*>(this->data.data()) + this->num_elements(),
                original.begin());

cpp/tests/column/compound_test.cu

@@ -66,7 +66,7 @@ struct checker_for_level2 {
 TEST_F(CompoundColumnTest, ChildrenLevel1)
 {
   rmm::device_uvector<int32_t> data(1000, cudf::default_stream_value);
-  thrust::sequence(rmm::exec_policy(), data.begin(), data.end(), 1);
+  thrust::sequence(rmm::exec_policy(cudf::default_stream_value), data.begin(), data.end(), 1);
 
   auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED);
   rmm::device_buffer data1{data.data() + 100, 100 * sizeof(int32_t), cudf::default_stream_value};
@@ -105,14 +105,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel1)
 
   {
     auto column = cudf::column_device_view::create(parent->view());
-    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(),
+    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(cudf::default_stream_value),
                                thrust::make_counting_iterator<int32_t>(0),
                                thrust::make_counting_iterator<int32_t>(100),
                                checker_for_level1<cudf::column_device_view>{*column}));
   }
   {
     auto column = cudf::mutable_column_device_view::create(parent->mutable_view());
-    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(),
+    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(cudf::default_stream_value),
                                thrust::make_counting_iterator<int32_t>(0),
                                thrust::make_counting_iterator<int32_t>(100),
                                checker_for_level1<cudf::mutable_column_device_view>{*column}));
@@ -122,7 +122,7 @@ TEST_F(CompoundColumnTest, ChildrenLevel1)
 TEST_F(CompoundColumnTest, ChildrenLevel2)
 {
   rmm::device_uvector<int32_t> data(1000, cudf::default_stream_value);
-  thrust::sequence(rmm::exec_policy(), data.begin(), data.end(), 1);
+  thrust::sequence(rmm::exec_policy(cudf::default_stream_value), data.begin(), data.end(), 1);
 
   auto null_mask = cudf::create_null_mask(100, cudf::mask_state::UNALLOCATED);
   rmm::device_buffer data11{data.data() + 100, 100 * sizeof(int32_t), cudf::default_stream_value};
@@ -202,14 +202,14 @@ TEST_F(CompoundColumnTest, ChildrenLevel2)
 
   {
     auto column = cudf::column_device_view::create(parent->view());
-    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(),
+    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(cudf::default_stream_value),
                                thrust::make_counting_iterator<int32_t>(0),
                                thrust::make_counting_iterator<int32_t>(100),
                                checker_for_level2<cudf::column_device_view>{*column}));
   }
   {
     auto column = cudf::mutable_column_device_view::create(parent->mutable_view());
-    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(),
+    EXPECT_TRUE(thrust::any_of(rmm::exec_policy(cudf::default_stream_value),
                                thrust::make_counting_iterator<int32_t>(0),
                                thrust::make_counting_iterator<int32_t>(100),
                                checker_for_level2<cudf::mutable_column_device_view>{*column}));

cpp/tests/copying/concatenate_tests.cu

@@ -522,11 +522,11 @@ TEST_F(OverflowTest, Presliced)
 
     // try and concatenate 4 string columns of with ~1/2 billion chars in each
     auto offsets = cudf::make_fixed_width_column(data_type{type_id::INT32}, num_rows + 1);
-    thrust::fill(rmm::exec_policy(),
+    thrust::fill(rmm::exec_policy(cudf::default_stream_value),
                  offsets->mutable_view().begin<offset_type>(),
                  offsets->mutable_view().end<offset_type>(),
                  string_size);
-    thrust::exclusive_scan(rmm::exec_policy(),
+    thrust::exclusive_scan(rmm::exec_policy(cudf::default_stream_value),
                            offsets->view().begin<offset_type>(),
                            offsets->view().end<offset_type>(),
                            offsets->mutable_view().begin<offset_type>());
@@ -596,11 +596,11 @@ TEST_F(OverflowTest, Presliced)
 
     // try and concatenate 4 struct columns of with ~1/2 billion elements in each
     auto offsets = cudf::make_fixed_width_column(data_type{type_id::INT32}, num_rows + 1);
-    thrust::fill(rmm::exec_policy(),
+    thrust::fill(rmm::exec_policy(cudf::default_stream_value),
                  offsets->mutable_view().begin<offset_type>(),
                  offsets->mutable_view().end<offset_type>(),
                  list_size);
-    thrust::exclusive_scan(rmm::exec_policy(),
+    thrust::exclusive_scan(rmm::exec_policy(cudf::default_stream_value),
                            offsets->view().begin<offset_type>(),
                            offsets->view().end<offset_type>(),
                            offsets->mutable_view().begin<offset_type>());
@@ -688,11 +688,11 @@ TEST_F(OverflowTest, BigColumnsSmallSlices)
     constexpr size_type string_size = inner_size / num_rows;
 
     auto offsets = cudf::make_fixed_width_column(data_type{type_id::INT32}, num_rows + 1);
-    thrust::fill(rmm::exec_policy(),
+    thrust::fill(rmm::exec_policy(cudf::default_stream_value),
                  offsets->mutable_view().begin<offset_type>(),
                  offsets->mutable_view().end<offset_type>(),
                  string_size);
-    thrust::exclusive_scan(rmm::exec_policy(),
+    thrust::exclusive_scan(rmm::exec_policy(cudf::default_stream_value),
                            offsets->view().begin<offset_type>(),
                            offsets->view().end<offset_type>(),
                            offsets->mutable_view().begin<offset_type>());
@@ -715,11 +715,11 @@ TEST_F(OverflowTest, BigColumnsSmallSlices)
     constexpr size_type list_size = inner_size / num_rows;
 
     auto offsets = cudf::make_fixed_width_column(data_type{type_id::INT32}, num_rows + 1);
-    thrust::fill(rmm::exec_policy(),
+    thrust::fill(rmm::exec_policy(cudf::default_stream_value),
                  offsets->mutable_view().begin<offset_type>(),
                  offsets->mutable_view().end<offset_type>(),
                  list_size);
-    thrust::exclusive_scan(rmm::exec_policy(),
+    thrust::exclusive_scan(rmm::exec_policy(cudf::default_stream_value),
                            offsets->view().begin<offset_type>(),
                            offsets->view().end<offset_type>(),
                            offsets->mutable_view().begin<offset_type>());
@@ -742,11 +742,11 @@ TEST_F(OverflowTest, BigColumnsSmallSlices)
     constexpr size_type list_size = inner_size / num_rows;
 
     auto offsets = cudf::make_fixed_width_column(data_type{type_id::INT32}, num_rows + 1);
-    thrust::fill(rmm::exec_policy(),
+    thrust::fill(rmm::exec_policy(cudf::default_stream_value),
                  offsets->mutable_view().begin<offset_type>(),
                  offsets->mutable_view().end<offset_type>(),
                  list_size);
-    thrust::exclusive_scan(rmm::exec_policy(),
+    thrust::exclusive_scan(rmm::exec_policy(cudf::default_stream_value),
                            offsets->view().begin<offset_type>(),
                            offsets->view().end<offset_type>(),
                            offsets->mutable_view().begin<offset_type>());

cpp/tests/copying/detail_gather_tests.cu

@@ -48,7 +48,8 @@ TYPED_TEST(GatherTest, GatherDetailDeviceVectorTest)
 {
   constexpr cudf::size_type source_size{1000};
   rmm::device_uvector<cudf::size_type> gather_map(source_size, cudf::default_stream_value);
-  thrust::sequence(thrust::device, gather_map.begin(), gather_map.end());
+  thrust::sequence(
+    rmm::exec_policy(cudf::default_stream_value), gather_map.begin(), gather_map.end());
 
   auto data = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i; });
   cudf::test::fixed_width_column_wrapper<TypeParam> source_column(data, data + source_size);

cpp/tests/device_atomics/device_atomics_test.cu

@@ -147,9 +147,11 @@ struct AtomicsTest : public cudf::test::BaseFixture {
     if (block_size == 0) { block_size = vec_size; }
 
     if (is_cas_test) {
-      gpu_atomicCAS_test<<<grid_size, block_size>>>(dev_result.data(), dev_data.data(), vec_size);
+      gpu_atomicCAS_test<<<grid_size, block_size, 0, cudf::default_stream_value.value()>>>(
+        dev_result.data(), dev_data.data(), vec_size);
     } else {
-      gpu_atomic_test<<<grid_size, block_size>>>(dev_result.data(), dev_data.data(), vec_size);
+      gpu_atomic_test<<<grid_size, block_size, 0, cudf::default_stream_value.value()>>>(
+        dev_result.data(), dev_data.data(), vec_size);
     }
 
     auto host_result = cudf::detail::make_host_vector_sync(dev_result);
@@ -296,7 +298,7 @@ struct AtomicsBitwiseOpTest : public cudf::test::BaseFixture {
 
     if (block_size == 0) { block_size = vec_size; }
 
-    gpu_atomic_bitwiseOp_test<T><<<grid_size, block_size>>>(
+    gpu_atomic_bitwiseOp_test<T><<<grid_size, block_size, 0, cudf::default_stream_value.value()>>>(
       reinterpret_cast<T*>(dev_result.data()), reinterpret_cast<T*>(dev_data.data()), vec_size);
 
     auto host_result = cudf::detail::make_host_vector_sync(dev_result);

cpp/tests/error/error_handling_test.cu

@@ -17,6 +17,7 @@
 #include <cudf_test/base_fixture.hpp>
 
 #include <cudf/filling.hpp>
+#include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/error.hpp>
 
 #include <rmm/cuda_stream.hpp>
@@ -89,7 +90,7 @@ TEST(DeathTest, CudaFatalError)
 {
   testing::FLAGS_gtest_death_test_style = "threadsafe";
   auto call_kernel                      = []() {
-    kernel<<<1, 1>>>();
+    kernel<<<1, 1, 0, cudf::default_stream_value.value()>>>();
     try {
       CUDF_CUDA_TRY(cudaDeviceSynchronize());
     } catch (const cudf::fatal_cuda_error& fe) {

cpp/tests/fixed_point/fixed_point_tests.cu

@@ -85,8 +85,10 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice)
   std::vector<decimal32> vec1(1000, decimal32{1, scale_type{-2}});
   auto d_vec1 = cudf::detail::make_device_uvector_sync(vec1);
 
-  auto const sum = thrust::reduce(
-    rmm::exec_policy(), std::cbegin(d_vec1), std::cend(d_vec1), decimal32{0, scale_type{-2}});
+  auto const sum = thrust::reduce(rmm::exec_policy(cudf::default_stream_value),
+                                  std::cbegin(d_vec1),
+                                  std::cend(d_vec1),
+                                  decimal32{0, scale_type{-2}});
 
   EXPECT_EQ(static_cast<int32_t>(sum), 1000);
 
@@ -99,16 +101,18 @@ TEST_F(FixedPointTest, DecimalXXThrustOnDevice)
   std::vector<int32_t> vec2(1000);
   std::iota(std::begin(vec2), std::end(vec2), 1);
 
-  auto const res1 = thrust::reduce(
-    rmm::exec_policy(), std::cbegin(d_vec1), std::cend(d_vec1), decimal32{0, scale_type{-2}});
+  auto const res1 = thrust::reduce(rmm::exec_policy(cudf::default_stream_value),
+                                   std::cbegin(d_vec1),
+                                   std::cend(d_vec1),
+                                   decimal32{0, scale_type{-2}});
 
   auto const res2 = std::accumulate(std::cbegin(vec2), std::cend(vec2), 0);
 
   EXPECT_EQ(static_cast<int32_t>(res1), res2);
 
   rmm::device_uvector<int32_t> d_vec3(1000, cudf::default_stream_value);
 
-  thrust::transform(rmm::exec_policy(),
+  thrust::transform(rmm::exec_policy(cudf::default_stream_value),
                     std::cbegin(d_vec1),
                     std::cend(d_vec1),
                     std::begin(d_vec3),

cpp/tests/groupby/lists_tests.cu

@@ -28,6 +28,8 @@
 #include <cudf/types.hpp>
 #include <cudf/utilities/default_stream.hpp>
 
+#include <rmm/exec_policy.hpp>
+
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/logical.h>
 
@@ -122,7 +124,8 @@ inline void test_hash_based_sum_agg(column_view const& keys,
 
   // For each row in expected table `t[0, num_rows)`, there must be a match
   // in the resulting table `t[num_rows, 2 * num_rows)`
-  EXPECT_TRUE(thrust::all_of(thrust::make_counting_iterator<cudf::size_type>(0),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
+                             thrust::make_counting_iterator<cudf::size_type>(0),
                              thrust::make_counting_iterator<cudf::size_type>(num_rows),
                              func));
 }

cpp/tests/hash_map/map_test.cu

@@ -140,16 +140,18 @@ TYPED_TEST(InsertTest, UniqueKeysUniqueValues)
 {
   using map_type  = typename TypeParam::map_type;
   using pair_type = typename TypeParam::pair_type;
-  thrust::tabulate(
-    rmm::exec_policy(), this->pairs.begin(), this->pairs.end(), unique_pair_generator<pair_type>{});
+  thrust::tabulate(rmm::exec_policy(cudf::default_stream_value),
+                   this->pairs.begin(),
+                   this->pairs.end(),
+                   unique_pair_generator<pair_type>{});
   // All pairs should be new inserts
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.end(),
                              insert_pair<map_type, pair_type>{*this->map}));
 
   // All pairs should be present in the map
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.end(),
                              find_pair<map_type, pair_type>{*this->map}));
@@ -159,23 +161,23 @@ TYPED_TEST(InsertTest, IdenticalKeysIdenticalValues)
 {
   using map_type  = typename TypeParam::map_type;
   using pair_type = typename TypeParam::pair_type;
-  thrust::tabulate(rmm::exec_policy(),
+  thrust::tabulate(rmm::exec_policy(cudf::default_stream_value),
                    this->pairs.begin(),
                    this->pairs.end(),
                    identical_pair_generator<pair_type>{});
   // Insert a single pair
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.begin() + 1,
                              insert_pair<map_type, pair_type>{*this->map}));
   // Identical inserts should all return false (no new insert)
-  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                               this->pairs.begin(),
                               this->pairs.end(),
                               insert_pair<map_type, pair_type>{*this->map}));
 
   // All pairs should be present in the map
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.end(),
                              find_pair<map_type, pair_type>{*this->map}));
@@ -185,30 +187,30 @@ TYPED_TEST(InsertTest, IdenticalKeysUniqueValues)
 {
   using map_type  = typename TypeParam::map_type;
   using pair_type = typename TypeParam::pair_type;
-  thrust::tabulate(rmm::exec_policy(),
+  thrust::tabulate(rmm::exec_policy(cudf::default_stream_value),
                    this->pairs.begin(),
                    this->pairs.end(),
                    identical_key_generator<pair_type>{});
 
   // Insert a single pair
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.begin() + 1,
                              insert_pair<map_type, pair_type>{*this->map}));
 
   // Identical key inserts should all return false (no new insert)
-  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                               this->pairs.begin() + 1,
                               this->pairs.end(),
                               insert_pair<map_type, pair_type>{*this->map}));
 
   // Only first pair is present in map
-  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_TRUE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                              this->pairs.begin(),
                              this->pairs.begin() + 1,
                              find_pair<map_type, pair_type>{*this->map}));
 
-  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(),
+  EXPECT_FALSE(thrust::all_of(rmm::exec_policy(cudf::default_stream_value),
                               this->pairs.begin() + 1,
                               this->pairs.end(),
                               find_pair<map_type, pair_type>{*this->map}));

cpp/tests/io/json_type_cast_test.cu

@@ -70,7 +70,7 @@ TEST_F(JSONTypeCastTest, String)
 
   auto d_column = cudf::column_device_view::create(input);
   rmm::device_uvector<thrust::pair<const char*, cudf::size_type>> svs(d_column->size(), stream);
-  thrust::transform(thrust::device,
+  thrust::transform(rmm::exec_policy(cudf::default_stream_value),
                     d_column->pair_begin<cudf::string_view, false>(),
                     d_column->pair_end<cudf::string_view, false>(),
                     svs.begin(),
@@ -100,7 +100,7 @@ TEST_F(JSONTypeCastTest, Int)
   cudf::test::strings_column_wrapper data({"1", "null", "3", "true", "5", "false"});
   auto d_column = cudf::column_device_view::create(data);
   rmm::device_uvector<thrust::pair<const char*, cudf::size_type>> svs(d_column->size(), stream);
-  thrust::transform(thrust::device,
+  thrust::transform(rmm::exec_policy(cudf::default_stream_value),
                     d_column->pair_begin<cudf::string_view, false>(),
                     d_column->pair_end<cudf::string_view, false>(),
                     svs.begin(),
@@ -137,7 +137,7 @@ TEST_F(JSONTypeCastTest, StringEscapes)
   });
   auto d_column = cudf::column_device_view::create(data);
   rmm::device_uvector<thrust::pair<const char*, cudf::size_type>> svs(d_column->size(), stream);
-  thrust::transform(thrust::device,
+  thrust::transform(rmm::exec_policy(cudf::default_stream_value),
                     d_column->pair_begin<cudf::string_view, false>(),
                     d_column->pair_end<cudf::string_view, false>(),
                     svs.begin(),