Add per-edge triplet neighbor-intersection transform-reduce primitive - #5551
Add per-edge triplet neighbor-intersection transform-reduce primitive#5551jnke2016 wants to merge 38 commits into
Conversation
…C++ tests for the primitives
…ings, and per-partition setup)
…nts_by_e invoke the operator once per common neighbor r
…f_e_endpoints_by_e
…(operator overload)
…intersection_of_e_endpoints_by_e
…oc/FIXME and readability cleanups
| VertexPairIterator vertex_pair_first, | ||
| VertexPairIterator vertex_pair_last, |
There was a problem hiding this comment.
If we are performing reduction "by_e", this means vertex pairs should be edges. Then, we should better use an edge bucket than general vertex pairs.
See
https://github.com/rapidsai/cugraph/blob/main/cpp/include/cugraph/prims/transform_gather_e.cuh#L151
…ction counting pass
| edge_t& nbr_offset, | ||
| edge_t& nbr_degree) |
There was a problem hiding this comment.
Better return cuda::std::tuple<vertex_t const*, edge_t, edge_t> rather than taking nbr_indices, nbr_offset and nbr_degree as L-value references. Returning a value is more functional and generally preferable unless there is a strong reason to take L-value references.
There was a problem hiding this comment.
And you can actually return just nbr_indices + nbr_offset and nbr_degree.
| // whole function k times and emit k times over. | ||
| if (s > edge_t{0} && short_indices[p - 1] == w) { return; } | ||
|
|
||
| // Binary search for w in the longer list. Could equivalently use the sequential thrust version: |
There was a problem hiding this comment.
Why aren't you using thrust::lower_bound(thrust::seq, ...)?
There was a problem hiding this comment.
And first, you can skip all the complexity if the graph is not a multi-graph (graph_view.is_multigraph()), and you may be able to update the caller site to call this only once per unique v in the shorter neighbor list (assuming that there is very few edges with a huge multiplicity value).
You can run thrust::lower_bound first and just scan from there.
| v1_off, | ||
| v1_deg); | ||
|
|
||
| // Scan the shorter neighbor list, binary-search the longer one. |
There was a problem hiding this comment.
You may use
__ballot_sync (set to 1 if its the first in the run) followed by __popc then __fns to find n'th unique neighbor in the short neighbor list.
There can be a trade-off. If multi-edges are infrequent, you may just let few threads in a warp. If multi-edges are common, we may allow few edges to process more than one neighbor in the short neighbor list.
Or if we assume that invoking intersection_op is expensive, we can separate finding common neighbor indices in the long and short arrays (each thread may process more than one neighbor in the short neighbor list) and calling intersection_op.
We use all threads in a warp to update the array indices. Then, for the identified indices, call intersection_op in the second round. This may lead to faster & simpler code.
2 participants
No description provided.