rapidsai · BradReesWork · Aug 12, 2021 · Aug 11, 2021
diff --git a/hornet/include/Core/BatchUpdate/BatchUpdate.cuh b/hornet/include/Core/BatchUpdate/BatchUpdate.cuh
@@ -50,7 +50,8 @@
 #include "BatchUpdateKernels.cuh"
 #include "../Static/Static.cuh"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Core/BatchUpdate/BatchUpdate.i.cuh b/hornet/include/Core/BatchUpdate/BatchUpdate.i.cuh
@@ -35,7 +35,9 @@
  */
 #include "Host/Metaprogramming.hpp"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
+
 using namespace rmm;
 
 template <typename T>
@@ -267,7 +269,7 @@ remove_duplicates_edges_only(
     cudaStream_t stream{nullptr};
     auto end_ptr =
         thrust::unique_copy(
-                rmm::exec_policy(stream)->on(stream),
+                rmm::exec_policy(stream),
                 begin_in_tuple, begin_in_tuple + nE,
                 begin_out_tuple,
                 IsSrcDstEqual());
@@ -304,7 +306,7 @@ remove_duplicates(
     cudaStream_t stream{nullptr};
     auto end_ptr =
         thrust::unique_copy(
-                rmm::exec_policy(stream)->on(stream),
+                rmm::exec_policy(stream),
                 begin_in_tuple, begin_in_tuple + nE,
                 begin_out_tuple,
                 IsSrcDstEqual());
@@ -332,7 +334,7 @@ remove_duplicates(
     cudaStream_t stream{nullptr};
     auto end_ptr =
         thrust::unique_copy(
-                rmm::exec_policy(stream)->on(stream),
+                rmm::exec_policy(stream),
                 begin_in_tuple, begin_in_tuple + nE,
                 begin_out_tuple,
                 IsSrcDstEqual());
@@ -694,7 +696,7 @@ locateEdgesToBeErased(
                 batch_src_out, destination_edges.begin()));
                 //realloc_sources.begin(), destination_edges.begin()));
     cudaStream_t stream{nullptr};
-    _nE = thrust::copy_if(rmm::exec_policy(stream)->on(stream),
+    _nE = thrust::copy_if(rmm::exec_policy(stream),
             ptr_tuple, ptr_tuple + _nE,
             batch_erase_flag.begin(),
             out_ptr_tuple,

diff --git a/hornet/include/Core/BatchUpdate/BatchUpdateKernels.cuh b/hornet/include/Core/BatchUpdate/BatchUpdateKernels.cuh
@@ -1,6 +1,7 @@
 #include "../Conf/EdgeOperations.cuh"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Core/HornetInitialize/HornetInitialize.i.cuh b/hornet/include/Core/HornetInitialize/HornetInitialize.i.cuh
@@ -35,7 +35,8 @@
  */
 #include "../SoA/SoAData.cuh"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Core/HornetInitialize/HornetStaticInitialize.i.cuh b/hornet/include/Core/HornetInitialize/HornetStaticInitialize.i.cuh
@@ -1,6 +1,7 @@
 #include "../SoA/SoAData.cuh"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 
@@ -119,7 +120,7 @@ HORNETSTATIC::
 max_degree_id() const noexcept {
     auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
     cudaStream_t stream{nullptr};    
-    auto* iter = thrust::max_element(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV);
+    auto* iter = thrust::max_element(rmm::exec_policy(stream), start_ptr, start_ptr + _nV);
     if (iter == start_ptr + _nV) {
         return static_cast<vid_t>(-1);
     } else {
@@ -134,7 +135,7 @@ max_degree() const noexcept {
     auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
     cudaStream_t stream{nullptr};
 
-    auto* iter = thrust::max_element(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV);
+    auto* iter = thrust::max_element(rmm::exec_policy(stream), start_ptr, start_ptr + _nV);
     if (iter == start_ptr + _nV) {
         return static_cast<degree_t>(0);
     } else {

diff --git a/hornet/include/Core/HornetOperations/HornetQuery.i.cuh b/hornet/include/Core/HornetOperations/HornetQuery.i.cuh
@@ -5,7 +5,8 @@
 #include <thrust/execution_policy.h>
 #include "../SoA/SoAData.cuh"
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 
@@ -18,7 +19,7 @@ namespace gpu {
   max_degree_id() const noexcept {
       auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
       cudaStream_t stream{nullptr};
-      auto* iter = thrust::max_element(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV);
+      auto* iter = thrust::max_element(rmm::exec_policy(stream), start_ptr, start_ptr + _nV);
       if (iter == start_ptr + _nV) {
           return static_cast<vid_t>(-1);
       } else {
@@ -32,7 +33,7 @@ namespace gpu {
   max_degree() const noexcept {
       auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
       cudaStream_t stream{nullptr};
-      auto* iter = thrust::max_element(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV);
+      auto* iter = thrust::max_element(rmm::exec_policy(stream), start_ptr, start_ptr + _nV);
       if (iter == start_ptr + _nV) {
           return static_cast<degree_t>(0);
       } else {
@@ -88,8 +89,8 @@ namespace gpu {
     auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
 
     cudaStream_t stream{nullptr};
-    thrust::copy(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV, offset.begin());
-    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream), offset.begin(), offset.end(), offset.begin());
+    thrust::copy(rmm::exec_policy(stream), start_ptr, start_ptr + _nV, offset.begin());
+    thrust::exclusive_scan(rmm::exec_policy(stream), offset.begin(), offset.end(), offset.begin());
 
     HornetDeviceT hornet_device = device();
     const int BLOCK_SIZE = 256;
@@ -136,8 +137,8 @@ namespace gpu {
     auto start_ptr = _vertex_data.get_soa_ptr().template get<0>();
     cudaStream_t stream{nullptr};
 
-    thrust::copy(rmm::exec_policy(stream)->on(stream), start_ptr, start_ptr + _nV, degree.begin());
-    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream), degree.begin(), degree.end(), degree.begin());
+    thrust::copy(rmm::exec_policy(stream), start_ptr, start_ptr + _nV, degree.begin());
+    thrust::exclusive_scan(rmm::exec_policy(stream), degree.begin(), degree.end(), degree.begin());
 
     HornetDeviceT hornet_device = device();
     const int BLOCK_SIZE = 256;

diff --git a/hornet/include/Core/HornetOperations/HornetSort.i.cuh b/hornet/include/Core/HornetOperations/HornetSort.i.cuh
@@ -15,6 +15,8 @@
  */
 #include <limits>
 
+#include <rmm/device_vector.hpp>
+
 namespace hornet {
 
 namespace gpu {
@@ -60,13 +62,13 @@ sort(void) {
 
   rmm::device_vector<degree_t> offsets(_nV + 1);
   degree_t * vertex_degrees = _vertex_data.get_soa_ptr().template get<0>();
-  thrust::transform(rmm::exec_policy(stream)->on(stream),
+  thrust::transform(rmm::exec_policy(stream),
       vertex_degrees, vertex_degrees + _nV,
       offsets.begin(),
       InvalidEdgeCount<degree_t>());
     CHECK_CUDA_ERROR
 
-  thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
       offsets.begin(), offsets.end(), offsets.begin());
     CHECK_CUDA_ERROR
 

diff --git a/hornet/include/Core/MemoryManager/lrb.cuh b/hornet/include/Core/MemoryManager/lrb.cuh
@@ -1,7 +1,8 @@
 #ifndef LRB_CUH
 #define LRB_CUH
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 #include <vector>
 
 ////////////////////////////////////////////////////////////////

diff --git a/hornet/include/Core/SoA/SoAData.cuh b/hornet/include/Core/SoA/SoAData.cuh
@@ -44,7 +44,8 @@
 #include <thrust/gather.h>
 #include <vector>
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Core/SoA/impl/SoAData.i.cuh b/hornet/include/Core/SoA/impl/SoAData.i.cuh
@@ -34,7 +34,8 @@
  * </blockquote>}
  */
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Core/SoA/impl/SoADataSort.i.cuh b/hornet/include/Core/SoA/impl/SoADataSort.i.cuh
@@ -3,6 +3,7 @@
 
 #include <cub/cub.cuh>
 #include <rmm/device_buffer.hpp>
+#include <rmm/device_vector.hpp>
 
 namespace hornet {
 
@@ -267,7 +268,7 @@ cub_block_segmented_sort(CSoAPtr<EdgeTypes...> &soa, degree_t capacity, degree_t
 
   cudaStream_t stream{nullptr};
   rmm::device_vector<degree_t> index(capacity);
-  thrust::sequence(rmm::exec_policy(stream)->on(stream), index.begin(), index.end());
+  thrust::sequence(rmm::exec_policy(stream), index.begin(), index.end());
   using T0 = typename xlib::SelectType<0, EdgeTypes...>::type;
   T0 * key = temp_soa.template get<0>();
   using T1 = degree_t;
@@ -299,7 +300,7 @@ cub_segmented_sort(CSoAPtr<EdgeTypes...> &soa, degree_t capacity, degree_t segme
 
   degree_t offset_count = capacity/segment_length;
   rmm::device_vector<degree_t> offsets(offset_count + 1);
-  thrust::transform(rmm::exec_policy(stream)->on(stream),
+  thrust::transform(rmm::exec_policy(stream),
       offsets.begin(), offsets.end(),
       thrust::make_constant_iterator(segment_length),
       offsets.begin(),
@@ -333,7 +334,7 @@ cub_segmented_sort(CSoAPtr<EdgeTypes...> &soa, degree_t capacity, degree_t segme
 
   degree_t offset_count = capacity/segment_length;
   rmm::device_vector<degree_t> offsets(offset_count + 1);
-  thrust::transform(rmm::exec_policy(stream)->on(stream),
+  thrust::transform(rmm::exec_policy(stream),
       offsets.begin(), offsets.end(),
       thrust::make_constant_iterator(segment_length),
       offsets.begin(),
@@ -371,7 +372,7 @@ cub_segmented_sort(CSoAPtr<EdgeTypes...> &soa, degree_t capacity, degree_t segme
 
   degree_t offset_count = capacity/segment_length;
   rmm::device_vector<degree_t> offsets(offset_count + 1);
-  thrust::transform(rmm::exec_policy(stream)->on(stream),
+  thrust::transform(rmm::exec_policy(stream),
       offsets.begin(), offsets.end(),
       thrust::make_constant_iterator(segment_length),
       offsets.begin(),

diff --git a/hornet/include/Core/SoA/impl/SoAPtr.i.cuh b/hornet/include/Core/SoA/impl/SoAPtr.i.cuh
@@ -34,7 +34,8 @@
  * </blockquote>}
  */
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 #include <thrust/host_vector.h>
 
 using namespace rmm;
@@ -277,7 +278,7 @@ struct RecursiveGather {
         if (N >= SIZE) { return; }
         cudaStream_t stream{nullptr};
         thrust::gather(
-                rmm::exec_policy(stream)->on(stream),
+                rmm::exec_policy(stream),
                 map.begin(), map.begin() + nE,
                 src.template get<N>(),
                 dst.template get<N>());
@@ -595,11 +596,11 @@ sort_edges(Ptr<EdgeTypes...> ptr, const degree_t nE) {
     cudaStream_t stream{nullptr};
 
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             ptr.template get<1>(), ptr.template get<1>() + nE,
             ptr.template get<0>());
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             ptr.template get<0>(), ptr.template get<0>() + nE,
             ptr.template get<1>());
 }
@@ -619,11 +620,11 @@ sort_batch(Ptr<EdgeTypes...> in_ptr, const degree_t nE, rmm::device_vector<degre
     cudaStream_t stream{nullptr};
 
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             in_ptr.template get<1>(), in_ptr.template get<1>() + nE,
             thrust::make_zip_iterator(thrust::make_tuple(in_ptr.template get<0>(), in_ptr.template get<2>())) );
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             in_ptr.template get<0>(), in_ptr.template get<0>() + nE,
             thrust::make_zip_iterator(thrust::make_tuple(in_ptr.template get<1>(), in_ptr.template get<2>())) );
     return false;
@@ -638,11 +639,11 @@ sort_batch(Ptr<EdgeTypes...> in_ptr, const degree_t nE, rmm::device_vector<degre
     cudaStream_t stream{nullptr};
 
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             in_ptr.template get<1>(), in_ptr.template get<1>() + nE,
             thrust::make_zip_iterator(thrust::make_tuple(in_ptr.template get<0>(), range.begin())) );
     thrust::sort_by_key(
-            rmm::exec_policy(stream)->on(stream),
+            rmm::exec_policy(stream),
             in_ptr.template get<0>(), in_ptr.template get<0>() + nE,
             thrust::make_zip_iterator(thrust::make_tuple(in_ptr.template get<1>(), range.begin())) );
     //FIXME : Check correctness of RecursiveCopy and RecursiveGather template parameters

diff --git a/hornet/include/Core/Static/Static.cuh b/hornet/include/Core/Static/Static.cuh
@@ -10,7 +10,8 @@
 #include "../Hornet.cuh"
 #include <map>
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
+#include <rmm/device_vector.hpp>
 
 using namespace rmm;
 

diff --git a/hornet/include/Util/RandomGraphData.cuh b/hornet/include/Util/RandomGraphData.cuh
@@ -9,6 +9,8 @@
 #include <thrust/random/linear_congruential_engine.h>
 #include <thrust/random/uniform_int_distribution.h>
 
+#include <rmm/device_vector.hpp>
+
 namespace hornet {
 
 template <typename T>

diff --git a/hornetsnest/include/Static/KTruss/KTruss.impl.cuh b/hornetsnest/include/Static/KTruss/KTruss.impl.cuh
@@ -34,7 +34,7 @@
 #include <iostream>
 #include <Device/Util/Timer.cuh>
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
 
 using namespace std;
 using namespace rmm;
@@ -94,7 +94,7 @@ void KTruss::createOffSetArray(){
     forAllVertices(hornet, getVertexSizes {tempSize});
 
     cudaStream_t stream{nullptr};
-    thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream), tempSize, tempSize + originalNV, hd_data().offset_array+1);
+    thrust::inclusive_scan(rmm::exec_policy(stream), tempSize, tempSize + originalNV, hd_data().offset_array+1);
 }
 
 void KTruss::copyOffsetArrayHost(const vert_t* host_offset_array) {

diff --git a/hornetsnest/include/Static/KTruss/KTrussWeighted.impl.cuh b/hornetsnest/include/Static/KTruss/KTrussWeighted.impl.cuh
@@ -34,7 +34,7 @@
 #include <iostream>
 #include <Device/Util/Timer.cuh>
 
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
 
 using namespace std;
 using namespace rmm;
@@ -100,7 +100,7 @@ void KTrussWeighted<T>::createOffSetArray(){
     forAllVertices(hnt, getVertexSizes {tempSize});
 
     cudaStream_t stream{nullptr};
-    thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream), tempSize, tempSize + originalNV, hd_data().offset_array+1);
+    thrust::inclusive_scan(rmm::exec_policy(stream), tempSize, tempSize + originalNV, hd_data().offset_array+1);
 }
 
 template <typename T>

diff --git a/xlib/include/Device/Primitives/CubWrapper.cuh b/xlib/include/Device/Primitives/CubWrapper.cuh
@@ -50,8 +50,9 @@
 #include "Host/Numeric.hpp"
 
 #include <cub/cub.cuh>
-#include <rmm/thrust_rmm_allocator.h>
+#include <rmm/exec_policy.hpp>
 #include <rmm/device_buffer.hpp>
+#include <rmm/device_vector.hpp>
 
 namespace xlib {
 

diff --git a/xlib/include/Device/Util/impl/Basic.i.cuh b/xlib/include/Device/Util/impl/Basic.i.cuh
@@ -204,7 +204,7 @@ unsigned discontinuity_mask(const T& value1, const T& value2, bool& lane_bit,
 template<unsigned WARP_SZ>
 constexpr unsigned warp_segmask() {
     unsigned value = 0;
-    for (int i = 0; i < xlib::WARP_SIZE; i += WARP_SZ) {
+    for (size_t i = 0; i < xlib::WARP_SIZE; i += WARP_SZ) {
         value <<= WARP_SZ;
         value |= 1;
     }