diff --git a/hornet/include/Core/BatchUpdate/BatchUpdate.cuh b/hornet/include/Core/BatchUpdate/BatchUpdate.cuh index b4549776..54801467 100644 --- 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 +#include +#include using namespace rmm; diff --git a/hornet/include/Core/BatchUpdate/BatchUpdate.i.cuh b/hornet/include/Core/BatchUpdate/BatchUpdate.i.cuh index a156efdd..8d59cc2d 100644 --- 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 +#include +#include + using namespace rmm; template @@ -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 index 0413f3ba..294831d5 100644 --- a/hornet/include/Core/BatchUpdate/BatchUpdateKernels.cuh +++ b/hornet/include/Core/BatchUpdate/BatchUpdateKernels.cuh @@ -1,6 +1,7 @@ #include "../Conf/EdgeOperations.cuh" -#include +#include +#include using namespace rmm; diff --git a/hornet/include/Core/HornetInitialize/HornetInitialize.i.cuh b/hornet/include/Core/HornetInitialize/HornetInitialize.i.cuh index 657fd6da..af6237df 100644 --- 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 +#include +#include using namespace rmm; diff --git a/hornet/include/Core/HornetInitialize/HornetStaticInitialize.i.cuh b/hornet/include/Core/HornetInitialize/HornetStaticInitialize.i.cuh index f2ac516f..f8fa2777 100644 --- 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 +#include +#include 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(-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(0); } else { diff --git a/hornet/include/Core/HornetOperations/HornetQuery.i.cuh b/hornet/include/Core/HornetOperations/HornetQuery.i.cuh index 74d6161b..952773c5 100644 --- a/hornet/include/Core/HornetOperations/HornetQuery.i.cuh +++ b/hornet/include/Core/HornetOperations/HornetQuery.i.cuh @@ -5,7 +5,8 @@ #include #include "../SoA/SoAData.cuh" -#include +#include +#include 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(-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(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 index 77fae31e..ec13e820 100644 --- a/hornet/include/Core/HornetOperations/HornetSort.i.cuh +++ b/hornet/include/Core/HornetOperations/HornetSort.i.cuh @@ -15,6 +15,8 @@ */ #include +#include + namespace hornet { namespace gpu { @@ -60,13 +62,13 @@ sort(void) { rmm::device_vector 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()); 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 index 07ad7edf..e8fba27f 100644 --- 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 +#include +#include #include //////////////////////////////////////////////////////////////// diff --git a/hornet/include/Core/SoA/SoAData.cuh b/hornet/include/Core/SoA/SoAData.cuh index e731f53f..46e29576 100644 --- a/hornet/include/Core/SoA/SoAData.cuh +++ b/hornet/include/Core/SoA/SoAData.cuh @@ -44,7 +44,8 @@ #include #include -#include +#include +#include using namespace rmm; diff --git a/hornet/include/Core/SoA/impl/SoAData.i.cuh b/hornet/include/Core/SoA/impl/SoAData.i.cuh index bd28597e..b34de32a 100644 --- a/hornet/include/Core/SoA/impl/SoAData.i.cuh +++ b/hornet/include/Core/SoA/impl/SoAData.i.cuh @@ -34,7 +34,8 @@ * } */ -#include +#include +#include using namespace rmm; diff --git a/hornet/include/Core/SoA/impl/SoADataSort.i.cuh b/hornet/include/Core/SoA/impl/SoADataSort.i.cuh index 05bf9043..8f64b125 100644 --- a/hornet/include/Core/SoA/impl/SoADataSort.i.cuh +++ b/hornet/include/Core/SoA/impl/SoADataSort.i.cuh @@ -3,6 +3,7 @@ #include #include +#include namespace hornet { @@ -267,7 +268,7 @@ cub_block_segmented_sort(CSoAPtr &soa, degree_t capacity, degree_t cudaStream_t stream{nullptr}; rmm::device_vector 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 &soa, degree_t capacity, degree_t segme degree_t offset_count = capacity/segment_length; rmm::device_vector 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 &soa, degree_t capacity, degree_t segme degree_t offset_count = capacity/segment_length; rmm::device_vector 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 &soa, degree_t capacity, degree_t segme degree_t offset_count = capacity/segment_length; rmm::device_vector 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 index fe6b81f7..6204ed3b 100644 --- a/hornet/include/Core/SoA/impl/SoAPtr.i.cuh +++ b/hornet/include/Core/SoA/impl/SoAPtr.i.cuh @@ -34,7 +34,8 @@ * } */ -#include +#include +#include #include 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(), dst.template get()); @@ -595,11 +596,11 @@ sort_edges(Ptr 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 in_ptr, const degree_t nE, rmm::device_vectoron(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 in_ptr, const degree_t nE, rmm::device_vectoron(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 index 76e469f0..e9ddceea 100644 --- a/hornet/include/Core/Static/Static.cuh +++ b/hornet/include/Core/Static/Static.cuh @@ -10,7 +10,8 @@ #include "../Hornet.cuh" #include -#include +#include +#include using namespace rmm; diff --git a/hornet/include/Util/RandomGraphData.cuh b/hornet/include/Util/RandomGraphData.cuh index bf4f6c85..9f5839eb 100644 --- a/hornet/include/Util/RandomGraphData.cuh +++ b/hornet/include/Util/RandomGraphData.cuh @@ -9,6 +9,8 @@ #include #include +#include + namespace hornet { template diff --git a/hornetsnest/include/Static/KTruss/KTruss.impl.cuh b/hornetsnest/include/Static/KTruss/KTruss.impl.cuh index 15fb0433..8c8f8a28 100644 --- a/hornetsnest/include/Static/KTruss/KTruss.impl.cuh +++ b/hornetsnest/include/Static/KTruss/KTruss.impl.cuh @@ -34,7 +34,7 @@ #include #include -#include +#include 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 index 8de19d09..1ea97c15 100644 --- a/hornetsnest/include/Static/KTruss/KTrussWeighted.impl.cuh +++ b/hornetsnest/include/Static/KTruss/KTrussWeighted.impl.cuh @@ -34,7 +34,7 @@ #include #include -#include +#include using namespace std; using namespace rmm; @@ -100,7 +100,7 @@ void KTrussWeighted::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 diff --git a/xlib/include/Device/Primitives/CubWrapper.cuh b/xlib/include/Device/Primitives/CubWrapper.cuh index 8d3bf5bc..8ffed2d4 100644 --- a/xlib/include/Device/Primitives/CubWrapper.cuh +++ b/xlib/include/Device/Primitives/CubWrapper.cuh @@ -50,8 +50,9 @@ #include "Host/Numeric.hpp" #include -#include +#include #include +#include namespace xlib { diff --git a/xlib/include/Device/Util/impl/Basic.i.cuh b/xlib/include/Device/Util/impl/Basic.i.cuh index f4b4c7a6..cfbc5d88 100644 --- 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 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; }