From 6272ef1073e7d600e6d49588846e916155b9b245 Mon Sep 17 00:00:00 2001
From: Paul Taylor <paul.e.taylor@me.com>
Date: Tue, 18 May 2021 16:21:21 -0500
Subject: [PATCH] Fix Thrust 1.12 compile errors (#231)

Initialize `rmm::exec_policy` with `rmm::cuda_stream_view` to fix Thrust 1.12 compile errors.

Authors:
  - Paul Taylor (https://github.com/trxcllnt)

Approvers:
  - Dante Gama Dessavre (https://github.com/dantegd)

URL: https://github.com/rapidsai/raft/pull/231
---
 .../raft/sparse/hierarchy/detail/agglomerative.cuh     |  2 +-
 .../raft/sparse/hierarchy/detail/connectivities.cuh    |  4 ++--
 cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh  | 10 +++++-----
 cpp/include/raft/sparse/op/reduce.cuh                  |  2 +-
 4 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh
index 598df519aa..69fd8d4dc9 100644
--- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh
+++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh
@@ -278,7 +278,7 @@ void extract_flattened_clusters(const raft::handle_t &handle, value_idx *labels,
                                 size_t n_leaves) {
   auto d_alloc = handle.get_device_allocator();
   auto stream = handle.get_stream();
-  auto thrust_policy = rmm::exec_policy(stream);
+  auto thrust_policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
 
   // Handle special case where n_clusters == 1
   if (n_clusters == 1) {
diff --git a/cpp/include/raft/sparse/hierarchy/detail/connectivities.cuh b/cpp/include/raft/sparse/hierarchy/detail/connectivities.cuh
index 286ea386c4..7cf959dda6 100644
--- a/cpp/include/raft/sparse/hierarchy/detail/connectivities.cuh
+++ b/cpp/include/raft/sparse/hierarchy/detail/connectivities.cuh
@@ -62,7 +62,7 @@ struct distance_graph_impl<raft::hierarchy::LinkageDistance::KNN_GRAPH,
            rmm::device_uvector<value_t> &data, int c) {
     auto d_alloc = handle.get_device_allocator();
     auto stream = handle.get_stream();
-    auto exec_policy = rmm::exec_policy(stream);
+    auto exec_policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
 
     // Need to symmetrize knn into undirected graph
     raft::sparse::COO<value_t, value_idx> knn_graph_coo(d_alloc, stream);
@@ -133,4 +133,4 @@ void get_distance_graph(const raft::handle_t &handle, const value_t *X,
 
 };  // namespace detail
 };  // namespace hierarchy
-};  // namespace raft
\ No newline at end of file
+};  // namespace raft
diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh
index 4fe40975bc..c12c4a8d02 100644
--- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh
+++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh
@@ -83,7 +83,7 @@ MST_solver<vertex_t, edge_t, weight_t>::MST_solver(
   sm_count = handle_.get_device_properties().multiProcessorCount;
 
   //Initially, color holds the vertex id as color
-  auto policy = rmm::exec_policy(stream);
+  auto policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
   if (initialize_colors_) {
     thrust::sequence(policy, color.begin(), color.end(), 0);
     thrust::sequence(policy, color_index, color_index + v, 0);
@@ -212,7 +212,7 @@ struct alteration_functor {
 // Compute the uper bound for the alteration
 template <typename vertex_t, typename edge_t, typename weight_t>
 weight_t MST_solver<vertex_t, edge_t, weight_t>::alteration_max() {
-  auto policy = rmm::exec_policy(stream);
+  auto policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
   rmm::device_vector<weight_t> tmp(e);
   thrust::device_ptr<const weight_t> weights_ptr(weights);
   thrust::copy(policy, weights_ptr, weights_ptr + e, tmp.begin());
@@ -308,7 +308,7 @@ void MST_solver<vertex_t, edge_t, weight_t>::label_prop(vertex_t* mst_src,
 // Finds the minimum edge from each vertex to the lowest color
 template <typename vertex_t, typename edge_t, typename weight_t>
 void MST_solver<vertex_t, edge_t, weight_t>::min_edge_per_vertex() {
-  auto policy = rmm::exec_policy(stream);
+  auto policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
   thrust::fill(policy, min_edge_color.begin(), min_edge_color.end(),
                std::numeric_limits<weight_t>::max());
   thrust::fill(policy, new_mst_edge.begin(), new_mst_edge.end(),
@@ -333,7 +333,7 @@ void MST_solver<vertex_t, edge_t, weight_t>::min_edge_per_supervertex() {
   auto nthreads = std::min(v, max_threads);
   auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks);
 
-  auto policy = rmm::exec_policy(stream);
+  auto policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
   thrust::fill(policy, temp_src.begin(), temp_src.end(),
                std::numeric_limits<vertex_t>::max());
 
@@ -388,7 +388,7 @@ struct new_edges_functor {
 template <typename vertex_t, typename edge_t, typename weight_t>
 void MST_solver<vertex_t, edge_t, weight_t>::append_src_dst_pair(
   vertex_t* mst_src, vertex_t* mst_dst, weight_t* mst_weights) {
-  auto policy = rmm::exec_policy(stream);
+  auto policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
 
   auto curr_mst_edge_count = prev_mst_edge_count[0];
 
diff --git a/cpp/include/raft/sparse/op/reduce.cuh b/cpp/include/raft/sparse/op/reduce.cuh
index a86a1acb86..5856f8b1d8 100644
--- a/cpp/include/raft/sparse/op/reduce.cuh
+++ b/cpp/include/raft/sparse/op/reduce.cuh
@@ -129,7 +129,7 @@ void max_duplicates(const raft::handle_t &handle,
   auto d_alloc = handle.get_device_allocator();
   auto stream = handle.get_stream();
 
-  auto exec_policy = rmm::exec_policy(stream);
+  auto exec_policy = rmm::exec_policy(rmm::cuda_stream_view{stream});
 
   // compute diffs & take exclusive scan
   rmm::device_uvector<value_idx> diff(nnz + 1, stream);