From 00563022426232238fd11874930f2dcbaf968922 Mon Sep 17 00:00:00 2001 From: Ray Bell Date: Mon, 20 May 2024 17:09:36 -0400 Subject: [PATCH 01/23] 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 02/23] 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 03/23] 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 04/23] 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 05/23] 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 06/23] 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 07/23] 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 08/23] 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 From ddfaacf1bd9d305b31e2c55b7ae954ba13f21899 Mon Sep 17 00:00:00 2001 From: Ray Bell Date: Tue, 21 May 2024 21:24:18 -0400 Subject: [PATCH 09/23] DOC: doc-update-link-for-cugraphops (#4279) Fixes a broken link https://github.com/rapidsai/cugraph-ops/blob/branch-23.04/README.md -> https://github.com/rapidsai/cugraph/blob/branch-24.04/readme_pages/cugraph_ops.md Authors: - Ray Bell (https://github.com/raybellwaves) - Alex Barghi (https://github.com/alexbarghi-nv) - Rick Ratzel (https://github.com/rlratzel) Approvers: - Don Acosta (https://github.com/acostadon) URL: https://github.com/rapidsai/cugraph/pull/4279 --- docs/cugraph/source/graph_support/cugraphops_support.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/cugraph/source/graph_support/cugraphops_support.rst b/docs/cugraph/source/graph_support/cugraphops_support.rst index fd79564f849..96b13f62a9c 100644 --- a/docs/cugraph/source/graph_support/cugraphops_support.rst +++ b/docs/cugraph/source/graph_support/cugraphops_support.rst @@ -7,4 +7,4 @@ cugraph-ops aims to be a low-level, framework agnostic library providing commonl .. toctree:: :maxdepth: 3 - https://github.com/rapidsai/cugraph-ops/blob/branch-23.04/README.md + https://github.com/rapidsai/cugraph/blob/branch-24.06/readme_pages/cugraph_ops.md From 0d64b729568ea5240ea32d990542a9edf471f349 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Thu, 23 May 2024 17:12:32 -0400 Subject: [PATCH 10/23] Fix stream synchronization in MTMG graph construction (#4275) Restructure to pass stream instead of handle and synchronize appropriately. Closes #4236 Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Seunghwa Kang (https://github.com/seunghwak) URL: https://github.com/rapidsai/cugraph/pull/4275 --- .../mtmg/detail/per_device_edgelist.hpp | 90 +++++++++---------- cpp/include/cugraph/mtmg/edge_property.hpp | 3 +- .../cugraph/mtmg/edge_property_view.hpp | 3 +- cpp/include/cugraph/mtmg/edgelist.hpp | 10 ++- cpp/include/cugraph/mtmg/handle.hpp | 7 +- .../cugraph/mtmg/per_thread_edgelist.hpp | 45 +++++----- cpp/tests/mtmg/multi_node_threaded_test.cu | 6 +- cpp/tests/mtmg/threaded_test.cu | 6 +- cpp/tests/mtmg/threaded_test_jaccard.cu | 6 +- cpp/tests/mtmg/threaded_test_louvain.cu | 6 +- 10 files changed, 96 insertions(+), 86 deletions(-) diff --git a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp index 7fd5bb726e6..63d7fd9685e 100644 --- a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp +++ b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -62,17 +62,17 @@ class per_device_edgelist_t { /** * @brief Construct a new per device edgelist t object * - * @param handle MTMG resource handle - used to identify GPU resources * @param device_buffer_size Number of edges to store in each device buffer * @param use_weight Whether or not the edgelist will have weights * @param use_edge_id Whether or not the edgelist will have edge ids * @param use_edge_type Whether or not the edgelist will have edge types + * @param stream_view CUDA stream view */ - per_device_edgelist_t(cugraph::mtmg::handle_t const& handle, - size_t device_buffer_size, + per_device_edgelist_t(size_t device_buffer_size, bool use_weight, bool use_edge_id, - bool use_edge_type) + bool use_edge_type, + rmm::cuda_stream_view stream_view) : device_buffer_size_{device_buffer_size}, current_pos_{0}, src_{}, @@ -89,7 +89,7 @@ class per_device_edgelist_t { edge_type_ = std::make_optional(std::vector>()); } - create_new_buffers(handle); + create_new_buffers(stream_view); } /** @@ -111,19 +111,19 @@ class per_device_edgelist_t { /** * @brief Append a list of edges to the edge list * - * @param handle The resource handle - * @param src Source vertex id - * @param dst Destination vertex id - * @param wgt Edge weight - * @param edge_id Edge id - * @param edge_type Edge type + * @param src Source vertex id + * @param dst Destination vertex id + * @param wgt Edge weight + * @param edge_id Edge id + * @param edge_type Edge type + * @param stream_view CUDA stream view */ - void append(handle_t const& handle, - raft::host_span src, + void append(raft::host_span src, raft::host_span dst, std::optional> wgt, std::optional> edge_id, - std::optional> edge_type) + std::optional> edge_type, + rmm::cuda_stream_view stream_view) { std::vector> copy_positions; @@ -142,13 +142,13 @@ class per_device_edgelist_t { pos += copy_count; current_pos_ += copy_count; - if (current_pos_ == src_.back().size()) { create_new_buffers(handle); } + if (current_pos_ == src_.back().size()) { create_new_buffers(stream_view); } } } std::for_each(copy_positions.begin(), copy_positions.end(), - [&handle, + [&stream_view, &this_src = src_, &src, &this_dst = dst_, @@ -164,47 +164,45 @@ class per_device_edgelist_t { raft::update_device(this_src[buffer_idx].begin() + buffer_pos, src.begin() + input_pos, copy_count, - handle.get_stream()); + stream_view); raft::update_device(this_dst[buffer_idx].begin() + buffer_pos, dst.begin() + input_pos, copy_count, - handle.get_stream()); + stream_view); if (this_wgt) raft::update_device((*this_wgt)[buffer_idx].begin() + buffer_pos, wgt->begin() + input_pos, copy_count, - handle.get_stream()); + stream_view); if (this_edge_id) raft::update_device((*this_edge_id)[buffer_idx].begin() + buffer_pos, edge_id->begin() + input_pos, copy_count, - handle.get_stream()); + stream_view); if (this_edge_type) raft::update_device((*this_edge_type)[buffer_idx].begin() + buffer_pos, edge_type->begin() + input_pos, copy_count, - handle.get_stream()); + stream_view); }); - - handle.sync_stream(); } /** * @brief Mark the edgelist as ready for reading (all writes are complete) * - * @param handle The resource handle + * @param stream_view CUDA stream view */ - void finalize_buffer(handle_t const& handle) + void finalize_buffer(rmm::cuda_stream_view stream_view) { - src_.back().resize(current_pos_, handle.get_stream()); - dst_.back().resize(current_pos_, handle.get_stream()); - if (wgt_) wgt_->back().resize(current_pos_, handle.get_stream()); - if (edge_id_) edge_id_->back().resize(current_pos_, handle.get_stream()); - if (edge_type_) edge_type_->back().resize(current_pos_, handle.get_stream()); + src_.back().resize(current_pos_, stream_view); + dst_.back().resize(current_pos_, stream_view); + if (wgt_) wgt_->back().resize(current_pos_, stream_view); + if (edge_id_) edge_id_->back().resize(current_pos_, stream_view); + if (edge_type_) edge_type_->back().resize(current_pos_, stream_view); } bool use_weight() const { return wgt_.has_value(); } @@ -230,16 +228,18 @@ class per_device_edgelist_t { void consolidate_and_shuffle(cugraph::mtmg::handle_t const& handle, bool store_transposed) { if (src_.size() > 1) { + auto stream = handle.raft_handle().get_stream(); + size_t total_size = std::transform_reduce( src_.begin(), src_.end(), size_t{0}, std::plus(), [](auto& d_vector) { return d_vector.size(); }); - resize_and_copy_buffers(handle.get_stream(), src_, total_size); - resize_and_copy_buffers(handle.get_stream(), dst_, total_size); - if (wgt_) resize_and_copy_buffers(handle.get_stream(), *wgt_, total_size); - if (edge_id_) resize_and_copy_buffers(handle.get_stream(), *edge_id_, total_size); - if (edge_type_) resize_and_copy_buffers(handle.get_stream(), *edge_type_, total_size); + resize_and_copy_buffers(src_, total_size, stream); + resize_and_copy_buffers(dst_, total_size, stream); + if (wgt_) resize_and_copy_buffers(*wgt_, total_size, stream); + if (edge_id_) resize_and_copy_buffers(*edge_id_, total_size, stream); + if (edge_type_) resize_and_copy_buffers(*edge_type_, total_size, stream); } auto tmp_wgt = wgt_ ? std::make_optional(std::move((*wgt_)[0])) : std::nullopt; @@ -267,9 +267,9 @@ class per_device_edgelist_t { private: template - void resize_and_copy_buffers(rmm::cuda_stream_view stream, - std::vector>& buffer, - size_t total_size) + void resize_and_copy_buffers(std::vector>& buffer, + size_t total_size, + rmm::cuda_stream_view stream) { size_t pos = buffer[0].size(); buffer[0].resize(total_size, stream); @@ -286,16 +286,16 @@ class per_device_edgelist_t { buffer = std::move(new_buffer); } - void create_new_buffers(cugraph::mtmg::handle_t const& handle) + void create_new_buffers(rmm::cuda_stream_view stream_view) { - src_.emplace_back(device_buffer_size_, handle.get_stream()); - dst_.emplace_back(device_buffer_size_, handle.get_stream()); + src_.emplace_back(device_buffer_size_, stream_view); + dst_.emplace_back(device_buffer_size_, stream_view); - if (wgt_) { wgt_->emplace_back(device_buffer_size_, handle.get_stream()); } + if (wgt_) { wgt_->emplace_back(device_buffer_size_, stream_view); } - if (edge_id_) { edge_id_->emplace_back(device_buffer_size_, handle.get_stream()); } + if (edge_id_) { edge_id_->emplace_back(device_buffer_size_, stream_view); } - if (edge_type_) { edge_type_->emplace_back(device_buffer_size_, handle.get_stream()); } + if (edge_type_) { edge_type_->emplace_back(device_buffer_size_, stream_view); } current_pos_ = 0; } diff --git a/cpp/include/cugraph/mtmg/edge_property.hpp b/cpp/include/cugraph/mtmg/edge_property.hpp index afa72492b9a..0b27ca85e46 100644 --- a/cpp/include/cugraph/mtmg/edge_property.hpp +++ b/cpp/include/cugraph/mtmg/edge_property.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -18,7 +18,6 @@ #include #include -#include namespace cugraph { namespace mtmg { diff --git a/cpp/include/cugraph/mtmg/edge_property_view.hpp b/cpp/include/cugraph/mtmg/edge_property_view.hpp index c84a6458e1d..6416ea382ef 100644 --- a/cpp/include/cugraph/mtmg/edge_property_view.hpp +++ b/cpp/include/cugraph/mtmg/edge_property_view.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -17,7 +17,6 @@ #pragma once #include -#include namespace cugraph { namespace mtmg { diff --git a/cpp/include/cugraph/mtmg/edgelist.hpp b/cpp/include/cugraph/mtmg/edgelist.hpp index 90c53dfbb64..d5d2bd2bca7 100644 --- a/cpp/include/cugraph/mtmg/edgelist.hpp +++ b/cpp/include/cugraph/mtmg/edgelist.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -39,7 +39,7 @@ class edgelist_t : public detail::device_shared_wrapper_t< bool use_edge_type) { detail::per_device_edgelist_t tmp( - handle, device_buffer_size, use_weight, use_edge_id, use_edge_type); + device_buffer_size, use_weight, use_edge_id, use_edge_type, handle.get_stream()); detail::device_shared_wrapper_t< detail::per_device_edgelist_t>::set(handle, @@ -49,7 +49,11 @@ class edgelist_t : public detail::device_shared_wrapper_t< /** * @brief Stop inserting edges into this edgelist so we can use the edges */ - void finalize_buffer(handle_t const& handle) { this->get(handle).finalize_buffer(handle); } + void finalize_buffer(handle_t const& handle) + { + handle.sync_stream_pool(); + this->get(handle).finalize_buffer(handle.get_stream()); + } /** * @brief Consolidate for the edgelist edges into a single edgelist and then diff --git a/cpp/include/cugraph/mtmg/handle.hpp b/cpp/include/cugraph/mtmg/handle.hpp index 0b02091a3cc..26c283f6acf 100644 --- a/cpp/include/cugraph/mtmg/handle.hpp +++ b/cpp/include/cugraph/mtmg/handle.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -79,6 +79,11 @@ class handle_t { */ void sync_stream() const { sync_stream(get_stream()); } + /** + * @brief Sync all streams in the stream pool + */ + void sync_stream_pool() const { raft::resource::sync_stream_pool(raft_handle_); } + /** * @brief get thrust policy for the stream * diff --git a/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp b/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp index b672db48719..73d69fdd5a7 100644 --- a/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp +++ b/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -18,7 +18,6 @@ #include #include -#include namespace cugraph { namespace mtmg { @@ -70,21 +69,21 @@ class per_thread_edgelist_t { /** * @brief Append an edge to the edge list * - * @param handle The resource handle - * @param src Source vertex id - * @param dst Destination vertex id - * @param wgt Edge weight - * @param edge_id Edge id - * @param edge_type Edge type + * @param src Source vertex id + * @param dst Destination vertex id + * @param wgt Edge weight + * @param edge_id Edge id + * @param edge_type Edge type + * @param stream_view The cuda stream */ - void append(handle_t const& handle, - vertex_t src, + void append(vertex_t src, vertex_t dst, std::optional wgt, std::optional edge_id, - std::optional edge_type) + std::optional edge_type, + rmm::cuda_stream_view stream_view) { - if (current_pos_ == src_.size()) { flush(handle); } + if (current_pos_ == src_.size()) { flush(stream_view); } src_[current_pos_] = src; dst_[current_pos_] = dst; @@ -98,19 +97,19 @@ class per_thread_edgelist_t { /** * @brief Append a list of edges to the edge list * - * @param handle The resource handle * @param src Source vertex id * @param dst Destination vertex id * @param wgt Edge weight * @param edge_id Edge id * @param edge_type Edge type + * @param stream_view The cuda stream */ - void append(handle_t const& handle, - raft::host_span src, + void append(raft::host_span src, raft::host_span dst, std::optional> wgt, std::optional> edge_id, - std::optional> edge_type) + std::optional> edge_type, + rmm::cuda_stream_view stream_view) { size_t count = src.size(); size_t pos = 0; @@ -131,7 +130,7 @@ class per_thread_edgelist_t { edge_type.begin() + pos + copy_count, edge_type_->begin() + current_pos_); - if (current_pos_ == src_.size()) { flush(handle); } + if (current_pos_ == src_.size()) { flush(stream_view); } count -= copy_count; pos += copy_count; @@ -141,12 +140,13 @@ class per_thread_edgelist_t { /** * @brief Flush thread data from host to GPU memory * - * @param handle The resource handle + * @param stream_view The cuda stream + * @param sync If true, synchronize the asynchronous copy of data; + * defaults to false. */ - void flush(handle_t const& handle) + void flush(rmm::cuda_stream_view stream_view, bool sync = false) { edgelist_.append( - handle, raft::host_span{src_.data(), current_pos_}, raft::host_span{dst_.data(), current_pos_}, wgt_ ? std::make_optional(raft::host_span{wgt_->data(), current_pos_}) @@ -155,9 +155,12 @@ class per_thread_edgelist_t { : std::nullopt, edge_type_ ? std::make_optional(raft::host_span{edge_type_->data(), current_pos_}) - : std::nullopt); + : std::nullopt, + stream_view); current_pos_ = 0; + + if (sync) stream_view.synchronize(); } private: diff --git a/cpp/tests/mtmg/multi_node_threaded_test.cu b/cpp/tests/mtmg/multi_node_threaded_test.cu index 1ad83761d51..24852562b86 100644 --- a/cpp/tests/mtmg/multi_node_threaded_test.cu +++ b/cpp/tests/mtmg/multi_node_threaded_test.cu @@ -175,15 +175,15 @@ class Tests_Multithreaded for (size_t j = starting_edge_offset; j < h_src_v.size(); j += stride) { per_thread_edgelist.append( - thread_handle, h_src_v[j], h_dst_v[j], h_weights_v ? std::make_optional((*h_weights_v)[j]) : std::nullopt, std::nullopt, - std::nullopt); + std::nullopt, + thread_handle.get_stream()); } - per_thread_edgelist.flush(thread_handle); + per_thread_edgelist.flush(thread_handle.get_stream()); }); } diff --git a/cpp/tests/mtmg/threaded_test.cu b/cpp/tests/mtmg/threaded_test.cu index f55a102ea67..df5a9e079df 100644 --- a/cpp/tests/mtmg/threaded_test.cu +++ b/cpp/tests/mtmg/threaded_test.cu @@ -191,15 +191,15 @@ class Tests_Multithreaded for (size_t j = i; j < h_src_v.size(); j += num_threads) { per_thread_edgelist.append( - thread_handle, h_src_v[j], h_dst_v[j], h_weights_v ? std::make_optional((*h_weights_v)[j]) : std::nullopt, std::nullopt, - std::nullopt); + std::nullopt, + thread_handle.get_stream()); } - per_thread_edgelist.flush(thread_handle); + per_thread_edgelist.flush(thread_handle.get_stream()); }); } diff --git a/cpp/tests/mtmg/threaded_test_jaccard.cu b/cpp/tests/mtmg/threaded_test_jaccard.cu index a64cc8ee1fa..0f531796cff 100644 --- a/cpp/tests/mtmg/threaded_test_jaccard.cu +++ b/cpp/tests/mtmg/threaded_test_jaccard.cu @@ -184,15 +184,15 @@ class Tests_Multithreaded for (size_t j = i; j < h_src_v.size(); j += num_threads) { per_thread_edgelist.append( - thread_handle, h_src_v[j], h_dst_v[j], h_weights_v ? std::make_optional((*h_weights_v)[j]) : std::nullopt, std::nullopt, - std::nullopt); + std::nullopt, + thread_handle.get_stream()); } - per_thread_edgelist.flush(thread_handle); + per_thread_edgelist.flush(thread_handle.get_stream()); }); } diff --git a/cpp/tests/mtmg/threaded_test_louvain.cu b/cpp/tests/mtmg/threaded_test_louvain.cu index c8faf33dae2..ab51d701b57 100644 --- a/cpp/tests/mtmg/threaded_test_louvain.cu +++ b/cpp/tests/mtmg/threaded_test_louvain.cu @@ -191,15 +191,15 @@ class Tests_Multithreaded for (size_t j = i; j < h_src_v.size(); j += num_threads) { per_thread_edgelist.append( - thread_handle, h_src_v[j], h_dst_v[j], h_weights_v ? std::make_optional((*h_weights_v)[j]) : std::nullopt, std::nullopt, - std::nullopt); + std::nullopt, + thread_handle.get_stream()); } - per_thread_edgelist.flush(thread_handle); + per_thread_edgelist.flush(thread_handle.get_stream()); }); } From e6c842fff4c88358f39d75cdcc01a689a6eea912 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang <45857425+seunghwak@users.noreply.github.com> Date: Thu, 23 May 2024 21:46:20 -0700 Subject: [PATCH 11/23] Biased sampling primitive (#4430) This PR restructures the current sampling primitive implementation and adds biased sampling support. Closes #4288 Authors: - Seunghwa Kang (https://github.com/seunghwak) - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/4430 --- .../cugraph/edge_partition_device_view.cuh | 70 +- .../cugraph/utilities/thrust_tuple_utils.hpp | 62 +- .../cugraph/vertex_partition_device_view.cuh | 3 +- .../detail/extract_transform_v_frontier_e.cuh | 57 +- cpp/src/prims/detail/partition_v_frontier.cuh | 91 + cpp/src/prims/detail/prim_functors.cuh | 7 +- .../sample_and_compute_local_nbr_indices.cuh | 2441 +++++++++++++++++ .../prims/detail/transform_v_frontier_e.cuh | 627 +++++ ...r_v_random_select_transform_outgoing_e.cuh | 1235 +-------- .../transform_reduce_e_by_src_dst_key.cuh | 26 +- ...rm_reduce_v_frontier_outgoing_e_by_dst.cuh | 17 +- cpp/src/prims/update_v_frontier.cuh | 16 +- ...er_v_random_select_transform_outgoing_e.cu | 130 +- cpp/tests/utilities/debug_utilities_mg.cpp | 12 +- 14 files changed, 3547 insertions(+), 1247 deletions(-) create mode 100644 cpp/src/prims/detail/partition_v_frontier.cuh create mode 100644 cpp/src/prims/detail/sample_and_compute_local_nbr_indices.cuh create mode 100644 cpp/src/prims/detail/transform_v_frontier_e.cuh diff --git a/cpp/include/cugraph/edge_partition_device_view.cuh b/cpp/include/cugraph/edge_partition_device_view.cuh index fc19a8f68dd..583b0a37214 100644 --- a/cpp/include/cugraph/edge_partition_device_view.cuh +++ b/cpp/include/cugraph/edge_partition_device_view.cuh @@ -214,9 +214,9 @@ class edge_partition_device_view_t - size_t compute_number_of_edges(MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ size_t compute_number_of_edges(MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { return dcs_nzd_vertices_ ? thrust::transform_reduce( rmm::exec_policy(stream), @@ -250,7 +250,7 @@ class edge_partition_device_view_t()); } - rmm::device_uvector compute_local_degrees(rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees(rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(this->major_range_size(), stream); if (dcs_nzd_vertices_) { @@ -277,9 +277,9 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees(MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees(MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(thrust::distance(major_first, major_last), stream); if (dcs_nzd_vertices_) { @@ -306,10 +306,10 @@ class edge_partition_device_view_t - size_t compute_number_of_edges_with_mask(MaskIterator mask_first, - MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ size_t compute_number_of_edges_with_mask(MaskIterator mask_first, + MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { return dcs_nzd_vertices_ ? thrust::transform_reduce( rmm::exec_policy(stream), @@ -348,8 +348,8 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees_with_mask(MaskIterator mask_first, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees_with_mask( + MaskIterator mask_first, rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(this->major_range_size(), stream); if (dcs_nzd_vertices_) { @@ -384,10 +384,11 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees_with_mask(MaskIterator mask_first, - MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees_with_mask( + MaskIterator mask_first, + MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(thrust::distance(major_first, major_last), stream); if (dcs_nzd_vertices_) { @@ -553,9 +554,9 @@ class edge_partition_device_view_t - size_t compute_number_of_edges(MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ size_t compute_number_of_edges(MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { return thrust::transform_reduce( rmm::exec_policy(stream), @@ -573,7 +574,7 @@ class edge_partition_device_view_t()); } - rmm::device_uvector compute_local_degrees(rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees(rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(this->major_range_size(), stream); thrust::transform(rmm::exec_policy(stream), @@ -589,9 +590,9 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees(MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees(MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(thrust::distance(major_first, major_last), stream); thrust::transform(rmm::exec_policy(stream), @@ -607,10 +608,10 @@ class edge_partition_device_view_t - size_t compute_number_of_edges_with_mask(MaskIterator mask_first, - MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ size_t compute_number_of_edges_with_mask(MaskIterator mask_first, + MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { return thrust::transform_reduce( rmm::exec_policy(stream), @@ -632,8 +633,8 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees_with_mask(MaskIterator mask_first, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees_with_mask( + MaskIterator mask_first, rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(this->major_range_size(), stream); thrust::transform( @@ -651,10 +652,11 @@ class edge_partition_device_view_t - rmm::device_uvector compute_local_degrees_with_mask(MaskIterator mask_first, - MajorIterator major_first, - MajorIterator major_last, - rmm::cuda_stream_view stream) const + __host__ rmm::device_uvector compute_local_degrees_with_mask( + MaskIterator mask_first, + MajorIterator major_first, + MajorIterator major_last, + rmm::cuda_stream_view stream) const { rmm::device_uvector local_degrees(thrust::distance(major_first, major_last), stream); thrust::transform( diff --git a/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp b/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp index d98754f51d1..304a5b94bd6 100644 --- a/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp +++ b/cpp/include/cugraph/utilities/thrust_tuple_utils.hpp @@ -1,5 +1,5 @@ /* - * 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. @@ -17,6 +17,7 @@ #include +#include #include #include @@ -30,7 +31,7 @@ template struct is_thrust_tuple_of_arithemetic_impl { constexpr bool evaluate() const { - if (!std::is_arithmetic::type>::value) { + if (!std::is_arithmetic_v::type>) { return false; } else { return is_thrust_tuple_of_arithemetic_impl().evaluate(); @@ -123,19 +124,19 @@ struct is_arithmetic_vector : std::false_type {}; template