From dd2eb361c97ed01bff2354ee7e57632edcc3cd69 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 11 Jan 2021 19:48:29 -0500 Subject: [PATCH 1/8] 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 2/8] 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 3/8] 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 4/8] 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 5/8] 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 6/8] 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 7/8] 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 8/8] 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);