From f458e3fa4b551396c58ec7708665644a0ee15ed0 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:38:51 -0500 Subject: [PATCH 1/8] fixed mem leak, cmake for limits, and refactors --- cpp/CMakeLists.txt | 2 +- cpp/include/raft/sparse/mst.cuh | 318 ------------------ .../raft/sparse/mst/detail/mst_kernels.cuh | 145 ++++++++ .../raft/sparse/mst/detail/mst_solver_inl.hpp | 156 +++++++++ cpp/include/raft/sparse/mst/detail/utils.cuh | 41 +++ cpp/include/raft/sparse/mst/mst.cuh | 36 ++ cpp/include/raft/sparse/mst/mst_solver.hpp | 70 ++++ cpp/test/mst.cu | 6 +- 8 files changed, 451 insertions(+), 323 deletions(-) delete mode 100644 cpp/include/raft/sparse/mst.cuh create mode 100644 cpp/include/raft/sparse/mst/detail/mst_kernels.cuh create mode 100644 cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp create mode 100644 cpp/include/raft/sparse/mst/detail/utils.cuh create mode 100644 cpp/include/raft/sparse/mst/mst.cuh create mode 100644 cpp/include/raft/sparse/mst/mst_solver.hpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 25d8641268..7db64f87cd 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -92,7 +92,7 @@ if(OPENMP_FOUND) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") endif(OPENMP_FOUND) -set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-extended-lambda") +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-extended-lambda --expt-relaxed-constexpr") if(${CMAKE_VERSION} VERSION_LESS "3.17.0") set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --std=c++14") diff --git a/cpp/include/raft/sparse/mst.cuh b/cpp/include/raft/sparse/mst.cuh deleted file mode 100644 index 73b048a5ad..0000000000 --- a/cpp/include/raft/sparse/mst.cuh +++ /dev/null @@ -1,318 +0,0 @@ - -/* - * Copyright (c) 2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include -#include -#include -#include -#include -#include -#include -#include - -namespace raft { -namespace mst { - -//FIXME this should live elswhere -template -void printv(rmm::device_vector& vec) { - std::cout.precision(15); - std::cout << "Size = " << vec.size() << std::endl; - thrust::copy(vec.begin(), vec.end(), - std::ostream_iterator(std::cout, " ")); - std::cout << std::endl; -} - -template -class MST_solver { - private: - raft::handle_t const& handle; - - //CSR - const vertex_t* offsets; - const vertex_t* indices; - const weight_t* weights; - const vertex_t v; - const vertex_t e; - - int max_blocks; - int max_threads; - int sm_count; - - rmm::device_vector color; // represent each supervertex as a color - rmm::device_vector next_color; //index of v color in color array - rmm::device_vector active_color; // track active supervertex color - //rmm::device_vector degree; // supervertices degrees - //rmm::device_vector cycle; // edges to be excluded from mst_edge - rmm::device_vector - successor; // current mst iteration. edge being added is (src=i, dst=successor[i]) - rmm::device_vector - mst_edge; // mst output - true if the edge belongs in mst - rmm::device_vector min_edge_color; // minimum incident edge per color - - void label_prop(); - - public: - MST_solver(const raft::handle_t& handle_, vertex_t const* offsets_, - vertex_t const* indices_, weight_t const* weights_, - vertex_t const v_, vertex_t const e_); - - void solve(rmm::device_vector& mst_src, - rmm::device_vector& mst_dst); - - ~MST_solver() {} -}; - -template -MST_solver::MST_solver( - const raft::handle_t& handle_, vertex_t const* offsets_, - vertex_t const* indices_, weight_t const* weights_, vertex_t const v_, - vertex_t const e_) - : handle(handle_), - offsets(offsets_), - indices(indices_), - weights(weights_), - v(v_), - e(e_), - color(v_), - next_color(v_), - active_color(v_), - successor(v_), - mst_edge(e_, false), - min_edge_color(v_, 100) { - max_blocks = handle_.get_device_properties().maxGridSize[0]; - max_threads = handle_.get_device_properties().maxThreadsPerBlock; - sm_count = handle_.get_device_properties().multiProcessorCount; - - //Initially, color holds the vertex id as color - thrust::sequence(color.begin(), color.end()); - //Initially, each next_color redirects to its own color - thrust::sequence(next_color.begin(), next_color.end()); - //Initially, each edge is not in the mst -} - -template -__global__ void kernel_min_edge_per_vertex(const vertex_t* offsets, - const edge_t* indices, - const weight_t* weights, - vertex_t* color, vertex_t* successor, - bool* mst_edge, const vertex_t v) { - edge_t tid = threadIdx.x + blockIdx.x * blockDim.x; - - unsigned warp_id = tid / 32; - unsigned lane_id = tid % 32; - - __shared__ edge_t min_edge_index[32]; - __shared__ weight_t min_edge_weight[32]; - __shared__ vertex_t min_color[32]; - - // min_edge_index[lane_id] = std::numeric_limits::max(); - // min_edge_weight[lane_id] = std::numeric_limits::max(); - // min_color[lane_id] = std::numeric_limits::max(); - - // TODO: Find a way to set limits - // Above does not work as it is host code - min_edge_index[lane_id] = 100; - min_edge_weight[lane_id] = 100; - min_color[lane_id] = 100; - __syncthreads(); - - vertex_t self_color = color[tid]; - - if (warp_id < v) { - // one row is associated with one warp - edge_t row_start = offsets[warp_id]; - edge_t row_end = offsets[warp_id + 1]; - - // assuming one warp per row - // find min for each thread in warp - for (edge_t e = row_start + lane_id; e < row_end; e += 32) { - weight_t curr_edge_weight = weights[e]; - vertex_t successor_color = color[indices[e]]; - - if (!mst_edge[e] && self_color != successor_color) { - if (curr_edge_weight < min_edge_weight[lane_id]) { - min_color[lane_id] = successor_color; - min_edge_weight[lane_id] = curr_edge_weight; - min_edge_index[lane_id] = e; - // theta = abs(curr_edge_weight - min_edge_weight[lane_id]); - } else if (curr_edge_weight == min_edge_weight[lane_id]) { - // tie break - if (min_color[lane_id] > successor_color) { - min_color[lane_id] = successor_color; - min_edge_weight[lane_id] = curr_edge_weight; - min_edge_index[lane_id] = e; - } - } - } - } - } - __syncthreads(); - - // reduce across threads in warp - for (int offset = 16; offset > 0; offset >>= 1) { - if (lane_id < offset) { - if (min_edge_weight[lane_id] > min_edge_weight[lane_id + offset]) { - min_color[lane_id] = min_color[lane_id + offset]; - min_edge_weight[lane_id] = min_edge_weight[lane_id + offset]; - min_edge_index[lane_id] = min_edge_index[lane_id + offset]; - } else if (min_edge_weight[lane_id] == - min_edge_weight[lane_id + offset]) { - if (min_color[lane_id] > min_color[lane_id + offset]) { - min_color[lane_id] = min_color[lane_id + offset]; - min_edge_weight[lane_id] = min_edge_weight[lane_id + offset]; - min_edge_index[lane_id] = min_edge_index[lane_id + offset]; - } - } - } - __syncthreads(); - } - - // min edge may now be found in first thread - if (lane_id == 0) { - if (min_edge_weight[0] != 100) { - successor[warp_id] = indices[min_edge_index[0]]; - } - } -} - -// TODO make this work in 64bit -__device__ int get_1D_idx() { return blockIdx.x * blockDim.x + threadIdx.x; } - -// executes for each vertex and updates the colors of both vertices to the lower color -template -__global__ void min_pair_colors(const vertex_t v, const vertex_t* successor, - vertex_t* color, vertex_t* next_color) { - int i = get_1D_idx(); - if (i < v) { - atomicMin(&next_color[i], color[successor[i]]); - atomicMin(&next_color[successor[i]], color[i]); - } -} - -template -__global__ void check_color_change(const vertex_t v, vertex_t* color, - vertex_t* next_color, bool* done) { - //This kernel works on the global_colors[] array - int i = get_1D_idx(); - if (i < v) { - if (color[i] > next_color[i]) { - //Termination for label propagation - done[0] = false; - color[i] = next_color[i]; - } - } - // Notice that some degree >1 and we run in parallel - // min_pair_colors kernel may result in pair color inconsitencies - // resolving here for next iteration - // TODO check experimentally - next_color[i] = color[i]; -} - -template -void MST_solver::label_prop() { - // update the colors of both ends its until there is no change in colors - int nthreads = std::min(v, max_threads); - int nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); - auto stream = handle.get_stream(); - - rmm::device_vector done(1, false); - vertex_t* color_ptr = thrust::raw_pointer_cast(color.data()); - vertex_t* next_color_ptr = thrust::raw_pointer_cast(next_color.data()); - vertex_t* successor_ptr = thrust::raw_pointer_cast(successor.data()); - - bool* done_ptr = thrust::raw_pointer_cast(done.data()); - - auto i = 0; - std::cout << "==================" << std::endl; - printv(color); - while (!done[0]) { - done[0] = true; - min_pair_colors<<>>( - v, successor_ptr, color_ptr, next_color_ptr); - printv(next_color); - check_color_change<<>>( - v, color_ptr, next_color_ptr, done_ptr); - printv(color); - i++; - } - std::cout << "Label prop iterations : " << i << std::endl; -} - -template -void MST_solver::solve( - rmm::device_vector& mst_src, - rmm::device_vector& mst_dst) { - RAFT_EXPECTS(v > 0, "0 vertices"); - RAFT_EXPECTS(e > 0, "0 edges"); - RAFT_EXPECTS(offsets != nullptr, "Null offsets."); - RAFT_EXPECTS(indices != nullptr, "Null indices."); - RAFT_EXPECTS(weights != nullptr, "Null weights."); - - auto stream = handle.get_stream(); - - kernel_min_edge_per_vertex<<>>( - offsets, indices, weights, thrust::raw_pointer_cast(color.data()), - thrust::raw_pointer_cast(successor.data()), - thrust::raw_pointer_cast(mst_edge.data()), v); - - printv(successor); - - label_prop(); - - printv(color); - - // Theorem : the minimum incident edge to any vertex has to be in the MST - // This is a segmented min scan/reduce - // cub::KeyValuePair* d_out = nullptr; - // void* cub_temp_storage = nullptr; - // size_t cub_temp_storage_bytes = 0; - // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, - // weights, d_out, v, offsets, offsets + 1); - // // FIXME RMM Allocate temporary storage - // cudaMalloc(&cub_temp_storage, cub_temp_storage_bytes); - // // Run argmin-reduction - // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, - // weights, d_out, v, offsets, offsets + 1); - // - // TODO: mst[offset[i]+key[i]]=true; (thrust)? - // Extract MST edge list by just filtering with the mask generated above? - - // bool mst_edge_found = true; - // Boruvka original formulation says "while more than 1 supervertex remains" - // Here we adjust it to support disconnected components (spanning forest) - // track completion with mst_edge_found status. - // should have max_iter ensure it always exits. - // for (auto i = 0; i < v; i++) { - // { - // updates colors of supervertices by propagating the lower color to the higher - // TODO - - // Finds the minimum outgoing edge from each supervertex to the lowest outgoing color - // by working at each vertex of the supervertex - // TODO - // segmented min with an extra check to discard edges leading to the same color - - // filter internal edges / remove cycles - // TODO - - // done - // if (!mst_edge_found) break; - // } - // } -} -} // namespace mst -} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh new file mode 100644 index 0000000000..061f7b4530 --- /dev/null +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -0,0 +1,145 @@ + +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "utils.cuh" + +#include + +namespace raft { +namespace mst { +namespace detail { + +template +__global__ void kernel_min_edge_per_color(const vertex_t* offsets, + const edge_t* indices, + const weight_t* weights, + vertex_t* color, vertex_t* successor, + bool* mst_edge, const vertex_t v) { + edge_t tid = threadIdx.x + blockIdx.x * blockDim.x; + + unsigned warp_id = tid / 32; + unsigned lane_id = tid % 32; + + __shared__ edge_t min_edge_index[32]; + __shared__ weight_t min_edge_weight[32]; + __shared__ vertex_t min_color[32]; + + min_edge_index[lane_id] = std::numeric_limits::max(); + min_edge_weight[lane_id] = std::numeric_limits::max(); + min_color[lane_id] = std::numeric_limits::max(); + + // TODO: Find a way to set limits + // Above does not work as it is host code + // min_edge_index[lane_id] = 100; + // min_edge_weight[lane_id] = 100; + // min_color[lane_id] = 100; + __syncthreads(); + + vertex_t self_color = color[warp_id]; + + if (warp_id < v) { + // one row is associated with one warp + edge_t row_start = offsets[warp_id]; + edge_t row_end = offsets[warp_id + 1]; + + // assuming one warp per row + // find min for each thread in warp + for (edge_t e = row_start + lane_id; e < row_end; e += 32) { + weight_t curr_edge_weight = weights[e]; + vertex_t successor_color = color[indices[e]]; + + if (!mst_edge[e] && self_color != successor_color) { + if (curr_edge_weight < min_edge_weight[lane_id]) { + min_color[lane_id] = successor_color; + min_edge_weight[lane_id] = curr_edge_weight; + min_edge_index[lane_id] = e; + // theta = abs(curr_edge_weight - min_edge_weight[lane_id]); + } else if (curr_edge_weight == min_edge_weight[lane_id]) { + // tie break + if (min_color[lane_id] > successor_color) { + min_color[lane_id] = successor_color; + min_edge_weight[lane_id] = curr_edge_weight; + min_edge_index[lane_id] = e; + } + } + } + } + } + __syncthreads(); + + // reduce across threads in warp + for (int offset = 16; offset > 0; offset >>= 1) { + if (lane_id < offset) { + if (min_edge_weight[lane_id] > min_edge_weight[lane_id + offset]) { + min_color[lane_id] = min_color[lane_id + offset]; + min_edge_weight[lane_id] = min_edge_weight[lane_id + offset]; + min_edge_index[lane_id] = min_edge_index[lane_id + offset]; + } else if (min_edge_weight[lane_id] == + min_edge_weight[lane_id + offset]) { + if (min_color[lane_id] > min_color[lane_id + offset]) { + min_color[lane_id] = min_color[lane_id + offset]; + min_edge_weight[lane_id] = min_edge_weight[lane_id + offset]; + min_edge_index[lane_id] = min_edge_index[lane_id + offset]; + } + } + } + __syncthreads(); + } + + // min edge may now be found in first thread + if (lane_id == 0) { + if (min_edge_weight[0] != std::numeric_limits::max()) { + successor[warp_id] = indices[min_edge_index[0]]; + } + } +} + +// executes for each vertex and updates the colors of both vertices to the lower color +template +__global__ void min_pair_colors(const vertex_t v, const vertex_t* successor, + vertex_t* color, vertex_t* next_color) { + int i = get_1D_idx(); + if (i < v) { + atomicMin(&next_color[i], color[successor[i]]); + atomicMin(&next_color[successor[i]], color[i]); + } +} + +template +__global__ void check_color_change(const vertex_t v, vertex_t* color, + vertex_t* next_color, bool* done) { + //This kernel works on the global_colors[] array + int i = get_1D_idx(); + if (i < v) { + if (color[i] > next_color[i]) { + //Termination for label propagation + done[0] = false; + color[i] = next_color[i]; + } + } + // Notice that some degree >1 and we run in parallel + // min_pair_colors kernel may result in pair color inconsitencies + // resolving here for next iteration + // TODO check experimentally + next_color[i] = color[i]; +} + +} // namespace detail +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp new file mode 100644 index 0000000000..d735b64f4e --- /dev/null +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp @@ -0,0 +1,156 @@ +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "mst_kernels.cuh" +#include "utils.cuh" + +namespace raft { +namespace mst { + +template +MST_solver::MST_solver( + const raft::handle_t& handle_, vertex_t const* offsets_, + vertex_t const* indices_, weight_t const* weights_, vertex_t const v_, + vertex_t const e_) + : handle(handle_), + offsets(offsets_), + indices(indices_), + weights(weights_), + v(v_), + e(e_), + color(v_), + next_color(v_), + active_color(v_), + successor(v_), + mst_edge(e_, false) { + max_blocks = handle_.get_device_properties().maxGridSize[0]; + max_threads = handle_.get_device_properties().maxThreadsPerBlock; + sm_count = handle_.get_device_properties().multiProcessorCount; + + //Initially, color holds the vertex id as color + thrust::sequence(color.begin(), color.end()); + //Initially, each next_color redirects to its own color + thrust::sequence(next_color.begin(), next_color.end()); + //Initially, each edge is not in the mst +} + +template +void MST_solver::solve() { + RAFT_EXPECTS(v > 0, "0 vertices"); + RAFT_EXPECTS(e > 0, "0 edges"); + RAFT_EXPECTS(offsets != nullptr, "Null offsets."); + RAFT_EXPECTS(indices != nullptr, "Null indices."); + RAFT_EXPECTS(weights != nullptr, "Null weights."); + + min_edge_per_color(); + + detail::printv(successor); + + label_prop(); + + detail::printv(color); + + // Theorem : the minimum incident edge to any vertex has to be in the MST + // This is a segmented min scan/reduce + // cub::KeyValuePair* d_out = nullptr; + // void* cub_temp_storage = nullptr; + // size_t cub_temp_storage_bytes = 0; + // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, + // weights, d_out, v, offsets, offsets + 1); + // // FIXME RMM Allocate temporary storage + // cudaMalloc(&cub_temp_storage, cub_temp_storage_bytes); + // // Run argmin-reduction + // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, + // weights, d_out, v, offsets, offsets + 1); + // + // TODO: mst[offset[i]+key[i]]=true; (thrust)? + // Extract MST edge list by just filtering with the mask generated above? + + // bool mst_edge_found = true; + // Boruvka original formulation says "while more than 1 supervertex remains" + // Here we adjust it to support disconnected components (spanning forest) + // track completion with mst_edge_found status. + // should have max_iter ensure it always exits. + // for (auto i = 0; i < v; i++) { + // { + // updates colors of supervertices by propagating the lower color to the higher + // TODO + + // Finds the minimum outgoing edge from each supervertex to the lowest outgoing color + // by working at each vertex of the supervertex + // TODO + // segmented min with an extra check to discard edges leading to the same color + + // filter internal edges / remove cycles + // TODO + + // done + // if (!mst_edge_found) break; + // } + // } +} + +template +void MST_solver::label_prop() { + // update the colors of both ends its until there is no change in colors + int nthreads = std::min(v, max_threads); + int nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); + auto stream = handle.get_stream(); + + rmm::device_vector done(1, false); + vertex_t* color_ptr = thrust::raw_pointer_cast(color.data()); + vertex_t* next_color_ptr = thrust::raw_pointer_cast(next_color.data()); + vertex_t* successor_ptr = thrust::raw_pointer_cast(successor.data()); + + bool* done_ptr = thrust::raw_pointer_cast(done.data()); + + auto i = 0; + std::cout << "==================" << std::endl; + detail::printv(color); + while (!done[0]) { + done[0] = true; + detail::min_pair_colors<<>>( + v, successor_ptr, color_ptr, next_color_ptr); + detail::printv(next_color); + detail::check_color_change<<>>( + v, color_ptr, next_color_ptr, done_ptr); + detail::printv(color); + i++; + } + std::cout << "Label prop iterations : " << i << std::endl; +} + +template +void MST_solver::min_edge_per_color() { + + auto stream = handle.get_stream(); + int n_threads = 32; + + vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); + vertex_t *successor_ptr = thrust::raw_pointer_cast(successor.data()); + bool *mst_edge_ptr = thrust::raw_pointer_cast(mst_edge.data()); + + detail::kernel_min_edge_per_color<<>>( + offsets, indices, weights, color_ptr, + successor_ptr, + mst_edge_ptr, v); + +} + +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/detail/utils.cuh b/cpp/include/raft/sparse/mst/detail/utils.cuh new file mode 100644 index 0000000000..aa5aec43ce --- /dev/null +++ b/cpp/include/raft/sparse/mst/detail/utils.cuh @@ -0,0 +1,41 @@ + +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace raft { +namespace mst { +namespace detail { + +// TODO make this work in 64bit +__device__ int get_1D_idx() { return blockIdx.x * blockDim.x + threadIdx.x; } + +//FIXME this should live elswhere +template +void printv(rmm::device_vector& vec) { + std::cout.precision(15); + std::cout << "Size = " << vec.size() << std::endl; + thrust::copy(vec.begin(), vec.end(), + std::ostream_iterator(std::cout, " ")); + std::cout << std::endl; +} + +} // namespace detail +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh new file mode 100644 index 0000000000..cd2deb81c1 --- /dev/null +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -0,0 +1,36 @@ + +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "mst_solver.hpp" + +namespace raft { +namespace mst { + +template +void mst(const raft::handle_t& handle, vertex_t const* offsets, + vertex_t const* indices, weight_t const* weights, + vertex_t const v, vertex_t const e) { + + MST_solver mst_solver(handle, offsets, indices, weights, v, e); + mst_solver.solve(); + +} + +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/mst_solver.hpp b/cpp/include/raft/sparse/mst/mst_solver.hpp new file mode 100644 index 0000000000..d39c0863ca --- /dev/null +++ b/cpp/include/raft/sparse/mst/mst_solver.hpp @@ -0,0 +1,70 @@ + +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +namespace raft { +namespace mst { + +template +class MST_solver { + + public: + MST_solver(const raft::handle_t& handle_, vertex_t const* offsets_, + vertex_t const* indices_, weight_t const* weights_, + vertex_t const v_, vertex_t const e_); + + void solve(); + + ~MST_solver() {} + + private: + raft::handle_t const& handle; + + //CSR + const vertex_t* offsets; + const vertex_t* indices; + const weight_t* weights; + const vertex_t v; + const vertex_t e; + + int max_blocks; + int max_threads; + int sm_count; + + rmm::device_vector color; // represent each supervertex as a color + rmm::device_vector next_color; //index of v color in color array + rmm::device_vector active_color; // track active supervertex color + //rmm::device_vector degree; // supervertices degrees + //rmm::device_vector cycle; // edges to be excluded from mst_edge + rmm::device_vector + successor; // current mst iteration. edge being added is (src=i, dst=successor[i]) + rmm::device_vector + mst_edge; // mst output - true if the edge belongs in mst + rmm::device_vector min_edge_color; // minimum incident edge per color + + void label_prop(); + void min_edge_per_color(); +}; + +} // namespace mst +} // namespace raft + +#include "detail/mst_solver_inl.hpp" \ No newline at end of file diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 912c8dcf19..d5d1af79a3 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -25,7 +25,7 @@ #include #include -#include +#include template struct CSRHost { @@ -125,10 +125,8 @@ class MSTTest auto v = static_cast((csr_d.offsets.size() / sizeof(value_t)) - 1); auto e = static_cast(csr_d.indices.size() / sizeof(edge_t)); - MST_solver solver(handle, offsets, indices, weights, v, e); + mst(handle, offsets, indices, weights, v, e); - //nullptr expected to trigger exceptions - solver.solve(mst_src, mst_dst); } void SetUp() override { From a4c5fbf1be5db7d10442c20bbe97ce6f35265615 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:41:03 -0500 Subject: [PATCH 2/8] style refactors --- .../raft/sparse/mst/detail/mst_kernels.cuh | 16 ++++++++-------- .../raft/sparse/mst/detail/mst_solver_inl.hpp | 16 ++++++---------- cpp/include/raft/sparse/mst/detail/utils.cuh | 6 +++--- cpp/include/raft/sparse/mst/mst.cuh | 9 ++++----- cpp/include/raft/sparse/mst/mst_solver.hpp | 5 ++--- cpp/test/mst.cu | 10 +++++----- 6 files changed, 28 insertions(+), 34 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 061f7b4530..563bf94da9 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -20,17 +20,17 @@ #include "utils.cuh" #include - + namespace raft { namespace mst { namespace detail { template __global__ void kernel_min_edge_per_color(const vertex_t* offsets, - const edge_t* indices, - const weight_t* weights, - vertex_t* color, vertex_t* successor, - bool* mst_edge, const vertex_t v) { + const edge_t* indices, + const weight_t* weights, + vertex_t* color, vertex_t* successor, + bool* mst_edge, const vertex_t v) { edge_t tid = threadIdx.x + blockIdx.x * blockDim.x; unsigned warp_id = tid / 32; @@ -140,6 +140,6 @@ __global__ void check_color_change(const vertex_t v, vertex_t* color, next_color[i] = color[i]; } -} // namespace detail -} // namespace mst -} // namespace raft \ No newline at end of file +} // namespace detail +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp index d735b64f4e..80e3aaf70c 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp @@ -137,20 +137,16 @@ void MST_solver::label_prop() { template void MST_solver::min_edge_per_color() { - auto stream = handle.get_stream(); int n_threads = 32; - vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); - vertex_t *successor_ptr = thrust::raw_pointer_cast(successor.data()); - bool *mst_edge_ptr = thrust::raw_pointer_cast(mst_edge.data()); + vertex_t* color_ptr = thrust::raw_pointer_cast(color.data()); + vertex_t* successor_ptr = thrust::raw_pointer_cast(successor.data()); + bool* mst_edge_ptr = thrust::raw_pointer_cast(mst_edge.data()); detail::kernel_min_edge_per_color<<>>( - offsets, indices, weights, color_ptr, - successor_ptr, - mst_edge_ptr, v); - + offsets, indices, weights, color_ptr, successor_ptr, mst_edge_ptr, v); } -} // namespace mst -} // namespace raft \ No newline at end of file +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/detail/utils.cuh b/cpp/include/raft/sparse/mst/detail/utils.cuh index aa5aec43ce..9d64e71d24 100644 --- a/cpp/include/raft/sparse/mst/detail/utils.cuh +++ b/cpp/include/raft/sparse/mst/detail/utils.cuh @@ -36,6 +36,6 @@ void printv(rmm::device_vector& vec) { std::cout << std::endl; } -} // namespace detail -} // namespace mst -} // namespace raft \ No newline at end of file +} // namespace detail +} // namespace mst +} // namespace raft \ No newline at end of file diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index cd2deb81c1..d8156c5a8f 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -24,12 +24,11 @@ namespace mst { template void mst(const raft::handle_t& handle, vertex_t const* offsets, - vertex_t const* indices, weight_t const* weights, - vertex_t const v, vertex_t const e) { - - MST_solver mst_solver(handle, offsets, indices, weights, v, e); + vertex_t const* indices, weight_t const* weights, vertex_t const v, + vertex_t const e) { + MST_solver mst_solver(handle, offsets, indices, + weights, v, e); mst_solver.solve(); - } } // namespace mst diff --git a/cpp/include/raft/sparse/mst/mst_solver.hpp b/cpp/include/raft/sparse/mst/mst_solver.hpp index d39c0863ca..f7e24d80d8 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.hpp +++ b/cpp/include/raft/sparse/mst/mst_solver.hpp @@ -25,7 +25,6 @@ namespace mst { template class MST_solver { - public: MST_solver(const raft::handle_t& handle_, vertex_t const* offsets_, vertex_t const* indices_, weight_t const* weights_, @@ -64,7 +63,7 @@ class MST_solver { void min_edge_per_color(); }; -} // namespace mst -} // namespace raft +} // namespace mst +} // namespace raft #include "detail/mst_solver_inl.hpp" \ No newline at end of file diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index d5d1af79a3..fa0540c6d9 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -118,15 +118,15 @@ class MSTTest rmm::device_vector mst_src; rmm::device_vector mst_dst; - vertex_t *offsets = static_cast(csr_d.offsets.data()); - edge_t *indices = static_cast(csr_d.indices.data()); - value_t *weights = static_cast(csr_d.weights.data()); + vertex_t *offsets = static_cast(csr_d.offsets.data()); + edge_t *indices = static_cast(csr_d.indices.data()); + value_t *weights = static_cast(csr_d.weights.data()); - auto v = static_cast((csr_d.offsets.size() / sizeof(value_t)) - 1); + auto v = + static_cast((csr_d.offsets.size() / sizeof(value_t)) - 1); auto e = static_cast(csr_d.indices.size() / sizeof(edge_t)); mst(handle, offsets, indices, weights, v, e); - } void SetUp() override { From b39ea589eca11f0e3f7d8edf28f79b0d0d72562c Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:42:18 -0500 Subject: [PATCH 3/8] extra line at EOF --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 2 +- cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp | 2 +- cpp/include/raft/sparse/mst/detail/utils.cuh | 2 +- cpp/include/raft/sparse/mst/mst.cuh | 2 +- cpp/include/raft/sparse/mst/mst_solver.hpp | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 563bf94da9..27fc29cea4 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -142,4 +142,4 @@ __global__ void check_color_change(const vertex_t v, vertex_t* color, } // namespace detail } // namespace mst -} // namespace raft \ No newline at end of file +} // namespace raft diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp index 80e3aaf70c..ec6a7310c7 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp @@ -149,4 +149,4 @@ void MST_solver::min_edge_per_color() { } } // namespace mst -} // namespace raft \ No newline at end of file +} // namespace raft diff --git a/cpp/include/raft/sparse/mst/detail/utils.cuh b/cpp/include/raft/sparse/mst/detail/utils.cuh index 9d64e71d24..7255f105e3 100644 --- a/cpp/include/raft/sparse/mst/detail/utils.cuh +++ b/cpp/include/raft/sparse/mst/detail/utils.cuh @@ -38,4 +38,4 @@ void printv(rmm::device_vector& vec) { } // namespace detail } // namespace mst -} // namespace raft \ No newline at end of file +} // namespace raft diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index d8156c5a8f..55f85f4996 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -32,4 +32,4 @@ void mst(const raft::handle_t& handle, vertex_t const* offsets, } } // namespace mst -} // namespace raft \ No newline at end of file +} // namespace raft diff --git a/cpp/include/raft/sparse/mst/mst_solver.hpp b/cpp/include/raft/sparse/mst/mst_solver.hpp index f7e24d80d8..cd4d20ed66 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.hpp +++ b/cpp/include/raft/sparse/mst/mst_solver.hpp @@ -66,4 +66,4 @@ class MST_solver { } // namespace mst } // namespace raft -#include "detail/mst_solver_inl.hpp" \ No newline at end of file +#include "detail/mst_solver_inl.hpp" From aa4fe7a10182b993fc8e222851842383cd11dbe2 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:45:39 -0500 Subject: [PATCH 4/8] consistency in template params --- cpp/test/mst.cu | 34 +++++++++++++++++----------------- 1 file changed, 17 insertions(+), 17 deletions(-) diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index fa0540c6d9..b04540a004 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -27,14 +27,14 @@ #include -template +template struct CSRHost { std::vector offsets; std::vector indices; - std::vector weights; + std::vector weights; }; -template +template struct CSRDevice { rmm::device_buffer offsets; rmm::device_buffer indices; @@ -46,13 +46,13 @@ namespace mst { // Sequential prims function // Returns total weight of MST -template -value_t prims(CSRHost &csr_h) { +template +weight_t prims(CSRHost &csr_h) { auto n_vertices = csr_h.offsets.size() - 1; bool active_vertex[n_vertices]; // bool mst_set[csr_h.n_edges]; - value_t curr_edge[n_vertices]; + weight_t curr_edge[n_vertices]; for (auto i = 0; i < n_vertices; i++) { active_vertex[i] = false; @@ -67,7 +67,7 @@ value_t prims(CSRHost &csr_h) { // function to pick next min vertex-edge auto min_vertex_edge = [](auto *curr_edge, auto *active_vertex, auto n_vertices) { - value_t min = INT_MAX; + weight_t min = INT_MAX; vertex_t min_vertex; for (auto v = 0; v < n_vertices; v++) { @@ -102,7 +102,7 @@ value_t prims(CSRHost &csr_h) { } // find sum of MST - value_t total_weight = 0; + weight_t total_weight = 0; for (auto v = 1; v < n_vertices; v++) { total_weight += curr_edge[v]; } @@ -110,9 +110,9 @@ value_t prims(CSRHost &csr_h) { return total_weight; } -template +template class MSTTest - : public ::testing::TestWithParam> { + : public ::testing::TestWithParam> { protected: void mst_sequential() { rmm::device_vector mst_src; @@ -120,32 +120,32 @@ class MSTTest vertex_t *offsets = static_cast(csr_d.offsets.data()); edge_t *indices = static_cast(csr_d.indices.data()); - value_t *weights = static_cast(csr_d.weights.data()); + weight_t *weights = static_cast(csr_d.weights.data()); auto v = - static_cast((csr_d.offsets.size() / sizeof(value_t)) - 1); + static_cast((csr_d.offsets.size() / sizeof(weight_t)) - 1); auto e = static_cast(csr_d.indices.size() / sizeof(edge_t)); - mst(handle, offsets, indices, weights, v, e); + mst(handle, offsets, indices, weights, v, e); } void SetUp() override { csr_h = - ::testing::TestWithParam>::GetParam(); + ::testing::TestWithParam>::GetParam(); csr_d.offsets = rmm::device_buffer(csr_h.offsets.data(), csr_h.offsets.size() * sizeof(vertex_t)); csr_d.indices = rmm::device_buffer(csr_h.indices.data(), csr_h.indices.size() * sizeof(edge_t)); csr_d.weights = rmm::device_buffer(csr_h.weights.data(), - csr_h.weights.size() * sizeof(value_t)); + csr_h.weights.size() * sizeof(weight_t)); } void TearDown() override {} protected: - CSRHost csr_h; - CSRDevice csr_d; + CSRHost csr_h; + CSRDevice csr_d; raft::handle_t handle; }; From fc77980950944eb8eb3e1749a89000358c29adc3 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:48:15 -0500 Subject: [PATCH 5/8] float weights in test --- cpp/test/mst.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index b04540a004..e760ca8776 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -159,11 +159,11 @@ Graph 1: */ -const std::vector> csr_in_h = { +const std::vector> csr_in_h = { // {nullptr, nullptr, nullptr, 0, 0}, {{0, 3, 5, 7, 8}, {1, 2, 3, 0, 3, 0, 0, 1}, {2, 3, 4, 2, 1, 3, 4, 1}}}; -typedef MSTTest MSTTestSequential; +typedef MSTTest MSTTestSequential; TEST_P(MSTTestSequential, Sequential) { mst_sequential(); From 0000c9436c958f85d9db1022bb56e2ac63f6ea28 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 27 Oct 2020 13:56:21 -0500 Subject: [PATCH 6/8] better naming --- cpp/include/raft/sparse/mst/detail/mst_kernels.cuh | 10 +++++----- cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp | 6 +++--- cpp/include/raft/sparse/mst/mst_solver.hpp | 2 +- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 27fc29cea4..c997c2819b 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -26,11 +26,11 @@ namespace mst { namespace detail { template -__global__ void kernel_min_edge_per_color(const vertex_t* offsets, - const edge_t* indices, - const weight_t* weights, - vertex_t* color, vertex_t* successor, - bool* mst_edge, const vertex_t v) { +__global__ void kernel_min_edge_per_vertex(const vertex_t* offsets, + const edge_t* indices, + const weight_t* weights, + vertex_t* color, vertex_t* successor, + bool* mst_edge, const vertex_t v) { edge_t tid = threadIdx.x + blockIdx.x * blockDim.x; unsigned warp_id = tid / 32; diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp index ec6a7310c7..fc8825b211 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp @@ -57,7 +57,7 @@ void MST_solver::solve() { RAFT_EXPECTS(indices != nullptr, "Null indices."); RAFT_EXPECTS(weights != nullptr, "Null weights."); - min_edge_per_color(); + min_edge_per_vertex(); detail::printv(successor); @@ -136,7 +136,7 @@ void MST_solver::label_prop() { } template -void MST_solver::min_edge_per_color() { +void MST_solver::min_edge_per_vertex() { auto stream = handle.get_stream(); int n_threads = 32; @@ -144,7 +144,7 @@ void MST_solver::min_edge_per_color() { vertex_t* successor_ptr = thrust::raw_pointer_cast(successor.data()); bool* mst_edge_ptr = thrust::raw_pointer_cast(mst_edge.data()); - detail::kernel_min_edge_per_color<<>>( + detail::kernel_min_edge_per_vertex<<>>( offsets, indices, weights, color_ptr, successor_ptr, mst_edge_ptr, v); } diff --git a/cpp/include/raft/sparse/mst/mst_solver.hpp b/cpp/include/raft/sparse/mst/mst_solver.hpp index cd4d20ed66..7071449047 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.hpp +++ b/cpp/include/raft/sparse/mst/mst_solver.hpp @@ -60,7 +60,7 @@ class MST_solver { rmm::device_vector min_edge_color; // minimum incident edge per color void label_prop(); - void min_edge_per_color(); + void min_edge_per_vertex(); }; } // namespace mst From 01c671a1890775ecf46f8e6619e604b9699e5838 Mon Sep 17 00:00:00 2001 From: Alex Fender Date: Tue, 27 Oct 2020 17:18:21 -0500 Subject: [PATCH 7/8] Update mst.cu --- cpp/test/mst.cu | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index e760ca8776..62bda18151 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -160,9 +160,14 @@ Graph 1: */ const std::vector> csr_in_h = { - // {nullptr, nullptr, nullptr, 0, 0}, {{0, 3, 5, 7, 8}, {1, 2, 3, 0, 3, 0, 0, 1}, {2, 3, 4, 2, 1, 3, 4, 1}}}; +const std::vector> csr_in2_h = { + {{0, 4, 6, 9, 12, 15, 17, 20}, + {2, 4, 5, 6, 3, 6, 0, 4, 5, 1, 4, 6, 0, 2, 3, 0, 2, 0, 1, 3}, + {5.0, 9.0, 1.0, 4.0, 8.0, 7.0, 5.0, 2.0, 6.0, 8.0, + 3.0, 4.0, 9.0, 2.0, 3.0, 1.0, 6.0, 4.0, 7.0, 10.0}}}; + typedef MSTTest MSTTestSequential; TEST_P(MSTTestSequential, Sequential) { mst_sequential(); From 1f547dc25fdd6b224a742b9b6403339a779d6ae8 Mon Sep 17 00:00:00 2001 From: afender Date: Tue, 27 Oct 2020 18:17:26 -0500 Subject: [PATCH 8/8] looping! --- .../raft/sparse/mst/detail/mst_solver_inl.hpp | 52 ++++--------------- cpp/test/mst.cu | 2 +- 2 files changed, 11 insertions(+), 43 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp index fc8825b211..43eb5a9aab 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.hpp @@ -57,52 +57,19 @@ void MST_solver::solve() { RAFT_EXPECTS(indices != nullptr, "Null indices."); RAFT_EXPECTS(weights != nullptr, "Null weights."); - min_edge_per_vertex(); - - detail::printv(successor); - - label_prop(); - - detail::printv(color); - - // Theorem : the minimum incident edge to any vertex has to be in the MST - // This is a segmented min scan/reduce - // cub::KeyValuePair* d_out = nullptr; - // void* cub_temp_storage = nullptr; - // size_t cub_temp_storage_bytes = 0; - // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, - // weights, d_out, v, offsets, offsets + 1); - // // FIXME RMM Allocate temporary storage - // cudaMalloc(&cub_temp_storage, cub_temp_storage_bytes); - // // Run argmin-reduction - // cub::DeviceSegmentedReduce::ArgMin(cub_temp_storage, cub_temp_storage_bytes, - // weights, d_out, v, offsets, offsets + 1); - // - // TODO: mst[offset[i]+key[i]]=true; (thrust)? - // Extract MST edge list by just filtering with the mask generated above? - - // bool mst_edge_found = true; // Boruvka original formulation says "while more than 1 supervertex remains" // Here we adjust it to support disconnected components (spanning forest) // track completion with mst_edge_found status. // should have max_iter ensure it always exits. - // for (auto i = 0; i < v; i++) { - // { - // updates colors of supervertices by propagating the lower color to the higher - // TODO - - // Finds the minimum outgoing edge from each supervertex to the lowest outgoing color - // by working at each vertex of the supervertex - // TODO - // segmented min with an extra check to discard edges leading to the same color - - // filter internal edges / remove cycles - // TODO - - // done - // if (!mst_edge_found) break; - // } - // } + for (auto i = 0; i < 2; i++) { + // Finds the minimum outgoing edge from each supervertex to the lowest outgoing color + // by working at each vertex of the supervertex + min_edge_per_vertex(); + detail::printv(successor); + // updates colors of supervertices by propagating the lower color to the higher + label_prop(); + detail::printv(color); + } } template @@ -133,6 +100,7 @@ void MST_solver::label_prop() { i++; } std::cout << "Label prop iterations : " << i << std::endl; + std::cout << "==================" << std::endl; } template diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 62bda18151..43bd62a220 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -178,7 +178,7 @@ TEST_P(MSTTestSequential, Sequential) { } INSTANTIATE_TEST_SUITE_P(MSTTests, MSTTestSequential, - ::testing::ValuesIn(csr_in_h)); + ::testing::ValuesIn(csr_in2_h)); } // namespace mst } // namespace raft