From 24d02a5b1c112a55eb40119a86ae58edd47c3172 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Fri, 12 Jan 2024 12:06:46 -0500 Subject: [PATCH] Fix OOB error, BFS C API should validate that the source vertex is a valid vertex (#4077) * Added error check to be sure that the source vertex is a valid vertex in the graph. * Updated `nx_cugraph.Graph` class to create PLC graphs using `vertices_array` in order to include isolated vertices. This is now needed since the error check added in this PR prevents NetworkX tests from passing if isolated vertices are treated as invalid, so this change prevents that. * This resolves the problem that required the test workarounds done [here](https://github.com/rapidsai/cugraph/pull/4029#discussion_r1433332010) in [4029](https://github.com/rapidsai/cugraph/pull/4029), so those workarounds have been removed in this PR. Closes #4067 Authors: - Chuck Hastings (https://github.com/ChuckHastings) - Rick Ratzel (https://github.com/rlratzel) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Ray Douglass (https://github.com/raydouglass) - Erik Welch (https://github.com/eriknw) URL: https://github.com/rapidsai/cugraph/pull/4077 --- ci/test_python.sh | 5 -- .../cugraph/detail/utility_wrappers.hpp | 17 +++++- cpp/src/c_api/abstract_functor.hpp | 12 +++-- cpp/src/c_api/bfs.cpp | 17 +++++- cpp/src/detail/utility_wrappers.cu | 19 ++++++- python/nx-cugraph/nx_cugraph/classes/graph.py | 17 +++++- python/nx-cugraph/nx_cugraph/interface.py | 13 +---- python/pylibcugraph/pylibcugraph/graphs.pyx | 53 ++++++++++--------- 8 files changed, 105 insertions(+), 48 deletions(-) diff --git a/ci/test_python.sh b/ci/test_python.sh index d8288758f3c..7eb5a08edc8 100755 --- a/ci/test_python.sh +++ b/ci/test_python.sh @@ -111,11 +111,6 @@ popd rapids-logger "pytest networkx using nx-cugraph backend" pushd python/nx-cugraph ./run_nx_tests.sh -# Individually run tests that are skipped above b/c they may run out of memory -PYTEST_NO_SKIP=True ./run_nx_tests.sh --cov-append -k "TestDAG and test_antichains" -PYTEST_NO_SKIP=True ./run_nx_tests.sh --cov-append -k "TestMultiDiGraph_DAGLCA and test_all_pairs_lca_pairs_without_lca" -PYTEST_NO_SKIP=True ./run_nx_tests.sh --cov-append -k "TestDAGLCA and test_all_pairs_lca_pairs_without_lca" -PYTEST_NO_SKIP=True ./run_nx_tests.sh --cov-append -k "TestEfficiency and test_using_ego_graph" # run_nx_tests.sh outputs coverage data, so check that total coverage is >0.0% # in case nx-cugraph failed to load but fallback mode allowed the run to pass. _coverage=$(coverage report|grep "^TOTAL") diff --git a/cpp/include/cugraph/detail/utility_wrappers.hpp b/cpp/include/cugraph/detail/utility_wrappers.hpp index faa0fbb841b..61ac1bd2804 100644 --- a/cpp/include/cugraph/detail/utility_wrappers.hpp +++ b/cpp/include/cugraph/detail/utility_wrappers.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -174,5 +174,20 @@ bool is_equal(raft::handle_t const& handle, raft::device_span span1, raft::device_span span2); +/** + * @brief Count the number of times a value appears in a span + * + * @tparam data_t type of data in span + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param span The span of data to compare + * @param value The value to count + * @return The count of how many instances of that value occur + */ +template +size_t count_values(raft::handle_t const& handle, + raft::device_span span, + data_t value); + } // namespace detail } // namespace cugraph diff --git a/cpp/src/c_api/abstract_functor.hpp b/cpp/src/c_api/abstract_functor.hpp index 7bff5b37380..72b433aa9af 100644 --- a/cpp/src/c_api/abstract_functor.hpp +++ b/cpp/src/c_api/abstract_functor.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -32,8 +32,14 @@ struct abstract_functor { void unsupported() { - error_code_ = CUGRAPH_UNSUPPORTED_TYPE_COMBINATION; - error_->error_message_ = "Type Dispatcher executing unsupported combination of types"; + mark_error(CUGRAPH_UNSUPPORTED_TYPE_COMBINATION, + "Type Dispatcher executing unsupported combination of types"); + } + + void mark_error(cugraph_error_code_t error_code, std::string const& error_message) + { + error_code_ = error_code; + error_->error_message_ = error_message; } }; diff --git a/cpp/src/c_api/bfs.cpp b/cpp/src/c_api/bfs.cpp index ae7667375d2..32841b2dd3c 100644 --- a/cpp/src/c_api/bfs.cpp +++ b/cpp/src/c_api/bfs.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -113,6 +113,21 @@ struct bfs_functor : public abstract_functor { graph_view.local_vertex_partition_range_last(), do_expensive_check_); + size_t invalid_count = cugraph::detail::count_values( + handle_, + raft::device_span{sources.data(), sources.size()}, + cugraph::invalid_vertex_id::value); + + if constexpr (multi_gpu) { + invalid_count = cugraph::host_scalar_allreduce( + handle_.get_comms(), invalid_count, raft::comms::op_t::SUM, handle_.get_stream()); + } + + if (invalid_count != 0) { + mark_error(CUGRAPH_INVALID_INPUT, "Found invalid vertex in the input sources"); + return; + } + cugraph::bfs( handle_, graph_view, diff --git a/cpp/src/detail/utility_wrappers.cu b/cpp/src/detail/utility_wrappers.cu index 2d5bf6215b1..9100ecbd5e1 100644 --- a/cpp/src/detail/utility_wrappers.cu +++ b/cpp/src/detail/utility_wrappers.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-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. @@ -15,11 +15,13 @@ */ #include #include +#include #include #include +#include #include #include #include @@ -227,5 +229,20 @@ template bool is_equal(raft::handle_t const& handle, raft::device_span span1, raft::device_span span2); +template +size_t count_values(raft::handle_t const& handle, + raft::device_span span, + data_t value) +{ + return thrust::count(handle.get_thrust_policy(), span.begin(), span.end(), value); +} + +template size_t count_values(raft::handle_t const& handle, + raft::device_span span, + int32_t value); +template size_t count_values(raft::handle_t const& handle, + raft::device_span span, + int64_t value); + } // namespace detail } // namespace cugraph diff --git a/python/nx-cugraph/nx_cugraph/classes/graph.py b/python/nx-cugraph/nx_cugraph/classes/graph.py index cdd3f744f24..cb6b4e7ae42 100644 --- a/python/nx-cugraph/nx_cugraph/classes/graph.py +++ b/python/nx-cugraph/nx_cugraph/classes/graph.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 @@ -65,6 +65,7 @@ class Graph: key_to_id: dict[NodeKey, IndexValue] | None _id_to_key: list[NodeKey] | None _N: int + _node_ids: cp.ndarray[IndexValue] | None # holds plc.SGGraph.vertices_array data # Used by graph._get_plc_graph _plc_type_map: ClassVar[dict[np.dtype, np.dtype]] = { @@ -116,6 +117,7 @@ def from_coo( new_graph.key_to_id = None if key_to_id is None else dict(key_to_id) new_graph._id_to_key = None if id_to_key is None else list(id_to_key) new_graph._N = op.index(N) # Ensure N is integral + new_graph._node_ids = None new_graph.graph = new_graph.graph_attr_dict_factory() new_graph.graph.update(attr) size = new_graph.src_indices.size @@ -157,6 +159,16 @@ def from_coo( f"(got {new_graph.dst_indices.dtype.name})." ) new_graph.dst_indices = dst_indices + + # If the graph contains isolates, plc.SGGraph() must be passed a value + # for vertices_array that contains every vertex ID, since the + # src/dst_indices arrays will not contain IDs for isolates. Create this + # only if needed. Like src/dst_indices, the _node_ids array must be + # maintained for the lifetime of the plc.SGGraph + isolates = nxcg.algorithms.isolate._isolates(new_graph) + if len(isolates) > 0: + new_graph._node_ids = cp.arange(new_graph._N, dtype=index_dtype) + return new_graph @classmethod @@ -405,6 +417,7 @@ def clear(self) -> None: self.src_indices = cp.empty(0, self.src_indices.dtype) self.dst_indices = cp.empty(0, self.dst_indices.dtype) self._N = 0 + self._node_ids = None self.key_to_id = None self._id_to_key = None @@ -637,6 +650,7 @@ def _get_plc_graph( dst_indices = self.dst_indices if switch_indices: src_indices, dst_indices = dst_indices, src_indices + return plc.SGGraph( resource_handle=plc.ResourceHandle(), graph_properties=plc.GraphProperties( @@ -649,6 +663,7 @@ def _get_plc_graph( store_transposed=store_transposed, renumber=False, do_expensive_check=False, + vertices_array=self._node_ids, ) def _sort_edge_indices(self, primary="src"): diff --git a/python/nx-cugraph/nx_cugraph/interface.py b/python/nx-cugraph/nx_cugraph/interface.py index 3f6449f571a..34eb5969869 100644 --- a/python/nx-cugraph/nx_cugraph/interface.py +++ b/python/nx-cugraph/nx_cugraph/interface.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 @@ -242,20 +242,9 @@ def key(testpath): ) too_slow = "Too slow to run" - maybe_oom = "out of memory in CI" skip = { key("test_tree_isomorphism.py:test_positive"): too_slow, key("test_tree_isomorphism.py:test_negative"): too_slow, - key("test_efficiency.py:TestEfficiency.test_using_ego_graph"): maybe_oom, - key("test_dag.py:TestDAG.test_antichains"): maybe_oom, - key( - "test_lowest_common_ancestors.py:" - "TestDAGLCA.test_all_pairs_lca_pairs_without_lca" - ): maybe_oom, - key( - "test_lowest_common_ancestors.py:" - "TestMultiDiGraph_DAGLCA.test_all_pairs_lca_pairs_without_lca" - ): maybe_oom, # These repeatedly call `bfs_layers`, which converts the graph every call key( "test_vf2pp.py:TestGraphISOVF2pp.test_custom_graph2_different_labels" diff --git a/python/pylibcugraph/pylibcugraph/graphs.pyx b/python/pylibcugraph/pylibcugraph/graphs.pyx index b3065fa0684..76ad7690840 100644 --- a/python/pylibcugraph/pylibcugraph/graphs.pyx +++ b/python/pylibcugraph/pylibcugraph/graphs.pyx @@ -1,4 +1,4 @@ -# Copyright (c) 2022-2023, NVIDIA CORPORATION. +# Copyright (c) 2022-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 @@ -64,7 +64,7 @@ cdef class SGGraph(_GPUGraph): Object defining intended properties for the graph. src_or_offset_array : device array type - Device array containing either the vertex identifiers of the source of + Device array containing either the vertex identifiers of the source of each directed edge if represented in COO format or the offset if CSR format. In the case of a COO, the order of the array corresponds to the ordering of the dst_or_index_array, where the ith item in @@ -77,9 +77,14 @@ cdef class SGGraph(_GPUGraph): CSR format. In the case of a COO, The order of the array corresponds to the ordering of the src_offset_array, where the ith item in src_offset_array and the ith item in dst_index_array define the ith edge of the graph. - + vertices_array : device array type - Device array containing the isolated vertices of the graph. + Device array containing all vertices of the graph. This array is + optional, but must be used if the graph contains isolated vertices + which cannot be represented in the src_or_offset_array and + dst_index_array arrays. If specified, this array must contain every + vertex identifier, including vertex identifiers that are already + included in the src_or_offset_array and dst_index_array arrays. weight_array : device array type Device array containing the weight values of each directed edge. The @@ -99,25 +104,25 @@ cdef class SGGraph(_GPUGraph): do_expensive_check : bool If True, performs more extensive tests on the inputs to ensure validitity, at the expense of increased run time. - + edge_id_array : device array type Device array containing the edge ids of each directed edge. Must match the ordering of the src/dst arrays. Optional (may be null). If provided, edge_type_array must also be provided. - + edge_type_array : device array type Device array containing the edge types of each directed edge. Must match the ordering of the src/dst/edge_id arrays. Optional (may be null). If provided, edge_id_array must be provided. - + input_array_format: str, optional (default='COO') Input representation used to construct a graph COO: arrays represent src_array and dst_array CSR: arrays represent offset_array and index_array - + drop_self_loops : bool, optional (default='False') If true, drop any self loops that exist in the provided edge list. - + drop_multi_edges: bool, optional (default='False') If true, drop any multi edges that exist in the provided edge list @@ -178,7 +183,7 @@ cdef class SGGraph(_GPUGraph): cdef cugraph_type_erased_device_array_view_t* srcs_or_offsets_view_ptr = \ create_cugraph_type_erased_device_array_view_from_py_obj( src_or_offset_array - ) + ) cdef cugraph_type_erased_device_array_view_t* dsts_or_indices_view_ptr = \ create_cugraph_type_erased_device_array_view_from_py_obj( dst_or_index_array @@ -192,7 +197,7 @@ cdef class SGGraph(_GPUGraph): ) self.edge_id_view_ptr = create_cugraph_type_erased_device_array_view_from_py_obj( edge_id_array - ) + ) cdef cugraph_type_erased_device_array_view_t* edge_type_view_ptr = \ create_cugraph_type_erased_device_array_view_from_py_obj( edge_type_array @@ -237,7 +242,7 @@ cdef class SGGraph(_GPUGraph): assert_success(error_code, error_ptr, "cugraph_sg_graph_create_from_csr()") - + else: raise ValueError("invalid 'input_array_format'. Only " "'COO' and 'CSR' format are supported." @@ -282,7 +287,7 @@ cdef class MGGraph(_GPUGraph): each directed edge. The order of the array corresponds to the ordering of the src_array, where the ith item in src_array and the ith item in dst_array define the ith edge of the graph. - + vertices_array : device array type Device array containing the isolated vertices of the graph. @@ -295,12 +300,12 @@ cdef class MGGraph(_GPUGraph): store_transposed : bool Set to True if the graph should be transposed. This is required for some algorithms, such as pagerank. - + num_arrays : size_t Number of arrays. - + If provided, all list of device arrays should be of the same size. - + do_expensive_check : bool If True, performs more extensive tests on the inputs to ensure validitity, at the expense of increased run time. @@ -309,15 +314,15 @@ cdef class MGGraph(_GPUGraph): Device array containing the edge ids of each directed edge. Must match the ordering of the src/dst arrays. Optional (may be null). If provided, edge_type_array must also be provided. - + edge_type_array : device array type Device array containing the edge types of each directed edge. Must match the ordering of the src/dst/edge_id arrays. Optional (may be null). If provided, edge_id_array must be provided. - + drop_self_loops : bool, optional (default='False') If true, drop any self loops that exist in the provided edge list. - + drop_multi_edges: bool, optional (default='False') If true, drop any multi edges that exist in the provided edge list """ @@ -357,12 +362,12 @@ cdef class MGGraph(_GPUGraph): dst_array = [dst_array] if not any(dst_array): dst_array = dst_array * num_arrays - + if not isinstance(weight_array, list): weight_array = [weight_array] if not any(weight_array): weight_array = weight_array * num_arrays - + if not isinstance(edge_id_array, list): edge_id_array = [edge_id_array] if not any(edge_id_array): @@ -372,7 +377,7 @@ cdef class MGGraph(_GPUGraph): edge_type_array = [edge_type_array] if not any(edge_type_array): edge_type_array = edge_type_array * num_arrays - + if not isinstance(vertices_array, list): vertices_array = [vertices_array] if not any(vertices_array): @@ -394,7 +399,7 @@ cdef class MGGraph(_GPUGraph): if edge_id_array is not None and len(edge_id_array[i]) != len(src_array[i]): raise ValueError('Edge id array must be same length as edgelist') - + assert_CAI_type(edge_type_array[i], "edge_type_array", True) if edge_type_array[i] is not None and len(edge_type_array[i]) != len(src_array[i]): raise ValueError('Edge type array must be same length as edgelist') @@ -421,7 +426,7 @@ cdef class MGGraph(_GPUGraph): malloc( num_arrays * sizeof(cugraph_type_erased_device_array_view_t*)) vertices_view_ptr_ptr[i] = \ - create_cugraph_type_erased_device_array_view_from_py_obj(vertices_array[i]) + create_cugraph_type_erased_device_array_view_from_py_obj(vertices_array[i]) if weight_array[i] is not None: if i == 0: