diff --git a/.github/labeler.yml b/.github/labeler.yml index c589fda6099..368bf328b99 100644 --- a/.github/labeler.yml +++ b/.github/labeler.yml @@ -9,17 +9,6 @@ python: benchmarks: - 'benchmarks/**' -doc: - - 'docs/**' - - '**/*.md' - - 'datasets/**' - - 'notebooks/**' - - '**/*.txt' - - '**/*.rst' - - '**/*.ipynb' - - '**/*.pdf' - - '**/*.png' - datasets: - 'datasets/**' diff --git a/github/workflows/labeler.yml b/.github/workflows/labeler.yml similarity index 83% rename from github/workflows/labeler.yml rename to .github/workflows/labeler.yml index 23956a02fbd..31e78f82a62 100644 --- a/github/workflows/labeler.yml +++ b/.github/workflows/labeler.yml @@ -6,6 +6,6 @@ jobs: triage: runs-on: ubuntu-latest steps: - - uses: actions/labeler@main + - uses: actions/labeler@v4 with: repo-token: "${{ secrets.GITHUB_TOKEN }}" diff --git a/CHANGELOG.md b/CHANGELOG.md index 33a5b2bc5e7..d165cd7efc4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,84 @@ +# cuGraph 23.12.00 (6 Dec 2023) + +## 🚨 Breaking Changes + +- [BUG] Restore the original default order of CSR, which does not reverse edges in cuGraph-PyG ([#3980](https://github.com/rapidsai/cugraph/pull/3980)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- `Resultset` and `Dataset` Refactors ([#3957](https://github.com/rapidsai/cugraph/pull/3957)) [@nv-rliu](https://github.com/nv-rliu) +- Moves more MG graph ETL to libcugraph and re-enables MG tests in CI ([#3941](https://github.com/rapidsai/cugraph/pull/3941)) [@jnke2016](https://github.com/jnke2016) + +## 🐛 Bug Fixes + +- Pin actions/labeler to v4 ([#4038](https://github.com/rapidsai/cugraph/pull/4038)) [@raydouglass](https://github.com/raydouglass) +- Find rmm before cuco ([#4011](https://github.com/rapidsai/cugraph/pull/4011)) [@vyasr](https://github.com/vyasr) +- Pin to minor versions of packages outside the cuGraph repository. ([#4004](https://github.com/rapidsai/cugraph/pull/4004)) [@bdice](https://github.com/bdice) +- Move MTMG_TEST to MG tests block ([#3993](https://github.com/rapidsai/cugraph/pull/3993)) [@naimnv](https://github.com/naimnv) +- Fix Leiden refinement phase ([#3990](https://github.com/rapidsai/cugraph/pull/3990)) [@naimnv](https://github.com/naimnv) +- [BUG] Fix Graph Construction From Pandas in cuGraph-PyG ([#3985](https://github.com/rapidsai/cugraph/pull/3985)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- [BUG] Restore the original default order of CSR, which does not reverse edges in cuGraph-PyG ([#3980](https://github.com/rapidsai/cugraph/pull/3980)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Fix eigenvector testing and HITS testing discrepancies ([#3979](https://github.com/rapidsai/cugraph/pull/3979)) [@ChuckHastings](https://github.com/ChuckHastings) +- [BUG] Fix Incorrect Edge Index, Directory Selection in cuGraph-PyG Loader ([#3978](https://github.com/rapidsai/cugraph/pull/3978)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- [BUG] Check if Dask has quit to avoid throwing an exception and triggering a segfault on ddp exit ([#3961](https://github.com/rapidsai/cugraph/pull/3961)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- nx-cugraph: xfail test_louvain.py:test_threshold in Python 3.9 ([#3944](https://github.com/rapidsai/cugraph/pull/3944)) [@eriknw](https://github.com/eriknw) + +## 📖 Documentation + +- [DOC]: Fix invalid links and add materials to notebooks ([#4002](https://github.com/rapidsai/cugraph/pull/4002)) [@huiyuxie](https://github.com/huiyuxie) +- Update Broken Links in README.md ([#3924](https://github.com/rapidsai/cugraph/pull/3924)) [@nv-rliu](https://github.com/nv-rliu) + +## 🚀 New Features + +- Implement the transform_e primitive (to update property values for all edges) ([#3917](https://github.com/rapidsai/cugraph/pull/3917)) [@seunghwak](https://github.com/seunghwak) +- Update the neighbor intersection primitive to support edge masking. ([#3550](https://github.com/rapidsai/cugraph/pull/3550)) [@seunghwak](https://github.com/seunghwak) + +## 🛠️ Improvements + +- Correct defect found in DLFW testing ([#4021](https://github.com/rapidsai/cugraph/pull/4021)) [@ChuckHastings](https://github.com/ChuckHastings) +- `nx-cugraph` README update: adds missing `connected_components` algo to table ([#4019](https://github.com/rapidsai/cugraph/pull/4019)) [@rlratzel](https://github.com/rlratzel) +- Build concurrency for nightly and merge triggers ([#4009](https://github.com/rapidsai/cugraph/pull/4009)) [@bdice](https://github.com/bdice) +- Support `drop_last` Argument in cuGraph-PyG Loader ([#3995](https://github.com/rapidsai/cugraph/pull/3995)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Adds `update-version.sh` support for recently added files containing RAPIDS versions ([#3994](https://github.com/rapidsai/cugraph/pull/3994)) [@rlratzel](https://github.com/rlratzel) +- Use new `rapids-dask-dependency` metapackage for managing `dask` versions ([#3991](https://github.com/rapidsai/cugraph/pull/3991)) [@galipremsagar](https://github.com/galipremsagar) +- Fixes to nx-cugraph README: fixes typos, updates link to NX backend docs ([#3989](https://github.com/rapidsai/cugraph/pull/3989)) [@rlratzel](https://github.com/rlratzel) +- Address FIXMEs ([#3988](https://github.com/rapidsai/cugraph/pull/3988)) [@seunghwak](https://github.com/seunghwak) +- Updates README file to include nx-cugraph user documentation, adds nx-cugraph to main README ([#3984](https://github.com/rapidsai/cugraph/pull/3984)) [@rlratzel](https://github.com/rlratzel) +- Update C API graph creation function signatures ([#3982](https://github.com/rapidsai/cugraph/pull/3982)) [@ChuckHastings](https://github.com/ChuckHastings) +- [REVIEW]Optimize cugraph-DGL csc codepath ([#3977](https://github.com/rapidsai/cugraph/pull/3977)) [@VibhuJawa](https://github.com/VibhuJawa) +- nx-cugraph: add SSSP (unweighted) ([#3976](https://github.com/rapidsai/cugraph/pull/3976)) [@eriknw](https://github.com/eriknw) +- CuGraph compatibility fixes ([#3973](https://github.com/rapidsai/cugraph/pull/3973)) [@brandon-b-miller](https://github.com/brandon-b-miller) +- Skip certain `cugraph-pyg` tests when torch-sparse is not available ([#3970](https://github.com/rapidsai/cugraph/pull/3970)) [@tingyu66](https://github.com/tingyu66) +- nx-cugraph: add `eigenvector_centrality`, `katz_centrality`, `hits`, `pagerank` ([#3968](https://github.com/rapidsai/cugraph/pull/3968)) [@eriknw](https://github.com/eriknw) +- Cut peak memory footprint in graph creation ([#3966](https://github.com/rapidsai/cugraph/pull/3966)) [@seunghwak](https://github.com/seunghwak) +- nx-cugraph: add CC for undirected graphs to fix k-truss ([#3965](https://github.com/rapidsai/cugraph/pull/3965)) [@eriknw](https://github.com/eriknw) +- Skip certain `cugraph-pyg` tests when `torch_sparse` is not available ([#3962](https://github.com/rapidsai/cugraph/pull/3962)) [@tingyu66](https://github.com/tingyu66) +- `Resultset` and `Dataset` Refactors ([#3957](https://github.com/rapidsai/cugraph/pull/3957)) [@nv-rliu](https://github.com/nv-rliu) +- Download `xml` docs artifact through CloudFront endpoint ([#3955](https://github.com/rapidsai/cugraph/pull/3955)) [@AyodeAwe](https://github.com/AyodeAwe) +- Add many graph generators to nx-cugraph ([#3954](https://github.com/rapidsai/cugraph/pull/3954)) [@eriknw](https://github.com/eriknw) +- Unpin `dask` and `distributed` for `23.12` development ([#3953](https://github.com/rapidsai/cugraph/pull/3953)) [@galipremsagar](https://github.com/galipremsagar) +- Errors compiling for DLFW on CUDA 12.3 ([#3952](https://github.com/rapidsai/cugraph/pull/3952)) [@ChuckHastings](https://github.com/ChuckHastings) +- nx-cugraph: add k_truss and degree centralities ([#3945](https://github.com/rapidsai/cugraph/pull/3945)) [@eriknw](https://github.com/eriknw) +- nx-cugraph: handle seed argument in edge_betweenness_centrality ([#3943](https://github.com/rapidsai/cugraph/pull/3943)) [@eriknw](https://github.com/eriknw) +- Moves more MG graph ETL to libcugraph and re-enables MG tests in CI ([#3941](https://github.com/rapidsai/cugraph/pull/3941)) [@jnke2016](https://github.com/jnke2016) +- Temporarily disable mg testing ([#3940](https://github.com/rapidsai/cugraph/pull/3940)) [@jnke2016](https://github.com/jnke2016) +- adding C/C++ API docs ([#3938](https://github.com/rapidsai/cugraph/pull/3938)) [@BradReesWork](https://github.com/BradReesWork) +- Add multigraph support to nx-cugraph ([#3934](https://github.com/rapidsai/cugraph/pull/3934)) [@eriknw](https://github.com/eriknw) +- Setup Consistent Nightly Versions for Pip and Conda ([#3933](https://github.com/rapidsai/cugraph/pull/3933)) [@divyegala](https://github.com/divyegala) +- MTMG multi node ([#3932](https://github.com/rapidsai/cugraph/pull/3932)) [@ChuckHastings](https://github.com/ChuckHastings) +- Use branch-23.12 workflows. ([#3928](https://github.com/rapidsai/cugraph/pull/3928)) [@bdice](https://github.com/bdice) +- Fix an issue occurring in the cuGraph-DGL example for "mixed" mode. ([#3927](https://github.com/rapidsai/cugraph/pull/3927)) [@drivanov](https://github.com/drivanov) +- Updating Docs ([#3923](https://github.com/rapidsai/cugraph/pull/3923)) [@BradReesWork](https://github.com/BradReesWork) +- Forward-merge branch-23.10 to branch-23.12 ([#3919](https://github.com/rapidsai/cugraph/pull/3919)) [@nv-rliu](https://github.com/nv-rliu) +- new build all option ([#3916](https://github.com/rapidsai/cugraph/pull/3916)) [@BradReesWork](https://github.com/BradReesWork) +- Silence spurious compiler warnings ([#3913](https://github.com/rapidsai/cugraph/pull/3913)) [@seunghwak](https://github.com/seunghwak) +- Link wholegrah and cugraphops XML docs ([#3906](https://github.com/rapidsai/cugraph/pull/3906)) [@AyodeAwe](https://github.com/AyodeAwe) +- Updates to 23.12 ([#3905](https://github.com/rapidsai/cugraph/pull/3905)) [@raydouglass](https://github.com/raydouglass) +- Forward-merge branch-23.10 to branch-23.12 ([#3904](https://github.com/rapidsai/cugraph/pull/3904)) [@GPUtester](https://github.com/GPUtester) +- Build CUDA 12.0 ARM conda packages. ([#3903](https://github.com/rapidsai/cugraph/pull/3903)) [@bdice](https://github.com/bdice) +- Merge branch-23.10 into branch-23.12 ([#3898](https://github.com/rapidsai/cugraph/pull/3898)) [@rlratzel](https://github.com/rlratzel) +- Some MTMG code cleanup and small optimizations ([#3894](https://github.com/rapidsai/cugraph/pull/3894)) [@ChuckHastings](https://github.com/ChuckHastings) +- Enable parallel mode ([#3875](https://github.com/rapidsai/cugraph/pull/3875)) [@jnke2016](https://github.com/jnke2016) +- Adds benchmarks for `nx-cugraph` ([#3854](https://github.com/rapidsai/cugraph/pull/3854)) [@rlratzel](https://github.com/rlratzel) +- Add nx-cugraph notebook for showing accelerated networkX APIs ([#3830](https://github.com/rapidsai/cugraph/pull/3830)) [@betochimas](https://github.com/betochimas) + # cuGraph 23.10.00 (11 Oct 2023) ## 🚨 Breaking Changes diff --git a/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp b/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp index c4cacb401af..3e4b2513a8d 100644 --- a/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp +++ b/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp @@ -57,10 +57,10 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos == objects_.end(), "Cannot overwrite wrapped object"); - objects_.insert(std::make_pair(handle.get_local_rank(), std::move(obj))); + objects_.insert(std::make_pair(handle.get_rank(), std::move(obj))); } /** @@ -90,7 +90,7 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos != objects_.end(), "Uninitialized wrapped object"); return pos->second; @@ -106,7 +106,7 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos != objects_.end(), "Uninitialized wrapped object"); diff --git a/cpp/include/cugraph/mtmg/handle.hpp b/cpp/include/cugraph/mtmg/handle.hpp index 6223de1781d..0b02091a3cc 100644 --- a/cpp/include/cugraph/mtmg/handle.hpp +++ b/cpp/include/cugraph/mtmg/handle.hpp @@ -32,18 +32,19 @@ namespace mtmg { * */ class handle_t { + handle_t(handle_t const&) = delete; + handle_t operator=(handle_t const&) = delete; + public: /** * @brief Constructor * * @param raft_handle Raft handle for the resources * @param thread_rank Rank for this thread + * @param device_id Device id for the device this handle operates on */ - handle_t(raft::handle_t const& raft_handle, int thread_rank, size_t device_id) - : raft_handle_(raft_handle), - thread_rank_(thread_rank), - local_rank_(raft_handle.get_comms().get_rank()), // FIXME: update for multi-node - device_id_(device_id) + handle_t(raft::handle_t const& raft_handle, int thread_rank, rmm::cuda_device_id device_id) + : raft_handle_(raft_handle), thread_rank_(thread_rank), device_id_raii_(device_id) { } @@ -118,18 +119,10 @@ class handle_t { */ int get_rank() const { return raft_handle_.get_comms().get_rank(); } - /** - * @brief Get local gpu rank - * - * @return local gpu rank - */ - int get_local_rank() const { return local_rank_; } - private: raft::handle_t const& raft_handle_; int thread_rank_; - int local_rank_; - size_t device_id_; + rmm::cuda_set_device_raii device_id_raii_; }; } // namespace mtmg diff --git a/cpp/include/cugraph/mtmg/instance_manager.hpp b/cpp/include/cugraph/mtmg/instance_manager.hpp index f819a5a0abe..f60063c4101 100644 --- a/cpp/include/cugraph/mtmg/instance_manager.hpp +++ b/cpp/include/cugraph/mtmg/instance_manager.hpp @@ -47,15 +47,10 @@ class instance_manager_t { ~instance_manager_t() { - int current_device{}; - RAFT_CUDA_TRY(cudaGetDevice(¤t_device)); - for (size_t i = 0; i < nccl_comms_.size(); ++i) { - RAFT_CUDA_TRY(cudaSetDevice(device_ids_[i].value())); + rmm::cuda_set_device_raii local_set_device(device_ids_[i]); RAFT_NCCL_TRY(ncclCommDestroy(*nccl_comms_[i])); } - - RAFT_CUDA_TRY(cudaSetDevice(current_device)); } /** @@ -75,8 +70,7 @@ class instance_manager_t { int gpu_id = local_id % raft_handle_.size(); int thread_id = local_id / raft_handle_.size(); - RAFT_CUDA_TRY(cudaSetDevice(device_ids_[gpu_id].value())); - return handle_t(*raft_handle_[gpu_id], thread_id, static_cast(gpu_id)); + return handle_t(*raft_handle_[gpu_id], thread_id, device_ids_[gpu_id]); } /** diff --git a/cpp/include/cugraph/mtmg/resource_manager.hpp b/cpp/include/cugraph/mtmg/resource_manager.hpp index 127944cf7ba..bc312c9ae77 100644 --- a/cpp/include/cugraph/mtmg/resource_manager.hpp +++ b/cpp/include/cugraph/mtmg/resource_manager.hpp @@ -89,7 +89,7 @@ class resource_manager_t { local_rank_map_.insert(std::pair(global_rank, local_device_id)); - RAFT_CUDA_TRY(cudaSetDevice(local_device_id.value())); + rmm::cuda_set_device_raii local_set_device(local_device_id); // FIXME: There is a bug in the cuda_memory_resource that results in a Hang. // using the pool resource as a work-around. @@ -182,14 +182,12 @@ class resource_manager_t { --gpu_row_comm_size; } - int current_device{}; - RAFT_CUDA_TRY(cudaGetDevice(¤t_device)); RAFT_NCCL_TRY(ncclGroupStart()); for (size_t i = 0; i < local_ranks_to_include.size(); ++i) { int rank = local_ranks_to_include[i]; auto pos = local_rank_map_.find(rank); - RAFT_CUDA_TRY(cudaSetDevice(pos->second.value())); + rmm::cuda_set_device_raii local_set_device(pos->second); nccl_comms.push_back(std::make_unique()); handles.push_back( @@ -204,7 +202,6 @@ class resource_manager_t { handles[i].get(), *nccl_comms[i], ranks_to_include.size(), rank); } RAFT_NCCL_TRY(ncclGroupEnd()); - RAFT_CUDA_TRY(cudaSetDevice(current_device)); std::vector running_threads; @@ -217,9 +214,7 @@ class resource_manager_t { &device_ids, &nccl_comms, &handles]() { - int rank = local_ranks_to_include[idx]; - RAFT_CUDA_TRY(cudaSetDevice(device_ids[idx].value())); - + rmm::cuda_set_device_raii local_set_device(device_ids[idx]); cugraph::partition_manager::init_subcomm(*handles[idx], gpu_row_comm_size); }); } diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index c49b62e4543..f0f729bce18 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -20,6 +20,7 @@ #include #include #include +#include #include #include @@ -524,34 +525,21 @@ std::tuple> mark_entries(raft::handle_t co return word; }); - size_t bit_count = thrust::transform_reduce( - handle.get_thrust_policy(), - marked_entries.begin(), - marked_entries.end(), - [] __device__(auto word) { return __popc(word); }, - size_t{0}, - thrust::plus()); + size_t bit_count = detail::count_set_bits(handle, marked_entries.begin(), num_entries); return std::make_tuple(bit_count, std::move(marked_entries)); } template -rmm::device_uvector remove_flagged_elements(raft::handle_t const& handle, - rmm::device_uvector&& vector, - raft::device_span remove_flags, - size_t remove_count) +rmm::device_uvector keep_flagged_elements(raft::handle_t const& handle, + rmm::device_uvector&& vector, + raft::device_span keep_flags, + size_t keep_count) { - rmm::device_uvector result(vector.size() - remove_count, handle.get_stream()); - - thrust::copy_if( - handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(vector.size()), - thrust::make_transform_output_iterator(result.begin(), - indirection_t{vector.data()}), - [remove_flags] __device__(size_t i) { - return !(remove_flags[cugraph::packed_bool_offset(i)] & cugraph::packed_bool_mask(i)); - }); + rmm::device_uvector result(keep_count, handle.get_stream()); + + detail::copy_if_mask_set( + handle, vector.begin(), vector.end(), keep_flags.begin(), result.begin()); return result; } diff --git a/cpp/src/structure/remove_multi_edges_impl.cuh b/cpp/src/structure/remove_multi_edges_impl.cuh index ab6b1fba8eb..fdd3059f874 100644 --- a/cpp/src/structure/remove_multi_edges_impl.cuh +++ b/cpp/src/structure/remove_multi_edges_impl.cuh @@ -254,50 +254,47 @@ remove_multi_edges(raft::handle_t const& handle, } } - auto [multi_edge_count, multi_edges_to_delete] = - detail::mark_entries(handle, - edgelist_srcs.size(), - [d_edgelist_srcs = edgelist_srcs.data(), - d_edgelist_dsts = edgelist_dsts.data()] __device__(auto idx) { - return (idx > 0) && (d_edgelist_srcs[idx - 1] == d_edgelist_srcs[idx]) && - (d_edgelist_dsts[idx - 1] == d_edgelist_dsts[idx]); - }); - - if (multi_edge_count > 0) { - edgelist_srcs = detail::remove_flagged_elements( + auto [keep_count, keep_flags] = detail::mark_entries( + handle, + edgelist_srcs.size(), + [d_edgelist_srcs = edgelist_srcs.data(), + d_edgelist_dsts = edgelist_dsts.data()] __device__(auto idx) { + return !((idx > 0) && (d_edgelist_srcs[idx - 1] == d_edgelist_srcs[idx]) && + (d_edgelist_dsts[idx - 1] == d_edgelist_dsts[idx])); + }); + + if (keep_count < edgelist_srcs.size()) { + edgelist_srcs = detail::keep_flagged_elements( handle, std::move(edgelist_srcs), - raft::device_span{multi_edges_to_delete.data(), multi_edges_to_delete.size()}, - multi_edge_count); - edgelist_dsts = detail::remove_flagged_elements( + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); + edgelist_dsts = detail::keep_flagged_elements( handle, std::move(edgelist_dsts), - raft::device_span{multi_edges_to_delete.data(), multi_edges_to_delete.size()}, - multi_edge_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_weights) - edgelist_weights = detail::remove_flagged_elements( + edgelist_weights = detail::keep_flagged_elements( handle, std::move(*edgelist_weights), - raft::device_span{multi_edges_to_delete.data(), - multi_edges_to_delete.size()}, - multi_edge_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_edge_ids) - edgelist_edge_ids = detail::remove_flagged_elements( + edgelist_edge_ids = detail::keep_flagged_elements( handle, std::move(*edgelist_edge_ids), - raft::device_span{multi_edges_to_delete.data(), - multi_edges_to_delete.size()}, - multi_edge_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_edge_types) - edgelist_edge_types = detail::remove_flagged_elements( + edgelist_edge_types = detail::keep_flagged_elements( handle, std::move(*edgelist_edge_types), - raft::device_span{multi_edges_to_delete.data(), - multi_edges_to_delete.size()}, - multi_edge_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); } return std::make_tuple(std::move(edgelist_srcs), diff --git a/cpp/src/structure/remove_self_loops_impl.cuh b/cpp/src/structure/remove_self_loops_impl.cuh index 161ffeae28e..dafe26cd1c5 100644 --- a/cpp/src/structure/remove_self_loops_impl.cuh +++ b/cpp/src/structure/remove_self_loops_impl.cuh @@ -44,44 +44,44 @@ remove_self_loops(raft::handle_t const& handle, std::optional>&& edgelist_edge_ids, std::optional>&& edgelist_edge_types) { - auto [self_loop_count, self_loops_to_delete] = + auto [keep_count, keep_flags] = detail::mark_entries(handle, edgelist_srcs.size(), [d_srcs = edgelist_srcs.data(), d_dsts = edgelist_dsts.data()] __device__( - size_t i) { return d_srcs[i] == d_dsts[i]; }); + size_t i) { return d_srcs[i] != d_dsts[i]; }); - if (self_loop_count > 0) { - edgelist_srcs = detail::remove_flagged_elements( + if (keep_count < edgelist_srcs.size()) { + edgelist_srcs = detail::keep_flagged_elements( handle, std::move(edgelist_srcs), - raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, - self_loop_count); - edgelist_dsts = detail::remove_flagged_elements( + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); + edgelist_dsts = detail::keep_flagged_elements( handle, std::move(edgelist_dsts), - raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, - self_loop_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_weights) - edgelist_weights = detail::remove_flagged_elements( + edgelist_weights = detail::keep_flagged_elements( handle, std::move(*edgelist_weights), - raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, - self_loop_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_edge_ids) - edgelist_edge_ids = detail::remove_flagged_elements( + edgelist_edge_ids = detail::keep_flagged_elements( handle, std::move(*edgelist_edge_ids), - raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, - self_loop_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); if (edgelist_edge_types) - edgelist_edge_types = detail::remove_flagged_elements( + edgelist_edge_types = detail::keep_flagged_elements( handle, std::move(*edgelist_edge_types), - raft::device_span{self_loops_to_delete.data(), self_loops_to_delete.size()}, - self_loop_count); + raft::device_span{keep_flags.data(), keep_flags.size()}, + keep_count); } return std::make_tuple(std::move(edgelist_srcs), diff --git a/cpp/tests/community/triangle_count_test.cpp b/cpp/tests/community/triangle_count_test.cpp index 836bab59457..592924c3c47 100644 --- a/cpp/tests/community/triangle_count_test.cpp +++ b/cpp/tests/community/triangle_count_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -232,7 +232,7 @@ class Tests_TriangleCount for (size_t i = 0; i < h_cugraph_vertices.size(); ++i) { auto v = h_cugraph_vertices[i]; auto count = h_cugraph_triangle_counts[i]; - ASSERT_TRUE(count == h_reference_triangle_counts[v]) + ASSERT_EQ(count, h_reference_triangle_counts[v]) << "Triangle count values do not match with the reference values."; } } diff --git a/cpp/tests/mtmg/threaded_test.cu b/cpp/tests/mtmg/threaded_test.cu index bc4d8cfef6a..1a6a17eaa18 100644 --- a/cpp/tests/mtmg/threaded_test.cu +++ b/cpp/tests/mtmg/threaded_test.cu @@ -155,10 +155,25 @@ class Tests_Multithreaded input_usecase.template construct_edgelist( handle, multithreaded_usecase.test_weighted, false, false); + rmm::device_uvector d_unique_vertices(2 * d_src_v.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), d_src_v.begin(), d_src_v.end(), d_unique_vertices.begin()); + thrust::copy(handle.get_thrust_policy(), + d_dst_v.begin(), + d_dst_v.end(), + d_unique_vertices.begin() + d_src_v.size()); + thrust::sort(handle.get_thrust_policy(), d_unique_vertices.begin(), d_unique_vertices.end()); + + d_unique_vertices.resize(thrust::distance(d_unique_vertices.begin(), + thrust::unique(handle.get_thrust_policy(), + d_unique_vertices.begin(), + d_unique_vertices.end())), + handle.get_stream()); + auto h_src_v = cugraph::test::to_host(handle, d_src_v); auto h_dst_v = cugraph::test::to_host(handle, d_dst_v); auto h_weights_v = cugraph::test::to_host(handle, d_weights_v); - auto unique_vertices = cugraph::test::to_host(handle, d_vertices_v); + auto unique_vertices = cugraph::test::to_host(handle, d_unique_vertices); // Load edgelist from different threads. We'll use more threads than GPUs here for (int i = 0; i < num_threads; ++i) { @@ -293,13 +308,13 @@ class Tests_Multithreaded num_threads]() { auto thread_handle = instance_manager->get_handle(); - auto number_of_vertices = unique_vertices->size(); + auto number_of_vertices = unique_vertices.size(); std::vector my_vertex_list; my_vertex_list.reserve((number_of_vertices + num_threads - 1) / num_threads); for (size_t j = i; j < number_of_vertices; j += num_threads) { - my_vertex_list.push_back((*unique_vertices)[j]); + my_vertex_list.push_back(unique_vertices[j]); } rmm::device_uvector d_my_vertex_list(my_vertex_list.size(), diff --git a/cpp/tests/structure/mg_select_random_vertices_test.cpp b/cpp/tests/structure/mg_select_random_vertices_test.cpp index 79c50301922..e49e1ebcb99 100644 --- a/cpp/tests/structure/mg_select_random_vertices_test.cpp +++ b/cpp/tests/structure/mg_select_random_vertices_test.cpp @@ -90,7 +90,7 @@ class Tests_MGSelectRandomVertices std::iota( h_given_set.begin(), h_given_set.end(), mg_graph_view.local_vertex_partition_range_first()); std::shuffle(h_given_set.begin(), h_given_set.end(), std::mt19937{std::random_device{}()}); - h_given_set.resize(std::rand() % mg_graph_view.local_vertex_partition_range_size() + 1); + h_given_set.resize(std::rand() % (mg_graph_view.local_vertex_partition_range_size() + 1)); // Compute size of the distributed vertex set int num_of_elements_in_given_set = static_cast(h_given_set.size()); @@ -105,7 +105,7 @@ class Tests_MGSelectRandomVertices size_t select_count = num_of_elements_in_given_set > select_random_vertices_usecase.select_count ? select_random_vertices_usecase.select_count - : std::rand() % num_of_elements_in_given_set + 1; + : std::rand() % (num_of_elements_in_given_set + 1); for (int idx = 0; idx < with_replacement_flags.size(); idx++) { bool with_replacement = with_replacement_flags[idx]; diff --git a/cpp/tests/utilities/test_graphs.hpp b/cpp/tests/utilities/test_graphs.hpp index 16c9d3ed145..8cc87b26f1d 100644 --- a/cpp/tests/utilities/test_graphs.hpp +++ b/cpp/tests/utilities/test_graphs.hpp @@ -621,9 +621,25 @@ construct_graph(raft::handle_t const& handle, CUGRAPH_EXPECTS(d_src_v.size() <= static_cast(std::numeric_limits::max()), "Invalid template parameter: edge_t overflow."); - if (drop_self_loops) { remove_self_loops(handle, d_src_v, d_dst_v, d_weights_v); } + if (drop_self_loops) { + std::tie(d_src_v, d_dst_v, d_weights_v, std::ignore, std::ignore) = + cugraph::remove_self_loops(handle, + std::move(d_src_v), + std::move(d_dst_v), + std::move(d_weights_v), + std::nullopt, + std::nullopt); + } - if (drop_multi_edges) { sort_and_remove_multi_edges(handle, d_src_v, d_dst_v, d_weights_v); } + if (drop_multi_edges) { + std::tie(d_src_v, d_dst_v, d_weights_v, std::ignore, std::ignore) = + cugraph::remove_multi_edges(handle, + std::move(d_src_v), + std::move(d_dst_v), + std::move(d_weights_v), + std::nullopt, + std::nullopt); + } graph_t graph(handle); std::optional< diff --git a/cpp/tests/utilities/thrust_wrapper.cu b/cpp/tests/utilities/thrust_wrapper.cu index cb7e6f1bd66..2daf250b4a2 100644 --- a/cpp/tests/utilities/thrust_wrapper.cu +++ b/cpp/tests/utilities/thrust_wrapper.cu @@ -206,131 +206,5 @@ template void populate_vertex_ids(raft::handle_t const& handle, rmm::device_uvector& d_vertices_v, int64_t vertex_id_offset); -template -void remove_self_loops(raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */) -{ - if (d_weight_v) { - auto edge_first = thrust::make_zip_iterator( - thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin())); - d_src_v.resize( - thrust::distance(edge_first, - thrust::remove_if( - handle.get_thrust_policy(), - edge_first, - edge_first + d_src_v.size(), - [] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })), - handle.get_stream()); - d_dst_v.resize(d_src_v.size(), handle.get_stream()); - (*d_weight_v).resize(d_src_v.size(), handle.get_stream()); - } else { - auto edge_first = - thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); - d_src_v.resize( - thrust::distance(edge_first, - thrust::remove_if( - handle.get_thrust_policy(), - edge_first, - edge_first + d_src_v.size(), - [] __device__(auto e) { return thrust::get<0>(e) == thrust::get<1>(e); })), - handle.get_stream()); - d_dst_v.resize(d_src_v.size(), handle.get_stream()); - } - - d_src_v.shrink_to_fit(handle.get_stream()); - d_dst_v.shrink_to_fit(handle.get_stream()); - if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); } -} - -template void remove_self_loops( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void remove_self_loops( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void remove_self_loops( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void remove_self_loops( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template -void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */) -{ - if (d_weight_v) { - auto edge_first = thrust::make_zip_iterator( - thrust::make_tuple(d_src_v.begin(), d_dst_v.begin(), (*d_weight_v).begin())); - thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size()); - d_src_v.resize( - thrust::distance(edge_first, - thrust::unique(handle.get_thrust_policy(), - edge_first, - edge_first + d_src_v.size(), - [] __device__(auto lhs, auto rhs) { - return (thrust::get<0>(lhs) == thrust::get<0>(rhs)) && - (thrust::get<1>(lhs) == thrust::get<1>(rhs)); - })), - handle.get_stream()); - d_dst_v.resize(d_src_v.size(), handle.get_stream()); - (*d_weight_v).resize(d_src_v.size(), handle.get_stream()); - } else { - auto edge_first = - thrust::make_zip_iterator(thrust::make_tuple(d_src_v.begin(), d_dst_v.begin())); - thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size()); - d_src_v.resize( - thrust::distance( - edge_first, - thrust::unique(handle.get_thrust_policy(), edge_first, edge_first + d_src_v.size())), - handle.get_stream()); - d_dst_v.resize(d_src_v.size(), handle.get_stream()); - } - - d_src_v.shrink_to_fit(handle.get_stream()); - d_dst_v.shrink_to_fit(handle.get_stream()); - if (d_weight_v) { (*d_weight_v).shrink_to_fit(handle.get_stream()); } -} - -template void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - } // namespace test } // namespace cugraph diff --git a/cpp/tests/utilities/thrust_wrapper.hpp b/cpp/tests/utilities/thrust_wrapper.hpp index eead4dc268f..fb82d781198 100644 --- a/cpp/tests/utilities/thrust_wrapper.hpp +++ b/cpp/tests/utilities/thrust_wrapper.hpp @@ -46,18 +46,5 @@ void populate_vertex_ids(raft::handle_t const& handle, rmm::device_uvector& d_vertices_v /* [INOUT] */, vertex_t vertex_id_offset); -template -void remove_self_loops(raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - -template -void sort_and_remove_multi_edges( - raft::handle_t const& handle, - rmm::device_uvector& d_src_v /* [INOUT] */, - rmm::device_uvector& d_dst_v /* [INOUT] */, - std::optional>& d_weight_v /* [INOUT] */); - } // namespace test } // namespace cugraph