From 33ff543194f62a3ddf99ad9ac7f230f23d4bb779 Mon Sep 17 00:00:00 2001 From: Ray Douglass Date: Wed, 7 Aug 2024 10:40:40 -0400 Subject: [PATCH 1/6] Update Changelog [skip ci] --- CHANGELOG.md | 85 ++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 85 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index f85c7d03f03..689a214751f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,88 @@ +# cugraph 24.08.00 (7 Aug 2024) + +## 🚨 Breaking Changes + +- Use MNMG version of ECG in python layer instead, and remove legacy ECG and Louvain ([#4514](https://github.com/rapidsai/cugraph/pull/4514)) [@naimnv](https://github.com/naimnv) + +## 🐛 Bug Fixes + +- add setuptools to host requirements for conda packages that need it ([#4582](https://github.com/rapidsai/cugraph/pull/4582)) [@jameslamb](https://github.com/jameslamb) +- Add pylibcugraph dependency on pylibraft. ([#4570](https://github.com/rapidsai/cugraph/pull/4570)) [@bdice](https://github.com/bdice) +- Fix build error with NO_CUGRAPH_OPS ([#4563](https://github.com/rapidsai/cugraph/pull/4563)) [@seunghwak](https://github.com/seunghwak) +- [BUG] Fix Failing WholeGraph Tests ([#4560](https://github.com/rapidsai/cugraph/pull/4560)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Temporarily Disable Feature Store Tests with WholeGraph ([#4559](https://github.com/rapidsai/cugraph/pull/4559)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Fix MG `katz_centrality`: Check if DataFrame Arg is Not None ([#4555](https://github.com/rapidsai/cugraph/pull/4555)) [@nv-rliu](https://github.com/nv-rliu) +- nx-cugraph: fix `from_pandas_edgekey` given edgekey but not edgeattr ([#4550](https://github.com/rapidsai/cugraph/pull/4550)) [@eriknw](https://github.com/eriknw) +- Fix triangle count test bug ([#4549](https://github.com/rapidsai/cugraph/pull/4549)) [@jnke2016](https://github.com/jnke2016) +- [BUG] Use the Correct WG Communicator ([#4548](https://github.com/rapidsai/cugraph/pull/4548)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Add Additional Check For SSSP Source Vertex & Fix SSSP Benchmark ([#4541](https://github.com/rapidsai/cugraph/pull/4541)) [@nv-rliu](https://github.com/nv-rliu) +- Fix OOM Bug for Jaccard, Sorensen, and Overlap benchmarks ([#4524](https://github.com/rapidsai/cugraph/pull/4524)) [@nv-rliu](https://github.com/nv-rliu) +- Distribute start_list across ranks ([#4519](https://github.com/rapidsai/cugraph/pull/4519)) [@jnke2016](https://github.com/jnke2016) +- [FIX] Skip Distributed Sampler Tests if PyTorch with CUDA is not Available ([#4518](https://github.com/rapidsai/cugraph/pull/4518)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- [BUG] Fix a hang issue in MG triangle counts (when invoked with a small number of vertices to update triangle counts) ([#4517](https://github.com/rapidsai/cugraph/pull/4517)) [@seunghwak](https://github.com/seunghwak) +- Update MG Benchmark List ([#4516](https://github.com/rapidsai/cugraph/pull/4516)) [@nv-rliu](https://github.com/nv-rliu) +- Fix TensorProductConv test and improve docs ([#4480](https://github.com/rapidsai/cugraph/pull/4480)) [@tingyu66](https://github.com/tingyu66) +- Test nx-cugraph package instead of editable install ([#4442](https://github.com/rapidsai/cugraph/pull/4442)) [@KyleFromNVIDIA](https://github.com/KyleFromNVIDIA) + +## 📖 Documentation + +- DOC: typo in nx_transition.rst ([#4491](https://github.com/rapidsai/cugraph/pull/4491)) [@raybellwaves](https://github.com/raybellwaves) +- Doc cleanup for nx-cugraph: fixed typos, cleaned up various descriptions, renamed notebook to match naming convetion. ([#4478](https://github.com/rapidsai/cugraph/pull/4478)) [@rlratzel](https://github.com/rlratzel) +- [DOC] Minor Improvements to cuGraph-PyG Documentation ([#4460](https://github.com/rapidsai/cugraph/pull/4460)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) + +## 🚀 New Features + +- Use MNMG version of ECG in python layer instead, and remove legacy ECG and Louvain ([#4514](https://github.com/rapidsai/cugraph/pull/4514)) [@naimnv](https://github.com/naimnv) +- c_api and plc binding for lookup src dst using edge ids and type(s) ([#4494](https://github.com/rapidsai/cugraph/pull/4494)) [@naimnv](https://github.com/naimnv) +- Forward merge branch-24.06 into branch-24.08 ([#4489](https://github.com/rapidsai/cugraph/pull/4489)) [@nv-rliu](https://github.com/nv-rliu) +- [FEA] New Graph Interface and Loaders for Distributed Sampling in DGL ([#4486](https://github.com/rapidsai/cugraph/pull/4486)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- compute cosine similarity for vertex pairs ([#4482](https://github.com/rapidsai/cugraph/pull/4482)) [@naimnv](https://github.com/naimnv) +- Define heterogeneous renumbering API ([#4463](https://github.com/rapidsai/cugraph/pull/4463)) [@seunghwak](https://github.com/seunghwak) +- Lookup edge src dst using edge id and type ([#4449](https://github.com/rapidsai/cugraph/pull/4449)) [@naimnv](https://github.com/naimnv) +- Biased sampling ([#4443](https://github.com/rapidsai/cugraph/pull/4443)) [@seunghwak](https://github.com/seunghwak) + +## 🛠️ Improvements + +- nx-cugraph: check networkx version ([#4571](https://github.com/rapidsai/cugraph/pull/4571)) [@eriknw](https://github.com/eriknw) +- nx-cugraph: add `G.__networkx_cache__` to enable graph conversion caching ([#4567](https://github.com/rapidsai/cugraph/pull/4567)) [@eriknw](https://github.com/eriknw) +- split up CUDA-suffixed dependencies in dependencies.yaml ([#4552](https://github.com/rapidsai/cugraph/pull/4552)) [@jameslamb](https://github.com/jameslamb) +- Use workflow branch 24.08 again ([#4544](https://github.com/rapidsai/cugraph/pull/4544)) [@KyleFromNVIDIA](https://github.com/KyleFromNVIDIA) +- Support non p2p configuration when initializing the comms ([#4543](https://github.com/rapidsai/cugraph/pull/4543)) [@jnke2016](https://github.com/jnke2016) +- Fix Warning from `simpleDistributedGraph.py` ([#4540](https://github.com/rapidsai/cugraph/pull/4540)) [@nv-rliu](https://github.com/nv-rliu) +- Create a graph from the edge list in multiple chunks ([#4539](https://github.com/rapidsai/cugraph/pull/4539)) [@seunghwak](https://github.com/seunghwak) +- nx-cugraph: add dijkstra sssp functions ([#4538](https://github.com/rapidsai/cugraph/pull/4538)) [@eriknw](https://github.com/eriknw) +- nx-cugraph: add `from_dict_of_lists` and `to_dict_of_lists` ([#4537](https://github.com/rapidsai/cugraph/pull/4537)) [@eriknw](https://github.com/eriknw) +- Ensure `get_test_data.sh` doesn't re-download datasets ([#4536](https://github.com/rapidsai/cugraph/pull/4536)) [@trxcllnt](https://github.com/trxcllnt) +- Define and Implement C API for biased sampling ([#4535](https://github.com/rapidsai/cugraph/pull/4535)) [@ChuckHastings](https://github.com/ChuckHastings) +- Build and test with CUDA 12.5.1 ([#4534](https://github.com/rapidsai/cugraph/pull/4534)) [@KyleFromNVIDIA](https://github.com/KyleFromNVIDIA) +- Refactor C++ unit tests to allow finer grained filtering ([#4533](https://github.com/rapidsai/cugraph/pull/4533)) [@ChuckHastings](https://github.com/ChuckHastings) +- [IMP] Set the Default WG Memory Type to 'distributed' for the MNMG PyG Example ([#4532](https://github.com/rapidsai/cugraph/pull/4532)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- nx-cugraph: add `relabel_nodes` and `convert_node_labels_to_integers` ([#4531](https://github.com/rapidsai/cugraph/pull/4531)) [@eriknw](https://github.com/eriknw) +- Add `-cuXX` suffixed versions of cugraph-service-client dependency to pyproject.toml's project.dependencies list ([#4530](https://github.com/rapidsai/cugraph/pull/4530)) [@trxcllnt](https://github.com/trxcllnt) +- Further optimize `from_pandas_edgelist` with cudf ([#4528](https://github.com/rapidsai/cugraph/pull/4528)) [@eriknw](https://github.com/eriknw) +- Performance optimize BFS (including direction optimizing BFS implementation, mainly for single-GPU) ([#4527](https://github.com/rapidsai/cugraph/pull/4527)) [@seunghwak](https://github.com/seunghwak) +- Add CUDA_STATIC_MATH_LIBRARIES ([#4526](https://github.com/rapidsai/cugraph/pull/4526)) [@KyleFromNVIDIA](https://github.com/KyleFromNVIDIA) +- Better handle cudf.pandas in `from_pandas_edgelist` ([#4525](https://github.com/rapidsai/cugraph/pull/4525)) [@eriknw](https://github.com/eriknw) +- Skip the benchmark ctests within CI ([#4522](https://github.com/rapidsai/cugraph/pull/4522)) [@ChuckHastings](https://github.com/ChuckHastings) +- remove thriftpy2 ceiling ([#4521](https://github.com/rapidsai/cugraph/pull/4521)) [@jameslamb](https://github.com/jameslamb) +- Avoid --find-links in wheel jobs ([#4509](https://github.com/rapidsai/cugraph/pull/4509)) [@jameslamb](https://github.com/jameslamb) +- Refactor code base to reduce memory requirement for building libcugraph ([#4506](https://github.com/rapidsai/cugraph/pull/4506)) [@naimnv](https://github.com/naimnv) +- Tweak rmm configuration for C++ unit tests ([#4503](https://github.com/rapidsai/cugraph/pull/4503)) [@ChuckHastings](https://github.com/ChuckHastings) +- Expose new all-pairs Similarity algorithms ([#4502](https://github.com/rapidsai/cugraph/pull/4502)) [@jnke2016](https://github.com/jnke2016) +- remove openmpi ceiling ([#4496](https://github.com/rapidsai/cugraph/pull/4496)) [@jameslamb](https://github.com/jameslamb) +- Cut peak memory footprint in per_v_transform_reduce_dst_key_aggregated_outgoing_e ([#4484](https://github.com/rapidsai/cugraph/pull/4484)) [@seunghwak](https://github.com/seunghwak) +- Skip MG `dgl_uniform_sampler` test in nightlies ([#4479](https://github.com/rapidsai/cugraph/pull/4479)) [@nv-rliu](https://github.com/nv-rliu) +- Remove text builds of documentation ([#4468](https://github.com/rapidsai/cugraph/pull/4468)) [@vyasr](https://github.com/vyasr) +- [IMP] Limit the Test Data Size when Running CI in `gcn_dist_sg.py` ([#4461](https://github.com/rapidsai/cugraph/pull/4461)) [@alexbarghi-nv](https://github.com/alexbarghi-nv) +- Forward Merge branch-24.06 into branch-24.08 ([#4454](https://github.com/rapidsai/cugraph/pull/4454)) [@nv-rliu](https://github.com/nv-rliu) +- Properly clean up python directories ([#4453](https://github.com/rapidsai/cugraph/pull/4453)) [@ChuckHastings](https://github.com/ChuckHastings) +- Fixes for On-Going MG Test Failures ([#4450](https://github.com/rapidsai/cugraph/pull/4450)) [@nv-rliu](https://github.com/nv-rliu) +- remove unnecessary 'setuptools' and 'wheel' dependencies ([#4448](https://github.com/rapidsai/cugraph/pull/4448)) [@jameslamb](https://github.com/jameslamb) +- MG Implementation K-Truss ([#4438](https://github.com/rapidsai/cugraph/pull/4438)) [@jnke2016](https://github.com/jnke2016) +- Overhaul ops-codeowners ([#4409](https://github.com/rapidsai/cugraph/pull/4409)) [@raydouglass](https://github.com/raydouglass) +- Use rapids-build-backend ([#4393](https://github.com/rapidsai/cugraph/pull/4393)) [@KyleFromNVIDIA](https://github.com/KyleFromNVIDIA) +- Optimize K-Truss ([#4375](https://github.com/rapidsai/cugraph/pull/4375)) [@jnke2016](https://github.com/jnke2016) + # cugraph 24.06.00 (5 Jun 2024) ## 🚨 Breaking Changes From 6372e20111d85577db31f23c322e193021efc257 Mon Sep 17 00:00:00 2001 From: Ray Douglass <3107146+raydouglass@users.noreply.github.com> Date: Wed, 7 Aug 2024 22:35:17 -0400 Subject: [PATCH 2/6] Fix ucx-py update-version script (#4603) Back ports the `update-version.sh` fixes from https://github.com/rapidsai/cugraph/pull/4562 --- ci/release/update-version.sh | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index ce2488ad0bf..08c22fca02e 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -69,18 +69,19 @@ DEPENDENCIES=( pyraft raft-dask rmm - ucx-py rapids-dask-dependency ) -for DEP in "${DEPENDENCIES[@]}"; do - for FILE in dependencies.yaml conda/environments/*.yaml python/cugraph-{pyg,dgl}/conda/*.yaml; do +for FILE in dependencies.yaml conda/environments/*.yaml python/cugraph-{pyg,dgl}/conda/*.yaml; do + for DEP in "${DEPENDENCIES[@]}"; do sed_runner "/-.* ${DEP}\(-cu[[:digit:]]\{2\}\)\{0,1\}==/ s/==.*/==${NEXT_SHORT_TAG_PEP440}.*,>=0.0.0a0/g" "${FILE}" - sed_runner "/-.* ucx-py==/ s/==.*/==${NEXT_UCX_PY_VERSION}.*,>=0.0.0a0/g" "${FILE}" done - for FILE in python/**/pyproject.toml python/**/**/pyproject.toml; do + sed_runner "/-.* ucx-py\(-cu[[:digit:]]\{2\}\)\{0,1\}==/ s/==.*/==${NEXT_UCX_PY_VERSION}.*,>=0.0.0a0/g" "${FILE}" +done +for FILE in python/**/pyproject.toml python/**/**/pyproject.toml; do + for DEP in "${DEPENDENCIES[@]}"; do sed_runner "/\"${DEP}\(-cu[[:digit:]]\{2\}\)\{0,1\}==/ s/==.*\"/==${NEXT_SHORT_TAG_PEP440}.*,>=0.0.0a0\"/g" "${FILE}" - sed_runner "/\"ucx-py==/ s/==.*\"/==${NEXT_UCX_PY_VERSION}.*,>=0.0.0a0\"/g" "${FILE}" done + sed_runner "/\"ucx-py\(-cu[[:digit:]]\{2\}\)\{0,1\}==/ s/==.*\"/==${NEXT_UCX_PY_VERSION}.*,>=0.0.0a0\"/g" "${FILE}" done # ucx-py version From d92c257acc88522e775850c2166cd723321caf69 Mon Sep 17 00:00:00 2001 From: Ralph Liu <137829296+nv-rliu@users.noreply.github.com> Date: Thu, 8 Aug 2024 16:46:22 -0400 Subject: [PATCH 3/6] Fix `test_property_graph_mg` Usage of Util Function (#4600) This PR fixes an issue that recently arose due to a change in constructing a `cudf.Series` object. `test_property_graph_mg.py` uses a function from `utilities/utils.py` to construct a Series object from 2d arrays. The function is now using the correct `cudf.Series._from_column` API. Authors: - Ralph Liu (https://github.com/nv-rliu) Approvers: - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/4600 --- python/cugraph/cugraph/utilities/utils.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/cugraph/cugraph/utilities/utils.py b/python/cugraph/cugraph/utilities/utils.py index 7a54a0bf2cf..d04322975c5 100644 --- a/python/cugraph/cugraph/utilities/utils.py +++ b/python/cugraph/cugraph/utilities/utils.py @@ -1,4 +1,4 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. # You may obtain a copy of the License at @@ -530,7 +530,7 @@ def create_list_series_from_2d_ar(ar, index): null_count=0, children=(offset_col, data), ) - return cudf.Series(lc, index=index) + return cudf.Series._from_column(lc, index=index) def create_directory_with_overwrite(directory): From cfda062631c18481dea3adfc89967e464f0cc0ef Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 8 Aug 2024 15:49:52 -0700 Subject: [PATCH 4/6] bug fix --- cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh b/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh index 43415ba6df4..c4c071b8f2c 100644 --- a/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh +++ b/cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh @@ -2041,7 +2041,7 @@ biased_sample_and_compute_local_nbr_indices( zero_bias_frontier_indices.resize(zero_bias_count_inclusive_sums.back(), handle.get_stream()); zero_bias_frontier_indices.shrink_to_fit(handle.get_stream()); - zero_bias_local_nbr_indices.resize(frontier_indices.size(), handle.get_stream()); + zero_bias_local_nbr_indices.resize(zero_bias_frontier_indices.size(), handle.get_stream()); zero_bias_local_nbr_indices.shrink_to_fit(handle.get_stream()); std::vector zero_bias_counts(zero_bias_count_inclusive_sums.size()); std::adjacent_difference(zero_bias_count_inclusive_sums.begin(), From 6091f71fe58e73228ec4c68a62cf263f731394a4 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 8 Aug 2024 21:50:52 -0700 Subject: [PATCH 5/6] udpate transform_v_frontier_e with edge mask to return the transformed values preserving the order in the local neighbor list --- .../prims/detail/transform_v_frontier_e.cuh | 55 ++++++++++--------- 1 file changed, 28 insertions(+), 27 deletions(-) diff --git a/cpp/src/prims/detail/transform_v_frontier_e.cuh b/cpp/src/prims/detail/transform_v_frontier_e.cuh index 7d8824849f0..5ebcddfe8da 100644 --- a/cpp/src/prims/detail/transform_v_frontier_e.cuh +++ b/cpp/src/prims/detail/transform_v_frontier_e.cuh @@ -209,9 +209,6 @@ __global__ static void transform_v_frontier_e_mid_degree( auto const lane_id = tid % raft::warp_size(); size_t idx = static_cast(tid / raft::warp_size()); - using WarpScan = cub::WarpScan; - __shared__ typename WarpScan::TempStorage temp_storage; - while (idx < static_cast(thrust::distance(edge_partition_frontier_key_index_first, edge_partition_frontier_key_index_last))) { auto key_idx = *(edge_partition_frontier_key_index_first + idx); @@ -224,16 +221,15 @@ __global__ static void transform_v_frontier_e_mid_degree( thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_offset); auto this_key_value_first = value_first + edge_partition_frontier_local_degree_offsets[key_idx]; if (edge_partition_e_mask) { - // FIXME: it might be faster to update in warp-sync way - edge_t counter{0}; - for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) { - if ((*edge_partition_e_mask).get(edge_offset + i)) { ++counter; } - } - edge_t offset_within_warp{}; - WarpScan(temp_storage).ExclusiveSum(counter, offset_within_warp); - counter = 0; - for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) { - if ((*edge_partition_e_mask).get(edge_offset + i)) { + auto rounded_up_local_degree = + ((static_cast(local_degree) + (raft::warp_size() - 1)) / raft::warp_size()) * + raft::warp_size(); + edge_t base_offset{0}; + for (edge_t i = lane_id; i < rounded_up_local_degree; i += raft::warp_size()) { + auto valid = (i < local_degree) && (*edge_partition_e_mask).get(edge_offset + i); + auto ballot = __ballot_sync(raft::warp_full_mask(), valid ? uint32_t{1} : uint32_t{0}); + if (valid) { + auto intra_warp_offset = __popc(ballot & ~(raft::warp_full_mask() << lane_id)); transform_v_frontier_e_update_buffer_element( edge_partition, key, @@ -244,9 +240,9 @@ __global__ static void transform_v_frontier_e_mid_degree( edge_partition_dst_value_input, edge_partition_e_value_input, e_op, - this_key_value_first + offset_within_warp + counter); - ++counter; + this_key_value_first + base_offset + intra_warp_offset); } + base_offset += __popc(ballot); } } else { for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) { @@ -300,6 +296,7 @@ __global__ static void transform_v_frontier_e_high_degree( using BlockScan = cub::BlockScan; __shared__ typename BlockScan::TempStorage temp_storage; + __shared__ edge_t increment; while (idx < static_cast(thrust::distance(edge_partition_frontier_key_index_first, edge_partition_frontier_key_index_last))) { @@ -313,16 +310,16 @@ __global__ static void transform_v_frontier_e_high_degree( thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_offset); auto this_key_value_first = value_first + edge_partition_frontier_local_degree_offsets[key_idx]; if (edge_partition_e_mask) { - // FIXME: it might be faster to update in block-sync way - edge_t counter{0}; - for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) { - if ((*edge_partition_e_mask).get(edge_offset + i)) { ++counter; } - } - edge_t offset_within_block{}; - BlockScan(temp_storage).ExclusiveSum(counter, offset_within_block); - counter = 0; - for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) { - if ((*edge_partition_e_mask).get(edge_offset + i)) { + auto rounded_up_local_degree = + ((static_cast(local_degree) + (transform_v_frontier_e_kernel_block_size - 1)) / + transform_v_frontier_e_kernel_block_size) * + transform_v_frontier_e_kernel_block_size; + edge_t base_offset{0}; + for (size_t i = threadIdx.x; i < rounded_up_local_degree; i += blockDim.x) { + auto valid = (i < local_degree) && (*edge_partition_e_mask).get(edge_offset + i); + edge_t intra_block_offset{}; + BlockScan(temp_storage).ExclusiveSum(valid ? edge_t{1} : edge_t{0}, intra_block_offset); + if (valid) { transform_v_frontier_e_update_buffer_element( edge_partition, key, @@ -333,9 +330,13 @@ __global__ static void transform_v_frontier_e_high_degree( edge_partition_dst_value_input, edge_partition_e_value_input, e_op, - this_key_value_first + offset_within_block + counter); - ++counter; + this_key_value_first + base_offset + intra_block_offset); + } + if (threadIdx.x == transform_v_frontier_e_kernel_block_size - 1) { + increment = intra_block_offset + (valid ? edge_t{1} : edge_t{0}); } + __syncthreads(); + base_offset += increment; } } else { for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) { From c574043eed3a3c747cc7b8af315fdf184a61547b Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Thu, 8 Aug 2024 22:59:24 -0700 Subject: [PATCH 6/6] add tests that include zero biases --- ...er_v_random_select_transform_outgoing_e.cu | 270 ++++++++++++------ 1 file changed, 185 insertions(+), 85 deletions(-) diff --git a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu index d77d8a7659e..f698701eb08 100644 --- a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu @@ -15,6 +15,7 @@ */ #include "prims/per_v_random_select_transform_outgoing_e.cuh" +#include "prims/transform_e.cuh" #include "prims/vertex_frontier.cuh" #include "utilities/base_fixture.hpp" #include "utilities/conversion_utilities.hpp" @@ -103,6 +104,7 @@ struct Prims_Usecase { bool with_replacement{false}; bool use_invalid_value{false}; bool use_weight_as_bias{false}; + bool inject_zero_bias{false}; // valid only when use_weight_as_bias is true bool edge_masking{false}; bool check_correctness{true}; }; @@ -159,6 +161,23 @@ class Tests_MGPerVRandomSelectTransformOutgoingE mg_graph_view.attach_edge_mask((*edge_mask).view()); } + if (mg_edge_weight_view && prims_usecase.inject_zero_bias) { + cugraph::transform_e( + *handle_, + mg_graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + *mg_edge_weight_view, + [] __device__(auto src, auto dst, auto, auto, auto w) { + if ((src % 2) == 0 && (dst % 2) == 0) { + return weight_t{0.0}; + } else { + return w; + } + }, + (*mg_edge_weights).mutable_view()); + } + // 2. run MG per_v_random_select_transform_outgoing_e primitive const int hash_bin_count = 5; @@ -324,11 +343,14 @@ class Tests_MGPerVRandomSelectTransformOutgoingE } cugraph::graph_t sg_graph(*handle_); - std::tie(sg_graph, std::ignore, std::ignore, std::ignore) = + std::optional< + cugraph::edge_property_t, weight_t>> + sg_edge_weights{std::nullopt}; + std::tie(sg_graph, sg_edge_weights, std::ignore, std::ignore) = cugraph::test::mg_graph_to_sg_graph( *handle_, mg_graph_view, - std::optional>{std::nullopt}, + mg_edge_weight_view, std::optional>{std::nullopt}, std::make_optional>((*mg_renumber_map).data(), (*mg_renumber_map).size()), @@ -347,6 +369,8 @@ class Tests_MGPerVRandomSelectTransformOutgoingE } auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; rmm::device_uvector sg_offsets(sg_graph_view.number_of_vertices() + vertex_t{1}, handle_->get_stream()); @@ -361,6 +385,17 @@ class Tests_MGPerVRandomSelectTransformOutgoingE sg_graph_view.local_edge_partition_view().indices().end(), sg_indices.begin()); + std::optional> sg_biases{std::nullopt}; + if (sg_edge_weight_view) { + auto firsts = (*sg_edge_weight_view).value_firsts(); + auto counts = (*sg_edge_weight_view).edge_counts(); + assert(firsts.size() == 1); + assert(counts.size() == 1); + sg_biases = rmm::device_uvector(counts[0], handle_->get_stream()); + thrust::copy( + handle_->get_thrust_policy(), firsts[0], firsts[0] + counts[0], (*sg_biases).begin()); + } + auto num_invalids = static_cast(thrust::count_if( handle_->get_thrust_policy(), thrust::make_counting_iterator(size_t{0}), @@ -371,9 +406,10 @@ class Tests_MGPerVRandomSelectTransformOutgoingE : thrust::nullopt, sample_e_op_result_first = cugraph::get_dataframe_buffer_begin(mg_aggregate_sample_e_op_results), - sg_offsets = sg_offsets.begin(), - sg_indices = sg_indices.begin(), - K = prims_usecase.K, + sg_offsets = sg_offsets.begin(), + sg_indices = sg_indices.begin(), + sg_biases = sg_biases ? thrust::make_optional((*sg_biases).begin()) : thrust::nullopt, + K = prims_usecase.K, with_replacement = prims_usecase.with_replacement, invalid_value = invalid_value ? thrust::make_optional(*invalid_value) : thrust::nullopt, @@ -402,6 +438,12 @@ class Tests_MGPerVRandomSelectTransformOutgoingE auto count = offset_last - offset_first; auto out_degree = *(sg_offsets + v + 1) - *(sg_offsets + v); + if (sg_biases) { + out_degree = thrust::count_if(thrust::seq, + *sg_biases + *(sg_offsets + v), + *sg_biases + *(sg_offsets + v + 1), + [] __device__(auto bias) { return bias > 0.0; }); + } if (with_replacement) { if ((out_degree > 0 && count != K) || (out_degree == 0 && count != 0)) { return true; @@ -418,12 +460,33 @@ class Tests_MGPerVRandomSelectTransformOutgoingE auto sg_dst = thrust::get<1>(e_op_result); auto sg_nbr_first = sg_indices + *(sg_offsets + sg_src); auto sg_nbr_last = sg_indices + *(sg_offsets + (sg_src + vertex_t{1})); - if (!thrust::binary_search(thrust::seq, - sg_nbr_first, - sg_nbr_last, - sg_dst)) { // assumed neighbor lists are sorted - return true; + auto sg_nbr_bias_first = + sg_biases ? thrust::make_optional((*sg_biases) + *(sg_offsets + sg_src)) + : thrust::nullopt; + if (sg_src != v) { return true; } + + if (sg_nbr_bias_first) { + auto lower_it = thrust::lower_bound(thrust::seq, sg_nbr_first, sg_nbr_last, sg_dst); + auto upper_it = thrust::upper_bound(thrust::seq, sg_nbr_first, sg_nbr_last, sg_dst); + bool found = false; + for (auto it = (*sg_nbr_bias_first + thrust::distance(sg_nbr_first, lower_it)); + it != (*sg_nbr_bias_first + thrust::distance(sg_nbr_first, upper_it)); + ++it) { + if (*it > 0.0) { + found = true; + break; + } + } + if (!found) { return true; } + } else { + if (!thrust::binary_search(thrust::seq, + sg_nbr_first, + sg_nbr_last, + sg_dst)) { // assumed neighbor lists are sorted + return true; + } } + property_t src_val{}; property_t dst_val{}; if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value) { @@ -443,20 +506,25 @@ class Tests_MGPerVRandomSelectTransformOutgoingE thrust::get<1>(sample_e_op_result_first.get_iterator_tuple()) + offset_first; auto sg_dst_last = thrust::get<1>(sample_e_op_result_first.get_iterator_tuple()) + offset_last; - auto dst_count = - thrust::count(thrust::seq, - sg_dst_first, - sg_dst_last, - sg_dst); // this could be inefficient for high-degree vertices, if - // we sort [sg_dst_first, sg_dst_last) we can use binary - // search but we may better not modify the sampling output - // and allow inefficiency as this is just for testing - auto multiplicity = thrust::distance( - thrust::lower_bound(thrust::seq, sg_nbr_first, sg_nbr_last, sg_dst), + auto dst_count = thrust::count(thrust::seq, sg_dst_first, sg_dst_last, sg_dst); + auto lower_it = + thrust::lower_bound(thrust::seq, + sg_nbr_first, + sg_nbr_last, + sg_dst); // this assumes neighbor lists are sorted + auto upper_it = thrust::upper_bound(thrust::seq, sg_nbr_first, sg_nbr_last, - sg_dst)); // this assumes neighbor lists are sorted + sg_dst); // this assumes neighbor lists are sorted + auto multiplicity = + sg_nbr_bias_first + ? thrust::count_if( + thrust::seq, + *sg_nbr_bias_first + thrust::distance(sg_nbr_first, lower_it), + *sg_nbr_bias_first + thrust::distance(sg_nbr_first, upper_it), + [] __device__(auto bias) { return bias > 0.0; }) + : thrust::distance(lower_it, upper_it); if (dst_count > multiplicity) { return true; } } } @@ -547,44 +615,60 @@ INSTANTIATE_TEST_SUITE_P( file_test, Tests_MGPerVRandomSelectTransformOutgoingE_File, ::testing::Combine( - ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true}), + ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); INSTANTIATE_TEST_SUITE_P( file_large_test, Tests_MGPerVRandomSelectTransformOutgoingE_File, ::testing::Combine( - ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true}), + ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), cugraph::test::File_Usecase("test/datasets/webbase-1M.mtx")))); @@ -593,22 +677,30 @@ INSTANTIATE_TEST_SUITE_P( rmat_small_test, Tests_MGPerVRandomSelectTransformOutgoingE_Rmat, ::testing::Combine( - ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, true}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false}, - Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true}), + ::testing::Values(Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, false, true, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, false, true, true, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, false, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, false}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, false, true}, + Prims_Usecase{size_t{1000}, size_t{4}, true, true, true, true, true}), ::testing::Values(cugraph::test::Rmat_Usecase(10, 16, 0.57, 0.19, 0.19, 0, false, false)))); INSTANTIATE_TEST_SUITE_P( @@ -620,22 +712,30 @@ INSTANTIATE_TEST_SUITE_P( Tests_MGPerVRandomSelectTransformOutgoingE_Rmat, ::testing::Combine( ::testing::Values( - Prims_Usecase{size_t{10000000}, size_t{25}, false, false, false, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, false, false, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, true, false, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, true, false, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, false, false, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, false, false, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, true, false, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, true, false, true, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, false, false}, - Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, true, false}), + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, false, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, false, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, true, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, false, true, true, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, false, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, false, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, true, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, false, true, true, true, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, false, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, false, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, true, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, false, true, true, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, false, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, false, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, false, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, true, false, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, false, true, false}, + Prims_Usecase{size_t{10000000}, size_t{25}, true, true, true, true, true, false}), ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); CUGRAPH_MG_TEST_PROGRAM_MAIN()