Skip to content

Commit

Permalink
Merge 38d63f8 into 5c75b5b
Browse files Browse the repository at this point in the history
  • Loading branch information
@ChuckHastings
ChuckHastings authored Mar 23, 2022
2 parents 5c75b5b + 38d63f8 commit 7f5cf33
Show file tree
Hide file tree
Showing 10 changed files with 395 additions and 47 deletions.
24 changes: 24 additions & 0 deletions cpp/include/cugraph/detail/shuffle_wrappers.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,5 +93,29 @@ rmm::device_uvector<size_t> groupby_and_count_edgelist_by_local_partition_id(
std::optional<rmm::device_uvector<weight_t>>& d_edgelist_weights,
bool groupby_and_count_local_partition_by_minor = false);

/**
* @brief Collect vertex values (represented as k/v pairs across cluster) and update the
* local value arrays on the GPU responsible for each vertex.
*
* Data will be shuffled and d_local_values[d_vertices[i]] = d_values[i]
*
* @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
* @tparam value_t Type of value associated with the vertex.
* @tparam bool multi_gpu flag
*
* @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator,
* and handles to various CUDA libraries) to run graph algorithms.
* @param[in/out] d_vertices Vertex IDs for the k/v pair
* @param[in/out] d_values Values for the k/v pair
* @param[out] d_local_values The device vector on each GPU that should be updated
* @param[in] local_vertex_first The first vertex id assigned to the local GPU
*/
template <typename vertex_t, typename value_t, bool multi_gpu>
void collect_vertex_values_to_local(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_vertices,
rmm::device_uvector<value_t>& d_values,
rmm::device_uvector<value_t>& d_local_values,
vertex_t local_vertex_first);

} // namespace detail
} // namespace cugraph
4 changes: 2 additions & 2 deletions cpp/include/cugraph_c/graph.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ typedef struct {
bool_t is_multigraph;
} cugraph_graph_properties_t;

// FIXME: Add support for specifying isloated vertices
// FIXME: Add support for specifying isolated vertices
/**
* @brief Construct an SG graph
*
Expand Down Expand Up @@ -74,7 +74,7 @@ cugraph_error_code_t cugraph_sg_graph_create(const cugraph_resource_handle_t* ha
// but didn't want to confuse with original cugraph_free_graph
void cugraph_sg_graph_free(cugraph_graph_t* graph);

// FIXME: Add support for specifying isloated vertices
// FIXME: Add support for specifying isolated vertices
/**
* @brief Construct an MG graph
*
Expand Down
65 changes: 36 additions & 29 deletions cpp/src/c_api/hits.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
#include <c_api/utils.hpp>

#include <cugraph/algorithms.hpp>
#include <cugraph/detail/shuffle_wrappers.hpp>
#include <cugraph/detail/utility_wrappers.hpp>
#include <cugraph/graph_functions.hpp>

Expand Down Expand Up @@ -113,40 +114,46 @@ struct hits_functor : public cugraph::c_api::abstract_functor {
weight_t hub_score_differences{0};
size_t number_of_iterations{0};

#if 0
// FIXME: Implementation will look something like this.

if (initial_hubs_guess_ != nullptr) {
//
// Need to renumber initial_hubs_guess_vertices
// Need to shuffle and populate hubs
//
// This is the original pagerank code, it will be sort of like this
renumber_ext_vertices<vertex_t, multi_gpu>(handle_,
personalization_vertices_->as_type<vertex_t>(),
personalization_vertices_->size_,
number_map->data(),
graph_view.get_local_vertex_first(),
graph_view.get_local_vertex_last(),
do_expensive_check_);
if (initial_hubs_guess_vertices_ != nullptr) {
rmm::device_uvector<vertex_t> guess_vertices(initial_hubs_guess_vertices_->size_,
handle_.get_stream());
rmm::device_uvector<weight_t> guess_values(initial_hubs_guess_values_->size_,
handle_.get_stream());

raft::copy(guess_vertices.data(),
initial_hubs_guess_vertices_->as_type<vertex_t>(),
guess_vertices.size(),
handle_.get_stream());
raft::copy(guess_values.data(),
initial_hubs_guess_values_->as_type<weight_t>(),
guess_values.size(),
handle_.get_stream());

cugraph::renumber_ext_vertices<vertex_t, multi_gpu>(handle_,
guess_vertices.data(),
guess_vertices.size(),
number_map->data(),
graph_view.get_local_vertex_first(),
graph_view.get_local_vertex_last(),
do_expensive_check_);

cugraph::detail::collect_vertex_values_to_local<vertex_t, weight_t, multi_gpu>(
handle_, guess_vertices, guess_values, hubs, graph_view.get_local_vertex_first());
}

// TODO: Add these to the result
std::tie(hub_score_differences, number_of_iterations) =
cugraph::hits<vertex_t, edge_t, weight_t, multi_gpu>(handle_,
graph_view,
hubs.data(),
authorities.data(),
epsilon_,
max_iterations_,
has_initial_hubs_guess,
normalize_,
do_expensive_check_);
cugraph::hits<vertex_t, edge_t, weight_t, multi_gpu>(
handle_,
graph_view,
hubs.data(),
authorities.data(),
epsilon_,
max_iterations_,
(initial_hubs_guess_vertices_ != nullptr),
normalize_,
do_expensive_check_);

raft::copy(vertex_ids.data(), number_map->data(), vertex_ids.size(), handle_.get_stream());
#else
unsupported();
#endif

result_ = new cugraph::c_api::cugraph_hits_result_t{
new cugraph::c_api::cugraph_type_erased_device_array_t(vertex_ids, graph_->vertex_type_),
Expand Down
98 changes: 98 additions & 0 deletions cpp/src/detail/shuffle_wrappers.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -335,5 +335,103 @@ template rmm::device_uvector<size_t> groupby_and_count_edgelist_by_local_partiti
std::optional<rmm::device_uvector<double>>& d_edgelist_weights,
bool groupby_and_counts_local_partition);

template <typename vertex_t, typename value_t, bool multi_gpu>
void collect_vertex_values_to_local(raft::handle_t const& handle,
rmm::device_uvector<vertex_t>& d_vertices,
rmm::device_uvector<value_t>& d_values,
rmm::device_uvector<value_t>& d_local_values,
vertex_t local_vertex_first)
{
if constexpr (multi_gpu) {
auto& comm = handle.get_comms();
auto const comm_size = comm.get_size();

rmm::device_uvector<vertex_t> d_rx_vertices(0, handle.get_stream());
rmm::device_uvector<value_t> d_rx_values(0, handle.get_stream());

std::tie(d_rx_vertices, d_rx_values, std::ignore) =
cugraph::groupby_gpu_id_and_shuffle_kv_pairs(
comm,
d_vertices.begin(),
d_vertices.end(),
d_values.begin(),
[key_func = cugraph::detail::compute_gpu_id_from_vertex_t<vertex_t>{comm_size}] __device__(
auto val) { return key_func(val); },
handle.get_stream());

auto vertex_iterator = thrust::make_transform_iterator(
d_rx_vertices.begin(),
[local_vertex_first] __device__(vertex_t v) { return v - local_vertex_first; });

thrust::scatter(handle.get_thrust_policy(),
d_rx_values.begin(),
d_rx_values.end(),
vertex_iterator,
d_local_values.begin());
} else {
thrust::scatter(handle.get_thrust_policy(),
d_values.begin(),
d_values.end(),
d_vertices.begin(),
d_local_values.begin());
}
}

template void collect_vertex_values_to_local<int32_t, float, false>(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_vertices,
rmm::device_uvector<float>& d_values,
rmm::device_uvector<float>& d_local_values,
int32_t local_vertex_first);

template void collect_vertex_values_to_local<int64_t, float, false>(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_vertices,
rmm::device_uvector<float>& d_values,
rmm::device_uvector<float>& d_local_values,
int64_t local_vertex_first);

template void collect_vertex_values_to_local<int32_t, double, false>(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_vertices,
rmm::device_uvector<double>& d_values,
rmm::device_uvector<double>& d_local_values,
int32_t local_vertex_first);

template void collect_vertex_values_to_local<int64_t, double, false>(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_vertices,
rmm::device_uvector<double>& d_values,
rmm::device_uvector<double>& d_local_values,
int64_t local_vertex_first);

template void collect_vertex_values_to_local<int32_t, float, true>(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_vertices,
rmm::device_uvector<float>& d_values,
rmm::device_uvector<float>& d_local_values,
int32_t local_vertex_first);

template void collect_vertex_values_to_local<int64_t, float, true>(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_vertices,
rmm::device_uvector<float>& d_values,
rmm::device_uvector<float>& d_local_values,
int64_t local_vertex_first);

template void collect_vertex_values_to_local<int32_t, double, true>(
raft::handle_t const& handle,
rmm::device_uvector<int32_t>& d_vertices,
rmm::device_uvector<double>& d_values,
rmm::device_uvector<double>& d_local_values,
int32_t local_vertex_first);

template void collect_vertex_values_to_local<int64_t, double, true>(
raft::handle_t const& handle,
rmm::device_uvector<int64_t>& d_vertices,
rmm::device_uvector<double>& d_values,
rmm::device_uvector<double>& d_local_values,
int64_t local_vertex_first);

} // namespace detail
} // namespace cugraph
1 change: 1 addition & 0 deletions cpp/tests/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -620,6 +620,7 @@ if(BUILD_CUGRAPH_MG_TESTS)
# - MG C API tests ------------------------------------------------------------------------
ConfigureCTestMG(MG_CAPI_CREATE_GRAPH c_api/mg_create_graph_test.c c_api/mg_test_utils.cpp)
ConfigureCTestMG(MG_CAPI_PAGERANK c_api/mg_pagerank_test.c c_api/mg_test_utils.cpp)
ConfigureCTestMG(MG_CAPI_HITS c_api/mg_hits_test.c c_api/mg_test_utils.cpp)
else()
message(FATAL_ERROR "OpenMPI NOT found, cannot build MG tests.")
endif()
Expand Down
17 changes: 4 additions & 13 deletions cpp/tests/c_api/hits_test.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,6 @@ int generic_hits_test(vertex_t* h_src,
size_t num_vertices,
size_t num_edges,
bool_t store_transposed,
double alpha,
double epsilon,
size_t max_iterations)
{
Expand All @@ -58,9 +57,6 @@ int generic_hits_test(vertex_t* h_src,
ret_code = cugraph_hits(
p_handle, p_graph, epsilon, max_iterations, NULL, NULL, FALSE, FALSE, &p_result, &ret_error);

TEST_ASSERT(test_ret_value, ret_code != CUGRAPH_SUCCESS, "cugraph_hits worked, but it's not implemented!!!");

#if 0
TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "cugraph_hits failed.");

cugraph_type_erased_device_array_view_t* vertices;
Expand Down Expand Up @@ -102,7 +98,6 @@ int generic_hits_test(vertex_t* h_src,
cugraph_sg_graph_free(p_graph);
cugraph_free_resource_handle(p_handle);
cugraph_error_free(ret_error);
#endif

return test_ret_value;
}
Expand All @@ -115,10 +110,9 @@ int test_hits()
vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4};
vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5};
weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f};
weight_t h_hubs[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237};
weight_t h_authorities[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237};
weight_t h_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0};
weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136};

double alpha = 0.95;
double epsilon = 0.0001;
size_t max_iterations = 20;

Expand All @@ -131,7 +125,6 @@ int test_hits()
num_vertices,
num_edges,
TRUE,
alpha,
epsilon,
max_iterations);
}
Expand All @@ -144,10 +137,9 @@ int test_hits_with_transpose()
vertex_t h_src[] = {0, 1, 1, 2, 2, 2, 3, 4};
vertex_t h_dst[] = {1, 3, 4, 0, 1, 3, 5, 5};
weight_t h_wgt[] = {0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f};
weight_t h_hubs[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237};
weight_t h_authorities[] = {0.0915528, 0.168382, 0.0656831, 0.191468, 0.120677, 0.362237};
weight_t h_hubs[] = {0.347296, 0.532089, 1, 0.00003608, 0.00003608, 0};
weight_t h_authorities[] = {0.652703, 0.879385, 0, 1, 0.347296, 0.00009136};

double alpha = 0.95;
double epsilon = 0.0001;
size_t max_iterations = 20;

Expand All @@ -162,7 +154,6 @@ int test_hits_with_transpose()
num_vertices,
num_edges,
FALSE,
alpha,
epsilon,
max_iterations);
}
Expand Down
Loading

0 comments on commit 7f5cf33

Please sign in to comment.