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: