Ktruss implementation

[jnke2016]

Implements SG and MG ktruss using graph primitives and drop cuHornet.

Closes #3447
Closes #3448
Closes #3449
Closes #3450
Closes #3451
Closes #3452
Closes #3453

[seunghwak]

Review part 1/2.

[seunghwak]

ktruss -> K-truss

[seunghwak]

auto [intersection_offsets, intersection_indices] = detail::nbr_intersection(...);

And you can cut your code size a bit. And better define r_nbr_intersection_property_values0 or 1 right before using them (minimize the scope of a variable).

[seunghwak]

We explicitly define them when we are taking returned variables with std::tie and use std::ignore as we can't use std::ignore with auto [...]. We are not using std::ignore here, so no need to explicitly define variables.

[jnke2016]

I am only using intersection_offsets and intersection_indices so I will remove r_nbr_intersection_property_values0 and r_nbr_intersection_property_values1

[seunghwak]

Please do not pre-allocate memory too much in advance. This unnecessarily increase peak memory usage.

[jnke2016]

in order to populate this vertex_pair_buffer with thrust::tabulate am I not required to know its size upfront?

[jnke2016]

auto vertex_pair_buffer_tmp = allocate_dataframe_buffer<thrust::tuple<vertex_t, vertex_t>>(
          intersection_indices.size() * 3, handle.get_stream());

The above pre-allocated memory for (p, q), (p, r) and (q, r) but I can do one group at a time and resize to reduce peak memory usage.

[seunghwak]

See the code from https://github.com/rapidsai/cugraph/blob/branch-24.02/cpp/src/sampling/sampling_post_processing_impl.cuh

    auto num_unique_labels = thrust::count_if(
      handle.get_thrust_policy(),
      thrust::make_counting_iterator(size_t{0}),
      thrust::make_counting_iterator((*renumber_map_label_indices).size()),
      detail::is_first_in_run_t<label_index_t const*>{(*renumber_map_label_indices).data()});
    rmm::device_uvector<label_index_t> unique_label_indices(num_unique_labels, handle.get_stream());
    rmm::device_uvector<vertex_t> vertex_counts(num_unique_labels, handle.get_stream());
    thrust::reduce_by_key(handle.get_thrust_policy(),
                          (*renumber_map_label_indices).begin(),
                          (*renumber_map_label_indices).end(),
                          thrust::make_constant_iterator(size_t{1}),
                          unique_label_indices.begin(),
                          vertex_counts.begin());

You can use the same logic to perform this with less memory.

[seunghwak]

You need to shuffle the resulting vertex pairs before reducing to be multi-GPU ready.

[seunghwak]

Now you can run thrust::partition here (Step 7 from #3446 (comment)) to find the edges to unroll.

[seunghwak]

@jnke2016

I thought more about implementing #3446 (comment); especially the steps 6, 7, and 8.

After step 6, sort the (edge source, edge destination, triangle count) triplets using edge destination as the primary key and edge source as the secondary key; this is to implement step 8-2.

7. Run thrust::stable_partition to place the edges with triangle counts smaller than K-2 at the end. These edges should be "unrolled" and removed (unrolled means undoing their contributions to triangle counts, see the Pierce & Sanders for additional details)<= run thrust::stable_partition instead of thrust::partition, so each partition is still sorted.
8. While (# edges to be remove > 0)
   8-1. Find (p, q, intersection of p&q's neighbor lists) for each edge to be removed. For each vertex r in the intersection, enumerate (p,r,-1) and (q,r,-1)<=call detail::nbr_intersection again but only for the edges to be removed. This covers the case the edge to be removed is a p->q edge. Note that we should eventually run step 6 in chunks as the aggregate size of the intersection indices can far exceed the number of edges. So, we cannot store all the intersection indices and we need to re-compute. This is a memory footprint vs compute trade-off.
   8-2. For each edge to be removed (say (q,r)), find all the incoming edges of r, then query the existence of (p,q) assuming p is an incoming neighbor of q. Accumulate (p,q,-1) and (p,r,-1) if (p,q) exists. <= This covers the case the edge to be removed is a q->r edge. We can run thrust::lower_bound & thrust::upper_bound on (s, d, triangle count) triplets to find all the incoming edges of q. Note that we need to run searches on both partitions of the (edge source, edge destination, triangle count) triplet array. In multi-GPU, we need to inspect all the GPUs that can possibly have r as an edge destination. This may sound a bit complicated unless you fully understand the partitioning and you may just add if constexpr (multi_gpu) { CUGRAPH_FAIL("unimplemented."); } for the initial SG implementation. We can come back to this once we validated the SG implementation.
   8-3. Similarly for each edge to be removed (say (p,r)), find all the incoming edges of r, then query the existence of (p,q) assuming q is an incoming neighbor of r. Accumulate (p,q,-1) and (q,r,-1). <= This covers the case the edge to be removed is a p->r edge. I assume there will be significant overlap in code for 8-2 & 8-3.
   8-4. Shuffle similar to 6-3 in multi-GPU.
   8-5. Update triangle counts based on triplets (similar to 6-4 & 6-5).
   8-6. Mask out the edges to be deleted. Also, resize the (edge source, edge destination, triangle count) triplet array so the deleted edges will no longer be considered.
   8-7. Identify the new set of edges to be deleted by running thrust::stable_partition again on the edge list based on newly update triangle counts.

Please read this and let me know when you are ready to sync again.

[ChuckHastings]

Will wait for @seunghwak to approve before merging, but LGTM.

[seunghwak]

I think this code is now nearly in mergable state. Let's quickly address few minor issues and let this PR pass the CI.

[seunghwak]

Don't forget to address this.

[seunghwak]

And you can provide thrust::make_transform_iterator(thrust::make_counting_iterator(size_t{0}), [query_vertices, num_edges = ..., sorted_veritces = ...]__device__(auto idx) { ... }); to thrust::exclusive_sum.

In the current code, you need to write dist values to prefix_sum; read dist values from prefix_sum; and compute prefix sum and write the results back to prefix_sum. (2 writes and 1 read). If you use transform_iterator, 1 write will be sufficient. This is a big cut in memory bandwidth and also reduces kernel launch overhead (2 kernel launches => 1 kernel launch).

[jnke2016]

Got it and thanks for the background information

[seunghwak]

(k+1)=>(k-1) and re-enable this. This shouldn't yield to incorrect results.

[jnke2016]

I tried and it still yield incorrect results. I tried even k-2

[jnke2016]

I added a fixme for this

[seunghwak]

edge_t num_invalid_edges{0};
      num_invalid_edges = ...

=>

auto num_invalid_edges = ...;

[seunghwak]

Don't forget to address this.

[seunghwak]

Should we really do this? Can we just compare cugraph COO output and reference COO output?

[jnke2016]

Ya but that means I will convert the reference ktruss to COO.

[jnke2016]

I think I will have to rely on csr comparison because my cugraph results are not sorted. I get the same concern I had last week regarding sorting edges.

[jnke2016]

Never mind. I extended the sorting utility functions in thrust_wrapper like you suggested.

[seunghwak]

LGTM and Thanks for all the hard work!!!

[seunghwak]

No need to address this in this PR, but just fyi,

You're passing h_offsets, h_indices, and h_values by value and based on the size of the vectors, this can be expensive. You can pass a const reference (const &) instead.

[jnke2016]

Right. I will resolve it in my next PR

[rlratzel]

/merge