diff --git a/cpp/examples/README.md b/cpp/examples/README.md index 103e752a597..682c72dbb22 100644 --- a/cpp/examples/README.md +++ b/cpp/examples/README.md @@ -22,7 +22,7 @@ For single_gpu `path_to_executable path_to_a_csv_graph_file [memory allocation mode]` -For multi_gpu and graph_partitioning +For multi_gpu, graph_partitioning and cugraph_operations `mpirun -np 2 path_to_executable path_to_a_csv_graph_file [memory allocation mode]` diff --git a/cpp/examples/cugraph_operations/graph_operations.cu b/cpp/examples/cugraph_operations/graph_operations.cu index 8fc077e4a2a..4440cb1cf47 100644 --- a/cpp/examples/cugraph_operations/graph_operations.cu +++ b/cpp/examples/cugraph_operations/graph_operations.cu @@ -104,7 +104,8 @@ std::unique_ptr initialize_mg_handle(std::string const& allocati template void perform_example_graph_operations(raft::handle_t const& handle, - std::string const& csv_graph_file_path) + std::string const& csv_graph_file_path, + const bool weighted = false) { auto const comm_rank = handle.get_comms().get_rank(); auto const comm_size = handle.get_comms().get_size(); @@ -118,7 +119,7 @@ void perform_example_graph_operations(raft::handle_t const& handle, auto [graph, edge_weights, renumber_map] = cugraph::test::read_graph_from_csv_file( - handle, csv_graph_file_path, true, renumber); + handle, csv_graph_file_path, weighted, renumber); // Non-owning view of the graph object auto graph_view = graph.view(); @@ -132,46 +133,99 @@ void perform_example_graph_operations(raft::handle_t const& handle, auto edge_weight_view = edge_weights ? std::make_optional((*edge_weights).view()) : std::nullopt; using graph_view_t = cugraph::graph_view_t; - rmm::device_uvector vertex_weights = - compute_out_weight_sums(handle, graph_view, *edge_weight_view); - - cugraph::edge_src_property_t src_vertex_weights_cache(handle, graph_view); - - cugraph::edge_dst_property_t dst_vertex_weights_cache(handle, graph_view); - - update_edge_src_property(handle, graph_view, vertex_weights.begin(), src_vertex_weights_cache); - - update_edge_dst_property(handle, graph_view, vertex_weights.begin(), dst_vertex_weights_cache); - - rmm::device_uvector outputs(size_of_the_vertex_partition_assigned_to_this_process, - handle.get_stream()); - - per_v_transform_reduce_incoming_e( - handle, - graph_view, - src_vertex_weights_cache.view(), - dst_vertex_weights_cache.view(), - *edge_weight_view, - [] __device__(auto src, auto dst, auto src_prop, auto dst_prop, auto edge_prop) { - printf("\n%d ---> %d : src_prop= %f dst_prop = %f edge_prop = %f \n", - static_cast(src), - static_cast(dst), - static_cast(src_prop), - static_cast(dst_prop), - static_cast(edge_prop)); - return dst_prop * edge_prop; - }, - weight_t{0}, - cugraph::reduce_op::plus{}, - outputs.begin()); - - auto outputs_title = std::string("outputs_").append(std::to_string(comm_rank)); - size_t max_nr_of_elements_to_print = 10; - RAFT_CUDA_TRY(cudaDeviceSynchronize()); - raft::print_device_vector(outputs_title.c_str(), - outputs.begin(), - std::min(outputs.size(), max_nr_of_elements_to_print), - std::cout); + + // + // As an example operation, compute the weighted average of the properties of + // neighboring vertices, weighted by the edge weights, if the input graph is weighted; + // Otherwise, compute the simple average. + // + if (weighted) { + using result_t = weight_t; + auto vertex_weights = compute_out_weight_sums(handle, graph_view, *edge_weight_view); + + cugraph::edge_src_property_t src_vertex_weights_cache(handle, + graph_view); + + cugraph::edge_dst_property_t dst_vertex_weights_cache(handle, + graph_view); + + update_edge_src_property(handle, graph_view, vertex_weights.begin(), src_vertex_weights_cache); + + update_edge_dst_property(handle, graph_view, vertex_weights.begin(), dst_vertex_weights_cache); + + rmm::device_uvector outputs(size_of_the_vertex_partition_assigned_to_this_process, + handle.get_stream()); + + per_v_transform_reduce_incoming_e( + handle, + graph_view, + src_vertex_weights_cache.view(), + dst_vertex_weights_cache.view(), + (*edge_weight_view), + [new_to_original_id_map = (*renumber_map).data()] __device__( + auto src, auto dst, auto src_prop, auto dst_prop, auto edge_prop) { + printf("\nsrc ---> %d dst = %d : src_prop = %f dst_prop = %f edge_prop = %f\n", + static_cast(new_to_original_id_map[src]), + static_cast(new_to_original_id_map[dst]), + static_cast(src_prop), + static_cast(dst_prop), + static_cast(edge_prop)); + return dst_prop * edge_prop; + }, + result_t{0}, + cugraph::reduce_op::plus{}, + outputs.begin()); + + auto outputs_title = std::string("outputs_").append(std::to_string(comm_rank)); + size_t max_nr_of_elements_to_print = 10; + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + raft::print_device_vector(outputs_title.c_str(), + outputs.begin(), + std::min(outputs.size(), max_nr_of_elements_to_print), + std::cout); + } else { + using result_t = edge_t; + auto vertex_weights = graph_view.compute_out_degrees(handle); + + cugraph::edge_src_property_t src_vertex_weights_cache(handle, + graph_view); + cugraph::edge_dst_property_t dst_vertex_weights_cache(handle, + graph_view); + + update_edge_src_property(handle, graph_view, vertex_weights.begin(), src_vertex_weights_cache); + + update_edge_dst_property(handle, graph_view, vertex_weights.begin(), dst_vertex_weights_cache); + + rmm::device_uvector outputs(size_of_the_vertex_partition_assigned_to_this_process, + handle.get_stream()); + + per_v_transform_reduce_incoming_e( + handle, + graph_view, + src_vertex_weights_cache.view(), + dst_vertex_weights_cache.view(), + cugraph::edge_dummy_property_t{}.view(), + [new_to_original_id_map = (*renumber_map).data()] __device__( + auto src, auto dst, auto src_prop, auto dst_prop, auto) { + printf("\nsrc ---> %d dst = %d : src_prop = %f dst_prop = %f\n", + static_cast(new_to_original_id_map[src]), + static_cast(new_to_original_id_map[dst]), + static_cast(src_prop), + static_cast(dst_prop)); + return dst_prop; + }, + result_t{0}, + cugraph::reduce_op::plus{}, + outputs.begin()); + + auto outputs_title = std::string("outputs_").append(std::to_string(comm_rank)); + size_t max_nr_of_elements_to_print = 10; + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + raft::print_device_vector(outputs_title.c_str(), + outputs.begin(), + std::min(outputs.size(), max_nr_of_elements_to_print), + std::cout); + } } int main(int argc, char** argv) @@ -196,6 +250,6 @@ int main(int argc, char** argv) using weight_t = float; constexpr bool multi_gpu = true; - perform_example_graph_operations(*handle, - csv_graph_file_path); + perform_example_graph_operations( + *handle, csv_graph_file_path, false); } diff --git a/cpp/examples/graph3.csv b/cpp/examples/graph3.csv new file mode 100644 index 00000000000..be50b88bd3d --- /dev/null +++ b/cpp/examples/graph3.csv @@ -0,0 +1,7 @@ +0 1 10 +0 2 20 +0 3 30 +0 4 40 +0 5 50 +6 0 6 + diff --git a/cpp/examples/graph4.csv b/cpp/examples/graph4.csv new file mode 100644 index 00000000000..fbefbcc8cb1 --- /dev/null +++ b/cpp/examples/graph4.csv @@ -0,0 +1,6 @@ +0 1 +0 2 +0 3 +0 4 +0 5 +6 0 diff --git a/cpp/examples/graph_partition/vertex_and_edge_partition.cu b/cpp/examples/graph_partition/vertex_and_edge_partition.cu index ba3df981b1f..edb78a1c290 100644 --- a/cpp/examples/graph_partition/vertex_and_edge_partition.cu +++ b/cpp/examples/graph_partition/vertex_and_edge_partition.cu @@ -98,7 +98,8 @@ std::unique_ptr initialize_mg_handle(std::string const& allocati template void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, - std::string const& csv_graph_file_path) + std::string const& csv_graph_file_path, + bool weighted = false) { auto const comm_rank = handle.get_comms().get_rank(); auto const comm_size = handle.get_comms().get_size(); @@ -112,7 +113,7 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, auto [graph, edge_weights, renumber_map] = cugraph::test::read_graph_from_csv_file( - handle, csv_graph_file_path, true, renumber); + handle, csv_graph_file_path, weighted, renumber); // Meta of the non-owning view of the graph object store vertex/edge partitioning map auto graph_view = graph.view(); @@ -160,35 +161,53 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, // vertex id. The new (aka renumbered) id of the first vertex, ie the vertex at position 0 // of `renumber_map`, assigned to this process - vertex_t renumber_vertex_id_of_local_first = graph_view.local_vertex_partition_range_first(); + vertex_t renumbered_vertex_id_of_local_first = graph_view.local_vertex_partition_range_first(); // The new (aka renumbered) id of the last vertex, ie the vertex at position // `size_of_the_vertex_partition_assigned_to_this_process` - 1 of `renumber_map`, // assigned to this process - vertex_t renumber_vertex_id_of_local_last = graph_view.local_vertex_partition_range_last(); + vertex_t renumbered_vertex_id_of_local_last = graph_view.local_vertex_partition_range_last(); // Print original vertex ids, new (aka renumbered) vertex ids and the ranks of the owner processes if (renumber_map) { - thrust::for_each( - thrust::host, - thrust::make_zip_iterator( - thrust::make_tuple(h_vertices_in_this_proces.begin(), - thrust::make_counting_iterator(renumber_vertex_id_of_local_first))), - thrust::make_zip_iterator( - thrust::make_tuple(h_vertices_in_this_proces.end(), - thrust::make_counting_iterator(renumber_vertex_id_of_local_last))), - [comm_rank](auto old_and_new_id_pair) { - auto old_id = thrust::get<0>(old_and_new_id_pair); - auto new_id = thrust::get<1>(old_and_new_id_pair); - printf("owner rank = %d, original vertex id %d -----> new (renumbered) vertex id %d\n", - comm_rank, - static_cast(old_id), - static_cast(new_id)); - }); + thrust::for_each(thrust::host, + thrust::make_zip_iterator(thrust::make_tuple( + h_vertices_in_this_proces.begin(), + thrust::make_counting_iterator(renumbered_vertex_id_of_local_first))), + thrust::make_zip_iterator(thrust::make_tuple( + h_vertices_in_this_proces.end(), + thrust::make_counting_iterator(renumbered_vertex_id_of_local_last))), + [comm_rank](auto old_and_new_id_pair) { + auto old_id = thrust::get<0>(old_and_new_id_pair); + auto new_id = thrust::get<1>(old_and_new_id_pair); + printf("owner rank = %d, original vertex id %d is renumbered to %d\n", + comm_rank, + static_cast(old_id), + static_cast(new_id)); + }); } + // if (renumber_map) { + // thrust::for_each( + // handle.get_thrust_policy(), + // thrust::make_zip_iterator( + // thrust::make_tuple((*renumber_map).begin(), + // thrust::make_counting_iterator(renumbered_vertex_id_of_local_first))), + // thrust::make_zip_iterator(thrust::make_tuple( + // (*renumber_map).end(), + // thrust::make_counting_iterator(renumbered_vertex_id_of_local_last))), + // [comm_rank](auto old_and_new_id_pair) { + // auto old_id = thrust::get<0>(old_and_new_id_pair); + // auto new_id = thrust::get<1>(old_and_new_id_pair); + // printf("owner rank = %d, original vertex id %d is renumbered to %d\n", + // comm_rank, + // static_cast(old_id), + // static_cast(new_id)); + // }); + // } + // // Look into edge partitions and their associated edge properties (if any) // @@ -239,7 +258,8 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, indices, major_range_first, is_weighted, - weights = weights_of_edges_stored_in_this_edge_partition.begin()] __device__(auto i) { + weights = weights_of_edges_stored_in_this_edge_partition.begin(), + new_to_original_id_map = (*renumber_map).data()] __device__(auto i) { auto v = major_range_first + i; auto deg_of_v_in_this_edge_partition = offsets[i + 1] - offsets[i]; @@ -247,15 +267,16 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, thrust::seq, thrust::make_counting_iterator(edge_t{offsets[i]}), thrust::make_counting_iterator(edge_t{offsets[i + 1]}), - [comm_rank, ep_idx, v, indices, is_weighted, weights] __device__(auto pos) { + [comm_rank, ep_idx, v, indices, new_to_original_id_map, is_weighted, weights] __device__( + auto pos) { if (is_weighted) { printf( "\n[comm_rank = %d local edge partition id = %d] edge: source = %d " "destination = %d weight = %f\n", static_cast(comm_rank), static_cast(ep_idx), - static_cast(v), - static_cast(indices[pos]), + static_cast(new_to_original_id_map[v]), + static_cast(new_to_original_id_map[indices[pos]]), static_cast(weights[pos])); } else { @@ -264,8 +285,8 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, "destination = %d\n", static_cast(comm_rank), static_cast(ep_idx), - static_cast(v), - static_cast(indices[pos])); + static_cast(new_to_original_id_map[v]), + static_cast(new_to_original_id_map[indices[pos]])); } }); }); @@ -283,6 +304,7 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, major_range_first, is_weighted, weights = weights_of_edges_stored_in_this_edge_partition.begin(), + new_to_original_id_map = (*renumber_map).data(), dcs_nzd_vertices = (*dcs_nzd_vertices), major_hypersparse_first = (*major_hypersparse_first)] __device__(auto i) { auto v = dcs_nzd_vertices[i]; @@ -293,15 +315,21 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, thrust::seq, thrust::make_counting_iterator(edge_t{offsets[major_idx]}), thrust::make_counting_iterator(edge_t{offsets[major_idx + 1]}), - [comm_rank, ep_idx, v, indices, is_weighted, weights] __device__(auto pos) { + [comm_rank, + ep_idx, + v, + indices, + new_to_original_id_map, + is_weighted, + weights] __device__(auto pos) { if (is_weighted) { printf( "\n[comm_rank = %d local edge partition id = %d] edge: source = %d " "destination = %d weight = %f\n", static_cast(comm_rank), static_cast(ep_idx), - static_cast(v), - static_cast(indices[pos]), + static_cast(new_to_original_id_map[v]), + static_cast(new_to_original_id_map[indices[pos]]), static_cast(weights[pos])); } else { @@ -310,8 +338,8 @@ void look_into_vertex_and_edge_partitions(raft::handle_t const& handle, "destination = %d\n", static_cast(comm_rank), static_cast(ep_idx), - static_cast(v), - static_cast(indices[pos])); + static_cast(new_to_original_id_map[v]), + static_cast(new_to_original_id_map[indices[pos]])); } }); }); @@ -360,6 +388,6 @@ int main(int argc, char** argv) using weight_t = float; constexpr bool multi_gpu = true; - look_into_vertex_and_edge_partitions(*handle, - csv_graph_file_path); + look_into_vertex_and_edge_partitions( + *handle, csv_graph_file_path, false); } diff --git a/cpp/examples/multi_gpu_application/mg_graph_algorithms.cpp b/cpp/examples/multi_gpu_application/mg_graph_algorithms.cpp index dacecbe10a9..d0d116e3139 100644 --- a/cpp/examples/multi_gpu_application/mg_graph_algorithms.cpp +++ b/cpp/examples/multi_gpu_application/mg_graph_algorithms.cpp @@ -87,13 +87,14 @@ std::unique_ptr initialize_mg_handle(std::string const& allocati return std::move(handle); } - /** * @brief This function reads graph from an input csv file and run BFS and Louvain on it. */ template -void run_graph_algos(raft::handle_t const& handle, std::string const& csv_graph_file_path) +void run_graph_algos(raft::handle_t const& handle, + std::string const& csv_graph_file_path, + bool weighted = false) { auto const comm_rank = handle.get_comms().get_rank(); auto const comm_size = handle.get_comms().get_size(); @@ -102,12 +103,14 @@ void run_graph_algos(raft::handle_t const& handle, std::string const& csv_graph_ bool renumber = true; // must be true for distributed graph. auto [graph, edge_weights, renumber_map] = cugraph::test::read_graph_from_csv_file( - handle, csv_graph_file_path, true, renumber); + handle, csv_graph_file_path, weighted, renumber); auto graph_view = graph.view(); auto edge_weight_view = edge_weights ? std::make_optional((*edge_weights).view()) : std::nullopt; - assert(graph_view.local_vertex_partition_range_size() == (*renumber_map).size()); + if (renumber_map.has_value()) { + assert(graph_view.local_vertex_partition_range_size() == (*renumber_map).size()); + } // run example algorithms // BFS @@ -200,5 +203,5 @@ int main(int argc, char** argv) using weight_t = float; constexpr bool multi_gpu = true; - run_graph_algos(*handle, csv_graph_file_path); + run_graph_algos(*handle, csv_graph_file_path, false); } diff --git a/cpp/examples/single_gpu_application/sg_graph_algorithms.cpp b/cpp/examples/single_gpu_application/sg_graph_algorithms.cpp index 9439788447f..980c847baad 100644 --- a/cpp/examples/single_gpu_application/sg_graph_algorithms.cpp +++ b/cpp/examples/single_gpu_application/sg_graph_algorithms.cpp @@ -57,17 +57,22 @@ std::unique_ptr initialize_sg_handle(std::string const& allocati */ template -void run_graph_algos(raft::handle_t const& handle, std::string const& csv_graph_file_path) +void run_graph_algos(raft::handle_t const& handle, + std::string const& csv_graph_file_path, + bool weighted = false) { std::cout << "Reading graph from " << csv_graph_file_path << std::endl; + bool renumber = false; // for single gpu, this can be true or false auto [graph, edge_weights, renumber_map] = cugraph::test::read_graph_from_csv_file( - handle, csv_graph_file_path, true, true); + handle, csv_graph_file_path, weighted, renumber); auto graph_view = graph.view(); auto edge_weight_view = edge_weights ? std::make_optional((*edge_weights).view()) : std::nullopt; - assert(graph_view.local_vertex_partition_range_size() == (*renumber_map).size()); + if (renumber_map.has_value()) { + assert(graph_view.local_vertex_partition_range_size() == (*renumber_map).size()); + } // run example algorithms // BFS @@ -150,5 +155,5 @@ int main(int argc, char** argv) using edge_t = int32_t; using weight_t = float; - run_graph_algos(*handle, argv[1]); + run_graph_algos(*handle, argv[1], false); } diff --git a/cpp/examples/ugraph.csv b/cpp/examples/ugraph.csv new file mode 100644 index 00000000000..eeb72d46c89 --- /dev/null +++ b/cpp/examples/ugraph.csv @@ -0,0 +1,6 @@ +0 1 +0 2 +1 0 +1 2 +2 0 +2 1