From f1490b0b3a6014bf8635560af6f745e15f9c4a77 Mon Sep 17 00:00:00 2001 From: Ralph Liu <137829296+nv-rliu@users.noreply.github.com> Date: Mon, 20 May 2024 10:47:48 -0400 Subject: [PATCH 01/10] Fix MG Betweenness Centrality Test Bug (#4429) Fixes bug seen in MG tests due to incorrect call to `enable_batch` Authors: - Ralph Liu (https://github.com/nv-rliu) Approvers: - Alex Barghi (https://github.com/alexbarghi-nv) URL: https://github.com/rapidsai/cugraph/pull/4429 --- .../centrality/test_betweenness_centrality.py | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/python/cugraph/cugraph/tests/centrality/test_betweenness_centrality.py b/python/cugraph/cugraph/tests/centrality/test_betweenness_centrality.py index ff8859a01b1..6d1f53f7fc3 100644 --- a/python/cugraph/cugraph/tests/centrality/test_betweenness_centrality.py +++ b/python/cugraph/cugraph/tests/centrality/test_betweenness_centrality.py @@ -111,18 +111,13 @@ def calc_betweenness_centrality( else: edge_attr = None - G = None + G = graph_file.get_graph( + download=True, + create_using=cugraph.Graph(directed=directed), + ignore_weights=not edgevals, + ) if multi_gpu_batch: - G = graph_file.get_dask_graph( - create_using=cugraph.Graph(directed=directed), ignore_weights=not edgevals - ) G.enable_batch() - else: - G = graph_file.get_graph( - download=True, - create_using=cugraph.Graph(directed=directed), - ignore_weights=not edgevals, - ) M = G.to_pandas_edgelist().rename( columns={"src": "0", "dst": "1", "wgt": edge_attr} From 30c5edb0a7c01de13587f6a0cc05fb82e59729c9 Mon Sep 17 00:00:00 2001 From: Matthew Roeschke <10647082+mroeschke@users.noreply.github.com> Date: Mon, 20 May 2024 07:06:22 -1000 Subject: [PATCH 02/10] Enable FutureWarnings and DeprecationWarnings as errors in cugraph (#4415) Supersedes https://github.com/rapidsai/cugraph/pull/4271 There are some tests and internal usages of APIs that emit these warnings still. I don't know the best way to handle these warnings, but this PR ensures that _new_ `FutureWarning`s and `DeprecationWarning`s will raise an error in the CI Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/4415 --- python/cugraph/pytest.ini | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/python/cugraph/pytest.ini b/python/cugraph/pytest.ini index ebbd6a7eaad..675a6cf8fde 100644 --- a/python/cugraph/pytest.ini +++ b/python/cugraph/pytest.ini @@ -59,5 +59,14 @@ python_functions = filterwarnings = error:::cudf + error::FutureWarning + error::DeprecationWarning + # TODO + ignore:Multi is deprecated and the removal of multi edges will no longer be supported:FutureWarning + ignore:The legacy column names:FutureWarning + ignore:The include_hop_column flag is deprecated and will be removed:FutureWarning + ignore:Calling uniform_neighbor_sample with the:FutureWarning + ignore:The with_edge_properties flag is deprecated and will be removed:FutureWarning + ignore:This function is deprecated. Batched support for multiple vertices:DeprecationWarning # Called via dask. Not obviously addressable in cugraph. - ignore:The behavior of array concatenation with empty entries is deprecated:FutureWarning:cudf + ignore:The behavior of array concatenation with empty entries is deprecated:FutureWarning From 00563022426232238fd11874930f2dcbaf968922 Mon Sep 17 00:00:00 2001 From: Ray Bell Date: Mon, 20 May 2024 17:09:36 -0400 Subject: [PATCH 03/10] Update operators.rst (#4339) Closes https://github.com/rapidsai/cugraph/issues/4337 Authors: - Ray Bell (https://github.com/raybellwaves) - Tingyu Wang (https://github.com/tingyu66) Approvers: - Tingyu Wang (https://github.com/tingyu66) - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/4339 --- .../api_docs/cugraph-ops/python/operators.rst | 24 +++++++++++++++---- 1 file changed, 20 insertions(+), 4 deletions(-) diff --git a/docs/cugraph/source/api_docs/cugraph-ops/python/operators.rst b/docs/cugraph/source/api_docs/cugraph-ops/python/operators.rst index 3e6664b2db5..8b5efd7aa36 100644 --- a/docs/cugraph/source/api_docs/cugraph-ops/python/operators.rst +++ b/docs/cugraph/source/api_docs/cugraph-ops/python/operators.rst @@ -47,10 +47,26 @@ Graph Attention (GATConv/GATv2Conv) .. autosummary:: :toctree: ../../api/ops - operators.mha_gat_n2n_fwd - operators.mha_gat_n2n_bwd - operators.mha_gat_n2n_efeat_fwd - operators.mha_gat_n2n_efeat_bwd + operators.mha_gat_n2n_fwd_bf16_fp32 + operators.mha_gat_n2n_fwd_fp16_fp32 + operators.mha_gat_n2n_fwd_fp32_fp32 + operators.mha_gat_n2n_bwd_bf16_bf16_bf16_fp32 + operators.mha_gat_n2n_bwd_bf16_bf16_fp32_fp32 + operators.mha_gat_n2n_bwd_bf16_fp32_fp32_fp32 + operators.mha_gat_n2n_bwd_fp16_fp16_fp16_fp32 + operators.mha_gat_n2n_bwd_fp16_fp16_fp32_fp32 + operators.mha_gat_n2n_bwd_fp16_fp32_fp32_fp32 + operators.mha_gat_n2n_bwd_fp32_fp32_fp32_fp32 + operators.mha_gat_n2n_efeat_fwd_bf16_fp32 + operators.mha_gat_n2n_efeat_fwd_fp16_fp32 + operators.mha_gat_n2n_efeat_fwd_fp32_fp32 + operators.mha_gat_n2n_efeat_bwd_bf16_bf16_bf16_fp32 + operators.mha_gat_n2n_efeat_bwd_bf16_bf16_fp32_fp32 + operators.mha_gat_n2n_efeat_bwd_bf16_fp32_fp32_fp32 + operators.mha_gat_n2n_efeat_bwd_fp16_fp16_fp16_fp32 + operators.mha_gat_n2n_efeat_bwd_fp16_fp16_fp32_fp32 + operators.mha_gat_n2n_efeat_bwd_fp16_fp32_fp32_fp32 + operators.mha_gat_n2n_efeat_bwd_fp32_fp32_fp32_fp32 operators.mha_gat_v2_n2n_fwd operators.mha_gat_v2_n2n_bwd From ed5f27fbacc5efc5b84dba1edf67eeb77ce4a1ad Mon Sep 17 00:00:00 2001 From: Paul Taylor <178183+trxcllnt@users.noreply.github.com> Date: Mon, 20 May 2024 14:11:36 -0700 Subject: [PATCH 04/10] Update pip devcontainers to UCX v1.15.0 (#4360) Authors: - Paul Taylor (https://github.com/trxcllnt) Approvers: - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cugraph/pull/4360 --- .devcontainer/cuda11.8-pip/devcontainer.json | 2 +- .devcontainer/cuda12.2-pip/devcontainer.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/.devcontainer/cuda11.8-pip/devcontainer.json b/.devcontainer/cuda11.8-pip/devcontainer.json index d225f15f755..851a992f5b9 100644 --- a/.devcontainer/cuda11.8-pip/devcontainer.json +++ b/.devcontainer/cuda11.8-pip/devcontainer.json @@ -11,7 +11,7 @@ "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/ucx:24.6": { - "version": "1.14.1" + "version": "1.15.0" }, "ghcr.io/rapidsai/devcontainers/features/cuda:24.6": { "version": "11.8", diff --git a/.devcontainer/cuda12.2-pip/devcontainer.json b/.devcontainer/cuda12.2-pip/devcontainer.json index e472f4621f9..c8654ded2ee 100644 --- a/.devcontainer/cuda12.2-pip/devcontainer.json +++ b/.devcontainer/cuda12.2-pip/devcontainer.json @@ -11,7 +11,7 @@ "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/ucx:24.6": { - "version": "1.14.1" + "version": "1.15.0" }, "ghcr.io/rapidsai/devcontainers/features/cuda:24.6": { "version": "12.2", From b9f6e8ca5f067b7c390dd13ed6251abe7ecc71d5 Mon Sep 17 00:00:00 2001 From: Paul Taylor <178183+trxcllnt@users.noreply.github.com> Date: Mon, 20 May 2024 14:13:23 -0700 Subject: [PATCH 05/10] add --rm and --name to devcontainer run args (#4361) * Remove the devcontainer when the VSCode window closes * Adds a descriptive name to the running container: ```shell $ docker ps -a CONTAINER ID IMAGE ... NAMES 0dbb364fe544 vsc-cugraph-... ... rapids-cugraph-24.06-cuda12.2-conda $ docker rm -f rapids-cugraph-24.06-cuda12.2-conda ``` Authors: - Paul Taylor (https://github.com/trxcllnt) - Ralph Liu (https://github.com/nv-rliu) Approvers: - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cugraph/pull/4361 --- .devcontainer/cuda11.8-conda/devcontainer.json | 5 +++++ .devcontainer/cuda11.8-pip/devcontainer.json | 5 +++++ .devcontainer/cuda12.2-conda/devcontainer.json | 5 +++++ .devcontainer/cuda12.2-pip/devcontainer.json | 5 +++++ ci/release/update-version.sh | 1 + 5 files changed, 21 insertions(+) diff --git a/.devcontainer/cuda11.8-conda/devcontainer.json b/.devcontainer/cuda11.8-conda/devcontainer.json index bab521f485d..7c9cd0258a4 100644 --- a/.devcontainer/cuda11.8-conda/devcontainer.json +++ b/.devcontainer/cuda11.8-conda/devcontainer.json @@ -8,6 +8,11 @@ "BASE": "rapidsai/devcontainers:24.06-cpp-cuda11.8-mambaforge-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda11.8-conda" + ], "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} diff --git a/.devcontainer/cuda11.8-pip/devcontainer.json b/.devcontainer/cuda11.8-pip/devcontainer.json index 851a992f5b9..a4dc168505b 100644 --- a/.devcontainer/cuda11.8-pip/devcontainer.json +++ b/.devcontainer/cuda11.8-pip/devcontainer.json @@ -8,6 +8,11 @@ "BASE": "rapidsai/devcontainers:24.06-cpp-cuda11.8-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda11.8-pip" + ], "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/ucx:24.6": { diff --git a/.devcontainer/cuda12.2-conda/devcontainer.json b/.devcontainer/cuda12.2-conda/devcontainer.json index bcaabab572b..eae4967f3b2 100644 --- a/.devcontainer/cuda12.2-conda/devcontainer.json +++ b/.devcontainer/cuda12.2-conda/devcontainer.json @@ -8,6 +8,11 @@ "BASE": "rapidsai/devcontainers:24.06-cpp-mambaforge-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda12.2-conda" + ], "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} diff --git a/.devcontainer/cuda12.2-pip/devcontainer.json b/.devcontainer/cuda12.2-pip/devcontainer.json index c8654ded2ee..393a5c63d23 100644 --- a/.devcontainer/cuda12.2-pip/devcontainer.json +++ b/.devcontainer/cuda12.2-pip/devcontainer.json @@ -8,6 +8,11 @@ "BASE": "rapidsai/devcontainers:24.06-cpp-cuda12.2-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda12.2-pip" + ], "hostRequirements": {"gpu": "optional"}, "features": { "ghcr.io/rapidsai/devcontainers/features/ucx:24.6": { diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 9a7324fb330..f5c14e8d315 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -105,6 +105,7 @@ find .devcontainer/ -type f -name devcontainer.json -print0 | while IFS= read -r sed_runner "s@rapidsai/devcontainers/features/ucx:[0-9.]*@rapidsai/devcontainers/features/ucx:${NEXT_SHORT_TAG_PEP440}@" "${filename}" sed_runner "s@rapidsai/devcontainers/features/cuda:[0-9.]*@rapidsai/devcontainers/features/cuda:${NEXT_SHORT_TAG_PEP440}@" "${filename}" sed_runner "s@rapidsai/devcontainers/features/rapids-build-utils:[0-9.]*@rapidsai/devcontainers/features/rapids-build-utils:${NEXT_SHORT_TAG_PEP440}@" "${filename}" + sed_runner "s@rapids-\${localWorkspaceFolderBasename}-[0-9.]*@rapids-\${localWorkspaceFolderBasename}-${NEXT_SHORT_TAG}@g" "${filename}" done sed_runner "s/:[0-9][0-9]\.[0-9][0-9]/:${NEXT_SHORT_TAG}/" ./notebooks/README.md From 624e961a91387580c788027dd4665b3efdf91f9b Mon Sep 17 00:00:00 2001 From: Erik Welch Date: Mon, 20 May 2024 19:24:47 -0500 Subject: [PATCH 06/10] nx-cugraph: add `ego_graph` (#4395) Authors: - Erik Welch (https://github.com/eriknw) Approvers: - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/4395 --- python/nx-cugraph/README.md | 2 + python/nx-cugraph/_nx_cugraph/__init__.py | 5 + python/nx-cugraph/lint.yaml | 8 +- python/nx-cugraph/nx_cugraph/convert.py | 9 +- .../nx_cugraph/generators/__init__.py | 3 +- .../nx-cugraph/nx_cugraph/generators/ego.py | 161 ++++++++++++++++++ .../nx_cugraph/tests/test_ego_graph.py | 81 +++++++++ python/nx-cugraph/pyproject.toml | 5 +- 8 files changed, 265 insertions(+), 9 deletions(-) create mode 100644 python/nx-cugraph/nx_cugraph/generators/ego.py create mode 100644 python/nx-cugraph/nx_cugraph/tests/test_ego_graph.py diff --git a/python/nx-cugraph/README.md b/python/nx-cugraph/README.md index 75b5c1c5aa9..27825585c28 100644 --- a/python/nx-cugraph/README.md +++ b/python/nx-cugraph/README.md @@ -216,6 +216,8 @@ Below is the list of algorithms that are currently supported in nx-cugraph. └─ wheel_graph community └─ caveman_graph +ego + └─ ego_graph small ├─ bull_graph ├─ chvatal_graph diff --git a/python/nx-cugraph/_nx_cugraph/__init__.py b/python/nx-cugraph/_nx_cugraph/__init__.py index edc96983b8f..f57b90eb402 100644 --- a/python/nx-cugraph/_nx_cugraph/__init__.py +++ b/python/nx-cugraph/_nx_cugraph/__init__.py @@ -77,6 +77,7 @@ "diamond_graph", "dodecahedral_graph", "edge_betweenness_centrality", + "ego_graph", "eigenvector_centrality", "empty_graph", "florentine_families_graph", @@ -163,6 +164,7 @@ "clustering": "Directed graphs and `weight` parameter are not yet supported.", "core_number": "Directed graphs are not yet supported.", "edge_betweenness_centrality": "`weight` parameter is not yet supported, and RNG with seed may be different.", + "ego_graph": "Weighted ego_graph with negative cycles is not yet supported. `NotImplementedError` will be raised if there are negative `distance` edge weights.", "eigenvector_centrality": "`nstart` parameter is not used, but it is checked for validity.", "from_pandas_edgelist": "cudf.DataFrame inputs also supported; value columns with str is unsuppported.", "generic_bfs_edges": "`neighbors` and `sort_neighbors` parameters are not yet supported.", @@ -191,6 +193,9 @@ "bellman_ford_path_length": { "dtype : dtype or None, optional": "The data type (np.float32, np.float64, or None) to use for the edge weights in the algorithm. If None, then dtype is determined by the edge values.", }, + "ego_graph": { + "dtype : dtype or None, optional": "The data type (np.float32, np.float64, or None) to use for the edge weights in the algorithm. If None, then dtype is determined by the edge values.", + }, "eigenvector_centrality": { "dtype : dtype or None, optional": "The data type (np.float32, np.float64, or None) to use for the edge weights in the algorithm. If None, then dtype is determined by the edge values.", }, diff --git a/python/nx-cugraph/lint.yaml b/python/nx-cugraph/lint.yaml index d220cb18df3..c4422ffb97d 100644 --- a/python/nx-cugraph/lint.yaml +++ b/python/nx-cugraph/lint.yaml @@ -26,7 +26,7 @@ repos: - id: mixed-line-ending - id: trailing-whitespace - repo: https://github.com/abravalheri/validate-pyproject - rev: v0.16 + rev: v0.17 hooks: - id: validate-pyproject name: Validate pyproject.toml @@ -50,7 +50,7 @@ repos: - id: black # - id: black-jupyter - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.4.2 + rev: v0.4.4 hooks: - id: ruff args: [--fix-only, --show-fixes] # --unsafe-fixes] @@ -62,7 +62,7 @@ repos: additional_dependencies: &flake8_dependencies # These versions need updated manually - flake8==7.0.0 - - flake8-bugbear==24.4.21 + - flake8-bugbear==24.4.26 - flake8-simplify==0.21.0 - repo: https://github.com/asottile/yesqa rev: v1.5.0 @@ -77,7 +77,7 @@ repos: additional_dependencies: [tomli] files: ^(nx_cugraph|docs)/ - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.4.2 + rev: v0.4.4 hooks: - id: ruff - repo: https://github.com/pre-commit/pre-commit-hooks diff --git a/python/nx-cugraph/nx_cugraph/convert.py b/python/nx-cugraph/nx_cugraph/convert.py index f265540a161..b34245d5031 100644 --- a/python/nx-cugraph/nx_cugraph/convert.py +++ b/python/nx-cugraph/nx_cugraph/convert.py @@ -1,4 +1,4 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. +# Copyright (c) 2023-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 @@ -562,7 +562,12 @@ def to_networkx(G: nxcg.Graph, *, sort_edges: bool = False) -> nx.Graph: dst_iter = map(id_to_key.__getitem__, dst_indices) if G.is_multigraph() and (G.edge_keys is not None or G.edge_indices is not None): if G.edge_keys is not None: - edge_keys = G.edge_keys + if not G.is_directed(): + edge_keys = [k for k, m in zip(G.edge_keys, mask.tolist()) if m] + else: + edge_keys = G.edge_keys + elif not G.is_directed(): + edge_keys = G.edge_indices[mask].tolist() else: edge_keys = G.edge_indices.tolist() if edge_values: diff --git a/python/nx-cugraph/nx_cugraph/generators/__init__.py b/python/nx-cugraph/nx_cugraph/generators/__init__.py index c1834a4dec7..60a9d92373a 100644 --- a/python/nx-cugraph/nx_cugraph/generators/__init__.py +++ b/python/nx-cugraph/nx_cugraph/generators/__init__.py @@ -1,4 +1,4 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. +# Copyright (c) 2023-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 @@ -12,5 +12,6 @@ # limitations under the License. from .classic import * from .community import * +from .ego import * from .small import * from .social import * diff --git a/python/nx-cugraph/nx_cugraph/generators/ego.py b/python/nx-cugraph/nx_cugraph/generators/ego.py new file mode 100644 index 00000000000..66c9c8b95ee --- /dev/null +++ b/python/nx-cugraph/nx_cugraph/generators/ego.py @@ -0,0 +1,161 @@ +# Copyright (c) 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 +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +import math + +import cupy as cp +import networkx as nx +import numpy as np +import pylibcugraph as plc + +import nx_cugraph as nxcg + +from ..utils import _dtype_param, _get_float_dtype, index_dtype, networkx_algorithm + +__all__ = ["ego_graph"] + + +@networkx_algorithm( + extra_params=_dtype_param, version_added="24.06", _plc={"bfs", "ego_graph", "sssp"} +) +def ego_graph( + G, n, radius=1, center=True, undirected=False, distance=None, *, dtype=None +): + """Weighted ego_graph with negative cycles is not yet supported. `NotImplementedError` will be raised if there are negative `distance` edge weights.""" # noqa: E501 + if isinstance(G, nx.Graph): + G = nxcg.from_networkx(G, preserve_all_attrs=True) + if n not in G: + if distance is None: + raise nx.NodeNotFound(f"Source {n} is not in G") + raise nx.NodeNotFound(f"Node {n} not found in graph") + src_index = n if G.key_to_id is None else G.key_to_id[n] + symmetrize = "union" if undirected and G.is_directed() else None + if distance is None or distance not in G.edge_values: + # Simple BFS to determine nodes + if radius is not None and radius <= 0: + if center: + node_ids = cp.array([src_index], dtype=index_dtype) + else: + node_ids = cp.empty(0, dtype=index_dtype) + node_mask = None + else: + if radius is None or np.isinf(radius): + radius = -1 + else: + radius = math.ceil(radius) + distances, unused_predecessors, node_ids = plc.bfs( + handle=plc.ResourceHandle(), + graph=G._get_plc_graph(symmetrize=symmetrize), + sources=cp.array([src_index], index_dtype), + direction_optimizing=False, # True for undirected only; what's best? + depth_limit=radius, + compute_predecessors=False, + do_expensive_check=False, + ) + node_mask = distances != np.iinfo(distances.dtype).max + else: + # SSSP to determine nodes + if callable(distance): + raise NotImplementedError("callable `distance` argument is not supported") + if symmetrize and G.is_multigraph(): + # G._get_plc_graph does not implement `symmetrize=True` w/ edge array + raise NotImplementedError( + "Weighted ego_graph with undirected=True not implemented" + ) + # Check for negative values since we don't support negative cycles + edge_vals = G.edge_values[distance] + if distance in G.edge_masks: + edge_vals = edge_vals[G.edge_masks[distance]] + if (edge_vals < 0).any(): + raise NotImplementedError( + "Negative edge weights not yet supported by ego_graph" + ) + # PERF: we could use BFS if all edges are equal + if radius is None: + radius = np.inf + dtype = _get_float_dtype(dtype, graph=G, weight=distance) + node_ids, distances, unused_predecessors = plc.sssp( + resource_handle=plc.ResourceHandle(), + graph=(G.to_undirected() if symmetrize else G)._get_plc_graph( + distance, 1, dtype + ), + source=src_index, + cutoff=np.nextafter(radius, np.inf, dtype=np.float64), + compute_predecessors=True, # TODO: False is not yet supported + do_expensive_check=False, + ) + node_mask = distances != np.finfo(distances.dtype).max + + if node_mask is not None: + if not center: + node_mask &= node_ids != src_index + node_ids = node_ids[node_mask] + if node_ids.size == G._N: + return G.copy() + # TODO: create renumbering helper function(s) + node_ids.sort() # TODO: is this ever necessary? Keep for safety + node_values = {key: val[node_ids] for key, val in G.node_values.items()} + node_masks = {key: val[node_ids] for key, val in G.node_masks.items()} + + G._sort_edge_indices() # TODO: is this ever necessary? Keep for safety + edge_mask = cp.isin(G.src_indices, node_ids) & cp.isin(G.dst_indices, node_ids) + src_indices = cp.searchsorted(node_ids, G.src_indices[edge_mask]).astype( + index_dtype + ) + dst_indices = cp.searchsorted(node_ids, G.dst_indices[edge_mask]).astype( + index_dtype + ) + edge_values = {key: val[edge_mask] for key, val in G.edge_values.items()} + edge_masks = {key: val[edge_mask] for key, val in G.edge_masks.items()} + + # Renumber nodes + if (id_to_key := G.id_to_key) is not None: + key_to_id = { + id_to_key[old_index]: new_index + for new_index, old_index in enumerate(node_ids.tolist()) + } + else: + key_to_id = { + old_index: new_index + for new_index, old_index in enumerate(node_ids.tolist()) + } + kwargs = { + "N": node_ids.size, + "src_indices": src_indices, + "dst_indices": dst_indices, + "edge_values": edge_values, + "edge_masks": edge_masks, + "node_values": node_values, + "node_masks": node_masks, + "key_to_id": key_to_id, + } + if G.is_multigraph(): + if G.edge_keys is not None: + kwargs["edge_keys"] = [ + x for x, m in zip(G.edge_keys, edge_mask.tolist()) if m + ] + if G.edge_indices is not None: + kwargs["edge_indices"] = G.edge_indices[edge_mask] + rv = G.__class__.from_coo(**kwargs) + rv.graph.update(G.graph) + return rv + + +@ego_graph._can_run +def _(G, n, radius=1, center=True, undirected=False, distance=None, *, dtype=None): + if distance is not None and undirected and G.is_directed() and G.is_multigraph(): + return "Weighted ego_graph with undirected=True not implemented" + if distance is not None and nx.is_negatively_weighted(G, weight=distance): + return "Weighted ego_graph with negative cycles not yet supported" + if callable(distance): + return "callable `distance` argument is not supported" + return True diff --git a/python/nx-cugraph/nx_cugraph/tests/test_ego_graph.py b/python/nx-cugraph/nx_cugraph/tests/test_ego_graph.py new file mode 100644 index 00000000000..5474f9d79e3 --- /dev/null +++ b/python/nx-cugraph/nx_cugraph/tests/test_ego_graph.py @@ -0,0 +1,81 @@ +# Copyright (c) 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 +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +import networkx as nx +import pytest +from packaging.version import parse + +import nx_cugraph as nxcg + +from .testing_utils import assert_graphs_equal + +nxver = parse(nx.__version__) + + +if nxver.major == 3 and nxver.minor < 2: + pytest.skip("Need NetworkX >=3.2 to test ego_graph", allow_module_level=True) + + +@pytest.mark.parametrize( + "create_using", [nx.Graph, nx.DiGraph, nx.MultiGraph, nx.MultiDiGraph] +) +@pytest.mark.parametrize("radius", [-1, 0, 1, 1.5, 2, float("inf"), None]) +@pytest.mark.parametrize("center", [True, False]) +@pytest.mark.parametrize("undirected", [False, True]) +@pytest.mark.parametrize("multiple_edges", [False, True]) +@pytest.mark.parametrize("n", [0, 3]) +def test_ego_graph_cycle_graph( + create_using, radius, center, undirected, multiple_edges, n +): + Gnx = nx.cycle_graph(7, create_using=create_using) + if multiple_edges: + # Test multigraph with multiple edges + if not Gnx.is_multigraph(): + return + Gnx.add_edges_from(nx.cycle_graph(7, create_using=nx.DiGraph).edges) + Gnx.add_edge(0, 1, 10) + Gcg = nxcg.from_networkx(Gnx, preserve_all_attrs=True) + assert_graphs_equal(Gnx, Gcg) # Sanity check + + kwargs = {"radius": radius, "center": center, "undirected": undirected} + Hnx = nx.ego_graph(Gnx, n, **kwargs) + Hcg = nx.ego_graph(Gnx, n, **kwargs, backend="cugraph") + assert_graphs_equal(Hnx, Hcg) + with pytest.raises(nx.NodeNotFound, match="not in G"): + nx.ego_graph(Gnx, -1, **kwargs) + with pytest.raises(nx.NodeNotFound, match="not in G"): + nx.ego_graph(Gnx, -1, **kwargs, backend="cugraph") + # Using sssp with default weight of 1 should give same answer as bfs + nx.set_edge_attributes(Gnx, 1, name="weight") + Gcg = nxcg.from_networkx(Gnx, preserve_all_attrs=True) + assert_graphs_equal(Gnx, Gcg) # Sanity check + + kwargs["distance"] = "weight" + H2nx = nx.ego_graph(Gnx, n, **kwargs) + is_nx32 = nxver.major == 3 and nxver.minor == 2 + if undirected and Gnx.is_directed() and Gnx.is_multigraph(): + if is_nx32: + # `should_run` was added in nx 3.3 + match = "Weighted ego_graph with undirected=True not implemented" + else: + match = "not implemented by cugraph" + with pytest.raises(RuntimeError, match=match): + nx.ego_graph(Gnx, n, **kwargs, backend="cugraph") + with pytest.raises(NotImplementedError, match="ego_graph"): + nx.ego_graph(Gcg, n, **kwargs) + else: + H2cg = nx.ego_graph(Gnx, n, **kwargs, backend="cugraph") + assert_graphs_equal(H2nx, H2cg) + with pytest.raises(nx.NodeNotFound, match="not found in graph"): + nx.ego_graph(Gnx, -1, **kwargs) + with pytest.raises(nx.NodeNotFound, match="not found in graph"): + nx.ego_graph(Gnx, -1, **kwargs, backend="cugraph") diff --git a/python/nx-cugraph/pyproject.toml b/python/nx-cugraph/pyproject.toml index a7daf01775b..477fe8bb493 100644 --- a/python/nx-cugraph/pyproject.toml +++ b/python/nx-cugraph/pyproject.toml @@ -1,4 +1,4 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. +# Copyright (c) 2023-2024, NVIDIA CORPORATION. [build-system] @@ -19,7 +19,7 @@ authors = [ license = { text = "Apache 2.0" } requires-python = ">=3.9" classifiers = [ - "Development Status :: 3 - Alpha", + "Development Status :: 4 - Beta", "License :: OSI Approved :: Apache Software License", "Programming Language :: Python", "Programming Language :: Python :: 3", @@ -233,6 +233,7 @@ ignore = [ "nx_cugraph/**/tests/*py" = ["S101", "S311", "T201", "D103", "D100"] "_nx_cugraph/__init__.py" = ["E501"] "nx_cugraph/algorithms/**/*py" = ["D205", "D401"] # Allow flexible docstrings for algorithms +"nx_cugraph/generators/**/*py" = ["D205", "D401"] # Allow flexible docstrings for generators "nx_cugraph/interface.py" = ["D401"] # Flexible docstrings "scripts/update_readme.py" = ["INP001"] # Not part of a package From 6bd08d245d82bb93d93f81b8875c04e31bb7450c Mon Sep 17 00:00:00 2001 From: Ray Bell Date: Mon, 20 May 2024 21:02:17 -0400 Subject: [PATCH 07/10] test sphinx mapping to networkx (#4323) Closes https://github.com/rapidsai/cugraph/issues/4285 I'll report back if it get to render locally. Ran out of mem building the library on my ec2 machine (g5.2xlarge: 32 Gb RAM) but I did just build cudf. I'll try again soon. ~~Also brainstorming here and I probably should upstream this to the rapids CI process. Would be nice to have /ok to test docs which just builds the docs for this PR~~ created https://github.com/nv-gha-runners/nvidia-runners/issues/25 I tried building the docs locally and got ``` WARNING: [autosummary] failed to import cugraph.jaccard_w. Possible hints: * AttributeError: module 'cugraph' has no attribute 'jaccard_w' * ImportError: * ModuleNotFoundError: No module named 'cugraph.jaccard_w' WARNING: [autosummary] failed to import cugraph.overlap_w. Possible hints: * ModuleNotFoundError: No module named 'cugraph.overlap_w' * ImportError: * AttributeError: module 'cugraph' has no attribute 'overlap_w' WARNING: [autosummary] failed to import cugraph.sorensen_w. Possible hints: * ModuleNotFoundError: No module named 'cugraph.sorensen_w' * ImportError: * AttributeError: module 'cugraph' has no attribute 'sorensen_w' ``` Think this comes from https://github.com/rapidsai/cugraph/blob/abe69c0419b67b567d3c8fce91ee1a062d53e385/docs/cugraph/source/api_docs/cugraph/link_prediction.rst#L14 but I may have messed up my build Authors: - Ray Bell (https://github.com/raybellwaves) - Rick Ratzel (https://github.com/rlratzel) Approvers: - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/4323 --- docs/cugraph/source/conf.py | 14 +++++++++++--- python/cugraph/cugraph/structure/convert_matrix.py | 4 +++- 2 files changed, 14 insertions(+), 4 deletions(-) diff --git a/docs/cugraph/source/conf.py b/docs/cugraph/source/conf.py index 952b962aca2..66bc3137fba 100644 --- a/docs/cugraph/source/conf.py +++ b/docs/cugraph/source/conf.py @@ -190,9 +190,17 @@ 'Miscellaneous'), ] -# Example configuration for intersphinx: refer to the Python standard library. -intersphinx_mapping = {'https://docs.python.org/': None} - +# Connect docs in other projects +intersphinx_mapping = { + "networkx": ( + "https://networkx.org/documentation/stable/", + "https://networkx.org/documentation/stable/objects.inv", + ), + "python": ( + "https://docs.python.org/3", + "https://docs.python.org/3/objects.inv", + ), +} # Config numpydoc numpydoc_show_inherited_class_members = False diff --git a/python/cugraph/cugraph/structure/convert_matrix.py b/python/cugraph/cugraph/structure/convert_matrix.py index ca8e93c482b..b9b9554b870 100644 --- a/python/cugraph/cugraph/structure/convert_matrix.py +++ b/python/cugraph/cugraph/structure/convert_matrix.py @@ -1,4 +1,4 @@ -# Copyright (c) 2019-2023, NVIDIA CORPORATION. +# Copyright (c) 2019-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 @@ -277,6 +277,8 @@ def from_pandas_edgelist( renumber=True, ): """ + See :func:`networkx.convert_matrix.from_pandas_edgelist`. + Initialize a graph from the edge list. It is an error to call this method on an initialized Graph object. Source argument is source column name and destination argument is destination column name. From 2779f3230feee01e1474c8374c688c8f8b021d65 Mon Sep 17 00:00:00 2001 From: Naim <110031745+naimnv@users.noreply.github.com> Date: Tue, 21 May 2024 03:44:59 +0200 Subject: [PATCH 08/10] MNMG Approximation Algorithm for the Weighted Matching Problem (#4315) MNMG [Approximation Algorithm for the Weighted Matching Problem](https://web.archive.org/web/20081031230449id_/http://www.ii.uib.no/~fredrikm/fredrik/papers/CP75.pdf) Authors: - Naim (https://github.com/naimnv) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) - Seunghwa Kang (https://github.com/seunghwak) URL: https://github.com/rapidsai/cugraph/pull/4315 --- cpp/CMakeLists.txt | 2 + cpp/include/cugraph/algorithms.hpp | 26 ++ .../approx_weighted_matching_impl.cuh | 392 ++++++++++++++++++ .../community/approx_weighted_matching_mg.cu | 50 +++ .../community/approx_weighted_matching_sg.cu | 50 +++ cpp/tests/CMakeLists.txt | 8 + .../community/mg_weighted_matching_test.cpp | 232 +++++++++++ .../community/weighted_matching_test.cpp | 182 ++++++++ 8 files changed, 942 insertions(+) create mode 100644 cpp/src/community/approx_weighted_matching_impl.cuh create mode 100644 cpp/src/community/approx_weighted_matching_mg.cu create mode 100644 cpp/src/community/approx_weighted_matching_sg.cu create mode 100644 cpp/tests/community/mg_weighted_matching_test.cpp create mode 100644 cpp/tests/community/weighted_matching_test.cpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d3dfdbd068c..57e0aa2d078 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -288,6 +288,8 @@ set(CUGRAPH_SOURCES src/structure/symmetrize_edgelist_mg.cu src/community/triangle_count_sg.cu src/community/triangle_count_mg.cu + src/community/approx_weighted_matching_sg.cu + src/community/approx_weighted_matching_mg.cu src/traversal/k_hop_nbrs_sg.cu src/traversal/k_hop_nbrs_mg.cu src/mtmg/vertex_result.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index 0caa151daac..7c4a978c4b4 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -2368,6 +2368,32 @@ rmm::device_uvector vertex_coloring( graph_view_t const& graph_view, raft::random::RngState& rng_state); +/* + * @brief Approximate Weighted Matching + * + * A matching in an undirected graph G = (V, E) is a pairing of adjacent vertices + * such that each vertex is matched with at most one other vertex, the objective + * being to match as many vertices as possible or to maximise the sum of the + * weights of the matched edges. Here we provide an implementation of an + * approximation algorithm to the weighted Maximum matching. See + * https://web.archive.org/web/20081031230449id_/http://www.ii.uib.no/~fredrikm/fredrik/papers/CP75.pdf + * for further information. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam edge_t Type of edge identifiers. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, + * and handles to various CUDA libraries) to run graph algorithms. + * @param[in] graph_view Graph view object. + * @param[in] edge_weight_view View object holding edge weights for @p graph_view. + * @return A tuple of device vector of matched vertex ids and sum of the weights of the matched + * edges. + */ +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); } // namespace cugraph /** diff --git a/cpp/src/community/approx_weighted_matching_impl.cuh b/cpp/src/community/approx_weighted_matching_impl.cuh new file mode 100644 index 00000000000..e693beee489 --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_impl.cuh @@ -0,0 +1,392 @@ +/* + * Copyright (c) 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#pragma once + +#include "prims/fill_edge_property.cuh" +#include "prims/reduce_op.cuh" +#include "prims/transform_e.cuh" +#include "prims/transform_reduce_e_by_src_dst_key.cuh" +#include "prims/update_edge_src_dst_property.cuh" +#include "utilities/collect_comm.cuh" + +#include +#include +#include + +#include + +#include + +namespace cugraph { + +namespace detail { + +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + edge_property_view_t edge_weight_view) +{ + CUGRAPH_EXPECTS(graph_view.is_symmetric(), + "Invalid input arguments: input graph for approximate_weighted_matching must " + "need to be symmetric"); + + using graph_view_t = cugraph::graph_view_t; + + graph_view_t current_graph_view(graph_view); + if (current_graph_view.has_edge_mask()) { current_graph_view.clear_edge_mask(); } + + cugraph::edge_property_t edge_masks_even(handle, current_graph_view); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_even); + cugraph::edge_property_t edge_masks_odd(handle, current_graph_view); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_odd); + + if (graph_view.has_edge_mask()) { + current_graph_view.attach_edge_mask(*(graph_view.edge_mask_view())); + } + // Mask out self-loop + cugraph::transform_e( + handle, + current_graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, thrust::nullopt_t, thrust::nullopt_t, thrust::nullopt_t) { + return !(src == dst); + }, + edge_masks_even.mutable_view()); + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + current_graph_view.attach_edge_mask(edge_masks_even.view()); + + auto constexpr invalid_partner = invalid_vertex_id::value; + rmm::device_uvector offers_from_partners( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + + rmm::device_uvector partners(current_graph_view.local_vertex_partition_range_size(), + handle.get_stream()); + + thrust::fill(handle.get_thrust_policy(), partners.begin(), partners.end(), invalid_partner); + thrust::fill(handle.get_thrust_policy(), + offers_from_partners.begin(), + offers_from_partners.end(), + weight_t{0.0}); + + rmm::device_uvector local_vertices( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + detail::sequence_fill(handle.get_stream(), + local_vertices.begin(), + local_vertices.size(), + current_graph_view.local_vertex_partition_range_first()); + + edge_src_property_t src_key_cache(handle); + cugraph::edge_src_property_t src_match_flags(handle); + cugraph::edge_dst_property_t dst_match_flags(handle); + + if constexpr (graph_view_t::is_multi_gpu) { + src_key_cache = edge_src_property_t(handle, current_graph_view); + + update_edge_src_property(handle, current_graph_view, local_vertices.begin(), src_key_cache); + + src_match_flags = cugraph::edge_src_property_t(handle, current_graph_view); + dst_match_flags = cugraph::edge_dst_property_t(handle, current_graph_view); + } + + vertex_t loop_counter = 0; + while (true) { + // + // For each candidate vertex, find the best possible target + // + + rmm::device_uvector candidates(0, handle.get_stream()); + rmm::device_uvector offers_from_candidates(0, handle.get_stream()); + rmm::device_uvector targets(0, handle.get_stream()); + + // FIXME: This can be implemented more efficiently if per_v_transform_reduce_incoming|outgoing_e + // is updated to support reduction on thrust::tuple. + std::forward_as_tuple(candidates, std::tie(offers_from_candidates, targets)) = + cugraph::transform_reduce_e_by_src_key( + handle, + current_graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + edge_weight_view, + graph_view_t::is_multi_gpu + ? src_key_cache.view() + : detail::edge_major_property_view_t(local_vertices.begin()), + [] __device__(auto, auto dst, thrust::nullopt_t, thrust::nullopt_t, auto wt) { + return thrust::make_tuple(wt, dst); + }, + thrust::make_tuple(weight_t{0.0}, invalid_partner), + reduce_op::maximum>{}, + true); + + // + // For each target, find the best offer + // + + if constexpr (graph_view_t::is_multi_gpu) { + auto vertex_partition_range_lasts = current_graph_view.vertex_partition_range_lasts(); + + rmm::device_uvector d_vertex_partition_range_lasts( + vertex_partition_range_lasts.size(), handle.get_stream()); + + raft::update_device(d_vertex_partition_range_lasts.data(), + vertex_partition_range_lasts.data(), + vertex_partition_range_lasts.size(), + handle.get_stream()); + + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + auto key_func = cugraph::detail::compute_gpu_id_from_int_vertex_t{ + raft::device_span(d_vertex_partition_range_lasts.data(), + d_vertex_partition_range_lasts.size()), + major_comm_size, + minor_comm_size}; + + std::forward_as_tuple(std::tie(candidates, offers_from_candidates, targets), std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator(thrust::make_tuple( + candidates.begin(), offers_from_candidates.begin(), targets.begin())), + thrust::make_zip_iterator( + thrust::make_tuple(candidates.end(), offers_from_candidates.end(), targets.end())), + [key_func] __device__(auto val) { return key_func(thrust::get<2>(val)); }, + handle.get_stream()); + } + + auto itr_to_tuples = thrust::make_zip_iterator( + thrust::make_tuple(offers_from_candidates.begin(), candidates.begin())); + + thrust::sort_by_key(handle.get_thrust_policy(), targets.begin(), targets.end(), itr_to_tuples); + + auto nr_unique_targets = thrust::count_if(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(targets.size()), + is_first_in_run_t{targets.data()}); + + rmm::device_uvector unique_targets(nr_unique_targets, handle.get_stream()); + rmm::device_uvector best_offers_to_targets(nr_unique_targets, handle.get_stream()); + rmm::device_uvector best_candidates(nr_unique_targets, handle.get_stream()); + + auto itr_to_reduced_tuples = thrust::make_zip_iterator( + thrust::make_tuple(best_offers_to_targets.begin(), best_candidates.begin())); + + auto new_end = thrust::reduce_by_key( + handle.get_thrust_policy(), + targets.begin(), + targets.end(), + itr_to_tuples, + unique_targets.begin(), + itr_to_reduced_tuples, + thrust::equal_to{}, + [] __device__(auto pair1, auto pair2) { return (pair1 > pair2) ? pair1 : pair2; }); + + vertex_t nr_reduces_tuples = + static_cast(thrust::distance(unique_targets.begin(), new_end.first)); + + targets = std::move(unique_targets); + offers_from_candidates = std::move(best_offers_to_targets); + candidates = std::move(best_candidates); + + // + // two vertex offer each other, that's a match + // + + kv_store_t target_candidate_map(targets.begin(), + targets.end(), + candidates.begin(), + invalid_vertex_id::value, + invalid_vertex_id::value, + handle.get_stream()); + + rmm::device_uvector candidates_of_candidates(0, handle.get_stream()); + + if (graph_view_t::is_multi_gpu) { + auto& comm = handle.get_comms(); + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + auto partitions_range_lasts = graph_view.vertex_partition_range_lasts(); + rmm::device_uvector d_partitions_range_lasts(partitions_range_lasts.size(), + handle.get_stream()); + + raft::update_device(d_partitions_range_lasts.data(), + partitions_range_lasts.data(), + partitions_range_lasts.size(), + handle.get_stream()); + + cugraph::detail::compute_gpu_id_from_int_vertex_t vertex_to_gpu_id_op{ + raft::device_span(d_partitions_range_lasts.data(), + d_partitions_range_lasts.size()), + major_comm_size, + minor_comm_size}; + + candidates_of_candidates = cugraph::collect_values_for_keys(handle, + target_candidate_map.view(), + candidates.begin(), + candidates.end(), + vertex_to_gpu_id_op); + } else { + candidates_of_candidates.resize(candidates.size(), handle.get_stream()); + + target_candidate_map.view().find(candidates.begin(), + candidates.end(), + candidates_of_candidates.begin(), + handle.get_stream()); + } + + // + // Mask out neighborhood of matched vertices + // + + rmm::device_uvector is_vertex_matched = rmm::device_uvector( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + thrust::fill( + handle.get_thrust_policy(), is_vertex_matched.begin(), is_vertex_matched.end(), bool{false}); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_zip_iterator(thrust::make_tuple(candidates_of_candidates.begin(), + targets.begin(), + candidates.begin(), + offers_from_candidates.begin())), + thrust::make_zip_iterator(thrust::make_tuple(candidates_of_candidates.end(), + targets.end(), + candidates.end(), + offers_from_candidates.end())), + [partners = partners.begin(), + offers_from_partners = offers_from_partners.begin(), + is_vertex_matched = + raft::device_span(is_vertex_matched.data(), is_vertex_matched.size()), + v_first = + current_graph_view.local_vertex_partition_range_first()] __device__(auto msrc_tgt) { + auto candidate_of_candidate = thrust::get<0>(msrc_tgt); + auto tgt = thrust::get<1>(msrc_tgt); + auto candiate = thrust::get<2>(msrc_tgt); + auto offer_value = thrust::get<3>(msrc_tgt); + + if (candidate_of_candidate != invalid_partner && candidate_of_candidate == tgt) { + auto tgt_offset = tgt - v_first; + is_vertex_matched[tgt_offset] = true; + partners[tgt_offset] = candiate; + offers_from_partners[tgt_offset] = offer_value; + } + }); + + if (current_graph_view.compute_number_of_edges(handle) == 0) { break; } + + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::update_edge_src_property( + handle, current_graph_view, is_vertex_matched.begin(), src_match_flags); + cugraph::update_edge_dst_property( + handle, current_graph_view, is_vertex_matched.begin(), dst_match_flags); + } + + if (loop_counter % 2 == 0) { + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::transform_e( + handle, + current_graph_view, + src_match_flags.view(), + dst_match_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_odd.mutable_view()); + } else { + cugraph::transform_e( + handle, + current_graph_view, + detail::edge_major_property_view_t(is_vertex_matched.begin()), + detail::edge_minor_property_view_t(is_vertex_matched.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_odd.mutable_view()); + } + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_even); + current_graph_view.attach_edge_mask(edge_masks_odd.view()); + } else { + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::transform_e( + handle, + current_graph_view, + src_match_flags.view(), + dst_match_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_even.mutable_view()); + } else { + cugraph::transform_e( + handle, + current_graph_view, + detail::edge_major_property_view_t(is_vertex_matched.begin()), + detail::edge_minor_property_view_t(is_vertex_matched.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_even.mutable_view()); + } + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_odd); + current_graph_view.attach_edge_mask(edge_masks_even.view()); + } + + loop_counter++; + } + + weight_t sum_matched_edge_weights = thrust::reduce( + handle.get_thrust_policy(), offers_from_partners.begin(), offers_from_partners.end()); + + if constexpr (graph_view_t::is_multi_gpu) { + sum_matched_edge_weights = host_scalar_allreduce( + handle.get_comms(), sum_matched_edge_weights, raft::comms::op_t::SUM, handle.get_stream()); + } + + return std::make_tuple(std::move(partners), sum_matched_edge_weights / 2.0); +} +} // namespace detail + +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view) +{ + return detail::approximate_weighted_matching(handle, graph_view, edge_weight_view); +} + +} // namespace cugraph diff --git a/cpp/src/community/approx_weighted_matching_mg.cu b/cpp/src/community/approx_weighted_matching_mg.cu new file mode 100644 index 00000000000..41d6c3d97e0 --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_mg.cu @@ -0,0 +1,50 @@ +/* + * Copyright (c) 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "approx_weighted_matching_impl.cuh" + +namespace cugraph { + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +} // namespace cugraph diff --git a/cpp/src/community/approx_weighted_matching_sg.cu b/cpp/src/community/approx_weighted_matching_sg.cu new file mode 100644 index 00000000000..418a43d51ae --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_sg.cu @@ -0,0 +1,50 @@ +/* + * Copyright (c) 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "approx_weighted_matching_impl.cuh" + +namespace cugraph { + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 19097add541..ced3b7bedb1 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -309,6 +309,10 @@ ConfigureTest(LOUVAIN_TEST community/louvain_test.cpp) # - LEIDEN tests ---------------------------------------------------------------------------------- ConfigureTest(LEIDEN_TEST community/leiden_test.cpp) +################################################################################################### +# - WEIGHTED MATCHING tests ---------------------------------------------------------------------------------- +ConfigureTest(WEIGHTED_MATCHING_TEST community/weighted_matching_test.cpp) + ################################################################################################### # - Legacy ECG tests ------------------------------------------------------------------------------------- ConfigureTest(LEGACY_ECG_TEST community/legacy_ecg_test.cpp) @@ -570,6 +574,10 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG LEIDEN tests -------------------------------------------------------------------------- ConfigureTestMG(MG_LEIDEN_TEST community/mg_leiden_test.cpp) + ############################################################################################### + # - MG WEIGHTED MATCHING tests -------------------------------------------------------------------------- + ConfigureTestMG(MG_WEIGHTED_MATCHING_TEST community/mg_weighted_matching_test.cpp) + ############################################################################################### # - MG ECG tests -------------------------------------------------------------------------- ConfigureTestMG(MG_ECG_TEST community/mg_ecg_test.cpp) diff --git a/cpp/tests/community/mg_weighted_matching_test.cpp b/cpp/tests/community/mg_weighted_matching_test.cpp new file mode 100644 index 00000000000..21963922ab1 --- /dev/null +++ b/cpp/tests/community/mg_weighted_matching_test.cpp @@ -0,0 +1,232 @@ +/* + * Copyright (c) 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct WeightedMatching_UseCase { + bool edge_masking{false}; + bool check_correctness{true}; +}; + +template +class Tests_MGWeightedMatching + : public ::testing::TestWithParam> { + public: + Tests_MGWeightedMatching() {} + + static void SetUpTestCase() { handle_ = cugraph::test::initialize_mg_handle(); } + static void TearDownTestCase() { handle_.reset(); } + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [weighted_matching_usecase, input_usecase] = param; + + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.start("MG Construct graph"); + } + + constexpr bool multi_gpu = true; + + bool test_weighted = true; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [mg_graph, mg_edge_weights, mg_renumber_map] = + cugraph::test::construct_graph( + *handle_, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(mg_graph, mg_edge_weights, mg_renumber_map) = cugraph::symmetrize_graph( + *handle_, + std::move(mg_graph), + std::move(mg_edge_weights), + mg_renumber_map ? std::optional>(std::move(*mg_renumber_map)) + : std::nullopt, + false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto mg_graph_view = mg_graph.view(); + auto mg_edge_weight_view = + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + if (weighted_matching_usecase.edge_masking) { + edge_mask = cugraph::test::generate::edge_property( + *handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + + rmm::device_uvector mg_partners(0, handle_->get_stream()); + weight_t mg_matching_weights; + + std::forward_as_tuple(mg_partners, mg_matching_weights) = + cugraph::approximate_weighted_matching( + *handle_, mg_graph_view, (*mg_edge_weights).view()); + + if (weighted_matching_usecase.check_correctness) { + auto h_mg_partners = cugraph::test::to_host(*handle_, mg_partners); + + auto constexpr invalid_partner = cugraph::invalid_vertex_id::value; + + rmm::device_uvector mg_aggregate_partners(0, handle_->get_stream()); + std::tie(std::ignore, mg_aggregate_partners) = + cugraph::test::mg_vertex_property_values_to_sg_vertex_property_values( + *handle_, + std::optional>{std::nullopt}, + mg_graph_view.local_vertex_partition_range(), + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + raft::device_span(mg_partners.data(), mg_partners.size())); + + cugraph::graph_t sg_graph(*handle_); + std::optional< + cugraph::edge_property_t, weight_t>> + sg_edge_weights{std::nullopt}; + std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>(std::nullopt), + false); + + if (handle_->get_comms().get_rank() == 0) { + auto sg_graph_view = sg_graph.view(); + + rmm::device_uvector sg_partners(0, handle_->get_stream()); + weight_t sg_matching_weights; + + std::forward_as_tuple(sg_partners, sg_matching_weights) = + cugraph::approximate_weighted_matching( + *handle_, sg_graph_view, (*sg_edge_weights).view()); + auto h_sg_partners = cugraph::test::to_host(*handle_, sg_partners); + auto h_mg_aggregate_partners = cugraph::test::to_host(*handle_, mg_aggregate_partners); + + ASSERT_FLOAT_EQ(mg_matching_weights, sg_matching_weights) + << "SG and MG matching weights are different"; + ASSERT_TRUE( + std::equal(h_sg_partners.begin(), h_sg_partners.end(), h_mg_aggregate_partners.begin())); + } + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGWeightedMatching::handle_ = nullptr; + +using Tests_MGWeightedMatching_File = Tests_MGWeightedMatching; +using Tests_MGWeightedMatching_Rmat = Tests_MGWeightedMatching; + +TEST_P(Tests_MGWeightedMatching_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_MGWeightedMatching_File, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGWeightedMatching_Rmat, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 2, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_MGWeightedMatching_Rmat, + ::testing::Combine( + ::testing::Values(WeightedMatching_UseCase{false, false}, + WeightedMatching_UseCase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, true, false)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/weighted_matching_test.cpp b/cpp/tests/community/weighted_matching_test.cpp new file mode 100644 index 00000000000..436273c3be3 --- /dev/null +++ b/cpp/tests/community/weighted_matching_test.cpp @@ -0,0 +1,182 @@ +/* + * Copyright (c) 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct WeightedMatching_UseCase { + bool edge_masking{false}; + bool check_correctness{true}; +}; + +template +class Tests_SGWeightedMatching + : public ::testing::TestWithParam> { + public: + Tests_SGWeightedMatching() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [weighted_matching_usecase, input_usecase] = param; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.start("Construct graph"); + } + + constexpr bool multi_gpu = false; + + bool test_weighted = true; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(sg_graph, sg_edge_weights, sg_renumber_map) = cugraph::symmetrize_graph( + handle, std::move(sg_graph), std::move(sg_edge_weights), std::move(sg_renumber_map), false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + if (weighted_matching_usecase.edge_masking) { + edge_mask = cugraph::test::generate::edge_property( + handle, sg_graph_view, 2); + sg_graph_view.attach_edge_mask((*edge_mask).view()); + } + + rmm::device_uvector d_partners(0, handle.get_stream()); + weight_t total_matching_weights; + + std::forward_as_tuple(d_partners, total_matching_weights) = + cugraph::approximate_weighted_matching( + handle, sg_graph_view, (*sg_edge_weights).view()); + + if (weighted_matching_usecase.check_correctness) { + auto h_partners = cugraph::test::to_host(handle, d_partners); + auto constexpr invalid_partner = cugraph::invalid_vertex_id::value; + + std::for_each(h_partners.begin(), h_partners.end(), [&invalid_partner, h_partners](auto& v) { + if (v != invalid_partner) ASSERT_TRUE(h_partners[h_partners[v]] == v); + }); + } + } +}; + +using Tests_SGWeightedMatching_File = Tests_SGWeightedMatching; +using Tests_SGWeightedMatching_Rmat = Tests_SGWeightedMatching; + +TEST_P(Tests_SGWeightedMatching_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGWeightedMatching_File, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_SGWeightedMatching_Rmat, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 3, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_SGWeightedMatching_Rmat, + ::testing::Combine( + ::testing::Values(WeightedMatching_UseCase{false, false}, + WeightedMatching_UseCase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN() From 5b72af3a8a8434201a017d659108701ff3a077f6 Mon Sep 17 00:00:00 2001 From: Naim <110031745+naimnv@users.noreply.github.com> Date: Tue, 21 May 2024 16:25:02 +0200 Subject: [PATCH 09/10] Fix a bug in kv_store_t implementation (#4434) Fix a bug in kv_store_t implementation Authors: - Naim (https://github.com/naimnv) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/4434 --- cpp/src/prims/kv_store.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/src/prims/kv_store.cuh b/cpp/src/prims/kv_store.cuh index 2cc7856d87a..76b64b5692b 100644 --- a/cpp/src/prims/kv_store.cuh +++ b/cpp/src/prims/kv_store.cuh @@ -526,6 +526,7 @@ class kv_cuco_store_t { std::conditional_t, value_t, void>>(0, stream)) { allocate(capacity, invalid_key, invalid_value, stream); + if constexpr (!std::is_arithmetic_v) { invalid_value_ = invalid_value; } capacity_ = capacity; size_ = 0; } From 8d8b4fde6922efbae88b46e840b9af7f089b30bb Mon Sep 17 00:00:00 2001 From: Ray Bell Date: Tue, 21 May 2024 12:18:07 -0400 Subject: [PATCH 10/10] Update DGL_support.md (#4327) Believe the path suggested previously is outdated. Authors: - Ray Bell (https://github.com/raybellwaves) - Don Acosta (https://github.com/acostadon) - Rick Ratzel (https://github.com/rlratzel) Approvers: - Don Acosta (https://github.com/acostadon) URL: https://github.com/rapidsai/cugraph/pull/4327 --- docs/cugraph/source/graph_support/DGL_support.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/cugraph/source/graph_support/DGL_support.md b/docs/cugraph/source/graph_support/DGL_support.md index dc4f66180ac..9df462155fd 100644 --- a/docs/cugraph/source/graph_support/DGL_support.md +++ b/docs/cugraph/source/graph_support/DGL_support.md @@ -17,7 +17,7 @@ mamba install cugraph-dgl -c rapidsai-nightly -c rapidsai -c pytorch -c conda-fo ### Create the conda development environment ``` -mamba env create -n cugraph_dgl_dev --file conda/cugraph_dgl_dev_11.6.yml +conda env create -n cugraph_dgl_dev --file conda/environments/all_cuda-122_arch-x86_64.yaml ``` ### Install in editable mode