Repository navigation
Conversation
seunghwak
left a comment
seunghwak
left a comment
There was a problem hiding this comment.
Thank you very much for the contribution!
I haven't reviewed "components/strongly_connected_components_ecl_impl.cuh" yet (will review soon) but want to discuss how should we handle the LICENSE file. Did you include the LICENSE file because you used any copyrighted code or you included the file because we used the published ECL-SCC algorithm?
| * @brief Finds (strongly-connected-)component IDs of each vertices in the input graph. | ||
| * | ||
| * Uses ECL vertex queues with global Manhattan scheduling for single-GPU graphs and | ||
| * forward-backward traversal for multi-GPU graphs. An attached edge mask is honored. |
There was a problem hiding this comment.
We generally don't expose implementation details in the public API. Better delete this.
| * The input graph may be asymmetric. Component IDs can be arbitrary integers (they can be | ||
| * non-consecutive and are not ordered by component size or any other criterion). | ||
| * | ||
| * Single-GPU graphs use ECL signature propagation; multi-GPU graphs use FWBW. |
There was a problem hiding this comment.
Same here. No need to say we are relying on ECL in a way we didn't mention we are using the FW-BW approach in the past.
There was a problem hiding this comment.
Why did you include this LICENSE file? Are we copying any code from the existing ECL-SCC implementation or you included this because we used the algorithm (but no code)?
And we have thirdpaty/LICENSES directory and it might be a more appropriate place to put this file.
There was a problem hiding this comment.
The current code is based on but very heavily modified from the original ECL-SCC implementation, though the basic algorithm logic is still consistent. I included the license to be safe. I can move it to the LICENSES subdirectory (apologies for missing that), or I can remove it entirely, if you do not think it would be necessary.
There was a problem hiding this comment.
Thanks for the update, and I will double check but I agree that it is safe to include the LICENSES. And better move this to the LICENSES subdirectory.
| */ | ||
| #pragma once | ||
|
|
||
| #include "components/strongly_connected_components_ecl_impl.cuh" |
There was a problem hiding this comment.
We can include "components/strongly_connected_components_ecl_impl.cuh" in strongly_connected_components_sg_v32_e32.cu and strongly_connected_components_sg_v64_e64.cu. Once we have both SG & MG implementations in ECL-SCC, later, we can replace the strongly_connected_components_impl.cuh with "components/strongly_connected_components_ecl_impl.cuh"
| include(CPack) | ||
|
|
||
| install(FILES src/components/ECL-SCC-LICENSE | ||
| DESTINATION share/licenses/cugraph) |
There was a problem hiding this comment.
Do we need this? We need to double check how should we handle the LICENSE file issue.
| - id: verify-copyright | ||
| name: verify-copyright-ecl | ||
| args: [--fix, --spdx, "--spdx-license-identifier=Apache-2.0 AND BSD-3-Clause"] | ||
| files: ^cpp/src/components/strongly_connected_components_ecl_impl[.]cuh$ |
There was a problem hiding this comment.
Not sure we really need to manually edit this file. We rarely modify this file on any algorithm's internal implementation changes.
seunghwak
added
improvement
seunghwak
left a comment
seunghwak
left a comment
There was a problem hiding this comment.
For cuGraph integration, do you want to go through multiple review rounds to make this code to fit into the cuGraph codebase coding styles/conventions, or I can take over from there, and ask a coding agent to convert this code with more specific instructions.
For example, we have code/function to create a reverse graph, transform_reduce_if_v_frontier_outgoing_e to propagate label, collect_values_for_int_vertices to implement the shortcut mechanism in the multi-GPU path, and so on.
We don't recommend application code to directly work withe CSR array or edge mask, and we want the MG path to be largely SG code + alpha rather than having two separate code paths.
And cuGraph coding style (mostly coming from its RAPIDS root) is to avoid using custom kernels unless the performance benefit is significant. A lot of this can be automated using a coding agent if the instruction is specific enough.
This code will work as a great starting point in either ways, and thank you very much for the contribution!
| template <typename GraphViewType> | ||
| std::unique_ptr< | ||
| graph_t<typename GraphViewType::vertex_type, typename GraphViewType::edge_type, false, false>> | ||
| build_ecl_reverse_graph( |
There was a problem hiding this comment.
You can refer to the old SCC code on creating a reverse graph (this code will work for both SG & MG)
| // Initialize our seed heuristic -> vertex with the largest in*out degree gets max label | ||
| auto masked_out_degrees = graph_view.has_edge_mask() ? graph_view.compute_out_degrees(handle) | ||
| : rmm::device_uvector<edge_t>(0, stream); | ||
| { | ||
| auto candidates = thrust::make_transform_iterator( | ||
| thrust::make_counting_iterator(vertex_t{0}), | ||
| degree_product<vertex_t, edge_t>{ | ||
| edges.offsets(), | ||
| reverse_edges.offsets(), | ||
| masked_out_degrees.size() ? masked_out_degrees.data() : nullptr}); | ||
| size_t scratch_bytes{}; | ||
| auto initial = degree_seed<vertex_t>{0, 0, std::numeric_limits<vertex_t>::max()}; | ||
| RAFT_CUDA_TRY(cub::DeviceReduce::Reduce(nullptr, | ||
| scratch_bytes, | ||
| candidates, | ||
| work.seed.data(), | ||
| n, | ||
| maximum_degree_seed<vertex_t>{}, | ||
| initial, | ||
| stream.get())); | ||
| work.scratch.resize(scratch_bytes, stream); | ||
| RAFT_CUDA_TRY(cub::DeviceReduce::Reduce(work.scratch.data(), | ||
| scratch_bytes, | ||
| candidates, | ||
| work.seed.data(), | ||
| n, | ||
| maximum_degree_seed<vertex_t>{}, | ||
| initial, | ||
| stream.get())); | ||
| } |
There was a problem hiding this comment.
cuGraph codebase generally prefers using thrust/cugraph primitives unless the performance difference is dramatic.
This picks the vertex with the highest in_degree * out_degree (and I assume this shouldn't be exact).
Can't we compute (double)in_degree * (double)out_degree instead of using 128bit integer (no more guarantee that we will pick the exact one but I assume this may not matter much in practice) and use cugraph::reduce_v instead? This code will work seamlessly in SG & MG.
|
|
||
| // Initialize our seed heuristic -> vertex with the largest in*out degree gets max label | ||
| auto masked_out_degrees = graph_view.has_edge_mask() ? graph_view.compute_out_degrees(handle) | ||
| : rmm::device_uvector<edge_t>(0, stream); |
There was a problem hiding this comment.
Can we just do out_degrees = graph_view.compute_out_degrees(hanlde)? I assume this won't impact performance significantly over computing from offsets on the go? This code will work in both SG & MG and simplify the codebase. And a single code path regardless of whether we are using edge mask or not.
| trim_singletons<<<blocks_for(n), block_size, 0, stream.get()>>>( | ||
| edges.offsets(), | ||
| reverse_edges.offsets(), | ||
| masked_out_degrees.size() ? masked_out_degrees.data() : nullptr, | ||
| n, | ||
| labels.data()); | ||
| RAFT_CUDA_TRY(cudaPeekAtLastError()); |
There was a problem hiding this comment.
We generally prefer thrust over custom kernels unless the performance difference is significant.
This sounds like thrust::transform_if. A predicate will check whether in_degree or out_degree is 0. If the predicate is satisfied, writes vertex ID.
| initialize_signatures<<<blocks_for(active_count), block_size, 0, stream.get()>>>( | ||
| active[active_buffer].data(), | ||
| active_count, | ||
| incoming.data(), | ||
| outgoing.data(), | ||
| n, | ||
| work.seed.data(), | ||
| first_round); |
There was a problem hiding this comment.
Similar here. I think thrust::scatter will be sufficient.
Something like
auto color_first = thrust::make_transform_iterator(
active_first,
[...]__device__(vertex_t v) {
auto olor = (first_round && v == seed->vertex) ? n : v;
return cuda::std::make_tuple(color, color);
}));
thrust::scatter(handle.get_thrust_policy(),
color_first, color_first + active_count,
active_first,
thrust::make_zip_iterator(incoming.begin(), outgoing.begin());
2 participants
Single-GPU strongly_connected_components now dispatches to a new ECL-SCC implementation through the existing C++ and C APIs. Multi-GPU graphs continue to use the FWBW implementation. This avoids the serialization and performance penalty that occurs using FWBW on graphs with a significant number of nontrivial strong components. A multi-GPU implementation of ECL-SCC is targeted for a future PR.
The implementation uses cuGraph graph views, RMM allocations, current stream APIs, and packed edge masks. It supports 32-bit and 64-bit types and uses a global Manhattan collapse (256 threads, 1024-edge tiles) for internal load balancing, a degree-product initial seed to target large components on the first propagation, initial trimming, and pointer jumping. Reverse CSR construction is internal to the call. All timing stays in external benchmark drivers.
The original ECL-SCC BSD notice is retained and installed with the library.
Validation
Performance
One first SCC invocation per fresh process, with no explicit CUDA module warmup. Input construction/upload and output copy are outside the timer; internal reverse CSR construction and allocations are inside. Control uses pristine-main SG FWBW with matching compiler flags and native dependencies. These are single samples, not statistical estimates. All eight cases produced matching canonical partitions relative to the existing FWBW algorithm. Correctness of a large suite of inputs was successfully tested against a reference serial algorithm.
Hardware/software: RTX PRO 5000 Blackwell Generation Laptop GPU, CUDA 12.9, GCC 14.4, Release sm_120a, cuGraph main 5d0bc45 (26.12), OpenMPI 4.1.6, NCCL 2.32.3.1. Test graphs were retrieved from the Stanford Large Network Dataset Collection, the SuiteSparse Matrix Collection, the ECL-SCC github, and cugraph::generate_rmat_edgelist. The R-MAT graphs used scale 26, edge factor 16, and seed 0 (67,108,864 vertices and 1,073,741,824 edges) in both renumbering modes.