From 046f703d2a4213feca9464bff3bf5cb3d56cb146 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 20 Apr 2021 17:15:17 -0400 Subject: [PATCH 01/31] Allowing epilogue in knn graph connection function --- .../sparse/hierarchy/detail/agglomerative.cuh | 4 ++ .../raft/sparse/hierarchy/detail/mst.cuh | 51 ++++++++----------- cpp/include/raft/sparse/op/sort.h | 23 +++++++++ 3 files changed, 49 insertions(+), 29 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index e94fdfb970..be326a83db 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -113,12 +113,14 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, std::vector children_h(n_edges * 2); std::vector out_size_h(n_edges); + std::vector out_delta_h(n_edges); UnionFind U(nnz + 1); for (value_idx i = 0; i < nnz; i++) { value_idx a = mst_src_h[i]; value_idx b = mst_dst_h[i]; + value_t delta = mst_weights_h[i]; value_idx aa = U.find(a); value_idx bb = U.find(b); @@ -127,6 +129,7 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, children_h[children_idx] = aa; children_h[children_idx + 1] = bb; + out_delta_h[i] = delta; out_size_h[i] = U.size[aa] + U.size[bb]; U.perform_union(aa, bb); @@ -138,6 +141,7 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, raft::update_device(children, children_h.data(), n_edges * 2, stream); raft::update_device(out_size.data(), out_size_h.data(), n_edges, stream); + raft::update_device(out_delta.data(), out_delta_h.data(), n_edges, stream); } /** diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 91c4e1642d..387b39fe14 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -21,6 +21,7 @@ #include #include +#include #include #include @@ -35,28 +36,7 @@ namespace raft { namespace hierarchy { namespace detail { -/** - * Sorts a COO by its weight - * @tparam value_idx - * @tparam value_t - * @param[inout] rows source edges - * @param[inout] cols dest edges - * @param[inout] data edge weights - * @param[in] nnz number of edges in edge list - * @param[in] stream cuda stream for which to order cuda operations - */ -template -void sort_coo_by_data(value_idx *rows, value_idx *cols, value_t *data, - value_idx nnz, cudaStream_t stream) { - thrust::device_ptr t_rows = thrust::device_pointer_cast(rows); - thrust::device_ptr t_cols = thrust::device_pointer_cast(cols); - thrust::device_ptr t_data = thrust::device_pointer_cast(data); - - auto first = thrust::make_zip_iterator(thrust::make_tuple(rows, cols)); - thrust::sort_by_key(thrust::cuda::par.on(stream), t_data, t_data + nnz, - first); -} template void merge_msts(raft::Graph_COO &coo1, @@ -82,6 +62,13 @@ void merge_msts(raft::Graph_COO &coo1, coo1.n_edges = final_nnz; } +template +struct MSTEpilogueNoOp { + + void operator()(raft::handle_t &handle, value_idx *coo_rows, + value_idx *coo_cols, value_t *coo_data) {} +} + /** * Connect an unconnected knn graph (one in which mst returns an msf). The * device buffers underlying the Graph_COO object are modified in-place. @@ -95,12 +82,13 @@ void merge_msts(raft::Graph_COO &coo1, * @param[inout] color the color labels array returned from the mst invocation * @return updated MST edge list */ -template +template void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, raft::distance::DistanceType metric = - raft::distance::DistanceType::L2SqrtExpanded) { + raft::distance::DistanceType::L2SqrtExpanded, + mst_epilogue_f mst_epilogue_func) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -120,6 +108,9 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, handle, indptr2.data(), connected_edges.cols(), connected_edges.vals(), m, connected_edges.nnz, color, stream, false, false); + distance_epilogue_func(handle, new_mst.src.data(), new_mst.dst.data(), + new_mst.weights.data(), new_mst.n_edges); + merge_msts(msf, new_mst, stream); } @@ -147,7 +138,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * @param[in] max_iter maximum iterations to run knn graph connection. This * argument is really just a safeguard against the potential for infinite loops. */ -template +template void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const value_idx *indptr, const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, @@ -157,7 +148,9 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const size_t nnz, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded, - int max_iter = 10) { + int max_iter = 10, + mst_epilogue_f epilogue_func = + MSTEpilogueNoOp()) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -174,7 +167,7 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, while (n_components > 1 && iters < max_iter) { connect_knn_graph(handle, X, mst_coo, m, n, - color.data()); + color.data(), epilogue_func); iters++; @@ -189,7 +182,7 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, * 1. There is a bug in this code somewhere * 2. Either the given KNN graph wasn't generated from X or the same metric is not being used * to generate the 1-nn (currently only L2SqrtExpanded is supported). - * 3. max_iter was not large enough to connect the graph. + * 3. max_iter was not large enough to connect the graph (less likely). * * Note that a KNN graph generated from 50 random isotropic balls (with significant overlap) * was able to be connected in a single iteration. @@ -207,7 +200,7 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, "MST or symmetrization stage.", m - 1, mst_coo.n_edges); - sort_coo_by_data(mst_coo.src.data(), mst_coo.dst.data(), + raft::sparse::op::coo_sort_by_weight(mst_coo.src.data(), mst_coo.dst.data(), mst_coo.weights.data(), mst_coo.n_edges, stream); // TODO: be nice if we could pass these directly into the MST @@ -223,4 +216,4 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, }; // namespace detail }; // namespace hierarchy -}; // namespace raft \ No newline at end of file +}; // namespace raft diff --git a/cpp/include/raft/sparse/op/sort.h b/cpp/include/raft/sparse/op/sort.h index 792983cc9b..3874b3b304 100644 --- a/cpp/include/raft/sparse/op/sort.h +++ b/cpp/include/raft/sparse/op/sort.h @@ -92,6 +92,29 @@ void coo_sort(COO *const in, coo_sort(in->n_rows, in->n_cols, in->nnz, in->rows(), in->cols(), in->vals(), d_alloc, stream); } + +/** + * Sorts a COO by its weight + * @tparam value_idx + * @tparam value_t + * @param[inout] rows source edges + * @param[inout] cols dest edges + * @param[inout] data edge weights + * @param[in] nnz number of edges in edge list + * @param[in] stream cuda stream for which to order cuda operations + */ +template +void coo_sort_by_weight(value_idx *rows, value_idx *cols, value_t *data, + value_idx nnz, cudaStream_t stream) { + thrust::device_ptr t_rows = thrust::device_pointer_cast(rows); + thrust::device_ptr t_cols = thrust::device_pointer_cast(cols); + thrust::device_ptr t_data = thrust::device_pointer_cast(data); + + auto first = thrust::make_zip_iterator(thrust::make_tuple(rows, cols)); + + thrust::sort_by_key(thrust::cuda::par.on(stream), t_data, t_data + nnz, + first); +} }; // namespace op }; // end NAMESPACE sparse }; // end NAMESPACE raft \ No newline at end of file From 086b08a4df6083d67af72ac1ea95151a69d97e4f Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 20 Apr 2021 17:16:51 -0400 Subject: [PATCH 02/31] Fixing style --- .../raft/sparse/hierarchy/detail/mst.cuh | 32 ++++++++----------- cpp/include/raft/sparse/op/sort.h | 2 +- 2 files changed, 15 insertions(+), 19 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 387b39fe14..47cbe9769b 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -19,9 +19,9 @@ #include #include +#include #include #include -#include #include #include @@ -36,8 +36,6 @@ namespace raft { namespace hierarchy { namespace detail { - - template void merge_msts(raft::Graph_COO &coo1, raft::Graph_COO &coo2, @@ -62,9 +60,8 @@ void merge_msts(raft::Graph_COO &coo1, coo1.n_edges = final_nnz; } -template +template struct MSTEpilogueNoOp { - void operator()(raft::handle_t &handle, value_idx *coo_rows, value_idx *coo_cols, value_t *coo_data) {} } @@ -139,18 +136,16 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * argument is really just a safeguard against the potential for infinite loops. */ template -void build_sorted_mst(const raft::handle_t &handle, const value_t *X, - const value_idx *indptr, const value_idx *indices, - const value_t *pw_dists, size_t m, size_t n, - rmm::device_uvector &mst_src, - rmm::device_uvector &mst_dst, - rmm::device_uvector &mst_weight, - const size_t nnz, - raft::distance::DistanceType metric = - raft::distance::DistanceType::L2SqrtExpanded, - int max_iter = 10, - mst_epilogue_f epilogue_func = - MSTEpilogueNoOp()) { +void build_sorted_mst( + const raft::handle_t &handle, const value_t *X, const value_idx *indptr, + const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, + rmm::device_uvector &mst_src, + rmm::device_uvector &mst_dst, + rmm::device_uvector &mst_weight, const size_t nnz, + raft::distance::DistanceType metric = + raft::distance::DistanceType::L2SqrtExpanded, + int max_iter = 10, + mst_epilogue_f epilogue_func = MSTEpilogueNoOp()) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -201,7 +196,8 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, m - 1, mst_coo.n_edges); raft::sparse::op::coo_sort_by_weight(mst_coo.src.data(), mst_coo.dst.data(), - mst_coo.weights.data(), mst_coo.n_edges, stream); + mst_coo.weights.data(), mst_coo.n_edges, + stream); // TODO: be nice if we could pass these directly into the MST mst_src.resize(mst_coo.n_edges, stream); diff --git a/cpp/include/raft/sparse/op/sort.h b/cpp/include/raft/sparse/op/sort.h index 3874b3b304..9dbe2b67c5 100644 --- a/cpp/include/raft/sparse/op/sort.h +++ b/cpp/include/raft/sparse/op/sort.h @@ -105,7 +105,7 @@ void coo_sort(COO *const in, */ template void coo_sort_by_weight(value_idx *rows, value_idx *cols, value_t *data, - value_idx nnz, cudaStream_t stream) { + value_idx nnz, cudaStream_t stream) { thrust::device_ptr t_rows = thrust::device_pointer_cast(rows); thrust::device_ptr t_cols = thrust::device_pointer_cast(cols); thrust::device_ptr t_data = thrust::device_pointer_cast(data); From 37d2e0dfb3154480ad4d63191b97eb3437bfb5ba Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 23 Apr 2021 14:25:22 -0400 Subject: [PATCH 03/31] Adding missing argument --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 47cbe9769b..d45dbb88ac 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -63,7 +63,7 @@ void merge_msts(raft::Graph_COO &coo1, template struct MSTEpilogueNoOp { void operator()(raft::handle_t &handle, value_idx *coo_rows, - value_idx *coo_cols, value_t *coo_data) {} + value_idx *coo_cols, value_t *coo_data, value_idx nnz) {} } /** From 709c04039ec1406863f23a7e7a2353a600d00ad6 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 23 Apr 2021 14:50:35 -0400 Subject: [PATCH 04/31] Fixing typename --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index d45dbb88ac..311e189451 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -60,7 +60,7 @@ void merge_msts(raft::Graph_COO &coo1, coo1.n_edges = final_nnz; } -template +template struct MSTEpilogueNoOp { void operator()(raft::handle_t &handle, value_idx *coo_rows, value_idx *coo_cols, value_t *coo_data, value_idx nnz) {} From b1fbc63dff09fbb251c35f343bb8bd179092196e Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 23 Apr 2021 15:06:44 -0400 Subject: [PATCH 05/31] Updates --- cpp/include/raft/sparse/hierarchy/common.h | 6 ++++++ .../sparse/hierarchy/detail/agglomerative.cuh | 2 +- .../raft/sparse/hierarchy/detail/mst.cuh | 19 +++++++------------ .../raft/sparse/hierarchy/single_linkage.hpp | 3 ++- 4 files changed, 16 insertions(+), 14 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/common.h b/cpp/include/raft/sparse/hierarchy/common.h index 29f541498b..da768dfc7c 100644 --- a/cpp/include/raft/sparse/hierarchy/common.h +++ b/cpp/include/raft/sparse/hierarchy/common.h @@ -21,6 +21,12 @@ namespace hierarchy { enum LinkageDistance { PAIRWISE = 0, KNN_GRAPH = 1 }; +template +struct MSTEpilogueNoOp { + void operator()(const raft::handle_t &handle, value_idx *coo_rows, + value_idx *coo_cols, value_t *coo_data, value_idx nnz) {} +}; + /** * Simple POCO for consolidating linkage results. This closely * mirrors the trained instance variables populated in diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index be326a83db..1a864a455f 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -113,7 +113,7 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, std::vector children_h(n_edges * 2); std::vector out_size_h(n_edges); - std::vector out_delta_h(n_edges); + std::vector out_delta_h(n_edges); UnionFind U(nnz + 1); diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 311e189451..497bbc499b 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -60,11 +60,6 @@ void merge_msts(raft::Graph_COO &coo1, coo1.n_edges = final_nnz; } -template -struct MSTEpilogueNoOp { - void operator()(raft::handle_t &handle, value_idx *coo_rows, - value_idx *coo_cols, value_t *coo_data, value_idx nnz) {} -} /** * Connect an unconnected knn graph (one in which mst returns an msf). The @@ -83,9 +78,9 @@ template void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, - raft::distance::DistanceType metric = - raft::distance::DistanceType::L2SqrtExpanded, - mst_epilogue_f mst_epilogue_func) { + mst_epilogue_f mst_epilogue_func, + raft::distance::DistanceType metric = + raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -105,7 +100,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, handle, indptr2.data(), connected_edges.cols(), connected_edges.vals(), m, connected_edges.nnz, color, stream, false, false); - distance_epilogue_func(handle, new_mst.src.data(), new_mst.dst.data(), + mst_epilogue_func(handle, new_mst.src.data(), new_mst.dst.data(), new_mst.weights.data(), new_mst.n_edges); merge_msts(msf, new_mst, stream); @@ -141,11 +136,11 @@ void build_sorted_mst( const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, rmm::device_uvector &mst_src, rmm::device_uvector &mst_dst, - rmm::device_uvector &mst_weight, const size_t nnz, + rmm::device_uvector &mst_weight, size_t nnz, + mst_epilogue_f epilogue_func, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded, - int max_iter = 10, - mst_epilogue_f epilogue_func = MSTEpilogueNoOp()) { + int max_iter = 10) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index 096e154e87..a3de078a24 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -79,7 +79,8 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, */ detail::build_sorted_mst( handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, mst_rows, - mst_cols, mst_data, indices.size(), metric); + mst_cols, mst_data, indices.size(), MSTEpilogueNoOp(), + metric); pw_dists.release(); From 81b630a698c010bcdd4ce7ac6f45f61eafeca078 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 23 Apr 2021 15:08:45 -0400 Subject: [PATCH 06/31] Fixing style --- .../raft/sparse/hierarchy/detail/mst.cuh | 25 +++++++++---------- .../raft/sparse/hierarchy/single_linkage.hpp | 2 +- 2 files changed, 13 insertions(+), 14 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 497bbc499b..a355acd883 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -60,7 +60,6 @@ void merge_msts(raft::Graph_COO &coo1, coo1.n_edges = final_nnz; } - /** * Connect an unconnected knn graph (one in which mst returns an msf). The * device buffers underlying the Graph_COO object are modified in-place. @@ -79,7 +78,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, mst_epilogue_f mst_epilogue_func, - raft::distance::DistanceType metric = + raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -101,7 +100,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, connected_edges.nnz, color, stream, false, false); mst_epilogue_func(handle, new_mst.src.data(), new_mst.dst.data(), - new_mst.weights.data(), new_mst.n_edges); + new_mst.weights.data(), new_mst.n_edges); merge_msts(msf, new_mst, stream); } @@ -131,16 +130,16 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * argument is really just a safeguard against the potential for infinite loops. */ template -void build_sorted_mst( - const raft::handle_t &handle, const value_t *X, const value_idx *indptr, - const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, - rmm::device_uvector &mst_src, - rmm::device_uvector &mst_dst, - rmm::device_uvector &mst_weight, size_t nnz, - mst_epilogue_f epilogue_func, - raft::distance::DistanceType metric = - raft::distance::DistanceType::L2SqrtExpanded, - int max_iter = 10) { +void build_sorted_mst(const raft::handle_t &handle, const value_t *X, + const value_idx *indptr, const value_idx *indices, + const value_t *pw_dists, size_t m, size_t n, + rmm::device_uvector &mst_src, + rmm::device_uvector &mst_dst, + rmm::device_uvector &mst_weight, size_t nnz, + mst_epilogue_f epilogue_func, + raft::distance::DistanceType metric = + raft::distance::DistanceType::L2SqrtExpanded, + int max_iter = 10) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index a3de078a24..3f4b279e69 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -80,7 +80,7 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, detail::build_sorted_mst( handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, mst_rows, mst_cols, mst_data, indices.size(), MSTEpilogueNoOp(), - metric); + metric); pw_dists.release(); From 6c525421ad9eadc7348a898ccabeaafbf9dc1bdf Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 23 Apr 2021 18:05:08 -0400 Subject: [PATCH 07/31] Updating to get hdbscan to compile --- cpp/include/raft/sparse/op/reduce.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/sparse/op/reduce.cuh b/cpp/include/raft/sparse/op/reduce.cuh index a86a1acb86..0a026718fa 100644 --- a/cpp/include/raft/sparse/op/reduce.cuh +++ b/cpp/include/raft/sparse/op/reduce.cuh @@ -136,7 +136,7 @@ void max_duplicates(const raft::handle_t &handle, compute_duplicates_mask(diff.data(), rows, cols, nnz, stream); - thrust::exclusive_scan(exec_policy, diff.data(), diff.data() + diff.size(), + thrust::exclusive_scan(thrust::cuda::par.on(stream), diff.data(), diff.data() + diff.size(), diff.data()); // compute final size From 4be3d248f20e063f57d7332a2d2bd93244317dfd Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 26 Apr 2021 16:19:37 -0400 Subject: [PATCH 08/31] Some updates --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index a355acd883..e72e30b618 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -93,15 +93,15 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, connected_edges.nnz, indptr2.data(), m + 1, d_alloc, stream); + mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), + connected_edges.vals(), connected_edges.nnz); + // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process auto new_mst = raft::mst::mst( handle, indptr2.data(), connected_edges.cols(), connected_edges.vals(), m, connected_edges.nnz, color, stream, false, false); - mst_epilogue_func(handle, new_mst.src.data(), new_mst.dst.data(), - new_mst.weights.data(), new_mst.n_edges); - merge_msts(msf, new_mst, stream); } @@ -183,12 +183,20 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); + if(mst_coo.n_edges != m-1) { + raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, std::cout); + raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, std::cout); + raft::print_device_vector("mst_weight", mst_coo.weights.data(), mst_coo.n_edges, std::cout); + } + + RAFT_EXPECTS(mst_coo.n_edges == m - 1, "n_edges should be %d but was %d. This" "could be an indication of duplicate edges returned from the" "MST or symmetrization stage.", m - 1, mst_coo.n_edges); + raft::sparse::op::coo_sort_by_weight(mst_coo.src.data(), mst_coo.dst.data(), mst_coo.weights.data(), mst_coo.n_edges, stream); From 1aaa7d5908aa20394e50136b87932a81e301b5f7 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Thu, 29 Apr 2021 18:25:46 -0400 Subject: [PATCH 09/31] changes --- cpp/include/raft/sparse/hierarchy/common.h | 2 ++ .../raft/sparse/hierarchy/detail/mst.cuh | 24 ++++++++++++++----- cpp/include/raft/sparse/op/reduce.cuh | 4 ++-- 3 files changed, 22 insertions(+), 8 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/common.h b/cpp/include/raft/sparse/hierarchy/common.h index da768dfc7c..f5a2b26efe 100644 --- a/cpp/include/raft/sparse/hierarchy/common.h +++ b/cpp/include/raft/sparse/hierarchy/common.h @@ -14,6 +14,8 @@ * limitations under the License. */ +#include + #pragma once namespace raft { diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index e72e30b618..46abbf8d35 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -88,6 +88,9 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::linkage::connect_components(handle, connected_edges, X, color, m, n); + raft::print_device_vector("connected_edges", connected_edges.vals(), + connected_edges.nnz, std::cout); + rmm::device_uvector indptr2(m + 1, stream); raft::sparse::convert::sorted_coo_to_csr(connected_edges.rows(), connected_edges.nnz, indptr2.data(), @@ -150,14 +153,22 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, false, true); + raft::print_device_vector("initial_dists", pw_dists, nnz, std::cout); + raft::print_device_vector("mst_weights", mst_coo.weights.data(), mst_coo.weights.size(), std::cout); + int iters = 1; int n_components = linkage::get_n_components(color.data(), m, d_alloc, stream); while (n_components > 1 && iters < max_iter) { + + printf("Didn't converge. trying again\n"); connect_knn_graph(handle, X, mst_coo, m, n, color.data(), epilogue_func); + + raft::print_device_vector("mst_weights", mst_coo.weights.data(), mst_coo.weights.size(), std::cout); + iters++; n_components = linkage::get_n_components(color.data(), m, d_alloc, stream); @@ -183,20 +194,21 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); - if(mst_coo.n_edges != m-1) { - raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, std::cout); - raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, std::cout); - raft::print_device_vector("mst_weight", mst_coo.weights.data(), mst_coo.n_edges, std::cout); + if (mst_coo.n_edges != m - 1) { + raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, + std::cout); + raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, + std::cout); + raft::print_device_vector("mst_weight", mst_coo.weights.data(), + mst_coo.n_edges, std::cout); } - RAFT_EXPECTS(mst_coo.n_edges == m - 1, "n_edges should be %d but was %d. This" "could be an indication of duplicate edges returned from the" "MST or symmetrization stage.", m - 1, mst_coo.n_edges); - raft::sparse::op::coo_sort_by_weight(mst_coo.src.data(), mst_coo.dst.data(), mst_coo.weights.data(), mst_coo.n_edges, stream); diff --git a/cpp/include/raft/sparse/op/reduce.cuh b/cpp/include/raft/sparse/op/reduce.cuh index 0a026718fa..c912e332b1 100644 --- a/cpp/include/raft/sparse/op/reduce.cuh +++ b/cpp/include/raft/sparse/op/reduce.cuh @@ -136,8 +136,8 @@ void max_duplicates(const raft::handle_t &handle, compute_duplicates_mask(diff.data(), rows, cols, nnz, stream); - thrust::exclusive_scan(thrust::cuda::par.on(stream), diff.data(), diff.data() + diff.size(), - diff.data()); + thrust::exclusive_scan(thrust::cuda::par.on(stream), diff.data(), + diff.data() + diff.size(), diff.data()); // compute final size value_idx size = 0; From dd6e53785e132957bfcee14b1a38a77794d30092 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 30 Apr 2021 15:45:21 -0400 Subject: [PATCH 10/31] agglomerative labeling to accept device arrays directly --- .../raft/sparse/hierarchy/detail/agglomerative.cuh | 10 ++++------ cpp/include/raft/sparse/hierarchy/single_linkage.hpp | 2 +- 2 files changed, 5 insertions(+), 7 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index 1a864a455f..0fd1268bcb 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -92,8 +92,8 @@ template void build_dendrogram_host(const handle_t &handle, const value_idx *rows, const value_idx *cols, const value_t *data, size_t nnz, value_idx *children, - rmm::device_uvector &out_delta, - rmm::device_uvector &out_size) { + value_t *out_delta, + value_idx *out_size) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); @@ -135,13 +135,11 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, U.perform_union(aa, bb); } - out_size.resize(n_edges, stream); - printf("Finished dendrogram\n"); raft::update_device(children, children_h.data(), n_edges * 2, stream); - raft::update_device(out_size.data(), out_size_h.data(), n_edges, stream); - raft::update_device(out_delta.data(), out_delta_h.data(), n_edges, stream); + raft::update_device(out_size, out_size_h.data(), n_edges, stream); + raft::update_device(out_delta, out_delta_h.data(), n_edges, stream); } /** diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index 3f4b279e69..cfe18c9f97 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -94,7 +94,7 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, // Create dendrogram detail::build_dendrogram_host( handle, mst_rows.data(), mst_cols.data(), mst_data.data(), n_edges, - out->children, out_delta, out_size); + out->children, out_delta.data(), out_size.data()); detail::extract_flattened_clusters(handle, out->labels, out->children, n_clusters, m); From e3085cb4065aaa24b1c21f3d2559c60d40ba3def Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 30 Apr 2021 16:48:14 -0400 Subject: [PATCH 11/31] Cleaning up inputs to some of the single linkage prims --- .../sparse/hierarchy/detail/agglomerative.cuh | 3 +-- .../raft/sparse/hierarchy/detail/mst.cuh | 25 +++++++------------ .../raft/sparse/hierarchy/single_linkage.hpp | 12 ++++----- 3 files changed, 16 insertions(+), 24 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index 0fd1268bcb..04affc2287 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -91,8 +91,7 @@ class UnionFind { template void build_dendrogram_host(const handle_t &handle, const value_idx *rows, const value_idx *cols, const value_t *data, - size_t nnz, value_idx *children, - value_t *out_delta, + size_t nnz, value_idx *children, value_t *out_delta, value_idx *out_size) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 46abbf8d35..6f534d77c2 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -136,9 +136,8 @@ template void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const value_idx *indptr, const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, - rmm::device_uvector &mst_src, - rmm::device_uvector &mst_dst, - rmm::device_uvector &mst_weight, size_t nnz, + value_idx *mst_src, value_idx *mst_dst, + value_t *mst_weight, size_t nnz, mst_epilogue_f epilogue_func, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded, @@ -154,20 +153,20 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, false, true); raft::print_device_vector("initial_dists", pw_dists, nnz, std::cout); - raft::print_device_vector("mst_weights", mst_coo.weights.data(), mst_coo.weights.size(), std::cout); + raft::print_device_vector("mst_weights", mst_coo.weights.data(), + mst_coo.weights.size(), std::cout); int iters = 1; int n_components = linkage::get_n_components(color.data(), m, d_alloc, stream); while (n_components > 1 && iters < max_iter) { - printf("Didn't converge. trying again\n"); connect_knn_graph(handle, X, mst_coo, m, n, color.data(), epilogue_func); - - raft::print_device_vector("mst_weights", mst_coo.weights.data(), mst_coo.weights.size(), std::cout); + raft::print_device_vector("mst_weights", mst_coo.weights.data(), + mst_coo.weights.size(), std::cout); iters++; @@ -213,15 +212,9 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, mst_coo.weights.data(), mst_coo.n_edges, stream); - // TODO: be nice if we could pass these directly into the MST - mst_src.resize(mst_coo.n_edges, stream); - mst_dst.resize(mst_coo.n_edges, stream); - mst_weight.resize(mst_coo.n_edges, stream); - - raft::copy_async(mst_src.data(), mst_coo.src.data(), mst_coo.n_edges, stream); - raft::copy_async(mst_dst.data(), mst_coo.dst.data(), mst_coo.n_edges, stream); - raft::copy_async(mst_weight.data(), mst_coo.weights.data(), mst_coo.n_edges, - stream); + raft::copy_async(mst_src, mst_coo.src.data(), mst_coo.n_edges, stream); + raft::copy_async(mst_dst, mst_coo.dst.data(), mst_coo.n_edges, stream); + raft::copy_async(mst_weight, mst_coo.weights.data(), mst_coo.n_edges, stream); } }; // namespace detail diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index cfe18c9f97..4a744bc048 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -70,17 +70,17 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, detail::get_distance_graph( handle, X, m, n, metric, indptr, indices, pw_dists, c); - rmm::device_uvector mst_rows(EMPTY, stream); - rmm::device_uvector mst_cols(EMPTY, stream); - rmm::device_uvector mst_data(EMPTY, stream); + rmm::device_uvector mst_rows(m - 1, stream); + rmm::device_uvector mst_cols(m - 1, stream); + rmm::device_uvector mst_data(m - 1, stream); /** * 2. Construct MST, sorted by weights */ detail::build_sorted_mst( - handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, mst_rows, - mst_cols, mst_data, indices.size(), MSTEpilogueNoOp(), - metric); + handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, + mst_rows.data(), mst_cols.data(), mst_data.data(), indices.size(), + MSTEpilogueNoOp(), metric); pw_dists.release(); From ec1cca59ab1076fdd083aac97b1b8a698570e984 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 3 May 2021 18:24:50 -0700 Subject: [PATCH 12/31] removing deprecated rmm policy usage --- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 29 +++++++++---------- cpp/include/raft/sparse/mst/detail/utils.cuh | 2 +- cpp/include/raft/sparse/mst/mst_solver.cuh | 2 +- cpp/test/mst.cu | 1 - 4 files changed, 16 insertions(+), 18 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 5ec7697606..9dbc41164c 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -82,12 +82,12 @@ MST_solver::MST_solver( //Initially, color holds the vertex id as color auto policy = rmm::exec_policy(stream); if (initialize_colors_) { - thrust::sequence(policy->on(stream), color.begin(), color.end(), 0); - thrust::sequence(policy->on(stream), color_index, color_index + v, 0); + thrust::sequence(policy, color.begin(), color.end(), 0); + thrust::sequence(policy, color_index, color_index + v, 0); } else { raft::copy(color.data().get(), color_index, v, stream); } - thrust::sequence(policy->on(stream), next_color.begin(), next_color.end(), 0); + thrust::sequence(policy, next_color.begin(), next_color.end(), 0); } template @@ -212,12 +212,12 @@ weight_t MST_solver::alteration_max() { auto policy = rmm::exec_policy(stream); rmm::device_vector tmp(e); thrust::device_ptr weights_ptr(weights); - thrust::copy(policy->on(stream), weights_ptr, weights_ptr + e, tmp.begin()); + thrust::copy(policy, weights_ptr, weights_ptr + e, tmp.begin()); //sort tmp weights - thrust::sort(policy->on(stream), tmp.begin(), tmp.end()); + thrust::sort(policy, tmp.begin(), tmp.end()); //remove duplicates - auto new_end = thrust::unique(policy->on(stream), tmp.begin(), tmp.end()); + auto new_end = thrust::unique(policy, tmp.begin(), tmp.end()); //min(a[i+1]-a[i])/2 auto begin = @@ -225,9 +225,9 @@ weight_t MST_solver::alteration_max() { auto end = thrust::make_zip_iterator(thrust::make_tuple(new_end - 1, new_end)); auto init = tmp[1] - tmp[0]; - auto max = thrust::transform_reduce(policy->on(stream), begin, end, - alteration_functor(), init, - thrust::minimum()); + auto max = + thrust::transform_reduce(policy, begin, end, alteration_functor(), + init, thrust::minimum()); return max / static_cast(2); } @@ -306,9 +306,9 @@ void MST_solver::label_prop(vertex_t* mst_src, template void MST_solver::min_edge_per_vertex() { auto policy = rmm::exec_policy(stream); - thrust::fill(policy->on(stream), min_edge_color.begin(), min_edge_color.end(), + thrust::fill(policy, min_edge_color.begin(), min_edge_color.end(), std::numeric_limits::max()); - thrust::fill(policy->on(stream), new_mst_edge.begin(), new_mst_edge.end(), + thrust::fill(policy, new_mst_edge.begin(), new_mst_edge.end(), std::numeric_limits::max()); int n_threads = 32; @@ -331,7 +331,7 @@ void MST_solver::min_edge_per_supervertex() { auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); auto policy = rmm::exec_policy(stream); - thrust::fill(policy->on(stream), temp_src.begin(), temp_src.end(), + thrust::fill(policy, temp_src.begin(), temp_src.end(), std::numeric_limits::max()); vertex_t* color_ptr = color.data().get(); @@ -401,9 +401,8 @@ void MST_solver::append_src_dst_pair( thrust::make_tuple(temp_src.end(), temp_dst.end(), temp_weights.end())); // copy new mst edges to final output - thrust::copy_if(policy->on(stream), temp_src_dst_zip_begin, - temp_src_dst_zip_end, src_dst_zip_end, - new_edges_functor()); + thrust::copy_if(policy, temp_src_dst_zip_begin, temp_src_dst_zip_end, + src_dst_zip_end, new_edges_functor()); } } // namespace mst diff --git a/cpp/include/raft/sparse/mst/detail/utils.cuh b/cpp/include/raft/sparse/mst/detail/utils.cuh index c4884479b4..8f755de459 100644 --- a/cpp/include/raft/sparse/mst/detail/utils.cuh +++ b/cpp/include/raft/sparse/mst/detail/utils.cuh @@ -17,8 +17,8 @@ #pragma once -#include #include +#include #define MST_TIME namespace raft { diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index af8604cdea..438dc79a49 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -17,9 +17,9 @@ #pragma once -#include #include #include +#include namespace raft { diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 3dbbdc40c9..d32c10d881 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -17,7 +17,6 @@ #include #include -#include #include #include #include From d32677dd1ee6cb45526cfde1af062288f3a3429f Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 5 May 2021 16:17:58 -0400 Subject: [PATCH 13/31] Changes --- .../raft/sparse/hierarchy/detail/agglomerative.cuh | 4 ---- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 13 ++----------- 2 files changed, 2 insertions(+), 15 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index 04affc2287..d67654f031 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -102,8 +102,6 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, std::vector mst_dst_h(n_edges); std::vector mst_weights_h(n_edges); - printf("n_edges: %d\n", n_edges); - update_host(mst_src_h.data(), rows, n_edges, stream); update_host(mst_dst_h.data(), cols, n_edges, stream); update_host(mst_weights_h.data(), data, n_edges, stream); @@ -134,8 +132,6 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, U.perform_union(aa, bb); } - printf("Finished dendrogram\n"); - raft::update_device(children, children_h.data(), n_edges * 2, stream); raft::update_device(out_size, out_size_h.data(), n_edges, stream); raft::update_device(out_delta, out_delta_h.data(), n_edges, stream); diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 6f534d77c2..d33e6c5447 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -88,17 +88,14 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::linkage::connect_components(handle, connected_edges, X, color, m, n); - raft::print_device_vector("connected_edges", connected_edges.vals(), - connected_edges.nnz, std::cout); + mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), + connected_edges.vals(), connected_edges.nnz); rmm::device_uvector indptr2(m + 1, stream); raft::sparse::convert::sorted_coo_to_csr(connected_edges.rows(), connected_edges.nnz, indptr2.data(), m + 1, d_alloc, stream); - mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), - connected_edges.vals(), connected_edges.nnz); - // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process auto new_mst = raft::mst::mst( @@ -152,9 +149,6 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, false, true); - raft::print_device_vector("initial_dists", pw_dists, nnz, std::cout); - raft::print_device_vector("mst_weights", mst_coo.weights.data(), - mst_coo.weights.size(), std::cout); int iters = 1; int n_components = @@ -165,9 +159,6 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, connect_knn_graph(handle, X, mst_coo, m, n, color.data(), epilogue_func); - raft::print_device_vector("mst_weights", mst_coo.weights.data(), - mst_coo.weights.size(), std::cout); - iters++; n_components = linkage::get_n_components(color.data(), m, d_alloc, stream); From 43f7cf8da00e106700c39b546d267d44408ef9f3 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 3 May 2021 18:24:50 -0700 Subject: [PATCH 14/31] removing deprecated rmm policy usage --- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 29 +++++++++---------- cpp/include/raft/sparse/mst/detail/utils.cuh | 2 +- cpp/include/raft/sparse/mst/mst_solver.cuh | 2 +- cpp/test/mst.cu | 1 - 4 files changed, 16 insertions(+), 18 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 3218f09c31..2847cf2077 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -83,12 +83,12 @@ MST_solver::MST_solver( //Initially, color holds the vertex id as color auto policy = rmm::exec_policy(stream); if (initialize_colors_) { - thrust::sequence(policy->on(stream), color.begin(), color.end(), 0); - thrust::sequence(policy->on(stream), color_index, color_index + v, 0); + thrust::sequence(policy, color.begin(), color.end(), 0); + thrust::sequence(policy, color_index, color_index + v, 0); } else { raft::copy(color.data().get(), color_index, v, stream); } - thrust::sequence(policy->on(stream), next_color.begin(), next_color.end(), 0); + thrust::sequence(policy, next_color.begin(), next_color.end(), 0); } template @@ -213,12 +213,12 @@ weight_t MST_solver::alteration_max() { auto policy = rmm::exec_policy(stream); rmm::device_vector tmp(e); thrust::device_ptr weights_ptr(weights); - thrust::copy(policy->on(stream), weights_ptr, weights_ptr + e, tmp.begin()); + thrust::copy(policy, weights_ptr, weights_ptr + e, tmp.begin()); //sort tmp weights - thrust::sort(policy->on(stream), tmp.begin(), tmp.end()); + thrust::sort(policy, tmp.begin(), tmp.end()); //remove duplicates - auto new_end = thrust::unique(policy->on(stream), tmp.begin(), tmp.end()); + auto new_end = thrust::unique(policy, tmp.begin(), tmp.end()); //min(a[i+1]-a[i])/2 auto begin = @@ -226,9 +226,9 @@ weight_t MST_solver::alteration_max() { auto end = thrust::make_zip_iterator(thrust::make_tuple(new_end - 1, new_end)); auto init = tmp[1] - tmp[0]; - auto max = thrust::transform_reduce(policy->on(stream), begin, end, - alteration_functor(), init, - thrust::minimum()); + auto max = + thrust::transform_reduce(policy, begin, end, alteration_functor(), + init, thrust::minimum()); return max / static_cast(2); } @@ -307,9 +307,9 @@ void MST_solver::label_prop(vertex_t* mst_src, template void MST_solver::min_edge_per_vertex() { auto policy = rmm::exec_policy(stream); - thrust::fill(policy->on(stream), min_edge_color.begin(), min_edge_color.end(), + thrust::fill(policy, min_edge_color.begin(), min_edge_color.end(), std::numeric_limits::max()); - thrust::fill(policy->on(stream), new_mst_edge.begin(), new_mst_edge.end(), + thrust::fill(policy, new_mst_edge.begin(), new_mst_edge.end(), std::numeric_limits::max()); int n_threads = 32; @@ -332,7 +332,7 @@ void MST_solver::min_edge_per_supervertex() { auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); auto policy = rmm::exec_policy(stream); - thrust::fill(policy->on(stream), temp_src.begin(), temp_src.end(), + thrust::fill(policy, temp_src.begin(), temp_src.end(), std::numeric_limits::max()); vertex_t* color_ptr = color.data().get(); @@ -402,9 +402,8 @@ void MST_solver::append_src_dst_pair( thrust::make_tuple(temp_src.end(), temp_dst.end(), temp_weights.end())); // copy new mst edges to final output - thrust::copy_if(policy->on(stream), temp_src_dst_zip_begin, - temp_src_dst_zip_end, src_dst_zip_end, - new_edges_functor()); + thrust::copy_if(policy, temp_src_dst_zip_begin, temp_src_dst_zip_end, + src_dst_zip_end, new_edges_functor()); } } // namespace mst diff --git a/cpp/include/raft/sparse/mst/detail/utils.cuh b/cpp/include/raft/sparse/mst/detail/utils.cuh index c4884479b4..8f755de459 100644 --- a/cpp/include/raft/sparse/mst/detail/utils.cuh +++ b/cpp/include/raft/sparse/mst/detail/utils.cuh @@ -17,8 +17,8 @@ #pragma once -#include #include +#include #define MST_TIME namespace raft { diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index af8604cdea..438dc79a49 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -17,9 +17,9 @@ #pragma once -#include #include #include +#include namespace raft { diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index 3dbbdc40c9..d32c10d881 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -17,7 +17,6 @@ #include #include -#include #include #include #include From e0ef5b33ca647bb9337e48ee4568f3759b0fe017 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 10 May 2021 11:34:05 -0700 Subject: [PATCH 15/31] alpha to weight alteration for precision --- .../raft/sparse/mst/detail/mst_kernels.cuh | 18 +++++++++++---- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 23 +++++++++++++------ cpp/include/raft/sparse/mst/mst.cuh | 2 +- cpp/include/raft/sparse/mst/mst_solver.cuh | 5 ++-- 4 files changed, 34 insertions(+), 14 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index d2f86d6dc8..8a44a2cc38 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -285,17 +285,27 @@ template __global__ void alteration_kernel(const vertex_t v, const edge_t e, const edge_t* offsets, const vertex_t* indices, - const weight_t* weights, weight_t max, + const weight_t* weights, double max, weight_t* random_values, - weight_t* altered_weights) { + weight_t* altered_weights, int alpha, + bool use_alpha) { auto row = get_1D_idx(); if (row < v) { auto row_begin = offsets[row]; auto row_end = offsets[row + 1]; for (auto i = row_begin; i < row_end; i++) { auto column = indices[i]; - altered_weights[i] = - weights[i] + max * (random_values[row] + random_values[column]); + // doing the later step explicity in double for precision + if (use_alpha) { + altered_weights[i] = + alpha * weights[i] + alpha * max * + (static_cast(random_values[row]) + + static_cast(random_values[column])); + } else { + altered_weights[i] = + weights[i] + max * (static_cast(random_values[row]) + + static_cast(random_values[column])); + } } } } 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 2847cf2077..bebef92027 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -53,7 +53,7 @@ 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_, - bool symmetrize_output_, bool initialize_colors_, int iterations_) + bool symmetrize_output_, bool initialize_colors_, int iterations_, int alpha_) : handle(handle_), offsets(offsets_), indices(indices_), @@ -75,7 +75,8 @@ MST_solver::MST_solver( stream(stream_), symmetrize_output(symmetrize_output_), initialize_colors(initialize_colors_), - iterations(iterations_) { + iterations(iterations_), + alpha(alpha_) { max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; sm_count = handle_.get_device_properties().multiProcessorCount; @@ -114,7 +115,9 @@ MST_solver::solve() { timer0 = duration_us(stop - start); #endif - Graph_COO mst_result(2 * v - 2, stream); + auto n_expected_edges = symmetrize_output ? 2 * v - 2 : v - 1; + + Graph_COO mst_result(n_expected_edges, stream); // Boruvka original formulation says "while more than 1 supervertex remains" // Here we adjust it to support disconnected components (spanning forest) @@ -150,6 +153,10 @@ MST_solver::solve() { timer3 += duration_us(stop - start); #endif + RAFT_EXPECTS(mst_edge_count[0] == n_expected_edges, + "Number of edges found by MST is invalid. This may be due to " + "loss in precision. Try increasing precision of weights.") + if (prev_mst_edge_count[0] == mst_edge_count[0]) { #ifdef MST_TIME std::cout << "Iterations: " << i << std::endl; @@ -209,7 +216,7 @@ struct alteration_functor { // Compute the uper bound for the alteration template -weight_t MST_solver::alteration_max() { +double MST_solver::alteration_max() { auto policy = rmm::exec_policy(stream); rmm::device_vector tmp(e); thrust::device_ptr weights_ptr(weights); @@ -229,7 +236,7 @@ weight_t MST_solver::alteration_max() { auto max = thrust::transform_reduce(policy, begin, end, alteration_functor(), init, thrust::minimum()); - return max / static_cast(2); + return max / static_cast(2); } // Compute the alteration to make all undirected edge weight unique @@ -240,7 +247,7 @@ void MST_solver::alteration() { auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); // maximum alteration that does not change realtive weights order - weight_t max = alteration_max(); + double max = alteration_max(); // pool of rand values rmm::device_vector rand_values(v); @@ -258,10 +265,12 @@ void MST_solver::alteration() { RAFT_EXPECTS(curand_status == CURAND_STATUS_SUCCESS, "MST: CURAND cleanup failed"); + bool use_alpha = max < 1e-3 && sizeof(weight_t) == 4; + //Alterate the weights, make all undirected edge weight unique while keeping Wuv == Wvu detail::alteration_kernel<<>>( v, e, offsets, indices, weights, max, rand_values.data().get(), - altered_weights.data().get()); + altered_weights.data().get(), alpha, use_alpha); } // updates colors of vertices by propagating the lower color to the higher diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index 4685431e7a..034f7a98b6 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -27,7 +27,7 @@ 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, bool symmetrize_output = true, - bool initialize_colors = true, int iterations = 0) { + bool initialize_colors = true, int iterations = 0, int alpha = 1e6) { MST_solver mst_solver( handle, offsets, indices, weights, v, e, color, stream, symmetrize_output, initialize_colors, iterations); diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index 438dc79a49..ab95c11744 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -43,7 +43,7 @@ class MST_solver { const vertex_t* indices_, const weight_t* weights_, const vertex_t v_, const edge_t e_, vertex_t* color_, cudaStream_t stream_, bool symmetrize_output_, - bool initialize_colors_, int iterations_); + bool initialize_colors_, int iterations_ int alpha_); raft::Graph_COO solve(); @@ -54,6 +54,7 @@ class MST_solver { cudaStream_t stream; bool symmetrize_output, initialize_colors; int iterations; + int alpha; //CSR const edge_t* offsets; @@ -90,7 +91,7 @@ class MST_solver { void min_edge_per_supervertex(); void check_termination(); void alteration(); - weight_t alteration_max(); + double alteration_max(); void append_src_dst_pair(vertex_t* mst_src, vertex_t* mst_dst, weight_t* mst_weights); }; From 1678c66c340778453757acff45a19906a6d9f1b4 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 10 May 2021 11:53:17 -0700 Subject: [PATCH 16/31] resolving errors --- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 12 +++++++----- cpp/include/raft/sparse/mst/mst.cuh | 2 +- cpp/include/raft/sparse/mst/mst_solver.cuh | 2 +- cpp/test/mst.cu | 10 +++++----- 4 files changed, 14 insertions(+), 12 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 bebef92027..4222c46b2a 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -22,6 +22,11 @@ #include "mst_kernels.cuh" #include "utils.cuh" + +#include +#include +#include + #include #include #include @@ -31,9 +36,6 @@ #include #include -#include -#include - namespace raft { namespace mst { typedef std::chrono::high_resolution_clock Clock; @@ -153,9 +155,9 @@ MST_solver::solve() { timer3 += duration_us(stop - start); #endif - RAFT_EXPECTS(mst_edge_count[0] == n_expected_edges, + RAFT_EXPECTS(mst_edge_count[0] <= n_expected_edges, "Number of edges found by MST is invalid. This may be due to " - "loss in precision. Try increasing precision of weights.") + "loss in precision. Try increasing precision of weights."); if (prev_mst_edge_count[0] == mst_edge_count[0]) { #ifdef MST_TIME diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index 034f7a98b6..abe74f3fe3 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -30,7 +30,7 @@ raft::Graph_COO mst( bool initialize_colors = true, int iterations = 0, int alpha = 1e6) { MST_solver mst_solver( handle, offsets, indices, weights, v, e, color, stream, symmetrize_output, - initialize_colors, iterations); + initialize_colors, iterations, alpha); 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 ab95c11744..e9e0749cdf 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -43,7 +43,7 @@ class MST_solver { const vertex_t* indices_, const weight_t* weights_, const vertex_t v_, const edge_t e_, vertex_t* color_, cudaStream_t stream_, bool symmetrize_output_, - bool initialize_colors_, int iterations_ int alpha_); + bool initialize_colors_, int iterations_, int alpha_); raft::Graph_COO solve(); diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index d32c10d881..a9119acbad 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -138,12 +138,12 @@ class MSTTest if (iterations == 0) { MST_solver symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, true, 0); + true, true, 0, 0); auto symmetric_result = symmetric_solver.solve(); MST_solver non_symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - false, true, 0); + false, true, 0, 0); auto non_symmetric_result = non_symmetric_solver.solve(); EXPECT_LE(symmetric_result.n_edges, 2 * v - 2); @@ -154,12 +154,12 @@ class MSTTest } else { MST_solver intermediate_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, true, iterations); + true, true, iterations, 0); auto intermediate_result = intermediate_solver.solve(); MST_solver symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, false, 0); + true, false, 0, 0); auto symmetric_result = symmetric_solver.solve(); // symmetric_result.n_edges += intermediate_result.n_edges; @@ -182,7 +182,7 @@ class MSTTest MST_solver non_symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - false, true, 0); + false, true, 0, 0); auto non_symmetric_result = non_symmetric_solver.solve(); EXPECT_LE(symmetric_result.n_edges, 2 * v - 2); From 5eda4aad36c1787ddf982621c84c9e4539ebfce2 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 10 May 2021 18:52:24 -0700 Subject: [PATCH 17/31] trying alpha for all --- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 4 +++- 1 file changed, 3 insertions(+), 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 80088ac00b..8f60bc0169 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -266,7 +266,9 @@ void MST_solver::alteration() { RAFT_EXPECTS(curand_status == CURAND_STATUS_SUCCESS, "MST: CURAND cleanup failed"); - bool use_alpha = max < 1e-3 && sizeof(weight_t) == 4; + std::cout << "Max: " << max << std::endl; + // bool use_alpha = max < 1e-3 && sizeof(weight_t) == 4; + bool use_alpha = true; //Alterate the weights, make all undirected edge weight unique while keeping Wuv == Wvu detail::alteration_kernel<<>>( From 030b3a92e9c74227c4607f99f58202ddfa743864 Mon Sep 17 00:00:00 2001 From: divyegala Date: Wed, 12 May 2021 09:47:57 -0700 Subject: [PATCH 18/31] double precision weight alteration --- .../raft/sparse/mst/detail/mst_kernels.cuh | 57 ++++++++++++------- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 35 +++++++++--- cpp/include/raft/sparse/mst/mst_solver.cuh | 4 +- 3 files changed, 64 insertions(+), 32 deletions(-) diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 8a44a2cc38..d9693583cc 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -27,22 +27,22 @@ namespace raft { namespace mst { namespace detail { -template +template __global__ void kernel_min_edge_per_vertex( - const edge_t* offsets, const vertex_t* indices, const weight_t* weights, + const edge_t* offsets, const vertex_t* indices, const double* weights, const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, - const bool* mst_edge, weight_t* min_edge_color, const vertex_t v) { + const bool* mst_edge, double* min_edge_color, 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__ double 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_edge_weight[lane_id] = std::numeric_limits::max(); min_color[lane_id] = std::numeric_limits::max(); __syncthreads(); @@ -61,7 +61,7 @@ __global__ void kernel_min_edge_per_vertex( // 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]; + auto curr_edge_weight = weights[e]; vertex_t successor_color_idx = color_index[indices[e]]; vertex_t successor_color = color[successor_color_idx]; @@ -92,7 +92,7 @@ __global__ void kernel_min_edge_per_vertex( // min edge may now be found in first thread if (lane_id == 0) { - if (min_edge_weight[0] != std::numeric_limits::max()) { + if (min_edge_weight[0] != std::numeric_limits::max()) { new_mst_edge[warp_id] = min_edge_index[0]; // atomically set min edge per color @@ -106,8 +106,8 @@ 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, + const double* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, + weight_t* temp_weights, const double* min_edge_color, const vertex_t v, bool symmetrize_output) { auto tid = get_1D_idx(); if (tid < v) { @@ -119,7 +119,7 @@ __global__ void min_edge_per_supervertex( // 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]; + double vertex_weight = altered_weights[edge_idx]; bool add_edge = false; if (min_edge_color[vertex_color] == vertex_weight) { @@ -286,8 +286,8 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, const edge_t* offsets, const vertex_t* indices, const weight_t* weights, double max, - weight_t* random_values, - weight_t* altered_weights, int alpha, + double* random_values, + double* altered_weights, int alpha, bool use_alpha) { auto row = get_1D_idx(); if (row < v) { @@ -296,16 +296,31 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, for (auto i = row_begin; i < row_end; i++) { auto column = indices[i]; // doing the later step explicity in double for precision - if (use_alpha) { - altered_weights[i] = - alpha * weights[i] + alpha * max * - (static_cast(random_values[row]) + - static_cast(random_values[column])); - } else { - altered_weights[i] = - weights[i] + max * (static_cast(random_values[row]) + - static_cast(random_values[column])); + altered_weights[i] = + weights[i] + max * (random_values[row] + random_values[column]); + + auto print = false; + if (row == 293 && column == 1276) print = true; + if (row == 3 && column == 1310) print = true; + if (row == 1507 && column == 594) print = true; + + if (row == 279 && column == 1310) print = true; + if (row == 1276 && column == 1686) print = true; + + if (print) { + printf("row: %d, col: %d, alt: %lf, weight: %lf, max: %lf, randr: %lf, randl: %lf\n", row, column, altered_weights[i], weights[i], max, random_values[row], random_values[column]); } + + // if (use_alpha) { + // altered_weights[i] = + // alpha * weights[i] + alpha * max * + // (static_cast(random_values[row]) + + // static_cast(random_values[column])); + // } else { + // altered_weights[i] = + // weights[i] + max * (static_cast(random_values[row]) + + // static_cast(random_values[column])); + // } } } } 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 8f60bc0169..693f44b033 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -106,6 +106,13 @@ MST_solver::solve() { auto start = Clock::now(); #endif + // int start_1507, end_1507; + // raft::update_host(&start_1507, offsets + 1507, 1, stream); + // raft::update_host(&end_1507, offsets + 1508, 1, stream); + + // raft::print_device_vector("1507 indices", indices + start_1507, end_1507 - start_1507, std::cout); + // raft::print_device_vector("1507 weights", weights + start_1507, end_1507 - start_1507, std::cout); + // Alterating the weights // this is done by identifying the lowest cost edge weight gap that is not 0, call this theta. // For each edge, add noise that is less than theta. That is, generate a random number in the range [0.0, theta) and add it to each edge weight. @@ -153,12 +160,22 @@ MST_solver::solve() { stop = Clock::now(); timer3 += duration_us(stop - start); #endif + // raft::print_device_vector("altered_weights", altered_weights.data().get(), e, std::cout); + // raft::print_device_vector("new_mst_edge", new_mst_edge.data().get(), v, std::cout); + // raft::print_device_vector("min_edge_color", min_edge_color.data().get(), v, std::cout); + + // raft::print_device_vector("temp_src", temp_src.data().get(), v, std::cout); + // raft::print_device_vector("temp_dst", temp_dst.data().get(), v, std::cout); + // raft::print_device_vector("temp_weights", temp_weights.data().get(), v, std::cout); + - RAFT_EXPECTS(mst_edge_count[0] <= n_expected_edges, + auto curr_mst_edge_count = mst_edge_count[0]; + std::cout << "edge count: " << curr_mst_edge_count << ", expected: " << n_expected_edges << std::endl; + RAFT_EXPECTS(curr_mst_edge_count <= n_expected_edges, "Number of edges found by MST is invalid. This may be due to " "loss in precision. Try increasing precision of weights."); - if (prev_mst_edge_count[0] == mst_edge_count[0]) { + if (prev_mst_edge_count[0] == curr_mst_edge_count) { #ifdef MST_TIME std::cout << "Iterations: " << i << std::endl; std::cout << timer0 << "," << timer1 << "," << timer2 << "," << timer3 @@ -251,7 +268,7 @@ void MST_solver::alteration() { double max = alteration_max(); // pool of rand values - rmm::device_vector rand_values(v); + rmm::device_vector rand_values(v); // Random number generator curandGenerator_t randGen; @@ -320,17 +337,17 @@ template void MST_solver::min_edge_per_vertex() { auto policy = rmm::exec_policy(stream); thrust::fill(policy, min_edge_color.begin(), min_edge_color.end(), - std::numeric_limits::max()); + std::numeric_limits::max()); thrust::fill(policy, new_mst_edge.begin(), new_mst_edge.end(), - std::numeric_limits::max()); + std::numeric_limits::max()); int n_threads = 32; vertex_t* color_ptr = color.data().get(); edge_t* new_mst_edge_ptr = new_mst_edge.data().get(); bool* mst_edge_ptr = mst_edge.data().get(); - weight_t* min_edge_color_ptr = min_edge_color.data().get(); - weight_t* altered_weights_ptr = altered_weights.data().get(); + double* min_edge_color_ptr = min_edge_color.data().get(); + double* altered_weights_ptr = altered_weights.data().get(); detail::kernel_min_edge_per_vertex<<>>( offsets, indices, altered_weights_ptr, color_ptr, color_index, @@ -350,8 +367,8 @@ void MST_solver::min_edge_per_supervertex() { vertex_t* color_ptr = color.data().get(); edge_t* new_mst_edge_ptr = new_mst_edge.data().get(); bool* mst_edge_ptr = mst_edge.data().get(); - weight_t* min_edge_color_ptr = min_edge_color.data().get(); - weight_t* altered_weights_ptr = altered_weights.data().get(); + double* min_edge_color_ptr = min_edge_color.data().get(); + double* altered_weights_ptr = altered_weights.data().get(); vertex_t* temp_src_ptr = temp_src.data().get(); vertex_t* temp_dst_ptr = temp_dst.data().get(); weight_t* temp_weights_ptr = temp_weights.data().get(); diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index e9e0749cdf..2bac221f31 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -68,10 +68,10 @@ class MST_solver { vertex_t sm_count; vertex_t* color_index; // represent each supervertex as a color - rmm::device_vector + rmm::device_vector min_edge_color; // minimum incident edge weight per color rmm::device_vector new_mst_edge; // new minimum edge per vertex - rmm::device_vector altered_weights; // weights to be used for mst + rmm::device_vector altered_weights; // weights to be used for mst rmm::device_vector mst_edge_count; // total number of edges added after every iteration rmm::device_vector From 635d018f687c82f571b7c511e7632400046fec6a Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Thu, 13 May 2021 20:52:55 -0400 Subject: [PATCH 19/31] Checking in --- .../raft/sparse/hierarchy/detail/mst.cuh | 40 ++++++++++--------- .../raft/sparse/hierarchy/single_linkage.hpp | 6 ++- .../sparse/selection/connect_components.cuh | 12 +++--- 3 files changed, 31 insertions(+), 27 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index d33e6c5447..19c21cb868 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -73,11 +73,11 @@ void merge_msts(raft::Graph_COO &coo1, * @param[inout] color the color labels array returned from the mst invocation * @return updated MST edge list */ -template +template void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, - mst_epilogue_f mst_epilogue_func, + mst_epilogue_f mst_epilogue_func, red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); @@ -86,7 +86,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::sparse::COO connected_edges(d_alloc, stream); raft::linkage::connect_components(handle, connected_edges, - X, color, m, n); + X, color, m, n, reduction_op); mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), connected_edges.vals(), connected_edges.nnz); @@ -96,6 +96,9 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, connected_edges.nnz, indptr2.data(), m + 1, d_alloc, stream); + raft::print_device_vector("new_rows", connected_edges.rows(), connected_edges.nnz, std::cout); + raft::print_device_vector("new_cols", connected_edges.cols(), connected_edges.nnz, std::cout); + raft::print_device_vector("new_dists", connected_edges.vals(), connected_edges.nnz, std::cout); // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process auto new_mst = raft::mst::mst( @@ -129,39 +132,38 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * @param[in] max_iter maximum iterations to run knn graph connection. This * argument is really just a safeguard against the potential for infinite loops. */ -template +template void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const value_idx *indptr, const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, value_idx *mst_src, value_idx *mst_dst, - value_t *mst_weight, size_t nnz, - mst_epilogue_f epilogue_func, + value_t *mst_weight, value_idx *color, size_t nnz, + mst_epilogue_f epilogue_func, red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded, int max_iter = 10) { auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); - rmm::device_uvector color(m, stream); // We want to have MST initialize colors on first call. auto mst_coo = raft::mst::mst( - handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, + handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, false, true); int iters = 1; int n_components = - linkage::get_n_components(color.data(), m, d_alloc, stream); + linkage::get_n_components(color, m, d_alloc, stream); while (n_components > 1 && iters < max_iter) { printf("Didn't converge. trying again\n"); connect_knn_graph(handle, X, mst_coo, m, n, - color.data(), epilogue_func); + color, epilogue_func, reduction_op); iters++; - n_components = linkage::get_n_components(color.data(), m, d_alloc, stream); + n_components = linkage::get_n_components(color, m, d_alloc, stream); } /** @@ -184,14 +186,14 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); - if (mst_coo.n_edges != m - 1) { - raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, - std::cout); - raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, - std::cout); - raft::print_device_vector("mst_weight", mst_coo.weights.data(), - mst_coo.n_edges, std::cout); - } +// if (mst_coo.n_edges != m - 1) { +// raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, +// std::cout); +// raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, +// std::cout); +// raft::print_device_vector("mst_weight", mst_coo.weights.data(), +// mst_coo.n_edges, std::cout); +// } RAFT_EXPECTS(mst_coo.n_edges == m - 1, "n_edges should be %d but was %d. This" diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index 4a744bc048..63f36b1ff8 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -77,10 +77,12 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, /** * 2. Construct MST, sorted by weights */ + rmm::device_uvector color(m, stream); + raft::linkage::FixConnectivitiesRedOp op(color.data(), m); detail::build_sorted_mst( handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, - mst_rows.data(), mst_cols.data(), mst_data.data(), indices.size(), - MSTEpilogueNoOp(), metric); + mst_rows.data(), mst_cols.data(), mst_data.data(), color.data(), indices.size(), + MSTEpilogueNoOp(), op, metric); pw_dists.release(); diff --git a/cpp/include/raft/sparse/selection/connect_components.cuh b/cpp/include/raft/sparse/selection/connect_components.cuh index 386f4f1830..6b76d210b1 100644 --- a/cpp/include/raft/sparse/selection/connect_components.cuh +++ b/cpp/include/raft/sparse/selection/connect_components.cuh @@ -200,23 +200,22 @@ struct LookupColorOp { * @param[in] d_alloc device allocator to use * @param[in] stream cuda stream for which to order cuda operations */ -template +template void perform_1nn(cub::KeyValuePair *kvp, value_idx *nn_colors, value_idx *colors, const value_t *X, size_t n_rows, size_t n_cols, std::shared_ptr d_alloc, - cudaStream_t stream) { + cudaStream_t stream, red_op reduction_op) { rmm::device_uvector workspace(n_rows, stream); rmm::device_uvector x_norm(n_rows, stream); raft::linalg::rowNorm(x_norm.data(), X, n_cols, n_rows, raft::linalg::L2Norm, true, stream); - FixConnectivitiesRedOp red_op(colors, n_rows); raft::distance::fusedL2NN, value_idx>(kvp, X, X, x_norm.data(), x_norm.data(), n_rows, n_rows, n_cols, workspace.data(), - red_op, red_op, true, true, stream); + reduction_op, reduction_op, true, true, stream); LookupColorOp extract_colors_op(colors); thrust::transform(thrust::cuda::par.on(stream), kvp, kvp + n_rows, nn_colors, @@ -318,11 +317,12 @@ void min_components_by_color(raft::sparse::COO &coo, * @param[in] n_rows number of rows in X * @param[in] n_cols number of cols in X */ -template +template void connect_components(const raft::handle_t &handle, raft::sparse::COO &out, const value_t *X, const value_idx *orig_colors, size_t n_rows, size_t n_cols, + red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); @@ -352,7 +352,7 @@ void connect_components(const raft::handle_t &handle, rmm::device_uvector src_indices(n_rows, stream); perform_1nn(temp_inds_dists.data(), nn_colors.data(), colors.data(), X, - n_rows, n_cols, d_alloc, stream); + n_rows, n_cols, d_alloc, stream, reduction_op); /** * Sort data points by color (neighbors are not sorted) From 26dff4d38a666b3ca13797548ce780f4dc299c07 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 19 May 2021 17:37:31 -0400 Subject: [PATCH 20/31] Fixing style --- .../raft/sparse/hierarchy/detail/mst.cuh | 48 ++++++++++--------- .../raft/sparse/hierarchy/single_linkage.hpp | 4 +- .../raft/sparse/mst/detail/mst_kernels.cuh | 19 ++++---- .../raft/sparse/mst/detail/mst_solver_inl.cuh | 4 +- cpp/include/raft/sparse/mst/mst_solver.cuh | 2 +- .../sparse/selection/connect_components.cuh | 9 ++-- 6 files changed, 44 insertions(+), 42 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 19c21cb868..1ec2d6716f 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -73,7 +73,8 @@ void merge_msts(raft::Graph_COO &coo1, * @param[inout] color the color labels array returned from the mst invocation * @return updated MST edge list */ -template +template void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, @@ -85,8 +86,8 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::sparse::COO connected_edges(d_alloc, stream); - raft::linkage::connect_components(handle, connected_edges, - X, color, m, n, reduction_op); + raft::linkage::connect_components( + handle, connected_edges, X, color, m, n, reduction_op); mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), connected_edges.vals(), connected_edges.nnz); @@ -96,9 +97,12 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, connected_edges.nnz, indptr2.data(), m + 1, d_alloc, stream); - raft::print_device_vector("new_rows", connected_edges.rows(), connected_edges.nnz, std::cout); - raft::print_device_vector("new_cols", connected_edges.cols(), connected_edges.nnz, std::cout); - raft::print_device_vector("new_dists", connected_edges.vals(), connected_edges.nnz, std::cout); + raft::print_device_vector("new_rows", connected_edges.rows(), + connected_edges.nnz, std::cout); + raft::print_device_vector("new_cols", connected_edges.cols(), + connected_edges.nnz, std::cout); + raft::print_device_vector("new_dists", connected_edges.vals(), + connected_edges.nnz, std::cout); // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process auto new_mst = raft::mst::mst( @@ -132,7 +136,8 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * @param[in] max_iter maximum iterations to run knn graph connection. This * argument is really just a safeguard against the potential for infinite loops. */ -template +template void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const value_idx *indptr, const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, @@ -145,21 +150,18 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, auto d_alloc = handle.get_device_allocator(); auto stream = handle.get_stream(); - // We want to have MST initialize colors on first call. auto mst_coo = raft::mst::mst( - handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, - false, true); - + handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, false, + true); int iters = 1; - int n_components = - linkage::get_n_components(color, m, d_alloc, stream); + int n_components = linkage::get_n_components(color, m, d_alloc, stream); while (n_components > 1 && iters < max_iter) { printf("Didn't converge. trying again\n"); - connect_knn_graph(handle, X, mst_coo, m, n, - color, epilogue_func, reduction_op); + connect_knn_graph(handle, X, mst_coo, m, n, color, + epilogue_func, reduction_op); iters++; @@ -186,14 +188,14 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); -// if (mst_coo.n_edges != m - 1) { -// raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, -// std::cout); -// raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, -// std::cout); -// raft::print_device_vector("mst_weight", mst_coo.weights.data(), -// mst_coo.n_edges, std::cout); -// } + // if (mst_coo.n_edges != m - 1) { + // raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, + // std::cout); + // raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, + // std::cout); + // raft::print_device_vector("mst_weight", mst_coo.weights.data(), + // mst_coo.n_edges, std::cout); + // } RAFT_EXPECTS(mst_coo.n_edges == m - 1, "n_edges should be %d but was %d. This" diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index 63f36b1ff8..5710853848 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -81,8 +81,8 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, raft::linkage::FixConnectivitiesRedOp op(color.data(), m); detail::build_sorted_mst( handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, - mst_rows.data(), mst_cols.data(), mst_data.data(), color.data(), indices.size(), - MSTEpilogueNoOp(), op, metric); + mst_rows.data(), mst_cols.data(), mst_data.data(), color.data(), + indices.size(), MSTEpilogueNoOp(), op, metric); pw_dists.release(); diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index d9693583cc..259cb8ae7f 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -282,13 +282,10 @@ __global__ void final_color_indices(const vertex_t v, const vertex_t* color, // Alterate the weights, make all undirected edge weight unique while keeping Wuv == Wvu // Consider using curand device API instead of precomputed random_values array template -__global__ void alteration_kernel(const vertex_t v, const edge_t e, - const edge_t* offsets, - const vertex_t* indices, - const weight_t* weights, double max, - double* random_values, - double* altered_weights, int alpha, - bool use_alpha) { +__global__ void alteration_kernel( + const vertex_t v, const edge_t e, const edge_t* offsets, + const vertex_t* indices, const weight_t* weights, double max, + double* random_values, double* altered_weights, int alpha, bool use_alpha) { auto row = get_1D_idx(); if (row < v) { auto row_begin = offsets[row]; @@ -297,7 +294,7 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, auto column = indices[i]; // doing the later step explicity in double for precision altered_weights[i] = - weights[i] + max * (random_values[row] + random_values[column]); + weights[i] + max * (random_values[row] + random_values[column]); auto print = false; if (row == 293 && column == 1276) print = true; @@ -308,7 +305,11 @@ __global__ void alteration_kernel(const vertex_t v, const edge_t e, if (row == 1276 && column == 1686) print = true; if (print) { - printf("row: %d, col: %d, alt: %lf, weight: %lf, max: %lf, randr: %lf, randl: %lf\n", row, column, altered_weights[i], weights[i], max, random_values[row], random_values[column]); + printf( + "row: %d, col: %d, alt: %lf, weight: %lf, max: %lf, randr: %lf, " + "randl: %lf\n", + row, column, altered_weights[i], weights[i], max, random_values[row], + random_values[column]); } // if (use_alpha) { 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 693f44b033..e3ee7df425 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -168,9 +168,9 @@ MST_solver::solve() { // raft::print_device_vector("temp_dst", temp_dst.data().get(), v, std::cout); // raft::print_device_vector("temp_weights", temp_weights.data().get(), v, std::cout); - auto curr_mst_edge_count = mst_edge_count[0]; - std::cout << "edge count: " << curr_mst_edge_count << ", expected: " << n_expected_edges << std::endl; + std::cout << "edge count: " << curr_mst_edge_count + << ", expected: " << n_expected_edges << std::endl; RAFT_EXPECTS(curr_mst_edge_count <= n_expected_edges, "Number of edges found by MST is invalid. This may be due to " "loss in precision. Try increasing precision of weights."); diff --git a/cpp/include/raft/sparse/mst/mst_solver.cuh b/cpp/include/raft/sparse/mst/mst_solver.cuh index 2bac221f31..dc47adf50c 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -70,7 +70,7 @@ class MST_solver { vertex_t* color_index; // represent each supervertex as a color rmm::device_vector min_edge_color; // minimum incident edge weight per color - rmm::device_vector new_mst_edge; // new minimum edge per vertex + rmm::device_vector new_mst_edge; // new minimum edge per vertex rmm::device_vector altered_weights; // weights to be used for mst rmm::device_vector mst_edge_count; // total number of edges added after every iteration diff --git a/cpp/include/raft/sparse/selection/connect_components.cuh b/cpp/include/raft/sparse/selection/connect_components.cuh index 6b76d210b1..8aae90f1d8 100644 --- a/cpp/include/raft/sparse/selection/connect_components.cuh +++ b/cpp/include/raft/sparse/selection/connect_components.cuh @@ -213,9 +213,9 @@ void perform_1nn(cub::KeyValuePair *kvp, true, stream); raft::distance::fusedL2NN, - value_idx>(kvp, X, X, x_norm.data(), x_norm.data(), - n_rows, n_rows, n_cols, workspace.data(), - reduction_op, reduction_op, true, true, stream); + value_idx>( + kvp, X, X, x_norm.data(), x_norm.data(), n_rows, n_rows, n_cols, + workspace.data(), reduction_op, reduction_op, true, true, stream); LookupColorOp extract_colors_op(colors); thrust::transform(thrust::cuda::par.on(stream), kvp, kvp + n_rows, nn_colors, @@ -321,8 +321,7 @@ template void connect_components(const raft::handle_t &handle, raft::sparse::COO &out, const value_t *X, const value_idx *orig_colors, - size_t n_rows, size_t n_cols, - red_op reduction_op, + size_t n_rows, size_t n_cols, red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); From 55f274c1e76ceecd9f4e6df5540772325e8984cb Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 19 May 2021 17:56:11 -0400 Subject: [PATCH 21/31] Removing unused epilogue from mst --- .../raft/sparse/hierarchy/detail/mst.cuh | 26 +++---------------- 1 file changed, 4 insertions(+), 22 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 1ec2d6716f..9e2026b363 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -73,12 +73,11 @@ void merge_msts(raft::Graph_COO &coo1, * @param[inout] color the color labels array returned from the mst invocation * @return updated MST edge list */ -template +template void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::Graph_COO &msf, size_t m, size_t n, value_idx *color, - mst_epilogue_f mst_epilogue_func, red_op reduction_op, + red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded) { auto d_alloc = handle.get_device_allocator(); @@ -97,12 +96,6 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, connected_edges.nnz, indptr2.data(), m + 1, d_alloc, stream); - raft::print_device_vector("new_rows", connected_edges.rows(), - connected_edges.nnz, std::cout); - raft::print_device_vector("new_cols", connected_edges.cols(), - connected_edges.nnz, std::cout); - raft::print_device_vector("new_dists", connected_edges.vals(), - connected_edges.nnz, std::cout); // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process auto new_mst = raft::mst::mst( @@ -136,14 +129,13 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, * @param[in] max_iter maximum iterations to run knn graph connection. This * argument is really just a safeguard against the potential for infinite loops. */ -template +template void build_sorted_mst(const raft::handle_t &handle, const value_t *X, const value_idx *indptr, const value_idx *indices, const value_t *pw_dists, size_t m, size_t n, value_idx *mst_src, value_idx *mst_dst, value_t *mst_weight, value_idx *color, size_t nnz, - mst_epilogue_f epilogue_func, red_op reduction_op, + red_op reduction_op, raft::distance::DistanceType metric = raft::distance::DistanceType::L2SqrtExpanded, int max_iter = 10) { @@ -159,7 +151,6 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, int n_components = linkage::get_n_components(color, m, d_alloc, stream); while (n_components > 1 && iters < max_iter) { - printf("Didn't converge. trying again\n"); connect_knn_graph(handle, X, mst_coo, m, n, color, epilogue_func, reduction_op); @@ -188,15 +179,6 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); - // if (mst_coo.n_edges != m - 1) { - // raft::print_device_vector("mst_src", mst_coo.src.data(), mst_coo.n_edges, - // std::cout); - // raft::print_device_vector("mst_dst", mst_coo.dst.data(), mst_coo.n_edges, - // std::cout); - // raft::print_device_vector("mst_weight", mst_coo.weights.data(), - // mst_coo.n_edges, std::cout); - // } - RAFT_EXPECTS(mst_coo.n_edges == m - 1, "n_edges should be %d but was %d. This" "could be an indication of duplicate edges returned from the" From eb69413dcd7340b440032be3ac96c981ae24f80a Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 19 May 2021 18:17:19 -0400 Subject: [PATCH 22/31] Removing mst epilogue functor --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 5 +---- cpp/include/raft/sparse/hierarchy/single_linkage.hpp | 2 +- 2 files changed, 2 insertions(+), 5 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 9e2026b363..441eff4570 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -88,9 +88,6 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, raft::linkage::connect_components( handle, connected_edges, X, color, m, n, reduction_op); - mst_epilogue_func(handle, connected_edges.rows(), connected_edges.cols(), - connected_edges.vals(), connected_edges.nnz); - rmm::device_uvector indptr2(m + 1, stream); raft::sparse::convert::sorted_coo_to_csr(connected_edges.rows(), connected_edges.nnz, indptr2.data(), @@ -152,7 +149,7 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, while (n_components > 1 && iters < max_iter) { connect_knn_graph(handle, X, mst_coo, m, n, color, - epilogue_func, reduction_op); + reduction_op); iters++; diff --git a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp index 5710853848..01a033945c 100644 --- a/cpp/include/raft/sparse/hierarchy/single_linkage.hpp +++ b/cpp/include/raft/sparse/hierarchy/single_linkage.hpp @@ -82,7 +82,7 @@ void single_linkage(const raft::handle_t &handle, const value_t *X, size_t m, detail::build_sorted_mst( handle, X, indptr.data(), indices.data(), pw_dists.data(), m, n, mst_rows.data(), mst_cols.data(), mst_data.data(), color.data(), - indices.size(), MSTEpilogueNoOp(), op, metric); + indices.size(), op, metric); pw_dists.release(); From 50d1cdc100bc31e649399ad1ab883b8334a94703 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 19 May 2021 22:27:36 -0400 Subject: [PATCH 23/31] Getting test to build --- cpp/test/sparse/connect_components.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/test/sparse/connect_components.cu b/cpp/test/sparse/connect_components.cu index 68db00374c..451b41e327 100644 --- a/cpp/test/sparse/connect_components.cu +++ b/cpp/test/sparse/connect_components.cu @@ -99,9 +99,10 @@ class ConnectComponentsTest : public ::testing::TestWithParam< /** * 3. connect_components to fix connectivities */ + raft::linkage::FixConnectivitiesRedOp red_op(colors.data(), params.n_row); raft::linkage::connect_components( handle, out_edges, data.data(), colors.data(), params.n_row, - params.n_col); + params.n_col, red_op); /** * Construct final edge list From c654156cbd0850d251ac1c5c64fd006732f4ea29 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 21 May 2021 20:08:24 -0400 Subject: [PATCH 24/31] Removing mstepiloguenoop since it's no longer being used --- cpp/include/raft/sparse/hierarchy/common.h | 8 -------- 1 file changed, 8 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/common.h b/cpp/include/raft/sparse/hierarchy/common.h index f5a2b26efe..29f541498b 100644 --- a/cpp/include/raft/sparse/hierarchy/common.h +++ b/cpp/include/raft/sparse/hierarchy/common.h @@ -14,8 +14,6 @@ * limitations under the License. */ -#include - #pragma once namespace raft { @@ -23,12 +21,6 @@ namespace hierarchy { enum LinkageDistance { PAIRWISE = 0, KNN_GRAPH = 1 }; -template -struct MSTEpilogueNoOp { - void operator()(const raft::handle_t &handle, value_idx *coo_rows, - value_idx *coo_cols, value_t *coo_data, value_idx nnz) {} -}; - /** * Simple POCO for consolidating linkage results. This closely * mirrors the trained instance variables populated in From 136529a91a439681c61dc381663240c1179cc29b Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 24 May 2021 14:00:42 -0700 Subject: [PATCH 25/31] another template param for weight alteration --- .../raft/sparse/hierarchy/detail/mst.cuh | 10 +-- .../raft/sparse/mst/detail/mst_kernels.cuh | 46 +++++------ .../raft/sparse/mst/detail/mst_solver_inl.cuh | 81 ++++++++++--------- cpp/include/raft/sparse/mst/mst.cuh | 9 ++- cpp/include/raft/sparse/mst/mst_solver.cuh | 15 ++-- cpp/test/mst.cu | 20 ++--- 6 files changed, 89 insertions(+), 92 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 91c4e1642d..8ffcfe0f2b 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -116,7 +116,7 @@ void connect_knn_graph(const raft::handle_t &handle, const value_t *X, // On the second call, we hand the MST the original colors // and the new set of edges and let it restart the optimization process - auto new_mst = raft::mst::mst( + auto new_mst = raft::mst::mst( handle, indptr2.data(), connected_edges.cols(), connected_edges.vals(), m, connected_edges.nnz, color, stream, false, false); @@ -164,7 +164,7 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, rmm::device_uvector color(m, stream); // We want to have MST initialize colors on first call. - auto mst_coo = raft::mst::mst( + auto mst_coo = raft::mst::mst( handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, false, true); @@ -201,12 +201,6 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); - RAFT_EXPECTS(mst_coo.n_edges == m - 1, - "n_edges should be %d but was %d. This" - "could be an indication of duplicate edges returned from the" - "MST or symmetrization stage.", - m - 1, mst_coo.n_edges); - sort_coo_by_data(mst_coo.src.data(), mst_coo.dst.data(), mst_coo.weights.data(), mst_coo.n_edges, stream); diff --git a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh index 8a44a2cc38..f0d30b0cb7 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_kernels.cuh @@ -27,22 +27,22 @@ namespace raft { namespace mst { namespace detail { -template +template __global__ void kernel_min_edge_per_vertex( - const edge_t* offsets, const vertex_t* indices, const weight_t* weights, + const edge_t* offsets, const vertex_t* indices, const alteration_t* weights, const vertex_t* color, const vertex_t* color_index, edge_t* new_mst_edge, - const bool* mst_edge, weight_t* min_edge_color, const vertex_t v) { + const bool* mst_edge, alteration_t* min_edge_color, 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__ alteration_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_edge_weight[lane_id] = std::numeric_limits::max(); min_color[lane_id] = std::numeric_limits::max(); __syncthreads(); @@ -61,7 +61,7 @@ __global__ void kernel_min_edge_per_vertex( // 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]; + alteration_t curr_edge_weight = weights[e]; vertex_t successor_color_idx = color_index[indices[e]]; vertex_t successor_color = color[successor_color_idx]; @@ -92,7 +92,7 @@ __global__ void kernel_min_edge_per_vertex( // min edge may now be found in first thread if (lane_id == 0) { - if (min_edge_weight[0] != std::numeric_limits::max()) { + if (min_edge_weight[0] != std::numeric_limits::max()) { new_mst_edge[warp_id] = min_edge_index[0]; // atomically set min edge per color @@ -102,12 +102,13 @@ __global__ void kernel_min_edge_per_vertex( } } -template +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, + const alteration_t* altered_weights, vertex_t* temp_src, vertex_t* temp_dst, + weight_t* temp_weights, const alteration_t* min_edge_color, const vertex_t v, bool symmetrize_output) { auto tid = get_1D_idx(); if (tid < v) { @@ -119,7 +120,7 @@ __global__ void min_edge_per_supervertex( // 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]; + alteration_t vertex_weight = altered_weights[edge_idx]; bool add_edge = false; if (min_edge_color[vertex_color] == vertex_weight) { @@ -281,31 +282,22 @@ __global__ void final_color_indices(const vertex_t v, const vertex_t* color, // Alterate the weights, make all undirected edge weight unique while keeping Wuv == Wvu // Consider using curand device API instead of precomputed random_values array -template +template __global__ void alteration_kernel(const vertex_t v, const edge_t e, const edge_t* offsets, const vertex_t* indices, - const weight_t* weights, double max, - weight_t* random_values, - weight_t* altered_weights, int alpha, - bool use_alpha) { + const weight_t* weights, alteration_t max, + alteration_t* random_values, + alteration_t* altered_weights) { auto row = get_1D_idx(); if (row < v) { auto row_begin = offsets[row]; auto row_end = offsets[row + 1]; for (auto i = row_begin; i < row_end; i++) { auto column = indices[i]; - // doing the later step explicity in double for precision - if (use_alpha) { - altered_weights[i] = - alpha * weights[i] + alpha * max * - (static_cast(random_values[row]) + - static_cast(random_values[column])); - } else { - altered_weights[i] = - weights[i] + max * (static_cast(random_values[row]) + - static_cast(random_values[column])); - } + altered_weights[i] = + weights[i] + max * (random_values[row] + random_values[column]); } } } 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 4222c46b2a..d68c89ba27 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -22,7 +22,6 @@ #include "mst_kernels.cuh" #include "utils.cuh" - #include #include #include @@ -50,12 +49,13 @@ inline curandStatus_t curand_generate_uniformX(curandGenerator_t generator, return curandGenerateUniformDouble(generator, outputPtr, n); } -template -MST_solver::MST_solver( +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_, - bool symmetrize_output_, bool initialize_colors_, int iterations_, int alpha_) + bool symmetrize_output_, bool initialize_colors_, int iterations_) : handle(handle_), offsets(offsets_), indices(indices_), @@ -77,8 +77,7 @@ MST_solver::MST_solver( stream(stream_), symmetrize_output(symmetrize_output_), initialize_colors(initialize_colors_), - iterations(iterations_), - alpha(alpha_) { + iterations(iterations_) { max_blocks = handle_.get_device_properties().maxGridSize[0]; max_threads = handle_.get_device_properties().maxThreadsPerBlock; sm_count = handle_.get_device_properties().multiProcessorCount; @@ -94,9 +93,10 @@ MST_solver::MST_solver( thrust::sequence(policy, next_color.begin(), next_color.end(), 0); } -template +template raft::Graph_COO -MST_solver::solve() { +MST_solver::solve() { RAFT_EXPECTS(v > 0, "0 vertices"); RAFT_EXPECTS(e > 0, "0 edges"); RAFT_EXPECTS(offsets != nullptr, "Null offsets."); @@ -155,11 +155,12 @@ MST_solver::solve() { timer3 += duration_us(stop - start); #endif - RAFT_EXPECTS(mst_edge_count[0] <= n_expected_edges, + auto curr_mst_edge_count = mst_edge_count[0]; + RAFT_EXPECTS(curr_mst_edge_count <= n_expected_edges, "Number of edges found by MST is invalid. This may be due to " "loss in precision. Try increasing precision of weights."); - if (prev_mst_edge_count[0] == mst_edge_count[0]) { + if (curr_mst_edge_count == prev_mst_edge_count[0]) { #ifdef MST_TIME std::cout << "Iterations: " << i << std::endl; std::cout << timer0 << "," << timer1 << "," << timer2 << "," << timer3 @@ -217,8 +218,10 @@ struct alteration_functor { }; // Compute the uper bound for the alteration -template -double MST_solver::alteration_max() { +template +alteration_t +MST_solver::alteration_max() { auto policy = rmm::exec_policy(stream); rmm::device_vector tmp(e); thrust::device_ptr weights_ptr(weights); @@ -238,21 +241,22 @@ double MST_solver::alteration_max() { auto max = thrust::transform_reduce(policy, begin, end, alteration_functor(), init, thrust::minimum()); - return max / static_cast(2); + return max / static_cast(2); } // Compute the alteration to make all undirected edge weight unique // Preserves weights order -template -void MST_solver::alteration() { +template +void MST_solver::alteration() { auto nthreads = std::min(v, max_threads); auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); // maximum alteration that does not change realtive weights order - double max = alteration_max(); + alteration_t max = alteration_max(); // pool of rand values - rmm::device_vector rand_values(v); + rmm::device_vector rand_values(v); // Random number generator curandGenerator_t randGen; @@ -267,18 +271,17 @@ void MST_solver::alteration() { RAFT_EXPECTS(curand_status == CURAND_STATUS_SUCCESS, "MST: CURAND cleanup failed"); - bool use_alpha = max < 1e-3 && sizeof(weight_t) == 4; - //Alterate the weights, make all undirected edge weight unique while keeping Wuv == Wvu detail::alteration_kernel<<>>( v, e, offsets, indices, weights, max, rand_values.data().get(), - altered_weights.data().get(), alpha, use_alpha); + altered_weights.data().get()); } // updates colors of vertices by propagating the lower color to the higher -template -void MST_solver::label_prop(vertex_t* mst_src, - vertex_t* mst_dst) { +template +void MST_solver::label_prop( + vertex_t* mst_src, vertex_t* mst_dst) { // update the colors of both ends its until there is no change in colors thrust::host_vector curr_mst_edge_count = mst_edge_count; @@ -315,11 +318,13 @@ void MST_solver::label_prop(vertex_t* mst_src, } // Finds the minimum edge from each vertex to the lowest color -template -void MST_solver::min_edge_per_vertex() { +template +void MST_solver::min_edge_per_vertex() { auto policy = rmm::exec_policy(stream); thrust::fill(policy, min_edge_color.begin(), min_edge_color.end(), - std::numeric_limits::max()); + std::numeric_limits::max()); thrust::fill(policy, new_mst_edge.begin(), new_mst_edge.end(), std::numeric_limits::max()); @@ -328,8 +333,8 @@ void MST_solver::min_edge_per_vertex() { vertex_t* color_ptr = color.data().get(); edge_t* new_mst_edge_ptr = new_mst_edge.data().get(); bool* mst_edge_ptr = mst_edge.data().get(); - weight_t* min_edge_color_ptr = min_edge_color.data().get(); - weight_t* altered_weights_ptr = altered_weights.data().get(); + alteration_t* min_edge_color_ptr = min_edge_color.data().get(); + alteration_t* altered_weights_ptr = altered_weights.data().get(); detail::kernel_min_edge_per_vertex<<>>( offsets, indices, altered_weights_ptr, color_ptr, color_index, @@ -337,8 +342,10 @@ 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() { +template +void MST_solver::min_edge_per_supervertex() { auto nthreads = std::min(v, max_threads); auto nblocks = std::min((v + nthreads - 1) / nthreads, max_blocks); @@ -349,8 +356,8 @@ void MST_solver::min_edge_per_supervertex() { vertex_t* color_ptr = color.data().get(); edge_t* new_mst_edge_ptr = new_mst_edge.data().get(); bool* mst_edge_ptr = mst_edge.data().get(); - weight_t* min_edge_color_ptr = min_edge_color.data().get(); - weight_t* altered_weights_ptr = altered_weights.data().get(); + alteration_t* min_edge_color_ptr = min_edge_color.data().get(); + alteration_t* altered_weights_ptr = altered_weights.data().get(); vertex_t* temp_src_ptr = temp_src.data().get(); vertex_t* temp_dst_ptr = temp_dst.data().get(); weight_t* temp_weights_ptr = temp_weights.data().get(); @@ -370,8 +377,9 @@ void MST_solver::min_edge_per_supervertex() { } } -template -void MST_solver::check_termination() { +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); @@ -394,8 +402,9 @@ struct new_edges_functor { } }; -template -void MST_solver::append_src_dst_pair( +template +void MST_solver::append_src_dst_pair( vertex_t* mst_src, vertex_t* mst_dst, weight_t* mst_weights) { auto policy = rmm::exec_policy(stream); diff --git a/cpp/include/raft/sparse/mst/mst.cuh b/cpp/include/raft/sparse/mst/mst.cuh index abe74f3fe3..10c981445e 100644 --- a/cpp/include/raft/sparse/mst/mst.cuh +++ b/cpp/include/raft/sparse/mst/mst.cuh @@ -22,15 +22,16 @@ namespace raft { namespace mst { -template +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, bool symmetrize_output = true, - bool initialize_colors = true, int iterations = 0, int alpha = 1e6) { - MST_solver mst_solver( + bool initialize_colors = true, int iterations = 0) { + MST_solver mst_solver( handle, offsets, indices, weights, v, e, color, stream, symmetrize_output, - initialize_colors, iterations, alpha); + initialize_colors, iterations); 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 e9e0749cdf..833882ea0d 100644 --- a/cpp/include/raft/sparse/mst/mst_solver.cuh +++ b/cpp/include/raft/sparse/mst/mst_solver.cuh @@ -36,14 +36,15 @@ struct Graph_COO { namespace mst { -template +template class MST_solver { public: 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_, bool symmetrize_output_, - bool initialize_colors_, int iterations_, int alpha_); + bool initialize_colors_, int iterations_); raft::Graph_COO solve(); @@ -54,7 +55,6 @@ class MST_solver { cudaStream_t stream; bool symmetrize_output, initialize_colors; int iterations; - int alpha; //CSR const edge_t* offsets; @@ -68,10 +68,11 @@ class MST_solver { vertex_t sm_count; vertex_t* color_index; // represent each supervertex as a color - rmm::device_vector + rmm::device_vector min_edge_color; // minimum incident edge weight per color - rmm::device_vector new_mst_edge; // new minimum edge per vertex - rmm::device_vector altered_weights; // weights to be used for mst + rmm::device_vector new_mst_edge; // new minimum edge per vertex + rmm::device_vector + altered_weights; // weights to be used for mst rmm::device_vector mst_edge_count; // total number of edges added after every iteration rmm::device_vector @@ -91,7 +92,7 @@ class MST_solver { void min_edge_per_supervertex(); void check_termination(); void alteration(); - double alteration_max(); + alteration_t alteration_max(); void append_src_dst_pair(vertex_t* mst_src, vertex_t* mst_dst, weight_t* mst_weights); }; diff --git a/cpp/test/mst.cu b/cpp/test/mst.cu index a9119acbad..215c6f6548 100644 --- a/cpp/test/mst.cu +++ b/cpp/test/mst.cu @@ -136,14 +136,14 @@ class MSTTest vertex_t *color_ptr = thrust::raw_pointer_cast(color.data()); if (iterations == 0) { - MST_solver symmetric_solver( + MST_solver symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, true, 0, 0); + true, true, 0); auto symmetric_result = symmetric_solver.solve(); - MST_solver non_symmetric_solver( + MST_solver non_symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - false, true, 0, 0); + false, true, 0); auto non_symmetric_result = non_symmetric_solver.solve(); EXPECT_LE(symmetric_result.n_edges, 2 * v - 2); @@ -152,14 +152,14 @@ class MSTTest return std::make_pair(std::move(symmetric_result), std::move(non_symmetric_result)); } else { - MST_solver intermediate_solver( + MST_solver intermediate_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, true, iterations, 0); + true, true, iterations); auto intermediate_result = intermediate_solver.solve(); - MST_solver symmetric_solver( + MST_solver symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - true, false, 0, 0); + true, false, 0); auto symmetric_result = symmetric_solver.solve(); // symmetric_result.n_edges += intermediate_result.n_edges; @@ -180,9 +180,9 @@ class MSTTest intermediate_result.n_edges, handle.get_stream()); symmetric_result.n_edges = total_edge_size; - MST_solver non_symmetric_solver( + MST_solver non_symmetric_solver( handle, offsets, indices, weights, v, e, color_ptr, handle.get_stream(), - false, true, 0, 0); + false, true, 0); auto non_symmetric_result = non_symmetric_solver.solve(); EXPECT_LE(symmetric_result.n_edges, 2 * v - 2); From d7e93d93e864a528e7a37f90318dcaef4c54b604 Mon Sep 17 00:00:00 2001 From: divyegala Date: Mon, 24 May 2021 14:30:25 -0700 Subject: [PATCH 26/31] renaming confusing variable name --- 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 ad1b2c6f8b..c5ba4fcb4f 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -122,9 +122,9 @@ MST_solver::solve() { timer0 = duration_us(stop - start); #endif - auto n_expected_edges = symmetrize_output ? 2 * v - 2 : v - 1; + auto max_mst_edges = symmetrize_output ? 2 * v - 2 : v - 1; - Graph_COO mst_result(n_expected_edges, stream); + Graph_COO mst_result(max_mst_edges, stream); // Boruvka original formulation says "while more than 1 supervertex remains" // Here we adjust it to support disconnected components (spanning forest) @@ -161,7 +161,7 @@ MST_solver::solve() { #endif auto curr_mst_edge_count = mst_edge_count[0]; - RAFT_EXPECTS(curr_mst_edge_count <= n_expected_edges, + RAFT_EXPECTS(curr_mst_edge_count <= max_mst_edges, "Number of edges found by MST is invalid. This may be due to " "loss in precision. Try increasing precision of weights."); From beea02053c6768f55515c62e53e5a1d8ed6fa879 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 25 May 2021 19:38:00 -0700 Subject: [PATCH 27/31] merging mst template PR --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 14 +++++--------- cpp/test/sparse/connect_components.cu | 7 ++++--- 2 files changed, 9 insertions(+), 12 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index 072561e494..765a5ad77f 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -141,8 +141,8 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, // We want to have MST initialize colors on first call. auto mst_coo = raft::mst::mst( - handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, - false, true); + handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, false, + true); int iters = 1; int n_components = linkage::get_n_components(color, m, d_alloc, stream); @@ -176,13 +176,9 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, " or increase 'max_iter'", max_iter); - sort_coo_by_data(mst_coo.src.data(), mst_coo.dst.data(), - mst_coo.weights.data(), mst_coo.n_edges, stream); - - // TODO: be nice if we could pass these directly into the MST - mst_src.resize(mst_coo.n_edges, stream); - mst_dst.resize(mst_coo.n_edges, stream); - mst_weight.resize(mst_coo.n_edges, stream); + raft::sparse::op::coo_sort_by_weight(mst_coo.src.data(), mst_coo.dst.data(), + mst_coo.weights.data(), mst_coo.n_edges, + stream); raft::copy_async(mst_src, mst_coo.src.data(), mst_coo.n_edges, stream); raft::copy_async(mst_dst, mst_coo.dst.data(), mst_coo.n_edges, stream); diff --git a/cpp/test/sparse/connect_components.cu b/cpp/test/sparse/connect_components.cu index 451b41e327..31173d8f61 100644 --- a/cpp/test/sparse/connect_components.cu +++ b/cpp/test/sparse/connect_components.cu @@ -99,10 +99,11 @@ class ConnectComponentsTest : public ::testing::TestWithParam< /** * 3. connect_components to fix connectivities */ - raft::linkage::FixConnectivitiesRedOp red_op(colors.data(), params.n_row); + raft::linkage::FixConnectivitiesRedOp red_op( + colors.data(), params.n_row); raft::linkage::connect_components( - handle, out_edges, data.data(), colors.data(), params.n_row, - params.n_col, red_op); + handle, out_edges, data.data(), colors.data(), params.n_row, params.n_col, + red_op); /** * Construct final edge list From 86cbf4239966246f61227cb0b24af1f926664f53 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 25 May 2021 19:39:58 -0700 Subject: [PATCH 28/31] removing unnecessary comments --- cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh | 7 ------- 1 file changed, 7 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 7a5df7197d..c5ba4fcb4f 100644 --- a/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh +++ b/cpp/include/raft/sparse/mst/detail/mst_solver_inl.cuh @@ -112,13 +112,6 @@ MST_solver::solve() { auto start = Clock::now(); #endif - // int start_1507, end_1507; - // raft::update_host(&start_1507, offsets + 1507, 1, stream); - // raft::update_host(&end_1507, offsets + 1508, 1, stream); - - // raft::print_device_vector("1507 indices", indices + start_1507, end_1507 - start_1507, std::cout); - // raft::print_device_vector("1507 weights", weights + start_1507, end_1507 - start_1507, std::cout); - // Alterating the weights // this is done by identifying the lowest cost edge weight gap that is not 0, call this theta. // For each edge, add noise that is less than theta. That is, generate a random number in the range [0.0, theta) and add it to each edge weight. From eb92e26e695662019ddaccf489219cdb7b6fee4d Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 26 May 2021 20:29:22 -0400 Subject: [PATCH 29/31] Review feedback --- .../sparse/hierarchy/detail/agglomerative.cuh | 66 +++---------------- 1 file changed, 9 insertions(+), 57 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index 486598a552..c5a53ac55f 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -73,9 +73,15 @@ class UnionFind { }; /** - * Standard single-threaded agglomerative labeling on host. This should work - * well for smaller sizes of m. This is a C++ port of the original reference - * implementation of HDBSCAN. + * Agglomerative labeling on host. This has not been found to be a bottleneck + * in the algorithm. A parallel version of this can be done using a parallel + * variant of Kruskal's MST algorithm + * (ref http://cucis.ece.northwestern.edu/publications/pdf/HenPat12.pdf), + * which breaks apart the sorted MST results into overlapping subsets and + * independently runs Kruskal's algorithm on each subset, merging them back + * together into a single hierarchy when complete. Unfortunately, + * this is nontrivial and the speedup wouldn't bit e useful until this + * becomes a bottleneck. * * @tparam value_idx * @tparam value_t @@ -137,60 +143,6 @@ void build_dendrogram_host(const handle_t &handle, const value_idx *rows, raft::update_device(out_delta, out_delta_h.data(), n_edges, stream); } -/** - * Parallel agglomerative labeling. This amounts to a parallel Kruskal's - * MST algorithm, which breaks apart the sorted MST results into overlapping - * subsets and independently runs Kruskal's algorithm on each subset, - * merging them back together into a single hierarchy when complete. - * - * This outputs the same format as the reference HDBSCAN, but as 4 separate - * arrays, rather than a single 2D array. - * - * Reference: http://cucis.ece.northwestern.edu/publications/pdf/HenPat12.pdf - * - * TODO: Investigate potential for the following end-to-end single-hierarchy batching: - * For each of k (independent) batches over the input: - * - Sample n elements from X - * - Compute mutual reachability graph of batch - * - Construct labels from batch - * - * The sampled datasets should have some overlap across batches. This will - * allow for the cluster hierarchies to be merged. Being able to batch - * will reduce the memory cost so that the full n^2 pairwise distances - * don't need to be materialized in memory all at once. - * - * @tparam value_idx - * @tparam value_t - * @param[in] handle the raft handle - * @param[in] rows src edges of the sorted MST - * @param[in] cols dst edges of the sorted MST - * @param[in] nnz the number of edges in the sorted MST - * @param[out] out_src parents of output - * @param[out] out_dst children of output - * @param[out] out_delta distances of output - * @param[out] out_size cluster sizes of output - * @param[in] k_folds number of folds for parallelizing label step - */ -template -void build_dendrogram_device(const handle_t &handle, const value_idx *rows, - const value_idx *cols, const value_t *data, - value_idx nnz, value_idx *children, - value_t *out_delta, value_idx *out_size, - value_idx k_folds) { - ASSERT(k_folds < nnz / 2, "k_folds must be < n_edges / 2"); - /** - * divide (sorted) mst coo into overlapping subsets. Easiest way to do this is to - * break it into k-folds and iterate through two folds at a time. - */ - - // 1. Generate ranges for the overlapping subsets - - // 2. Run union-find in parallel for each pair of folds - - // 3. Sort individual label hierarchies - - // 4. Merge label hierarchies together -} template __global__ void write_levels_kernel(const value_idx *children, From 8b1e344dd375dca737b7b601f2978581685b6a31 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Thu, 27 May 2021 10:57:51 -0400 Subject: [PATCH 30/31] Update cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh Co-authored-by: Dante Gama Dessavre --- cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh index 9c87cd147c..1ac075489a 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/agglomerative.cuh @@ -80,7 +80,7 @@ class UnionFind { * which breaks apart the sorted MST results into overlapping subsets and * independently runs Kruskal's algorithm on each subset, merging them back * together into a single hierarchy when complete. Unfortunately, - * this is nontrivial and the speedup wouldn't bit e useful until this + * this is nontrivial and the speedup wouldn't be useful until this * becomes a bottleneck. * * @tparam value_idx From 193352039ee7f1e87a8daed6c5db08356c49ea26 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Thu, 27 May 2021 11:13:55 -0400 Subject: [PATCH 31/31] Fixing bad merge --- cpp/include/raft/sparse/hierarchy/detail/mst.cuh | 4 ++-- cpp/test/sparse/connect_components.cu | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh index f34b5ac3ed..765a5ad77f 100644 --- a/cpp/include/raft/sparse/hierarchy/detail/mst.cuh +++ b/cpp/include/raft/sparse/hierarchy/detail/mst.cuh @@ -141,8 +141,8 @@ void build_sorted_mst(const raft::handle_t &handle, const value_t *X, // We want to have MST initialize colors on first call. auto mst_coo = raft::mst::mst( - handle, indptr, indices, pw_dists, (value_idx)m, nnz, color.data(), stream, - false, true); + handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, false, + true); int iters = 1; int n_components = linkage::get_n_components(color, m, d_alloc, stream); diff --git a/cpp/test/sparse/connect_components.cu b/cpp/test/sparse/connect_components.cu index 31173d8f61..d98f9de9c3 100644 --- a/cpp/test/sparse/connect_components.cu +++ b/cpp/test/sparse/connect_components.cu @@ -92,7 +92,7 @@ class ConnectComponentsTest : public ::testing::TestWithParam< */ rmm::device_uvector colors(params.n_row, stream); - auto mst_coo = raft::mst::mst( + auto mst_coo = raft::mst::mst( handle, indptr.data(), knn_graph_coo.cols(), knn_graph_coo.vals(), params.n_row, knn_graph_coo.nnz, colors.data(), stream, false, true);