From dd2eb361c97ed01bff2354ee7e57632edcc3cd69 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 11 Jan 2021 19:48:29 -0500 Subject: [PATCH 01/12] Turn off symmetrization in MST output --- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 2ae4d93113..87cb7b9a13 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -339,9 +339,9 @@ void MST_solver::min_edge_per_supervertex() { // the above kernel only adds directed mst edges in the case where // a pair of vertices don't pick the same min edge between them // so, now we add the reverse edge to make it undirected - detail::add_reverse_edge<<>>( - new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, - temp_weights_ptr, v); +// detail::add_reverse_edge<<>>( +// new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, +// temp_weights_ptr, v); } template From 1d3f4f4cf36b2806a9d5eb6116e566e4352b4d8b Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 12 Jan 2021 19:24:16 -0500 Subject: [PATCH 02/12] Some changes --- .../raft/sparse/mst/detail/mst_kernels.cuh | 63 ++++++++++++++++--- 1 file changed, 56 insertions(+), 7 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 895927d6e1..620f75fd26 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -102,6 +102,37 @@ __global__ void kernel_min_edge_per_vertex( } } +//template +//__global__ void min_edge_per_supervertex( +// const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, +// bool* mst_edge, const vertex_t* indices, const weight_t* weights, +// const weight_t* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, +// weight_t* temp_weights, const weight_t* min_edge_color, const vertex_t v) { +// auto tid = get_1D_idx(); +// +// if (tid < v) { +// vertex_t vertex_color_idx = color_index[tid]; +// vertex_t vertex_color = color[vertex_color_idx]; +// edge_t edge_idx = new_mst_edge[tid]; +// +// // check if valid outgoing edge was found +// // find minimum edge is same as minimum edge of whole supervertex +// // if yes, that is part of mst +// if (edge_idx != std::numeric_limits::max()) { +// weight_t vertex_weight = altered_weights[edge_idx]; +// if (min_edge_color[vertex_color] == vertex_weight) { +// temp_src[tid] = tid; +// temp_dst[tid] = indices[edge_idx]; +// temp_weights[tid] = weights[edge_idx]; +// +// mst_edge[edge_idx] = true; +// } else { +// new_mst_edge[tid] = std::numeric_limits::max(); +// } +// } +// } +//} + template __global__ void min_edge_per_supervertex( const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, @@ -109,23 +140,40 @@ __global__ void min_edge_per_supervertex( const weight_t* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, weight_t* temp_weights, const weight_t* min_edge_color, const vertex_t v) { auto tid = get_1D_idx(); - if (tid < v) { vertex_t vertex_color_idx = color_index[tid]; vertex_t vertex_color = color[vertex_color_idx]; edge_t edge_idx = new_mst_edge[tid]; - // check if valid outgoing edge was found // find minimum edge is same as minimum edge of whole supervertex // if yes, that is part of mst if (edge_idx != std::numeric_limits::max()) { weight_t vertex_weight = altered_weights[edge_idx]; if (min_edge_color[vertex_color] == vertex_weight) { - temp_src[tid] = tid; - temp_dst[tid] = indices[edge_idx]; - temp_weights[tid] = weights[edge_idx]; - - mst_edge[edge_idx] = true; + auto dst = indices[edge_idx]; + auto dst_edge_idx = new_mst_edge[dst]; + auto dst_color = color[color_index[dst]]; + // vertices added each other + // only if destination has found an edge + // the edge points back to source + // the edge is minimum edge found for dst color + if (dst_edge_idx != std::numeric_limits::max() && + indices[dst_edge_idx] == tid && + min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { + auto dst_src = indices[dst_edge_idx]; + if (vertex_color < dst_color) { + temp_src[tid] = tid; + temp_dst[tid] = dst; + temp_weights[tid] = weights[edge_idx]; + mst_edge[edge_idx] = true; + } + } + else { + temp_src[tid] = tid; + temp_dst[tid] = dst; + temp_weights[tid] = weights[edge_idx]; + mst_edge[edge_idx] = true; + } } else { new_mst_edge[tid] = std::numeric_limits::max(); } @@ -273,6 +321,7 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, } } + template __global__ void kernel_count_new_mst_edges(const vertex_t* mst_src, edge_t* mst_edge_count, From 3af9ff7f3b35d5e76992bd36aa35335f5a9c526b Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 20 Jan 2021 19:28:44 -0500 Subject: [PATCH 03/12] Updates to MST --- .../raft/sparse/mst/detail/mst_kernels.cuh | 32 ------------------- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 15 +++++---- cpp/include/raft/sparse/mst/mst.cuh | 3 ++ cpp/include/raft/sparse/mst/mst_solver.cuh | 9 +++--- 4 files changed, 17 insertions(+), 42 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 620f75fd26..d38c8e6b2a 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -102,37 +102,6 @@ __global__ void kernel_min_edge_per_vertex( } } -//template -//__global__ void min_edge_per_supervertex( -// const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, -// bool* mst_edge, const vertex_t* indices, const weight_t* weights, -// const weight_t* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, -// weight_t* temp_weights, const weight_t* min_edge_color, const vertex_t v) { -// auto tid = get_1D_idx(); -// -// if (tid < v) { -// vertex_t vertex_color_idx = color_index[tid]; -// vertex_t vertex_color = color[vertex_color_idx]; -// edge_t edge_idx = new_mst_edge[tid]; -// -// // check if valid outgoing edge was found -// // find minimum edge is same as minimum edge of whole supervertex -// // if yes, that is part of mst -// if (edge_idx != std::numeric_limits::max()) { -// weight_t vertex_weight = altered_weights[edge_idx]; -// if (min_edge_color[vertex_color] == vertex_weight) { -// temp_src[tid] = tid; -// temp_dst[tid] = indices[edge_idx]; -// temp_weights[tid] = weights[edge_idx]; -// -// mst_edge[edge_idx] = true; -// } else { -// new_mst_edge[tid] = std::numeric_limits::max(); -// } -// } -// } -//} - template __global__ void min_edge_per_supervertex( const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, @@ -160,7 +129,6 @@ __global__ void min_edge_per_supervertex( if (dst_edge_idx != std::numeric_limits::max() && indices[dst_edge_idx] == tid && min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { - auto dst_src = indices[dst_edge_idx]; if (vertex_color < dst_color) { temp_src[tid] = tid; temp_dst[tid] = dst; diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 87cb7b9a13..0920759324 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -71,6 +71,9 @@ MST_solver::MST_solver( mst_edge_count(1, 0), prev_mst_edge_count(1, 0), stream(stream_) { + + printf("GOT HERE!!!\n"); + max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; sm_count = handle_.get_device_properties().multiProcessorCount; @@ -262,9 +265,9 @@ void MST_solver::label_prop(vertex_t* mst_src, // update the colors of both ends its until there is no change in colors thrust::host_vector curr_mst_edge_count = mst_edge_count; - auto min_pair_nthreads = std::min(v, max_threads); + auto min_pair_nthreads = std::min(v, (vertex_t)max_threads); auto min_pair_nblocks = - std::min((v + min_pair_nthreads - 1) / min_pair_nthreads, max_blocks); + std::min((v + min_pair_nthreads - 1) / min_pair_nthreads, (vertex_t)max_blocks); rmm::device_vector done(1, false); @@ -316,8 +319,8 @@ void MST_solver::min_edge_per_vertex() { // Finds the minimum edge from each supervertex to the lowest color template void MST_solver::min_edge_per_supervertex() { - int nthreads = std::min(v, max_threads); - int nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); + auto nthreads = std::min(v, max_threads); + auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); thrust::fill(temp_src.begin(), temp_src.end(), std::numeric_limits::max()); @@ -346,8 +349,8 @@ void MST_solver::min_edge_per_supervertex() { template void MST_solver::check_termination() { - int nthreads = std::min(2 * v, max_threads); - int nblocks = std::min((2 * v + nthreads - 1) / nthreads, max_blocks); + vertex_t nthreads = std::min(2 * v, (vertex_t)max_threads); + vertex_t nblocks = std::min((2 * v + nthreads - 1) / nthreads, (vertex_t)max_blocks); // count number of new mst edges edge_t* mst_edge_count_ptr = mst_edge_count.data().get(); diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index d9caca3ba4..e49e118c32 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -27,6 +27,9 @@ raft::Graph_COO mst( const raft::handle_t& handle, edge_t const* offsets, vertex_t const* indices, weight_t const* weights, vertex_t const v, edge_t const e, vertex_t* color, cudaStream_t stream) { + + + MST_solver mst_solver( handle, offsets, indices, weights, v, e, color, stream); return mst_solver.solve(); diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index d747a32eaf..218869e20a 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -31,7 +31,8 @@ struct Graph_COO { edge_t n_edges; Graph_COO(vertex_t size, cudaStream_t stream) - : src(size, stream), dst(size, stream), weights(size, stream) {} + : src(size, stream), dst(size, stream), weights(size, stream) { + } }; namespace mst { @@ -59,9 +60,9 @@ class MST_solver { const vertex_t v; const edge_t e; - int max_blocks; - int max_threads; - int sm_count; + vertex_t max_blocks; + vertex_t max_threads; + vertex_t sm_count; vertex_t* color; // represent each supervertex as a color rmm::device_vector From 7737535a4f124b0c4a0c9bfe381aa25e5d5c9b5d Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 20 Jan 2021 19:29:08 -0500 Subject: [PATCH 04/12] Fixing style --- .../raft/sparse/mst/detail/mst_kernels.cuh | 4 +--- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 18 +++++++++--------- cpp/include/raft/sparse/mst/mst.cuh | 3 --- cpp/include/raft/sparse/mst/mst_solver.cuh | 3 +-- 4 files changed, 11 insertions(+), 17 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index d38c8e6b2a..32208f79f4 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -135,8 +135,7 @@ __global__ void min_edge_per_supervertex( temp_weights[tid] = weights[edge_idx]; mst_edge[edge_idx] = true; } - } - else { + } else { temp_src[tid] = tid; temp_dst[tid] = dst; temp_weights[tid] = weights[edge_idx]; @@ -289,7 +288,6 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, } } - template __global__ void kernel_count_new_mst_edges(const vertex_t* mst_src, edge_t* mst_edge_count, diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 0920759324..76192115ec 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -39,11 +39,11 @@ typedef std::chrono::high_resolution_clock Clock; // curand generator uniform inline curandStatus_t curand_generate_uniformX(curandGenerator_t generator, - float* outputPtr, size_t n) { + float* outputPtr, size_t n) { return curandGenerateUniform(generator, outputPtr, n); } inline curandStatus_t curand_generate_uniformX(curandGenerator_t generator, - double* outputPtr, size_t n) { + double* outputPtr, size_t n) { return curandGenerateUniformDouble(generator, outputPtr, n); } @@ -71,7 +71,6 @@ MST_solver::MST_solver( mst_edge_count(1, 0), prev_mst_edge_count(1, 0), stream(stream_) { - printf("GOT HERE!!!\n"); max_blocks = handle_.get_device_properties().maxGridSize[0]; @@ -266,8 +265,8 @@ void MST_solver::label_prop(vertex_t* mst_src, thrust::host_vector curr_mst_edge_count = mst_edge_count; auto min_pair_nthreads = std::min(v, (vertex_t)max_threads); - auto min_pair_nblocks = - std::min((v + min_pair_nthreads - 1) / min_pair_nthreads, (vertex_t)max_blocks); + auto min_pair_nblocks = std::min( + (v + min_pair_nthreads - 1) / min_pair_nthreads, (vertex_t)max_blocks); rmm::device_vector done(1, false); @@ -342,15 +341,16 @@ void MST_solver::min_edge_per_supervertex() { // the above kernel only adds directed mst edges in the case where // a pair of vertices don't pick the same min edge between them // so, now we add the reverse edge to make it undirected -// detail::add_reverse_edge<<>>( -// new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, -// temp_weights_ptr, v); + // detail::add_reverse_edge<<>>( + // new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, + // temp_weights_ptr, v); } template void MST_solver::check_termination() { vertex_t nthreads = std::min(2 * v, (vertex_t)max_threads); - vertex_t nblocks = std::min((2 * v + nthreads - 1) / nthreads, (vertex_t)max_blocks); + vertex_t nblocks = + std::min((2 * v + nthreads - 1) / nthreads, (vertex_t)max_blocks); // count number of new mst edges edge_t* mst_edge_count_ptr = mst_edge_count.data().get(); diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index e49e118c32..d9caca3ba4 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -27,9 +27,6 @@ raft::Graph_COO mst( const raft::handle_t& handle, edge_t const* offsets, vertex_t const* indices, weight_t const* weights, vertex_t const v, edge_t const e, vertex_t* color, cudaStream_t stream) { - - - MST_solver mst_solver( handle, offsets, indices, weights, v, e, color, stream); return mst_solver.solve(); diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index 218869e20a..c46baa9c27 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -31,8 +31,7 @@ struct Graph_COO { edge_t n_edges; Graph_COO(vertex_t size, cudaStream_t stream) - : src(size, stream), dst(size, stream), weights(size, stream) { - } + : src(size, stream), dst(size, stream), weights(size, stream) {} }; namespace mst { From 91091908454ba51da7b0158933d682a52a77b8c9 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 22 Feb 2021 13:37:45 -0500 Subject: [PATCH 05/12] remove debug print from mst solver --- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 76192115ec..7e579546eb 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -71,7 +71,6 @@ MST_solver::MST_solver( mst_edge_count(1, 0), prev_mst_edge_count(1, 0), stream(stream_) { - printf("GOT HERE!!!\n"); max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; From 31239c4ee42e4339dab7aa88c7536947fb7d0a18 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 22 Feb 2021 14:03:43 -0500 Subject: [PATCH 06/12] Adding `symmetrize_output` optional argument to MST --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 9 +++++---- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 14 +++++++++----- cpp/include/raft/sparse/mst/mst.cuh | 4 ++-- cpp/include/raft/sparse/mst/mst_solver.cuh | 4 +++- cpp/test/mst.cu | 2 +- 5 files changed, 20 insertions(+), 13 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 32208f79f4..d90264cb1f 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -126,16 +126,17 @@ __global__ void min_edge_per_supervertex( // only if destination has found an edge // the edge points back to source // the edge is minimum edge found for dst color + bool add_edge = false; if (dst_edge_idx != std::numeric_limits::max() && indices[dst_edge_idx] == tid && min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { if (vertex_color < dst_color) { - temp_src[tid] = tid; - temp_dst[tid] = dst; - temp_weights[tid] = weights[edge_idx]; - mst_edge[edge_idx] = true; + add_edge = true; } } else { + add_edge = true; + } + if (add_edge) { temp_src[tid] = tid; temp_dst[tid] = dst; temp_weights[tid] = weights[edge_idx]; diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 7e579546eb..f1d1f6bb2c 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -51,7 +51,8 @@ template MST_solver::MST_solver( const raft::handle_t& handle_, const edge_t* offsets_, const vertex_t* indices_, const weight_t* weights_, const vertex_t v_, - const edge_t e_, vertex_t* color_, cudaStream_t stream_) + const edge_t e_, vertex_t* color_, cudaStream_t stream_, + bool symmetrize_output_) : handle(handle_), offsets(offsets_), indices(indices_), @@ -70,7 +71,8 @@ MST_solver::MST_solver( temp_weights(2 * v_), mst_edge_count(1, 0), prev_mst_edge_count(1, 0), - stream(stream_) { + stream(stream_), + symmetrize_output(symmetrize_output_){ max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; @@ -340,9 +342,11 @@ void MST_solver::min_edge_per_supervertex() { // the above kernel only adds directed mst edges in the case where // a pair of vertices don't pick the same min edge between them // so, now we add the reverse edge to make it undirected - // detail::add_reverse_edge<<>>( - // new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, - // temp_weights_ptr, v); + if(symmetrize_output) { + detail::add_reverse_edge<<>>( + new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, + temp_weights_ptr, v); + } } template diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index d9caca3ba4..453fa9f1c1 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -26,9 +26,9 @@ template raft::Graph_COO mst( const raft::handle_t& handle, edge_t const* offsets, vertex_t const* indices, weight_t const* weights, vertex_t const v, edge_t const e, vertex_t* color, - cudaStream_t stream) { + cudaStream_t stream, bool symmetrize_output = true) { MST_solver mst_solver( - handle, offsets, indices, weights, v, e, color, stream); + handle, offsets, indices, weights, v, e, color, stream, symmetrize_output); return mst_solver.solve(); } diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index c46baa9c27..dedbe06370 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -42,7 +42,7 @@ class MST_solver { MST_solver(const raft::handle_t& handle_, const edge_t* offsets_, const vertex_t* indices_, const weight_t* weights_, const vertex_t v_, const edge_t e_, vertex_t* color_, - cudaStream_t stream_); + cudaStream_t stream_, bool symmetrize_output_); raft::Graph_COO solve(); @@ -52,6 +52,8 @@ class MST_solver { const raft::handle_t& handle; cudaStream_t stream; + bool symmetrize_output; + //CSR const edge_t* offsets; const vertex_t* indices; diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 4005238812..7e73102fcd 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -129,7 +129,7 @@ class MSTTest vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); MST_solver mst_solver( - handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream()); + handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), true); auto result = mst_solver.solve(); raft::print_device_vector("Final MST Src: ", result.src.data(), result.n_edges, std::cout); From b45d899c70b95415522849e93e3fa3e73f1a8ad6 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 22 Feb 2021 14:14:32 -0500 Subject: [PATCH 07/12] Pushing Divye's changes --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 10 +++++----- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 2 +- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index d90264cb1f..ee6055a7e6 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -107,7 +107,8 @@ __global__ void min_edge_per_supervertex( const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, bool* mst_edge, const vertex_t* indices, const weight_t* weights, const weight_t* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, - weight_t* temp_weights, const weight_t* min_edge_color, const vertex_t v) { + weight_t* temp_weights, const weight_t* min_edge_color, const vertex_t v, + bool symmetrize_output) { auto tid = get_1D_idx(); if (tid < v) { vertex_t vertex_color_idx = color_index[tid]; @@ -130,7 +131,7 @@ __global__ void min_edge_per_supervertex( if (dst_edge_idx != std::numeric_limits::max() && indices[dst_edge_idx] == tid && min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { - if (vertex_color < dst_color) { + if (symmetrize_output || vertex_color < dst_color) { add_edge = true; } } else { @@ -141,14 +142,13 @@ __global__ void min_edge_per_supervertex( temp_dst[tid] = dst; temp_weights[tid] = weights[edge_idx]; mst_edge[edge_idx] = true; + } else { + new_mst_edge[tid] = std::numeric_limits::max(); } - } else { - new_mst_edge[tid] = std::numeric_limits::max(); } } } } - template __global__ void add_reverse_edge(const edge_t* new_mst_edge, const vertex_t* indices, diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index f1d1f6bb2c..487d7874db 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -337,7 +337,7 @@ void MST_solver::min_edge_per_supervertex() { detail::min_edge_per_supervertex<<>>( color, color_index_ptr, new_mst_edge_ptr, mst_edge_ptr, indices, weights, altered_weights_ptr, temp_src_ptr, temp_dst_ptr, temp_weights_ptr, - min_edge_color_ptr, v); + min_edge_color_ptr, v, symmetrize_output); // the above kernel only adds directed mst edges in the case where // a pair of vertices don't pick the same min edge between them From c16d0e5689ea32d3cb077588b3c48846ad9d7f86 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 22 Feb 2021 14:16:50 -0500 Subject: [PATCH 08/12] Fixing style --- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 7 +++---- cpp/test/mst.cu | 3 ++- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 487d7874db..00fae46ecc 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -72,8 +72,7 @@ MST_solver::MST_solver( mst_edge_count(1, 0), prev_mst_edge_count(1, 0), stream(stream_), - symmetrize_output(symmetrize_output_){ - + symmetrize_output(symmetrize_output_) { max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; sm_count = handle_.get_device_properties().multiProcessorCount; @@ -342,10 +341,10 @@ void MST_solver::min_edge_per_supervertex() { // the above kernel only adds directed mst edges in the case where // a pair of vertices don't pick the same min edge between them // so, now we add the reverse edge to make it undirected - if(symmetrize_output) { + if (symmetrize_output) { detail::add_reverse_edge<<>>( new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, - temp_weights_ptr, v); + temp_weights_ptr, v); } } diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 7e73102fcd..59d9a80a86 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -129,7 +129,8 @@ class MSTTest vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); MST_solver mst_solver( - handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), true); + handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), + true); auto result = mst_solver.solve(); raft::print_device_vector("Final MST Src: ", result.src.data(), result.n_edges, std::cout); From e9c6e56276c6f63f0cf0752b64d5e10b39438b33 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 23 Feb 2021 20:09:42 -0800 Subject: [PATCH 09/12] debugging symmetrize output --- .../raft/sparse/mst/detail/mst_kernels.cuh | 59 +++++++++++-------- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 2 +- 2 files changed, 37 insertions(+), 24 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index ee6055a7e6..eec3d9edce 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -114,47 +114,58 @@ __global__ void min_edge_per_supervertex( vertex_t vertex_color_idx = color_index[tid]; vertex_t vertex_color = color[vertex_color_idx]; edge_t edge_idx = new_mst_edge[tid]; + // check if valid outgoing edge was found // find minimum edge is same as minimum edge of whole supervertex // if yes, that is part of mst if (edge_idx != std::numeric_limits::max()) { weight_t vertex_weight = altered_weights[edge_idx]; + + bool add_edge = false; if (min_edge_color[vertex_color] == vertex_weight) { + add_edge = true; + auto dst = indices[edge_idx]; - auto dst_edge_idx = new_mst_edge[dst]; - auto dst_color = color[color_index[dst]]; - // vertices added each other - // only if destination has found an edge - // the edge points back to source - // the edge is minimum edge found for dst color - bool add_edge = false; - if (dst_edge_idx != std::numeric_limits::max() && - indices[dst_edge_idx] == tid && - min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { - if (symmetrize_output || vertex_color < dst_color) { - add_edge = true; - } - } else { - add_edge = true; + if (!symmetrize_output) { + auto dst_edge_idx = new_mst_edge[dst]; + auto dst_color = color[color_index[dst]]; + + // vertices added each other + // only if destination has found an edge + // the edge points back to source + // the edge is minimum edge found for dst color + if (dst_edge_idx != std::numeric_limits::max() && + indices[dst_edge_idx] == tid && + min_edge_color[dst_color] == altered_weights[dst_edge_idx]) { + if (vertex_color > dst_color) { + add_edge = false; + } + } } + // else { + // add_edge = true; + // } if (add_edge) { temp_src[tid] = tid; temp_dst[tid] = dst; temp_weights[tid] = weights[edge_idx]; mst_edge[edge_idx] = true; - } else { - new_mst_edge[tid] = std::numeric_limits::max(); } } + + if (!add_edge) { + new_mst_edge[tid] = std::numeric_limits::max(); + } } } } + template __global__ void add_reverse_edge(const edge_t* new_mst_edge, const vertex_t* indices, const weight_t* weights, vertex_t* temp_src, vertex_t* temp_dst, weight_t* temp_weights, - const vertex_t v) { + const vertex_t v, bool symmetrize_output) { auto tid = get_1D_idx(); if (tid < v) { @@ -171,12 +182,14 @@ __global__ void add_reverse_edge(const edge_t* new_mst_edge, reverse_needed = true; } else { // check what vertex the neighbor vertex picked - vertex_t neighbor_vertex_neighbor = indices[neighbor_edge_idx]; + if (symmetrize_output) { + vertex_t neighbor_vertex_neighbor = indices[neighbor_edge_idx]; - // if vertices did not pick each other - // add a reverse edge - if (tid != neighbor_vertex_neighbor) { - reverse_needed = true; + // if vertices did not pick each other + // add a reverse edge + if (tid != neighbor_vertex_neighbor) { + reverse_needed = true; + } } } diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh index 00fae46ecc..d3a82ca711 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -344,7 +344,7 @@ void MST_solver::min_edge_per_supervertex() { if (symmetrize_output) { detail::add_reverse_edge<<>>( new_mst_edge_ptr, indices, weights, temp_src_ptr, temp_dst_ptr, - temp_weights_ptr, v); + temp_weights_ptr, v, symmetrize_output); } } From ac440d4b4886ac8013feb99405a211c1f58224b9 Mon Sep 17 00:00:00 2001 From: divyegala Date: Wed, 24 Feb 2021 09:20:09 -0800 Subject: [PATCH 10/12] style fixes --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index eec3d9edce..2bbcef19e6 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -114,7 +114,7 @@ __global__ void min_edge_per_supervertex( vertex_t vertex_color_idx = color_index[tid]; vertex_t vertex_color = color[vertex_color_idx]; edge_t edge_idx = new_mst_edge[tid]; - + // check if valid outgoing edge was found // find minimum edge is same as minimum edge of whole supervertex // if yes, that is part of mst @@ -140,7 +140,7 @@ __global__ void min_edge_per_supervertex( if (vertex_color > dst_color) { add_edge = false; } - } + } } // else { // add_edge = true; From d27fd4d3c84a00568423c4b15d3741d74c10ca7f Mon Sep 17 00:00:00 2001 From: divyegala Date: Wed, 24 Feb 2021 09:45:35 -0800 Subject: [PATCH 11/12] adding non symmetric tests --- cpp/test/mst.cu | 49 +++++++++++++++++++++++++++++-------------------- 1 file changed, 29 insertions(+), 20 deletions(-) diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 59d9a80a86..733d56a7b1 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -112,7 +112,9 @@ template class MSTTest : public ::testing::TestWithParam> { protected: - raft::Graph_COO mst_sequential() { + std::pair, + raft::Graph_COO> + mst_gpu() { edge_t *offsets = static_cast(csr_d.offsets.data()); vertex_t *indices = static_cast(csr_d.indices.data()); weight_t *weights = static_cast(csr_d.weights.data()); @@ -128,22 +130,22 @@ class MSTTest vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); - MST_solver mst_solver( + MST_solver symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), true); - auto result = mst_solver.solve(); - raft::print_device_vector("Final MST Src: ", result.src.data(), - result.n_edges, std::cout); - raft::print_device_vector("Final MST Dst: ", result.dst.data(), - result.n_edges, std::cout); - raft::print_device_vector("Final MST Weights: ", result.weights.data(), - result.n_edges, std::cout); - raft::print_device_vector("Final MST Colors: ", color_ptr, v, std::cout); - - std::cout << "number_of_MST_edges: " << result.n_edges << std::endl; - EXPECT_LE(result.n_edges, 2 * v - 2); - - return result; + auto symmetric_result = symmetric_solver.solve(); + + MST_solver non_symmetric_solver( + handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), + false); + auto non_symmetric_result = non_symmetric_solver.solve(); + + std::cout << "number_of_MST_edges: " << symmetric_result.n_edges + << std::endl; + EXPECT_LE(symmetric_result.n_edges, 2 * v - 2); + + return std::make_pair(std::move(symmetric_result), + std::move(non_symmetric_result)); } void SetUp() override { @@ -216,17 +218,24 @@ const std::vector> csr_in5_h = { typedef MSTTest MSTTestSequential; TEST_P(MSTTestSequential, Sequential) { - auto gpu_result = mst_sequential(); + auto results_pair = mst_gpu(); + auto &symmetric_result = results_pair.first; + auto &non_symmetric_result = results_pair.second; // do assertions here // in this case, running sequential MST auto prims_result = prims(csr_h); - auto parallel_mst_result = - thrust::reduce(thrust::device, gpu_result.weights.data(), - gpu_result.weights.data() + gpu_result.n_edges); + auto symmetric_sum = + thrust::reduce(thrust::device, symmetric_result.weights.data(), + symmetric_result.weights.data() + symmetric_result.n_edges); + auto non_symmetric_sum = thrust::reduce( + thrust::device, non_symmetric_result.weights.data(), + non_symmetric_result.weights.data() + non_symmetric_result.n_edges); - ASSERT_TRUE(raft::match(2 * prims_result, parallel_mst_result, + ASSERT_TRUE(raft::match(2 * prims_result, symmetric_sum, + raft::CompareApprox(0.1))); + ASSERT_TRUE(raft::match(prims_result, non_symmetric_sum, raft::CompareApprox(0.1))); } From bdd9fd41ca82c58e1362229169b6ede29b4f623f Mon Sep 17 00:00:00 2001 From: divyegala Date: Wed, 24 Feb 2021 09:47:18 -0800 Subject: [PATCH 12/12] removing trailing comment --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 2bbcef19e6..d2f86d6dc8 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -142,9 +142,7 @@ __global__ void min_edge_per_supervertex( } } } - // else { - // add_edge = true; - // } + if (add_edge) { temp_src[tid] = tid; temp_dst[tid] = dst;