diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 108cb0748a8..b0365c3cfd6 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -113,7 +113,6 @@ set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-extended-lambda --expt-relaxed- set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Werror=cross-execution-space-call -Wno-deprecated-declarations -Xptxas --disable-warnings") set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -Wall,-Wno-error=sign-compare,-Wno-error=unused-but-set-variable") - # Option to enable line info in CUDA device compilation to allow introspection when profiling / # memchecking option(CMAKE_CUDA_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" OFF) @@ -298,7 +297,8 @@ else(DEFINED ENV{RAFT_PATH}) FetchContent_Declare( raft GIT_REPOSITORY https://github.com/rapidsai/raft.git - GIT_TAG a3461b201ea1c9f61571f1927274f739e775d2d2 + GIT_TAG b055cf862a599fd45537d21a309edd8a6e06da4c + SOURCE_SUBDIR raft ) @@ -446,6 +446,10 @@ target_link_directories(cugraph # add_dependencies(cugraph gunrock_ext) +# Per-thread default stream option see https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html +# The per-thread default stream does not synchronize with other streams +target_compile_definitions(cugraph PUBLIC CUDA_API_PER_THREAD_DEFAULT_STREAM) + ################################################################################################### # - include paths --------------------------------------------------------------------------------- target_include_directories(cugraph diff --git a/cpp/include/algorithms.hpp b/cpp/include/algorithms.hpp index 0b8bd59587f..c3a4f3ec985 100644 --- a/cpp/include/algorithms.hpp +++ b/cpp/include/algorithms.hpp @@ -1167,7 +1167,7 @@ void katz_centrality(raft::handle_t const &handle, * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) * or multi-GPU (true). * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and - * handles to various CUDA libraries) to run graph algorithms. + * handles to various CUDA libraries) to run graph algorithms. Must have at least one worker stream. * @param graph_view Graph view object of, we extract induced egonet subgraphs from @p graph_view. * @param source_vertex Pointer to egonet center vertices (size == @p n_subgraphs). * @param n_subgraphs Number of induced EgoNet subgraphs to extract (ie. number of elements in @p diff --git a/cpp/src/community/egonet.cu b/cpp/src/community/egonet.cu index fa788aa307b..067d27f9a92 100644 --- a/cpp/src/community/egonet.cu +++ b/cpp/src/community/egonet.cu @@ -22,6 +22,9 @@ #include #include +#include +#include + #include #include @@ -34,6 +37,8 @@ #include #include +#include + namespace { /* @@ -61,58 +66,111 @@ extract( vertex_t n_subgraphs, vertex_t radius) { - auto v = csr_view.get_number_of_vertices(); - auto e = csr_view.get_number_of_edges(); - auto stream = handle.get_stream(); - float avg_degree = e / v; + auto v = csr_view.get_number_of_vertices(); + auto e = csr_view.get_number_of_edges(); + auto user_stream_view = handle.get_stream_view(); rmm::device_vector neighbors_offsets(n_subgraphs + 1); rmm::device_vector neighbors; - // It is the right thing to accept device memory for source_vertex - // FIXME consider adding a device API to BFS (ie. accept source on the device) std::vector h_source_vertex(n_subgraphs); - raft::update_host(&h_source_vertex[0], source_vertex, n_subgraphs, stream); + std::vector h_neighbors_offsets(n_subgraphs + 1); + + raft::update_host(&h_source_vertex[0], source_vertex, n_subgraphs, user_stream_view.value()); + + // Streams will allocate concurrently later + std::vector> reached{}; + reached.reserve(handle.get_num_internal_streams()); - // reserve some reasonable memory, but could grow larger than that - neighbors.reserve(v + avg_degree * n_subgraphs * radius); - neighbors_offsets[0] = 0; - // each source should be done concurently in the future + // h_source_vertex[i] is used by other streams in the for loop + user_stream_view.synchronize(); +#ifdef TIMING + HighResTimer hr_timer; + hr_timer.start("ego_neighbors"); +#endif for (vertex_t i = 0; i < n_subgraphs; i++) { + // get light handle from worker pool + raft::handle_t light_handle(handle, i); + auto worker_stream_view = light_handle.get_stream_view(); + + // Allocations and operations are attached to the worker stream + rmm::device_uvector local_reach(v, worker_stream_view); + reached.push_back(std::move(local_reach)); + // BFS with cutoff - rmm::device_vector reached(v); - rmm::device_vector predecessors(v); // not used + // consider adding a device API to BFS (ie. accept source on the device) + rmm::device_uvector predecessors(v, worker_stream_view); // not used bool direction_optimizing = false; - cugraph::experimental::bfs(handle, + thrust::fill(rmm::exec_policy(worker_stream_view), + reached[i].begin(), + reached[i].end(), + std::numeric_limits::max()); + thrust::fill( + rmm::exec_policy(worker_stream_view), reached[i].begin(), reached[i].begin() + 100, 1.0); + + cugraph::experimental::bfs(light_handle, csr_view, - reached.data().get(), - predecessors.data().get(), + reached[i].data(), + predecessors.data(), h_source_vertex[i], direction_optimizing, radius); // identify reached vertex ids from distance array - thrust::transform(rmm::exec_policy(stream)->on(stream), + thrust::transform(rmm::exec_policy(worker_stream_view), thrust::make_counting_iterator(vertex_t{0}), thrust::make_counting_iterator(v), - reached.begin(), - reached.begin(), + reached[i].begin(), + reached[i].begin(), [sentinel = std::numeric_limits::max()] __device__( auto id, auto val) { return val < sentinel ? id : sentinel; }); // removes unreached data - auto reached_end = thrust::remove(rmm::exec_policy(stream)->on(stream), - reached.begin(), - reached.end(), + auto reached_end = thrust::remove(rmm::exec_policy(worker_stream_view), + reached[i].begin(), + reached[i].end(), std::numeric_limits::max()); + // release temp storage + reached[i].resize(thrust::distance(reached[i].begin(), reached_end), worker_stream_view); + reached[i].shrink_to_fit(worker_stream_view); + } - // update extraction input - size_t n_reached = thrust::distance(reached.begin(), reached_end); - neighbors_offsets[i + 1] = neighbors_offsets[i] + n_reached; - if (neighbors_offsets[i + 1] > neighbors.capacity()) - neighbors.reserve(neighbors_offsets[i + 1] * 2); - neighbors.insert(neighbors.end(), reached.begin(), reached_end); + // wait on every one to identify their neighboors before proceeding to concatenation + handle.wait_on_internal_streams(); + + // Construct neighboors offsets (just a scan on neighborhod vector sizes) + h_neighbors_offsets[0] = 0; + for (vertex_t i = 0; i < n_subgraphs; i++) { + h_neighbors_offsets[i + 1] = h_neighbors_offsets[i] + reached[i].size(); + } + raft::update_device(neighbors_offsets.data().get(), + &h_neighbors_offsets[0], + n_subgraphs + 1, + user_stream_view.value()); + neighbors.resize(h_neighbors_offsets[n_subgraphs]); + user_stream_view.synchronize(); + + // Construct the neighboors list concurrently + for (vertex_t i = 0; i < n_subgraphs; i++) { + raft::handle_t light_handle(handle, i); + auto worker_stream_view = light_handle.get_stream_view(); + thrust::copy(rmm::exec_policy(worker_stream_view), + reached[i].begin(), + reached[i].end(), + neighbors.begin() + h_neighbors_offsets[i]); + + // reached info is not needed anymore + reached[i].resize(0, worker_stream_view); + reached[i].shrink_to_fit(worker_stream_view); } + // wait on every one before proceeding to grouped extraction + handle.wait_on_internal_streams(); + +#ifdef TIMING + hr_timer.stop(); + hr_timer.display(std::cout); +#endif + // extract return cugraph::experimental::extract_induced_subgraphs( handle, csr_view, neighbors_offsets.data().get(), neighbors.data().get(), n_subgraphs); @@ -207,4 +265,4 @@ extract_ego(raft::handle_t const &, int64_t, int64_t); } // namespace experimental -} // namespace cugraph +} // namespace cugraph \ No newline at end of file diff --git a/cpp/src/experimental/induced_subgraph.cu b/cpp/src/experimental/induced_subgraph.cu index a88adf76ef4..5cda36ad7e2 100644 --- a/cpp/src/experimental/induced_subgraph.cu +++ b/cpp/src/experimental/induced_subgraph.cu @@ -32,6 +32,8 @@ #include +#include + namespace cugraph { namespace experimental { @@ -52,6 +54,10 @@ extract_induced_subgraphs( size_t num_subgraphs, bool do_expensive_check) { +#ifdef TIMING + HighResTimer hr_timer; + hr_timer.start("extract_induced_subgraphs"); +#endif // FIXME: this code is inefficient for the vertices with their local degrees much larger than the // number of vertices in the subgraphs (in this case, searching that the subgraph vertices are // included in the local neighbors is more efficient than searching the local neighbors are @@ -244,7 +250,10 @@ extract_induced_subgraphs( subgraph_offsets + (num_subgraphs + 1), subgraph_vertex_output_offsets.begin(), subgraph_edge_offsets.begin()); - +#ifdef TIMING + hr_timer.stop(); + hr_timer.display(std::cout); +#endif return std::make_tuple(std::move(edge_majors), std::move(edge_minors), std::move(edge_weights), diff --git a/cpp/src/utilities/high_res_timer.hpp b/cpp/src/utilities/high_res_timer.hpp index f2d6bc6e13f..a731c5edc9d 100644 --- a/cpp/src/utilities/high_res_timer.hpp +++ b/cpp/src/utilities/high_res_timer.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,8 @@ #include #include +//#define TIMING + class HighResTimer { public: HighResTimer() : timers() {} diff --git a/cpp/tests/community/egonet_test.cu b/cpp/tests/community/egonet_test.cu index ef2699bd1d0..a9224b42bc1 100644 --- a/cpp/tests/community/egonet_test.cu +++ b/cpp/tests/community/egonet_test.cu @@ -15,6 +15,7 @@ */ #include +#include #include #include @@ -35,6 +36,8 @@ #include #include +#include + typedef struct InducedEgo_Usecase_t { std::string graph_file_full_path{}; std::vector ego_sources{}; @@ -67,7 +70,8 @@ class Tests_InducedEgo : public ::testing::TestWithParam { template void run_current_test(InducedEgo_Usecase const& configuration) { - raft::handle_t handle{}; + int n_streams = std::min(configuration.ego_sources.size(), static_cast(128)); + raft::handle_t handle(n_streams); cugraph::experimental::graph_t graph( handle); @@ -88,14 +92,18 @@ class Tests_InducedEgo : public ::testing::TestWithParam { rmm::device_uvector d_ego_edgelist_dst(0, handle.get_stream()); rmm::device_uvector d_ego_edgelist_weights(0, handle.get_stream()); rmm::device_uvector d_ego_edge_offsets(0, handle.get_stream()); - + HighResTimer hr_timer; + hr_timer.start("egonet"); + cudaProfilerStart(); std::tie(d_ego_edgelist_src, d_ego_edgelist_dst, d_ego_edgelist_weights, d_ego_edge_offsets) = cugraph::experimental::extract_ego(handle, graph_view, d_ego_sources.data(), static_cast(configuration.ego_sources.size()), configuration.radius); - + cudaProfilerStop(); + hr_timer.stop(); + hr_timer.display(std::cout); std::vector h_cugraph_ego_edge_offsets(d_ego_edge_offsets.size()); std::vector h_cugraph_ego_edgelist_src(d_ego_edgelist_src.size()); std::vector h_cugraph_ego_edgelist_dst(d_ego_edgelist_dst.size()); @@ -118,13 +126,11 @@ class Tests_InducedEgo : public ::testing::TestWithParam { ASSERT_TRUE(h_cugraph_ego_edge_offsets[configuration.ego_sources.size()] == d_ego_edgelist_src.size()); for (size_t i = 0; i < configuration.ego_sources.size(); i++) - ASSERT_TRUE(h_cugraph_ego_edge_offsets[i] < h_cugraph_ego_edge_offsets[i + 1]); + ASSERT_TRUE(h_cugraph_ego_edge_offsets[i] <= h_cugraph_ego_edge_offsets[i + 1]); auto n_vertices = graph_view.get_number_of_vertices(); for (size_t i = 0; i < d_ego_edgelist_src.size(); i++) { - ASSERT_TRUE(h_cugraph_ego_edgelist_src[i] >= 0); - ASSERT_TRUE(h_cugraph_ego_edgelist_src[i] < n_vertices); - ASSERT_TRUE(h_cugraph_ego_edgelist_dst[i] >= 0); - ASSERT_TRUE(h_cugraph_ego_edgelist_dst[i] < n_vertices); + ASSERT_TRUE(cugraph::test::is_valid_vertex(n_vertices, h_cugraph_ego_edgelist_src[i])); + ASSERT_TRUE(cugraph::test::is_valid_vertex(n_vertices, h_cugraph_ego_edgelist_dst[i])); } /* @@ -170,6 +176,156 @@ INSTANTIATE_TEST_CASE_P( InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{1}, 3, false), InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{10, 0, 5}, 2, false), InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{9, 3, 10}, 2, false), - InducedEgo_Usecase("test/datasets/karate.mtx", std::vector{5, 12, 13}, 2, true))); + InducedEgo_Usecase( + "test/datasets/karate.mtx", std::vector{5, 9, 3, 10, 12, 13}, 2, true))); +// For perf analysis +/* +INSTANTIATE_TEST_CASE_P( +simple_test, +Tests_InducedEgo, +::testing::Values( +InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 1, false), +InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 2, false), +InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 3, false), +InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 4, false), +InducedEgo_Usecase("test/datasets/soc-LiveJournal1.mtx", std::vector{0}, 5, false), +InducedEgo_Usecase( +"test/datasets/soc-LiveJournal1.mtx", std::vector{363617}, 2, false), +InducedEgo_Usecase( +"test/datasets/soc-LiveJournal1.mtx", +std::vector{ + 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755}, + 2, + false), + InducedEgo_Usecase( + "test/datasets/soc-LiveJournal1.mtx", + std::vector{ + 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755, + 2536834, 3070144, 3888415, 3131712, 2382526, 1040771, 2631543, 4607218, 4465829, + 3341686, 2772973, 2611175, 4526129, 2624421, 1220593, 2593137, 3270705, 1503899, + 1213033, 4840102, 4529036, 3421116, 4264831, 4089751, 4272322, 3486998, 2830318, + 320953, 2388331, 520808, 3023094, 1600294, 3631119, 1716614, 4829213, 1175844, 960680, + 847662, 3277365, 3957318, 3455123, 2454259, 670953, 4465677, 1027332, 2560721, 89061, + 1163406, 3109528, 3221856, 4714426, 2382774, 37828, 4433616, 3283229, 591911, + 4200188, 442522, 872207, 2437601, 741003, 266241, 914618, 3626195, 2021080, + 4679624, 777476, 2527796, 1114017, 640142, 49259, 4069879, 3869098, 1105040, + 4707804, 3208582, 3325885, 1450601, 4072548, 2037062, 2029646, 4575891, 1488598, 79105, + 4827273, 3795434, 4647518, 4733397, 3980718, 1184627}, + 2, + false), + InducedEgo_Usecase( + "test/datasets/soc-LiveJournal1.mtx", + std::vector{ + 363617, 722214, 2337449, 2510183, 2513389, 225853, 2035807, 3836330, 1865496, 28755, + 2536834, 3070144, 3888415, 3131712, 2382526, 1040771, 2631543, 4607218, 4465829, + 3341686, 2772973, 2611175, 4526129, 2624421, 1220593, 2593137, 3270705, 1503899, + 1213033, 4840102, 4529036, 3421116, 4264831, 4089751, 4272322, 3486998, 2830318, + 320953, 2388331, 520808, 3023094, 1600294, 3631119, 1716614, 4829213, 1175844, 960680, + 847662, 3277365, 3957318, 3455123, 2454259, 670953, 4465677, 1027332, 2560721, 89061, + 1163406, 3109528, 3221856, 4714426, 2382774, 37828, 4433616, 3283229, 591911, + 4200188, 442522, 872207, 2437601, 741003, 266241, 914618, 3626195, 2021080, + 4679624, 777476, 2527796, 1114017, 640142, 49259, 4069879, 3869098, 1105040, + 4707804, 3208582, 3325885, 1450601, 4072548, 2037062, 2029646, 4575891, 1488598, 79105, + 4827273, 3795434, 4647518, 4733397, 3980718, 1184627, 984983, 3114832, 1967741, + 1599818, 144593, 2698770, 2889449, 2495550, 1053813, 1193622, 686026, 3989015, + 2040719, 4693428, 3190376, 2926728, 3399030, 1664419, 662429, 4526841, 2186957, + 3752558, 2440046, 2930226, 3633006, 4058166, 3137060, 3499296, 2126343, 148971, + 2199672, 275811, 2813976, 2274536, 1189239, 1335942, 2465624, 2596042, 829684, 193400, + 2682845, 3691697, 4022437, 4051170, 4195175, 2876420, 3984220, 2174475, 326134, + 2606530, 2493046, 4706121, 1498980, 4576225, 1271339, 44832, 1875673, 4664940, + 134931, 736397, 4333554, 2751031, 2163610, 2879676, 3174153, 3317403, 2052464, + 1881883, 4757859, 3596257, 2358088, 2578758, 447504, 590720, 1717038, 1869795, + 1133885, 3027521, 840312, 2818881, 3654321, 2730947, 353585, 1134903, 2223378, + 1508824, 3662521, 1363776, 2712071, 288441, 1204581, 3502242, 4645567, 2767267, + 1514366, 3956099, 1422145, 1216608, 2253360, 189132, 4238225, 1345783, 451571, 1599442, + 3237284, 4711405, 929446, 1857675, 150759, 1277633, 761210, 138628, 1026833, + 2599544, 2464737, 989203, 3399615, 2144292, 216142, 637312, 2044964, 716256, 1660632, + 1762919, 4784357, 2213415, 2764769, 291806, 609772, 3264819, 1870953, 1516385, + 235647, 1045474, 2664957, 819095, 1824119, 4045271, 4448109, 1676788, 4285177, + 1580502, 3546548, 2771971, 3927086, 1339779, 3156204, 1730998, 1172522, 2433024, + 4533449, 479930, 2010695, 672994, 3542039, 3176455, 26352, 2137735, 866910, + 4410835, 2623982, 3603159, 2555625, 2765653, 267865, 2015523, 1009052, 4713994, + 1600667, 2176195, 3179631, 4570390, 2018424, 3356384, 1784287, 894861, 3622099, + 1647273, 3044136, 950354, 1491760, 3416929, 3757300, 2244912, 4129215, 1600848, + 3867343, 72329, 919189, 992521, 3445975, 4712557, 4680974, 188419, 2612093, + 1991268, 3566207, 2281468, 3859078, 2492806, 3398628, 763441, 2679107, 2554420, + 2130132, 4664374, 1182901, 3890770, 4714667, 4209303, 4013060, 3617653, 2040022, + 3296519, 4190671, 1693353, 2678411, 3788834, 2781815, 191965, 1083926, 503974, 3529226, + 1650522, 1900976, 542080, 3423929, 3418905, 878165, 4701703, 3022790, 4316365, 76365, + 4053672, 1358185, 3830478, 4445661, 3210024, 1895915, 4541133, 2938808, 562788, + 3920065, 1458776, 4052046, 2967475, 1092809, 3203538, 159626, 3399464, 214467, + 3343982, 1811854, 3189045, 4272117, 4701563, 424807, 4341116, 760545, 4674683, + 1538018, 386762, 194237, 2162719, 1694433, 943728, 2389036, 2196653, 3085571, + 1513424, 3689413, 3278747, 4197291, 3324063, 3651090, 1737936, 2768803, 2768889, + 3108096, 4311775, 3569480, 886705, 733256, 2477493, 1735412, 2960895, 1983781, + 1861797, 3566460, 4537673, 1164093, 3499764, 4553071, 3518985, 847658, 918948, + 2922351, 1056144, 652895, 1013195, 780505, 1702928, 3562838, 1432719, 2405207, + 1054920, 641647, 2240939, 3617702, 383165, 652641, 879593, 1810739, 2096385, + 4497865, 4768530, 1743968, 3582014, 1025009, 3002122, 2422190, 527647, 1251821, + 2571153, 4095874, 3705333, 3637407, 1385567, 4043855, 4041930, 2433139, 1710383, + 1127734, 4362316, 711588, 817839, 3214775, 910077, 1313768, 2382229, 16864, 2081770, + 3095420, 3195272, 548711, 2259860, 1167323, 2435974, 425238, 2085179, 2630042, + 2632881, 2867923, 3703565, 1037695, 226617, 4379130, 1541468, 3581937, 605965, + 1137674, 4655221, 4769963, 1394370, 4425315, 2990132, 2364485, 1561137, 2713384, + 481509, 2900382, 934766, 2986774, 1767669, 298593, 2502539, 139296, 3794229, + 4002180, 4718138, 2909238, 423691, 3023810, 2784924, 2760160, 1971980, 316683, + 3828090, 3253691, 4839313, 1203624, 584938, 3901482, 1747543, 1572737, 3533226, + 774708, 1691195, 1037110, 1557763, 225120, 4424243, 3524086, 1717663, 4332507, + 3513592, 4274932, 1232118, 873498, 1416042, 2488925, 111391, 4704545, 4492545, + 445317, 1584812, 2187737, 2471948, 3731678, 219255, 2282627, 2589971, 2372185, + 4609096, 3673961, 2524410, 12823, 2437155, 3015974, 4188352, 3184084, 3690756, + 1222341, 1278376, 3652030, 4162647, 326548, 3930062, 3926100, 1551222, 2722165, + 4526695, 3997534, 4815513, 3139056, 2547644, 3028915, 4149092, 3656554, 2691582, + 2676699, 1878842, 260174, 3129900, 4379993, 182347, 2189338, 3783616, 2616666, + 2596952, 243007, 4179282, 2730, 1939894, 2332032, 3335636, 182332, 3112260, + 2174584, 587481, 4527368, 3154106, 3403059, 673206, 2150292, 446521, 1600204, + 4819428, 2591357, 48490, 2917012, 2285923, 1072926, 2824281, 4364250, 956033, 311938, + 37251, 3729300, 2726300, 644966, 1623020, 1419070, 4646747, 2417222, 2680238, + 2561083, 1793801, 2349366, 339747, 611366, 4684147, 4356907, 1277161, 4510381, + 3218352, 4161658, 3200733, 1172372, 3997786, 3169266, 3353418, 2248955, 2875885, + 2365369, 498208, 2968066, 2681505, 2059048, 2097106, 3607540, 1121504, 2016789, + 1762605, 3138431, 866081, 3705757, 3833066, 2599788, 760816, 4046672, 1544367, + 2983906, 4842911, 209599, 1250954, 3333704, 561212, 4674336, 2831841, 3690724, + 2929360, 4830834, 1177524, 2487687, 3525137, 875283, 651241, 2110742, 1296646, + 1543739, 4349417, 2384725, 1931751, 1519208, 1520034, 3385008, 3219962, 734912, 170230, + 1741419, 729913, 2860117, 2362381, 1199807, 2424230, 177824, 125948, 2722701, + 4687548, 1140771, 3232742, 4522020, 4376360, 1125603, 590312, 2481884, 138951, + 4086775, 615155, 3395781, 4587272, 283209, 568470, 4296185, 4344150, 2454321, + 2672602, 838828, 4051647, 1709120, 3074610, 693235, 4356087, 3018806, 239410, + 2431497, 691186, 766276, 4462126, 859155, 2370304, 1571808, 1938673, 1694955, + 3871296, 4245059, 3987376, 301524, 2512461, 3410437, 3300380, 684922, 4581995, + 3599557, 683515, 1850634, 3704678, 1937490, 2035591, 3718533, 2065879, 3160765, + 1467884, 1912241, 2501509, 3668572, 3390469, 2501150, 612319, 713633, 1976262, 135946, + 3641535, 632083, 13414, 4217765, 4137712, 2550250, 3281035, 4179598, 961045, + 2020694, 4380006, 1345936, 289162, 1359035, 770872, 4509911, 3947317, 4719693, + 248568, 2625660, 1237232, 2153208, 4814282, 1259954, 3677369, 861222, 2883506, + 3339149, 3998335, 491017, 1609022, 2648112, 742132, 649609, 4206953, 3131106, + 3504814, 3344486, 611721, 3215620, 2856233, 4447505, 1949222, 1868345, 712710, 6966, + 4730666, 3181872, 2972889, 3038521, 3525444, 4385208, 1845613, 1124187, 2030476, + 4468651, 2478792, 3473580, 3783357, 1852991, 1648485, 871319, 1670723, 4458328, + 3218600, 1811100, 3443356, 2233873, 3035207, 2548692, 3337891, 3773674, 1552957, + 4782811, 3144712, 3523466, 1491315, 3955852, 1838410, 3164028, 1092543, 776459, + 2959379, 2541744, 4064418, 3908320, 2854145, 3960709, 1348188, 977678, 853619, + 1304291, 2848702, 1657913, 1319826, 3322665, 788037, 2913686, 4471279, 1766285, 348304, + 56570, 1892118, 4017244, 401006, 3524539, 4310134, 1624693, 4081113, 957511, 849400, + 129975, 2616130, 378537, 1556787, 3916162, 1039980, 4407778, 2027690, 4213675, + 839863, 683134, 75805, 2493150, 4215796, 81587, 751845, 1255588, 1947964, + 1950470, 859401, 3077088, 3931110, 2316256, 1523761, 4527477, 4237511, 1123513, + 4209796, 3584772, 4250563, 2091754, 1618766, 2139944, 4525352, 382159, 2955887, 41760, + 2313998, 496912, 3791570, 3904792, 3613654, 873959, 127076, 2537797, 2458107, + 4543265, 3661909, 26828, 271816, 17854, 2461269, 1776042, 1573899, 3409957, + 4335712, 4534313, 3392751, 1230124, 2159031, 4444015, 3373087, 3848014, 2026600, + 1382747, 3537242, 4536743, 4714155, 3788371, 3570849, 173741, 211962, 4377778, + 119369, 2856973, 2945854, 1508054, 4503932, 3141566, 1842177, 3448683, 3384614, + 2886508, 1573965, 990618, 3053734, 2918742, 4508753, 1032149, 60943, 4291620, + 722607, 2883224, 169359, 4356585, 3725543, 3678729, 341673, 3592828, 4077251, + 3382936, 3885685, 4630994, 1286698, 4449616, 1138430, 3113385, 4660578, 2539973, + 4562286, 4085089, 494737, 3967610, 2130702, 1823755, 1369324, 3796951, 956299, 141730, + 935144, 4381893, 4412545, 1382250, 3024476, 2364546, 3396164, 3573511, 314081, 577688, + 4154135, 1567018, 4047761, 2446220, 1148833, 4842497, 3967186, 1175290, 3749667, + 1209593, 3295627, 3169065, 2460328, 1838486, 1436923, 2843887, 3676426, 2079145, + 2975635, 535071, 4287509, 3281107, 39606, 3115500, 3204573, 722131, 3124073}, +2, +false)));*/ CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/python/cugraph/community/egonet_wrapper.pyx b/python/cugraph/community/egonet_wrapper.pyx index 7d2a1169e25..ff9f2b8b3de 100644 --- a/python/cugraph/community/egonet_wrapper.pyx +++ b/python/cugraph/community/egonet_wrapper.pyx @@ -56,9 +56,11 @@ def egonet(input_graph, vertices, radius=1): # Pointers for egonet cdef uintptr_t c_source_vertex_ptr = vertices.__cuda_array_interface__['data'][0] n_subgraphs = vertices.size - + n_streams = 1 + if n_subgraphs > 1 : + n_streams = min(n_subgraphs, 32) cdef unique_ptr[handle_t] handle_ptr - handle_ptr.reset(new handle_t()) + handle_ptr.reset(new handle_t(n_streams)) handle_ = handle_ptr.get(); cdef graph_container_t graph_container diff --git a/python/cugraph/structure/graph_primtypes.pxd b/python/cugraph/structure/graph_primtypes.pxd index 3bb42633ac1..1e0d9626727 100644 --- a/python/cugraph/structure/graph_primtypes.pxd +++ b/python/cugraph/structure/graph_primtypes.pxd @@ -23,7 +23,6 @@ from libcpp.vector cimport vector from cugraph.raft.common.handle cimport * from rmm._lib.device_buffer cimport device_buffer - cdef extern from "graph.hpp" namespace "cugraph": ctypedef enum PropType: