From 1635f874f26727a0b7bf3fb1d65511fec6ab8ee4 Mon Sep 17 00:00:00 2001 From: jnke2016 Date: Mon, 8 Jan 2024 10:38:46 -0800 Subject: [PATCH] update implementation --- cpp/src/community/ktruss_impl.cuh | 415 ++++++++++++++++++++++++++++-- 1 file changed, 397 insertions(+), 18 deletions(-) diff --git a/cpp/src/community/ktruss_impl.cuh b/cpp/src/community/ktruss_impl.cuh index d1c76e0a212..c6662c1f043 100644 --- a/cpp/src/community/ktruss_impl.cuh +++ b/cpp/src/community/ktruss_impl.cuh @@ -106,11 +106,18 @@ struct extract_p_q { __device__ thrust::tuple operator()(edge_t i) const { + printf("\n i = %d", i); auto itr = thrust::upper_bound(thrust::seq, intersection_offsets.begin()+1, intersection_offsets.end(), i); + auto idx = thrust::distance(intersection_offsets.begin()+1, itr); - thrust::tuple pair = *(vertex_pairs_begin + idx); + printf("\nthe idx = %d", idx); + + printf("\n major = %d and minor = %d", thrust::get<0>(*(vertex_pairs_begin + idx)), thrust::get<1>(*(vertex_pairs_begin + idx))); + printf("\nintersection = %d\n", intersection_offsets[i+1] - intersection_offsets[i]); + + thrust::tuple pair = *(vertex_pairs_begin + idx); - return pair; + return pair; } }; @@ -121,15 +128,20 @@ template struct extract_p_r { raft::device_span intersection_offsets{}; raft::device_span intersection_indices{}; + VertexPairIterator vertex_pairs_begin; + __device__ thrust::tuple operator()(edge_t i) const { auto itr = thrust::upper_bound(thrust::seq, intersection_offsets.begin()+1, intersection_offsets.end(), i); + auto idx = thrust::distance(intersection_offsets.begin()+1, itr); - auto pair = thrust::make_tuple(thrust::get<0>(*(vertex_pairs_begin + idx)), intersection_indices[i]); - return pair; + printf("\n ** (p, q) major = %d and minor = %d", thrust::get<0>(*(vertex_pairs_begin + idx)), intersection_indices[i]); + auto pair = thrust::make_tuple(thrust::get<0>(*(vertex_pairs_begin + idx)), intersection_indices[i]); + + return pair; } }; @@ -139,19 +151,86 @@ template struct extract_q_r { raft::device_span intersection_offsets{}; raft::device_span intersection_indices{}; + VertexPairIterator vertex_pairs_begin; + __device__ thrust::tuple operator()(edge_t i) const { + printf("\n i = %d", i); auto itr = thrust::upper_bound(thrust::seq, intersection_offsets.begin()+1, intersection_offsets.end(), i); + auto idx = thrust::distance(intersection_offsets.begin()+1, itr); - auto pair = thrust::make_tuple(thrust::get<1>(*(vertex_pairs_begin + idx)), intersection_indices[i]); - return pair; + printf("\n ** (p, q) major = %d and minor = %d", thrust::get<1>(*(vertex_pairs_begin + idx)), intersection_indices[i]); + + auto pair = thrust::make_tuple(thrust::get<1>(*(vertex_pairs_begin + idx)), intersection_indices[i]); + + return pair; } }; +template +struct triangle_counts_less_t { + raft::device_span num_triangles{}; + edge_t k; + +}; + +template +struct find_intersection { + size_t vertex_pair_buffer_size; + size_t num_invalid_edges; + raft::device_span intersection_offsets{}; + raft::device_span intersection_indices{}; + + VertexPairBuffer vertex_pair_buffer; + Iterator invalid_edge_first; + + __device__ size_t operator()(edge_t i) const + { + printf("\n ***** i = %d\n", i); + printf("\n src = %d, dst = %d\n", thrust::get<0>(*(vertex_pair_buffer+i)), thrust::get<1>(*(vertex_pair_buffer+i))); + printf("\nthe vertex pair buffer size = %d\n", vertex_pair_buffer_size); + //auto x = thrust::get<1>(*invalid_edge_first); + //printf("\nx = %d\n", x); + // Run binary search + //if (i < vertex_pair_buffer_size) { + if (i < num_invalid_edges) { + auto pair = thrust::get<0>(*(invalid_edge_first + i)); // FIXME only increment if it is less than 'invalid_edge_first.size()' + printf("\nsrc_pair = %d, dst_pair = %d\n", thrust::get<0>(pair), thrust::get<1>(pair)); + + auto itr = thrust::upper_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + vertex_pair_buffer_size, + //pair, + thrust::make_tuple(thrust::get<0>(pair), thrust::get<1>(pair)), + thrust::less>{}); + + printf("\n itr = %d\n", itr); + auto idx = thrust::distance(vertex_pair_buffer + 1, itr); + printf("\n idx = %d \n", idx); + + return 1; + } + else { + return 0; + } + + + return 0; + } +}; + + + + + + + + + template struct intersection_op_t { __device__ thrust::tuple operator()( @@ -386,6 +465,10 @@ void ktruss(raft::handle_t const& handle, auto vertex_pairs_begin = thrust::make_zip_iterator( edgelist_srcs.begin(), edgelist_dsts.begin()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + raft::print_device_vector("src ", edgelist_srcs.data(), edgelist_srcs.size(), std::cout); + raft::print_device_vector("dst ", edgelist_dsts.data(), edgelist_dsts.size(), std::cout); size_t num_vertex_pairs = edgelist_srcs.size(); auto out_degrees = cur_graph_view.compute_out_degrees(handle); @@ -404,10 +487,15 @@ void ktruss(raft::handle_t const& handle, vertex_pairs_begin + num_vertex_pairs, std::array{true, true}, do_expensive_check); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + raft::print_device_vector("intersection_offsets ", intersection_offsets.data(), intersection_offsets.size(), std::cout); + raft::print_device_vector("intersection_indices ", intersection_indices.data(), intersection_indices.size(), std::cout); auto vertex_pair_buffer = allocate_dataframe_buffer>( num_vertex_pairs, handle.get_stream()); + // stores all the pairs (p, q), (p, r) and (q, r) auto vertex_pair_buffer_ = allocate_dataframe_buffer>( intersection_indices.size() * 3, handle.get_stream()); @@ -452,6 +540,11 @@ void ktruss(raft::handle_t const& handle, vertex_pairs_begin }); + printf("\nbefore sorting\n"); + raft::print_device_vector("src", std::get<0>(vertex_pair_buffer_).data(), std::get<0>(vertex_pair_buffer_).size(), std::cout); + raft::print_device_vector("dst", std::get<1>(vertex_pair_buffer_).data(), std::get<1>(vertex_pair_buffer_).size(), std::cout); + + thrust::sort(handle.get_thrust_policy(), get_dataframe_buffer_begin(vertex_pair_buffer_), get_dataframe_buffer_end(vertex_pair_buffer_)); @@ -473,39 +566,325 @@ void ktruss(raft::handle_t const& handle, // Note: ensure 'edge_list' and 'cur_graph_view' have the same transpose flag + // one of the void cugraph::edge_bucket_tedge_list(handle); edge_list.insert(std::get<0>(vertex_pair_buffer).begin(), std::get<0>(vertex_pair_buffer).end(), + // num triangle here std::get<1>(vertex_pair_buffer).begin()); cugraph::edge_property_t edge_value_output(handle, cur_graph_view); - + thrust::sort(handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer), + get_dataframe_buffer_end(vertex_pair_buffer), + thrust::less>{}); + raft::print_device_vector("sorted - src", std::get<0>(vertex_pair_buffer).data(), std::get<0>(vertex_pair_buffer).size(), std::cout); + raft::print_device_vector("sorted - dst", std::get<1>(vertex_pair_buffer).data(), std::get<1>(vertex_pair_buffer).size(), std::cout); + raft::print_device_vector("new num triangles", num_triangles.data(), num_triangles.size(), std::cout); + handle.sync_stream(); cugraph::transform_e( handle, cur_graph_view, - edge_list, + edge_list, // assigning th e number of triangles as an edge property cugraph::edge_src_dummy_property_t{}.view(), cugraph::edge_dst_dummy_property_t{}.view(), cugraph::edge_dummy_property_t{}.view(), [num_triangles = num_triangles.data(), - vertex_pair_buffer = get_dataframe_buffer_begin(vertex_pair_buffer), - size=get_dataframe_buffer_end(vertex_pair_buffer) - get_dataframe_buffer_begin(vertex_pair_buffer)] - __device__(auto src, auto dst, thrust::nullopt_t, thrust::nullopt_t, thrust::nullopt_t) { + vertex_pair_buffer = get_dataframe_buffer_begin(vertex_pair_buffer), + size=get_dataframe_buffer_end(vertex_pair_buffer) - get_dataframe_buffer_begin(vertex_pair_buffer)] // FIXME: use 'size_dataframe_buffer(vertex_pair_buffer)' + // FIXME: Maybe: the last 'thrust::nullopt_t' for the edge prop containing the number of triangles. + // Take a look at the 'vertex_bucket_t' for context + __device__(auto src, auto dst, thrust::nullopt_t, thrust::nullopt_t, thrust::nullopt_t) { + /* + for ( int i =0 ; i < size; i++){ + //printf("\n i = %d major = %d and minor = %d", i, thrust::get<0>(*(vertex_pair_buffer+i)), thrust::get<1>(*(vertex_pair_buffer+i))); + printf("hello"); + } + */ + + + + + auto x = thrust::make_tuple(thrust::get<0>(*(vertex_pair_buffer)), thrust::get<1>(*(vertex_pair_buffer))); + printf("\nsrc = %d , dst = %d\n", src, dst); + //printf("\ny = %d\n", vertex_pair_buffer); + + auto it = thrust::lower_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + size, + thrust::make_tuple(src, dst)); + printf("\ntransform_e - itr = %d\n", it); + auto idx = thrust::distance(vertex_pair_buffer, it); + printf("\nthe distance = %d\n", idx); - auto it = thrust::upper_bound(thrust::seq, - vertex_pair_buffer, - vertex_pair_buffer + size, - thrust::make_tuple(src, dst), - thrust::equal_to>{}); + /* + auto it_ = thrust::lower_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + size, + //thrust::make_tuple(src, dst), + thrust::make_tuple(1, 3), + thrust::less>{}); + printf("\ntransform_e - itr____ = %d\n", it_); + auto idx_ = thrust::distance(vertex_pair_buffer, it_); + printf("\nthe distance = %d\n", idx_); + */ - auto idx = thrust::distance(vertex_pair_buffer, it); - return num_triangles[idx]; + //printf("\nitr = %d\n", *it); + return num_triangles[idx]; }, edge_value_output.mutable_view(), false); + + + /* + rmm::device_uvector intersection_size_(intersection_offsets.size() - 1, handle.get_stream()); // take the worst case where all edges are invalidated + // FIXME: might not be necessary to fill it with zeros + thrust::fill(handle.get_thrust_policy(), intersection_size_.begin(), intersection_size_.end(), size_t{0}); + printf("\n*******dummy******\n"); + thrust::tabulate(handle.get_thrust_policy(), + intersection_size_.begin(), + intersection_size_.end(), + [vertex_pair_buffer_size=size_dataframe_buffer(vertex_pair_buffer), + intersection_offsets = raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + intersection_indices = raft::device_span(intersection_indices.data(), intersection_indices.size()), + vertex_pair_buffer = get_dataframe_buffer_begin(vertex_pair_buffer) + ] + __device__(edge_t i){ + printf("\n ***** i = %d\n", i); + printf("\n src = %d, dst = %d\n", thrust::get<0>(*(vertex_pair_buffer+i)), thrust::get<1>(*(vertex_pair_buffer+i))); + printf("\nthe vertex pair buffer size = %d\n", vertex_pair_buffer_size); + auto num_invalid_edges = 6; + if (i < num_invalid_edges) { + //auto pair = thrust::get<0>(*(invalid_edge_first + i)); // FIXME only increment if it is less than 'invalid_edge_first.size()' + //printf("\nsrc_pair = %d, dst_pair = %d\n", thrust::get<0>(pair), thrust::get<1>(pair)); + + auto itr = thrust::lower_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + vertex_pair_buffer_size, + //pair, + //thrust::make_tuple(thrust::get<0>(pair), thrust::get<1>(pair)), + thrust::make_tuple(2, 3), + thrust::less>{}); + + printf("\n itr = %d\n", itr); + auto idx = thrust::distance(vertex_pair_buffer, itr); + printf("\n idx = %d \n", idx); + + return 1; + } + else { + return 0; + } + + }); + printf("\n*******Done dummy******\n"); + */ + + + + + + + + + + + + auto value_firsts = edge_value_output.view().value_firsts(); + auto edge_counts = edge_value_output.view().edge_counts(); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + raft::print_device_vector( + "num_triangles_:", value_firsts[0], edge_counts[0], std::cout); + + vertex_pairs_begin = thrust::make_zip_iterator(std::get<0>(vertex_pair_buffer).begin(), std::get<1>(vertex_pair_buffer).begin()); // FIXME: not used below but used up + printf("\nafter sorting\n"); + raft::print_device_vector("src", std::get<0>(vertex_pair_buffer).data(), std::get<0>(vertex_pair_buffer).size(), std::cout); + raft::print_device_vector("dst", std::get<1>(vertex_pair_buffer).data(), std::get<1>(vertex_pair_buffer).size(), std::cout); + + // FIXME: shuffle before sorting and reducing - Bring this part up + if constexpr (multi_gpu) { + thrust::copy(handle.get_thrust_policy(), + std::get<0>(vertex_pair_buffer).data(), + std::get<0>(vertex_pair_buffer).data() + std::get<0>(vertex_pair_buffer).size(), + edgelist_srcs.begin()); + + thrust::copy(handle.get_thrust_policy(), + std::get<1>(vertex_pair_buffer).data(), + std::get<1>(vertex_pair_buffer).data() + std::get<1>(vertex_pair_buffer).size(), + edgelist_dsts.begin()); + + edgelist_srcs.resize(std::get<0>(vertex_pair_buffer).size(), handle.get_stream()); + edgelist_dsts.resize(std::get<0>(vertex_pair_buffer).size(), handle.get_stream()); + + std::tie(edgelist_srcs, edgelist_dsts, std::ignore, std::ignore, std::ignore) = + detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( + handle, std::move(edgelist_srcs), std::move(edgelist_dsts), std::nullopt, std::nullopt, std::nullopt); + } + + // Run thrust::partition to place the edges with triangle counts smaller than K-2 at the end + // zip iterator for the pair + auto edges_to_num_triangles = thrust::make_zip_iterator( + get_dataframe_buffer_begin(vertex_pair_buffer), num_triangles.begin()); + + // FIXME: rename it to edge_to_remove_first + auto invalid_edge_first = thrust::partition( + handle.get_thrust_policy(), + edges_to_num_triangles, + edges_to_num_triangles + num_triangles.size(), + [k]__device__(auto e){ + auto num_triangles = thrust::get<1>(e); + auto is_in_k_truss = num_triangles == 2; + printf("\n is in k_truss = %d\n", is_in_k_truss); + return num_triangles > k; // FIXME k-2 + + } + ); + + // FIXME: rename it num_edges_to_remove + size_t num_invalid_edges{0}; + num_invalid_edges = + static_cast(thrust::distance(invalid_edge_first, edges_to_num_triangles + num_triangles.size())); + + printf("\nthe number of invalid edges = %d\n", num_invalid_edges); + /* + void resize_dataframe_buffer(BufferType& buffer, + size_t new_buffer_size, + rmm::cuda_stream_view stream_view) + */ + + if (num_invalid_edges != 0) { // unroll and remove/mask edges + // Find (p, q, intersection of p&q's neighbor lists) + // before that, determine the interection size of each edge to be removed + rmm::device_uvector intersection_size(intersection_offsets.size() - 1, handle.get_stream()); // take the worst case where all edges are invalidated + // FIXME: might not be necessary to fill it with zeros + thrust::fill(handle.get_thrust_policy(), intersection_size.begin(), intersection_size.end(), size_t{0}); + /* + thrust::tabulate(handle.get_thrust_policy(), + intersection_size.begin(), + intersection_size.end(), + find_intersection{ + //find_intersection{ + size_dataframe_buffer(vertex_pair_buffer), // FIXME: should we just do this inside the functor? + num_invalid_edges, + raft::device_span( + intersection_offsets.data(), intersection_offsets.size()), + raft::device_span( + intersection_indices.data(), intersection_indices.size()), + get_dataframe_buffer_begin(vertex_pair_buffer_), // FIXME: rename to vertex_pairs_buffer? maybe I should pass a pointer here 'get_dataframe_buffer_begin(vertex_pair_buffer)' + //vertex_pairs_begin, + invalid_edge_first + }); + */ + + thrust::tabulate(handle.get_thrust_policy(), + intersection_size.begin(), + intersection_size.end(), + [vertex_pair_buffer_size=size_dataframe_buffer(vertex_pair_buffer), + num_invalid_edges=num_invalid_edges, + intersection_offsets = raft::device_span(intersection_offsets.data(), intersection_offsets.size()), + intersection_indices = raft::device_span(intersection_indices.data(), intersection_indices.size()), + vertex_pair_buffer = get_dataframe_buffer_begin(vertex_pair_buffer), + invalid_edge_first = invalid_edge_first + ] + __device__(edge_t i){ + + printf("\n ***** i = %d\n", i); + printf("\n src = %d, dst = %d\n", thrust::get<0>(*(vertex_pair_buffer+i)), thrust::get<1>(*(vertex_pair_buffer+i))); + printf("\nthe vertex pair buffer size = %d\n", vertex_pair_buffer_size); + if (i < num_invalid_edges) { + //auto pair = thrust::get<0>(*(invalid_edge_first + i)); // FIXME only increment if it is less than 'invalid_edge_first.size()' + vertex_t src = thrust::get<0>(thrust::get<0>(*(invalid_edge_first + i))); + vertex_t dst = thrust::get<1>(thrust::get<0>(*(invalid_edge_first + i))); + //printf("\nsrc_pair = %d, dst_pair = %d, pair = %d\n", thrust::get<0>(pair), thrust::get<1>(pair), pair); + printf("\nsrc_pair = %d, dst_pair = %d", src, dst); + + auto itr = thrust::lower_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + vertex_pair_buffer_size, + //pair, + //thrust::make_tuple(thrust::get<0>(pair), thrust::get<1>(pair)), + //thrust::make_tuple(1, 3), + thrust::make_tuple(src, dst), + thrust::less>{}); + + printf("\n itr = %d\n", itr); + auto idx = thrust::distance(vertex_pair_buffer, itr); + printf("\n distance = %d \n", idx); + + return 1; // dummy + } + else { + return 0; //dummy + } + + }); + + + + + + + + + } + + /* + thrust::partition( + handle.get_thrust_policy(), + get_dataframe_buffer_begin(vertex_pair_buffer), + get_dataframe_buffer_end(vertex_pair_buffer), + [num_triangles = num_triangles.data(), + vertex_pair_buffer = get_dataframe_buffer_begin(vertex_pair_buffer), + //size=get_dataframe_buffer_end(vertex_pair_buffer) - get_dataframe_buffer_begin(vertex_pair_buffer) + size = size_dataframe_buffer(vertex_pair_buffer)] //size_dataframe_buffer + __device__(auto e){ + //printf("src = %d, dst = %d\n", std::get<0>(e), std::get<1>(e)); + auto[src, dst] = e; + auto x = thrust::make_tuple(src, dst); + printf("src_ = %d, dst_ = %d\n", thrust::get<0>(x), thrust::get<1>(x)); + //vertex_t src_ = thrust::get<0>(x); + //vertex_t dst_ = thrust::get<1>(x); + vertex_t dst_ = thrust::get<1>(e); + //vertex_t dst__ = e; + + + //auto x = thrust::make_tuple(src, dst); + //printf("src = %d, dst = %d\n", src, dst); + //printf("src_ = %d, dst_ = %d\n", thrust::get<0>(x), thrust::get<1>(x)); + //vertex_t src_ = thrust::get<0>(x); + //printf("src___ = %d\n", src_); + //vertex_t src_ = thrust::get<0>(x); + //vertex_t dst_ = thrust::get<1>(x); // error: no suitable conversion function from "thrust::detail::cons" to "int32_t" exists + + + + // sort by key. key num of triangles and the values =vertex pairs. thrust::greater + printf("the isze of the buffer = %d\n", size); + auto it = thrust::lower_bound(thrust::seq, + vertex_pair_buffer, + vertex_pair_buffer + size, + e, + thrust::less>{}); + + + auto idx = thrust::distance(vertex_pair_buffer, it); + printf("\nthe distance = %d\n", idx); + + return 0; + + } + ); + */ + + + } }