diff --git a/cpp/include/cuml/cluster/hdbscan.hpp b/cpp/include/cuml/cluster/hdbscan.hpp index e3184726bb..52fc1dc855 100644 --- a/cpp/include/cuml/cluster/hdbscan.hpp +++ b/cpp/include/cuml/cluster/hdbscan.hpp @@ -311,7 +311,7 @@ class hdbscan_output : public robust_single_linkage_output { CondensedHierarchy condensed_tree; }; -template class CondensedHierarchy; +template class CondensedHierarchy; /** * Container object for computing and storing intermediate information needed later for computing @@ -387,14 +387,14 @@ class PredictionData { rmm::device_uvector index_into_children; }; -template class PredictionData; +template class PredictionData; void generate_prediction_data(const raft::handle_t& handle, - CondensedHierarchy& condensed_tree, - int* labels, - int* inverse_label_map, + CondensedHierarchy& condensed_tree, + int64_t* labels, + int64_t* inverse_label_map, int n_selected_clusters, - PredictionData& prediction_data); + PredictionData& prediction_data); }; // namespace Common }; // namespace HDBSCAN @@ -427,43 +427,43 @@ void hdbscan(const raft::handle_t& handle, size_t n, ML::distance::DistanceType metric, HDBSCAN::Common::HDBSCANParams& params, - HDBSCAN::Common::hdbscan_output& out, + HDBSCAN::Common::hdbscan_output& out, float* core_dists); void build_condensed_hierarchy(const raft::handle_t& handle, - const int* children, + const int64_t* children, const float* delta, - const int* sizes, + const int64_t* sizes, int min_cluster_size, int n_leaves, - HDBSCAN::Common::CondensedHierarchy& condensed_tree); + HDBSCAN::Common::CondensedHierarchy& condensed_tree); void _extract_clusters(const raft::handle_t& handle, size_t n_leaves, int n_edges, - int* parents, - int* children, + int64_t* parents, + int64_t* children, float* lambdas, - int* sizes, - int* labels, + int64_t* sizes, + int64_t* labels, float* probabilities, HDBSCAN::Common::CLUSTER_SELECTION_METHOD cluster_selection_method, bool allow_single_cluster, - int max_cluster_size, + int64_t max_cluster_size, float cluster_selection_epsilon); void compute_all_points_membership_vectors( const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, ML::distance::DistanceType metric, float* membership_vec, size_t batch_size = 4096); void compute_membership_vector(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, const float* points_to_predict, size_t n_prediction_points, @@ -473,15 +473,15 @@ void compute_membership_vector(const raft::handle_t& handle, size_t batch_size = 4096); void out_of_sample_predict(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, - int* labels, + int64_t* labels, const float* points_to_predict, size_t n_prediction_points, ML::distance::DistanceType metric, int min_samples, - int* out_labels, + int64_t* out_labels, float* out_probabilities); namespace HDBSCAN::HELPER { @@ -519,12 +519,12 @@ void compute_core_dists(const raft::handle_t& handle, * @param[in] cluster_selection_epsilon cluster selection epsilon */ void compute_inverse_label_map(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, size_t n_leaves, HDBSCAN::Common::CLUSTER_SELECTION_METHOD cluster_selection_method, - rmm::device_uvector& inverse_label_map, + rmm::device_uvector& inverse_label_map, bool allow_single_cluster, - int max_cluster_size, + int64_t max_cluster_size, float cluster_selection_epsilon); } // namespace HDBSCAN::HELPER diff --git a/cpp/src/hdbscan/hdbscan.cu b/cpp/src/hdbscan/hdbscan.cu index cf025ae141..919c0fca30 100644 --- a/cpp/src/hdbscan/hdbscan.cu +++ b/cpp/src/hdbscan/hdbscan.cu @@ -32,20 +32,20 @@ void hdbscan(const raft::handle_t& handle, size_t n, ML::distance::DistanceType metric, HDBSCAN::Common::HDBSCANParams& params, - HDBSCAN::Common::hdbscan_output& out, + HDBSCAN::Common::hdbscan_output& out, float* core_dists) { - rmm::device_uvector labels(m, handle.get_stream()); + rmm::device_uvector labels(m, handle.get_stream()); HDBSCAN::_fit_hdbscan(handle, X, m, n, metric, params, labels.data(), core_dists, out); } void build_condensed_hierarchy(const raft::handle_t& handle, - const int* children, + const int64_t* children, const float* delta, - const int* sizes, + const int64_t* sizes, int min_cluster_size, int n_leaves, - HDBSCAN::Common::CondensedHierarchy& condensed_tree) + HDBSCAN::Common::CondensedHierarchy& condensed_tree) { HDBSCAN::detail::Condense::build_condensed_hierarchy( handle, children, delta, sizes, min_cluster_size, n_leaves, condensed_tree); @@ -54,23 +54,23 @@ void build_condensed_hierarchy(const raft::handle_t& handle, void _extract_clusters(const raft::handle_t& handle, size_t n_leaves, int n_edges, - int* parents, - int* children, + int64_t* parents, + int64_t* children, float* lambdas, - int* sizes, - int* labels, + int64_t* sizes, + int64_t* labels, float* probabilities, HDBSCAN::Common::CLUSTER_SELECTION_METHOD cluster_selection_method, bool allow_single_cluster, - int max_cluster_size, + int64_t max_cluster_size, float cluster_selection_epsilon) { HDBSCAN::Common::CondensedHierarchy condensed_tree( handle, n_leaves, n_edges, parents, children, lambdas, sizes); rmm::device_uvector stabilities(condensed_tree.get_n_clusters(), handle.get_stream()); - rmm::device_uvector label_map(condensed_tree.get_n_clusters(), handle.get_stream()); - rmm::device_uvector inverse_label_map(0, handle.get_stream()); + rmm::device_uvector label_map(condensed_tree.get_n_clusters(), handle.get_stream()); + rmm::device_uvector inverse_label_map(0, handle.get_stream()); HDBSCAN::detail::Extract::extract_clusters(handle, condensed_tree, @@ -88,8 +88,8 @@ void _extract_clusters(const raft::handle_t& handle, void compute_all_points_membership_vectors( const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, ML::distance::DistanceType metric, float* membership_vec, @@ -100,8 +100,8 @@ void compute_all_points_membership_vectors( } void compute_membership_vector(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, const float* points_to_predict, size_t n_prediction_points, @@ -125,15 +125,15 @@ void compute_membership_vector(const raft::handle_t& handle, } void out_of_sample_predict(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, - HDBSCAN::Common::PredictionData& prediction_data, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::PredictionData& prediction_data, const float* X, - int* labels, + int64_t* labels, const float* points_to_predict, size_t n_prediction_points, ML::distance::DistanceType metric, int min_samples, - int* out_labels, + int64_t* out_labels, float* out_probabilities) { // Note that (min_samples+1) is parsed to the approximate_predict function. This was done for the @@ -161,17 +161,17 @@ void compute_core_dists(const raft::handle_t& handle, ML::distance::DistanceType metric, int min_samples) { - HDBSCAN::detail::Reachability::_compute_core_dists( + HDBSCAN::detail::Reachability::_compute_core_dists( handle, X, core_dists, m, n, metric, min_samples); } void compute_inverse_label_map(const raft::handle_t& handle, - HDBSCAN::Common::CondensedHierarchy& condensed_tree, + HDBSCAN::Common::CondensedHierarchy& condensed_tree, size_t n_leaves, HDBSCAN::Common::CLUSTER_SELECTION_METHOD cluster_selection_method, - rmm::device_uvector& inverse_label_map, + rmm::device_uvector& inverse_label_map, bool allow_single_cluster, - int max_cluster_size, + int64_t max_cluster_size, float cluster_selection_epsilon) { HDBSCAN::detail::Extract::_compute_inverse_label_map(handle, diff --git a/cpp/src/hdbscan/prediction_data.cu b/cpp/src/hdbscan/prediction_data.cu index afe9e0ef55..d68b64b9bd 100644 --- a/cpp/src/hdbscan/prediction_data.cu +++ b/cpp/src/hdbscan/prediction_data.cu @@ -100,11 +100,11 @@ void build_index_into_children(const raft::handle_t& handle, * @param[in] prediction_data PreditionData object */ void generate_prediction_data(const raft::handle_t& handle, - CondensedHierarchy& condensed_tree, - int* labels, - int* inverse_label_map, + CondensedHierarchy& condensed_tree, + int64_t* labels, + int64_t* inverse_label_map, int n_selected_clusters, - PredictionData& prediction_data) + PredictionData& prediction_data) { auto stream = handle.get_stream(); auto exec_policy = handle.get_thrust_policy(); @@ -120,10 +120,10 @@ void generate_prediction_data(const raft::handle_t& handle, auto sizes = condensed_tree.get_sizes(); // first compute the death of each cluster in the condensed hierarchy - rmm::device_uvector sorted_parents(n_edges, stream); + rmm::device_uvector sorted_parents(n_edges, stream); raft::copy_async(sorted_parents.data(), parents, n_edges, stream); - rmm::device_uvector sorted_parents_offsets(n_clusters + 1, stream); + rmm::device_uvector sorted_parents_offsets(n_clusters + 1, stream); detail::Utils::parent_csr( handle, condensed_tree, sorted_parents.data(), sorted_parents_offsets.data()); @@ -135,8 +135,8 @@ void generate_prediction_data(const raft::handle_t& handle, const float* d_in, float* d_out, int num_segments, - const int* d_begin_offsets, - const int* d_end_offsets, + const int64_t* d_begin_offsets, + const int64_t* d_end_offsets, cudaStream_t stream = 0) -> cudaError_t { return cub::DeviceSegmentedReduce::Max(d_temp_storage, temp_storage_bytes, @@ -201,7 +201,7 @@ void generate_prediction_data(const raft::handle_t& handle, [is_exemplar = is_exemplar.data()] __device__(auto idx) { return is_exemplar[idx]; }); // use the exemplar labels to fetch the set of selected clusters from the condensed hierarchy - rmm::device_uvector exemplar_labels(n_exemplars, stream); + rmm::device_uvector exemplar_labels(n_exemplars, stream); // this uses the original, pre-normalized label by // using the inverse label_map to lookup the original labels from final labels @@ -212,7 +212,7 @@ void generate_prediction_data(const raft::handle_t& handle, [labels, inverse_label_map] __device__(auto idx) { auto label = labels[idx]; if (label != -1) { return inverse_label_map[label]; } - return -1; + return static_cast(-1); }); thrust::sort_by_key(exec_policy, diff --git a/cpp/tests/sg/hdbscan_inputs.hpp b/cpp/tests/sg/hdbscan_inputs.hpp index 9933c6e8a9..54d2b2c3ce 100644 --- a/cpp/tests/sg/hdbscan_inputs.hpp +++ b/cpp/tests/sg/hdbscan_inputs.hpp @@ -135,7 +135,7 @@ struct MembershipVectorInputs { std::vector expected_probabilities; }; -const std::vector> hdbscan_inputsf2 = { +const std::vector> hdbscan_inputsf2 = { // Test n_clusters == n_points {10, 5, @@ -256,7 +256,7 @@ const std::vector> hdbscan_inputsf2 = { -1, -1, -1, -1, -1, -1, -1, -1, 3, 0, 4, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, 4, -1, -1, -1, -1, -1, -1, 1, -1, 1, -1, -1, -1, -1, -1, -1, -1, -1, 2, 6, 5, -1, -1, -1}}}; -const std::vector> cluster_condensing_inputs = { +const std::vector> cluster_condensing_inputs = { {9, 3, {0, 2, 4, 6, 7, 1, 8, 8}, @@ -868,8 +868,8 @@ const std::vector> cluster_condensing_inputs -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1}}}; namespace Iris { -constexpr int n_row = 150; -const std::vector parents = { +constexpr int n_row = 150; +const std::vector parents = { 150, 150, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, 152, 151, @@ -880,7 +880,7 @@ const std::vector parents = { 154, 154, 154, 154, 153, 153, 153, 153, 153, 153, 153, 153, 153, 153, 154, 154, 154, 154, 154, 154, 154, 154, 154, 154, 154, 154, 154, 154}; -const std::vector children = { +const std::vector children = { 151, 152, 41, 131, 15, 117, 14, 118, 22, 106, 18, 98, 13, 109, 33, 57, 44, 60, 93, 32, 129, 24, 68, 43, 122, 16, 135, 5, 134, 23, 119, 20, 125, 8, 114, 36, 108, 31, 148, 10, 87, 46, 100, 35, 105, 6, 62, 19, 107, 42, 113, 25, 130, 11, @@ -913,7 +913,7 @@ const std::vector lambdas = { 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242, 2.67261242}; -const std::vector sizes = { +const std::vector sizes = { 50, 100, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, @@ -924,7 +924,7 @@ const std::vector sizes = { namespace Digits { constexpr int n_row = 1797; -const std::vector parents = { +const std::vector parents = { 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, 1797, @@ -1041,7 +1041,7 @@ const std::vector parents = { 1828, 1828, 1825, 1834, 1825, 1834, 1825, 1834, 1834, 1834, 1834, 1834, 1834, 1834, 1834, 1834, 1834, 1825, 1825, 1825, 1825, 1825, 1825, 1825, 1825, 1825, 1825}; -const std::vector children = { +const std::vector children = { 1113, 1149, 1572, 1660, 1595, 985, 1024, 77, 1727, 1551, 757, 1562, 891, 1581, 673, 1552, 1729, 421, 1274, 1195, 1154, 1150, 1690, 1152, 1264, 792, 9, 1038, 1593, 1580, 1628, 1712, 1685, 905, 341, 1100, 769, 1574, 502, 87, 1210, 1165, 1657, 1571, 1057, 1118, 538, 678, @@ -1390,7 +1390,7 @@ const std::vector lambdas = { 0.06917145, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813, 0.05902813}; -const std::vector sizes = { +const std::vector sizes = { 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, @@ -1481,7 +1481,7 @@ const std::vector sizes = { 1, 1, 1, 1, 1, 1, 1, 1}; }; // namespace Digits -const std::vector> cluster_selection_inputs = { +const std::vector> cluster_selection_inputs = { {150, 5, 10, @@ -2325,7 +2325,7 @@ const std::vector> cluster_selection_inputs = -1, -1, -1, -1, 11, 13, -1, 15, 13, -1, -1, -1, -1, -1, 13, -1, 1, 16, 11, -1, -1, 18, 13, -1, -1, -1, -1, -1, 3, -1, 3, 3, -1, -1, -1, -1, -1, -1, -1, -1, 15, 1, -1, -1, -1}}}; -const std::vector> +const std::vector> all_points_membership_vectors_inputs = { {MLCommon::Datasets::Digits::n_samples, MLCommon::Datasets::Digits::n_features, @@ -6191,7 +6191,7 @@ const std::vector> 0.0501381, 0.053290438, 0.047078118, 0.040903572, 0.0667583, 0.056616243, 0.07065173, 0.0443147, 0.04315245, 0.056318313, 0.049550712, 0.042784892}}}; -const std::vector> approximate_predict_inputs = { +const std::vector> approximate_predict_inputs = { {1000, 15, 200, @@ -9544,7 +9544,7 @@ const std::vector> approximate_predict_inpu 0.9011331, 0.77075315, 0.9326203, 0.8810431, 0.90232337, 0.8286241, 1.0, 0.83256954, 0.83346933, 0.83974075, 0.6880056}}}; -const std::vector> membership_vector_inputs = { +const std::vector> membership_vector_inputs = { {1000, 15, 200, diff --git a/cpp/tests/sg/hdbscan_test.cu b/cpp/tests/sg/hdbscan_test.cu index 2e985f999d..733ff732cb 100644 --- a/cpp/tests/sg/hdbscan_test.cu +++ b/cpp/tests/sg/hdbscan_test.cu @@ -36,6 +36,7 @@ #include #include +#include #include #include #include @@ -133,7 +134,7 @@ class HDBSCANTest : public ::testing::TestWithParam> { double score; }; -typedef HDBSCANTest HDBSCANTestF_Int; +typedef HDBSCANTest HDBSCANTestF_Int; TEST_P(HDBSCANTestF_Int, Result) { EXPECT_TRUE(score >= 0.85); } INSTANTIATE_TEST_CASE_P(HDBSCANTest, HDBSCANTestF_Int, ::testing::ValuesIn(hdbscan_inputsf2)); @@ -239,7 +240,7 @@ class ClusterCondensingTest : public ::testing::TestWithParam ClusterCondensingTestF_Int; +typedef ClusterCondensingTest ClusterCondensingTestF_Int; TEST_P(ClusterCondensingTestF_Int, Result) { EXPECT_TRUE(score == 1.0); } // This will be reactivate in 21.08 with better, contrived examples to @@ -313,7 +314,7 @@ class ClusterSelectionTest : public ::testing::TestWithParam(0), params.cluster_selection_epsilon); handle.sync_stream(handle.get_stream()); @@ -340,7 +341,7 @@ class ClusterSelectionTest : public ::testing::TestWithParam ClusterSelectionTestF_Int; +typedef ClusterSelectionTest ClusterSelectionTestF_Int; TEST_P(ClusterSelectionTestF_Int, Result) { EXPECT_TRUE(score == 1.0); } INSTANTIATE_TEST_CASE_P(ClusterSelectionTest, @@ -353,7 +354,7 @@ void transformLabels(const raft::handle_t& handle, IdxT* labels, IdxT* label_map thrust::transform( handle.get_thrust_policy(), labels, labels + m, labels, [label_map] __device__(IdxT label) { if (label != -1) return label_map[label]; - return -1; + return static_cast(-1); }); } @@ -441,7 +442,7 @@ class AllPointsMembershipVectorsTest params.cluster_selection_method, inverse_label_map, params.allow_single_cluster, - 0, + static_cast(0), params.cluster_selection_epsilon); rmm::device_uvector membership_vec(params.n_row * n_selected_clusters, handle.get_stream()); @@ -481,7 +482,7 @@ class AllPointsMembershipVectorsTest // T score; }; -typedef AllPointsMembershipVectorsTest AllPointsMembershipVectorsTestF_Int; +typedef AllPointsMembershipVectorsTest AllPointsMembershipVectorsTestF_Int; TEST_P(AllPointsMembershipVectorsTestF_Int, Result) { EXPECT_TRUE(true); } INSTANTIATE_TEST_CASE_P(AllPointsMembershipVectorsTest, @@ -491,7 +492,6 @@ INSTANTIATE_TEST_CASE_P(AllPointsMembershipVectorsTest, template class ApproximatePredictTest : public ::testing::TestWithParam> { public: - protected: void basicTest() { raft::handle_t handle; @@ -562,7 +562,7 @@ class ApproximatePredictTest : public ::testing::TestWithParam(0), params.cluster_selection_epsilon); rmm::device_uvector core_dists{static_cast(params.n_row), handle.get_stream()}; @@ -574,16 +574,64 @@ class ApproximatePredictTest : public ::testing::TestWithParam mutual_reachability_coo(stream, (params.min_samples + 1) * params.n_row * 2); - cuvs::neighbors::reachability::mutual_reachability_graph( + auto exec_policy = raft::resource::get_thrust_policy(handle); + auto new_min_samples = params.min_samples + 1; + auto inds = raft::make_device_matrix(handle, params.n_row, new_min_samples); + auto dists = raft::make_device_matrix(handle, params.n_row, new_min_samples); + + cuvs::neighbors::all_neighbors::all_neighbors_params all_neigh_p; + all_neigh_p.metric = cuvs::distance::DistanceType::L2SqrtExpanded; + auto brute_force_p = cuvs::neighbors::graph_build_params::brute_force_params{}; + brute_force_p.build_params.metric = cuvs::distance::DistanceType::L2SqrtExpanded; + all_neigh_p.graph_build_params = brute_force_p; + cuvs::neighbors::all_neighbors::build( handle, - raft::make_device_matrix_view(data.data(), params.n_row, params.n_col), - params.min_samples + 1, - raft::make_device_vector_view(mutual_reachability_indptr.data(), params.n_row + 1), + all_neigh_p, + raft::make_device_matrix_view(data.data(), params.n_row, params.n_col), + inds.view(), + dists.view(), raft::make_device_vector_view(core_dists.data(), params.n_row), - mutual_reachability_coo, - cuvs::distance::DistanceType::L2SqrtExpanded, 1.0); + // self-loops get max distance + rmm::device_uvector coo_rows(new_min_samples * params.n_row, stream); + auto coo_rows_counting_itr = thrust::make_counting_iterator(0); + thrust::transform(exec_policy, + coo_rows_counting_itr, + coo_rows_counting_itr + (params.n_row * new_min_samples), + coo_rows.data(), + [new_min_samples] __device__(IdxT c) -> IdxT { return c / new_min_samples; }); + + raft::sparse::linalg::symmetrize(handle, + coo_rows.data(), + inds.data_handle(), + dists.data_handle(), + static_cast(params.n_row), + static_cast(params.n_row), + static_cast(new_min_samples * params.n_row), + mutual_reachability_coo); + + raft::sparse::convert::sorted_coo_to_csr(mutual_reachability_coo.rows(), + mutual_reachability_coo.nnz, + mutual_reachability_indptr.data(), + (int)params.n_row + 1, + stream); + + auto transform_in = + thrust::make_zip_iterator(thrust::make_tuple(mutual_reachability_coo.rows(), + mutual_reachability_coo.cols(), + mutual_reachability_coo.vals())); + + thrust::transform(exec_policy, + transform_in, + transform_in + mutual_reachability_coo.nnz, + mutual_reachability_coo.vals(), + [=] __device__(const thrust::tuple& tup) { + return thrust::get<0>(tup) == thrust::get<1>(tup) + ? std::numeric_limits::max() + : thrust::get<2>(tup); + }); + transformLabels(handle, labels.data(), label_map.data(), params.n_row); ML::HDBSCAN::Common::generate_prediction_data(handle, condensed_tree, @@ -633,7 +681,7 @@ class ApproximatePredictTest : public ::testing::TestWithParam ApproximatePredictTestF_Int; +typedef ApproximatePredictTest ApproximatePredictTestF_Int; TEST_P(ApproximatePredictTestF_Int, Result) { EXPECT_TRUE(true); } INSTANTIATE_TEST_CASE_P(ApproximatePredictTest, @@ -642,7 +690,7 @@ INSTANTIATE_TEST_CASE_P(ApproximatePredictTest, template class MembershipVectorTest : public ::testing::TestWithParam> { - protected: + public: void basicTest() { raft::handle_t handle; @@ -713,7 +761,7 @@ class MembershipVectorTest : public ::testing::TestWithParam(0), params.cluster_selection_epsilon); rmm::device_uvector membership_vec(params.n_points_to_predict * n_selected_clusters, @@ -728,16 +776,64 @@ class MembershipVectorTest : public ::testing::TestWithParam mutual_reachability_coo(stream, (params.min_samples + 1) * params.n_row * 2); - cuvs::neighbors::reachability::mutual_reachability_graph( + auto exec_policy = raft::resource::get_thrust_policy(handle); + auto new_min_samples = params.min_samples + 1; + auto inds = raft::make_device_matrix(handle, params.n_row, new_min_samples); + auto dists = raft::make_device_matrix(handle, params.n_row, new_min_samples); + + cuvs::neighbors::all_neighbors::all_neighbors_params all_neigh_p; + all_neigh_p.metric = cuvs::distance::DistanceType::L2SqrtExpanded; + auto brute_force_p = cuvs::neighbors::graph_build_params::brute_force_params{}; + brute_force_p.build_params.metric = cuvs::distance::DistanceType::L2SqrtExpanded; + all_neigh_p.graph_build_params = brute_force_p; + cuvs::neighbors::all_neighbors::build( handle, - raft::make_device_matrix_view(data.data(), params.n_row, params.n_col), - params.min_samples + 1, - raft::make_device_vector_view(mutual_reachability_indptr.data(), params.n_row + 1), + all_neigh_p, + raft::make_device_matrix_view(data.data(), params.n_row, params.n_col), + inds.view(), + dists.view(), raft::make_device_vector_view(core_dists.data(), params.n_row), - mutual_reachability_coo, - cuvs::distance::DistanceType::L2SqrtExpanded, 1.0); + // self-loops get max distance + rmm::device_uvector coo_rows(new_min_samples * params.n_row, stream); + auto coo_rows_counting_itr = thrust::make_counting_iterator(0); + thrust::transform(exec_policy, + coo_rows_counting_itr, + coo_rows_counting_itr + (params.n_row * new_min_samples), + coo_rows.data(), + [new_min_samples] __device__(IdxT c) -> IdxT { return c / new_min_samples; }); + + raft::sparse::linalg::symmetrize(handle, + coo_rows.data(), + inds.data_handle(), + dists.data_handle(), + static_cast(params.n_row), + static_cast(params.n_row), + static_cast(new_min_samples * params.n_row), + mutual_reachability_coo); + + raft::sparse::convert::sorted_coo_to_csr(mutual_reachability_coo.rows(), + mutual_reachability_coo.nnz, + mutual_reachability_indptr.data(), + (int)params.n_row + 1, + stream); + + auto transform_in = + thrust::make_zip_iterator(thrust::make_tuple(mutual_reachability_coo.rows(), + mutual_reachability_coo.cols(), + mutual_reachability_coo.vals())); + + thrust::transform(exec_policy, + transform_in, + transform_in + mutual_reachability_coo.nnz, + mutual_reachability_coo.vals(), + [=] __device__(const thrust::tuple& tup) { + return thrust::get<0>(tup) == thrust::get<1>(tup) + ? std::numeric_limits::max() + : thrust::get<2>(tup); + }); + transformLabels(handle, labels.data(), label_map.data(), params.n_row); ML::HDBSCAN::Common::generate_prediction_data(handle, @@ -773,7 +869,7 @@ class MembershipVectorTest : public ::testing::TestWithParam MembershipVectorTestF_Int; +typedef MembershipVectorTest MembershipVectorTestF_Int; TEST_P(MembershipVectorTestF_Int, Result) { EXPECT_TRUE(true); } INSTANTIATE_TEST_CASE_P(MembershipVectorTest, diff --git a/python/cuml/cuml/cluster/hdbscan/hdbscan.pyx b/python/cuml/cuml/cluster/hdbscan/hdbscan.pyx index c7579141c9..44d0634730 100644 --- a/python/cuml/cuml/cluster/hdbscan/hdbscan.pyx +++ b/python/cuml/cuml/cluster/hdbscan/hdbscan.pyx @@ -34,7 +34,7 @@ from cuml.internals.interop import ( from cuml.internals.mixins import ClusterMixin, CMajorInputTagMixin from cython.operator cimport dereference as deref -from libc.stdint cimport uintptr_t +from libc.stdint cimport int64_t, uintptr_t from libcpp cimport bool from pylibraft.common.handle cimport handle_t from rmm.librmm.device_uvector cimport device_uvector @@ -98,10 +98,10 @@ cdef class _HDBSCANState: # A pointer to a `CondensedHierarchy`, or `NULL` if this state was # initialized through a `fit` call. - cdef lib.CondensedHierarchy[int, float] *condensed_tree + cdef lib.CondensedHierarchy[int64_t, float] *condensed_tree # The generated PredictionData, or NULL if prediction data was not yet generated. - cdef lib.PredictionData[int, float] *prediction_data + cdef lib.PredictionData[int64_t, float] *prediction_data # The number of clusters cdef public int n_clusters @@ -146,21 +146,21 @@ cdef class _HDBSCANState: """Shared helper for initializing a `CondensedHierarchy` from a condensed_tree array""" self.cached_condensed_tree = tree - parents = np.ascontiguousarray(tree["parent"], dtype=np.int32) - children = np.ascontiguousarray(tree["child"], dtype=np.int32) + parents = np.ascontiguousarray(tree["parent"], dtype=np.int64) + children = np.ascontiguousarray(tree["child"], dtype=np.int64) lambdas = np.ascontiguousarray(tree["lambda_val"], dtype=np.float32) - sizes = np.ascontiguousarray(tree["child_size"], dtype=np.int32) + sizes = np.ascontiguousarray(tree["child_size"], dtype=np.int64) cdef int n_edges = len(tree) cdef handle_t *handle_ = handle.getHandle() - self.condensed_tree = new lib.CondensedHierarchy[int, float]( + self.condensed_tree = new lib.CondensedHierarchy[int64_t, float]( handle_[0], n_leaves, n_edges, - (parents.ctypes.data), - (children.ctypes.data), + (parents.ctypes.data), + (children.ctypes.data), (lambdas.ctypes.data), - (sizes.ctypes.data), + (sizes.ctypes.data), ) @staticmethod @@ -196,10 +196,10 @@ cdef class _HDBSCANState: cdef float* X_ptr = X.ptr cdef float* core_dists_ptr = self.core_dists.ptr cdef bool allow_single_cluster = model.allow_single_cluster - cdef int max_cluster_size = model.max_cluster_size + cdef int64_t max_cluster_size = model.max_cluster_size cdef float cluster_selection_epsilon = model.cluster_selection_epsilon cdef int min_samples = model.min_samples or model.min_cluster_size - cdef device_uvector[int] *temp_buffer = new device_uvector[int]( + cdef device_uvector[int64_t] *temp_buffer = new device_uvector[int64_t]( 0, handle_[0].get_stream(), ) @@ -233,18 +233,18 @@ cdef class _HDBSCANState: data=_cupy_array_from_ptr( temp_buffer.data(), (self.n_clusters,), - np.int32, + np.int64, self ).copy() ) else: - self.inverse_label_map = CumlArray.empty((0,), dtype=np.int32) + self.inverse_label_map = CumlArray.empty((0,), dtype=np.int64) del temp_buffer return self - cdef lib.CondensedHierarchy[int, float]* get_condensed_tree(self) nogil: + cdef lib.CondensedHierarchy[int64_t, float]* get_condensed_tree(self) nogil: if self.hdbscan_output != NULL: return &(self.hdbscan_output.get_condensed_tree()) return self.condensed_tree @@ -266,8 +266,8 @@ cdef class _HDBSCANState: children = input_to_cuml_array( dendrogram[:, 0:2], order='C', - check_dtype=[np.int32], - convert_to_dtype=np.int32, + check_dtype=[np.int64], + convert_to_dtype=np.int64, )[0] lambdas = input_to_cuml_array( @@ -280,8 +280,8 @@ cdef class _HDBSCANState: sizes = input_to_cuml_array( dendrogram[:, 3], order='C', - check_dtype=[np.int32], - convert_to_dtype=np.int32, + check_dtype=[np.int64], + convert_to_dtype=np.int64, )[0] cdef size_t n_leaves = dendrogram.shape[0] + 1 @@ -289,10 +289,10 @@ cdef class _HDBSCANState: handle = Handle() cdef handle_t *handle_ = handle.getHandle() - self.condensed_tree = new lib.CondensedHierarchy[int, float](handle_[0], n_leaves) - cdef int* children_ptr = children.ptr + self.condensed_tree = new lib.CondensedHierarchy[int64_t, float](handle_[0], n_leaves) + cdef int64_t* children_ptr = children.ptr cdef float* lambdas_ptr = lambdas.ptr - cdef int* sizes_ptr = sizes.ptr + cdef int64_t* sizes_ptr = sizes.ptr with nogil: lib.build_condensed_hierarchy( handle_[0], @@ -320,15 +320,15 @@ cdef class _HDBSCANState: cdef int n_cols = X.shape[1] # Allocate output structures - labels = CumlArray.empty(n_rows, dtype="int32", index=X.index) + labels = CumlArray.empty(n_rows, dtype="int64", index=X.index) probabilities = CumlArray.empty(n_rows, dtype="float32") - children = CumlArray.empty((2, n_rows), dtype="int32") + children = CumlArray.empty((2, n_rows), dtype="int64") lambdas = CumlArray.empty(n_rows, dtype="float32") - sizes = CumlArray.empty(n_rows, dtype="int32") + sizes = CumlArray.empty(n_rows, dtype="int64") - mst_src = CumlArray.empty(n_rows - 1, dtype="int32") - mst_dst = CumlArray.empty(n_rows - 1, dtype="int32") + mst_src = CumlArray.empty(n_rows - 1, dtype="int64") + mst_dst = CumlArray.empty(n_rows - 1, dtype="int64") mst_weights = CumlArray.empty(n_rows - 1, dtype="float32") core_dists = CumlArray.empty(n_rows, dtype="float32") @@ -340,13 +340,13 @@ cdef class _HDBSCANState: self.hdbscan_output = new lib.hdbscan_output( handle_[0], n_rows, - (labels.ptr), + (labels.ptr), (probabilities.ptr), - (children.ptr), - (sizes.ptr), + (children.ptr), + (sizes.ptr), (lambdas.ptr), - (mst_src.ptr), - (mst_dst.ptr), + (mst_src.ptr), + (mst_dst.ptr), (mst_weights.ptr) ) @@ -371,12 +371,12 @@ cdef class _HDBSCANState: data=_cupy_array_from_ptr( self.hdbscan_output.get_inverse_label_map(), (self.n_clusters,), - np.int32, + np.int64, self ) ) else: - self.inverse_label_map = CumlArray.empty((0,), dtype=np.int32) + self.inverse_label_map = CumlArray.empty((0,), dtype=np.int64) self.core_dists = core_dists # Extract and prepare results @@ -428,18 +428,18 @@ cdef class _HDBSCANState: cdef int n_rows = X.shape[0] cdef int n_cols = X.shape[1] - cdef int* labels_ptr = labels.ptr - cdef int* inverse_label_map_ptr = self.inverse_label_map.ptr + cdef int64_t* labels_ptr = labels.ptr + cdef int64_t* inverse_label_map_ptr = self.inverse_label_map.ptr cdef handle_t* handle_ = handle.getHandle() - self.prediction_data = new lib.PredictionData[int, float]( + self.prediction_data = new lib.PredictionData[int64_t, float]( handle_[0], n_rows, n_cols, (self.core_dists.ptr), ) - cdef lib.CondensedHierarchy[int, float] *condensed_tree = self.get_condensed_tree() + cdef lib.CondensedHierarchy[int64_t, float] *condensed_tree = self.get_condensed_tree() with nogil: lib.generate_prediction_data( @@ -458,7 +458,7 @@ cdef class _HDBSCANState: # Cached, return the same result return self.cached_condensed_tree - cdef lib.CondensedHierarchy[int, float]* condensed_tree = self.get_condensed_tree() + cdef lib.CondensedHierarchy[int64_t, float]* condensed_tree = self.get_condensed_tree() n_condensed_tree_edges = condensed_tree.get_n_edges() @@ -471,14 +471,14 @@ cdef class _HDBSCANState: parents = _cupy_array_from_ptr( condensed_tree.get_parents(), (n_condensed_tree_edges,), - np.int32, + np.int64, self, ) children = _cupy_array_from_ptr( condensed_tree.get_children(), (n_condensed_tree_edges,), - np.int32, + np.int64, self, ) @@ -492,7 +492,7 @@ cdef class _HDBSCANState: sizes = _cupy_array_from_ptr( condensed_tree.get_sizes(), (n_condensed_tree_edges,), - np.int32, + np.int64, self, ) @@ -720,7 +720,7 @@ class HDBSCAN(Base, InteropMixin, ClusterMixin, CMajorInputTagMixin): raise UnsupportedOnGPU("Sparse inputs are not supported") raw_data = to_gpu(raw_data_cpu, order="C", dtype="float32") - labels = to_gpu(model.labels_, order="C", dtype="int32") + labels = to_gpu(model.labels_, order="C", dtype="int64") state = _HDBSCANState.from_sklearn(self.handle, model, raw_data) if model._prediction_data is not None: state.generate_prediction_data(self.handle, raw_data, labels) @@ -1218,7 +1218,7 @@ def approximate_predict(clusterer, points_to_predict, convert_dtype=True): prediction_labels = CumlArray.empty( (n_prediction_points,), - dtype="int32", + dtype="int64", index=points_to_predict_m.index, ) prediction_probs = CumlArray.empty( @@ -1232,9 +1232,9 @@ def approximate_predict(clusterer, points_to_predict, convert_dtype=True): cdef _HDBSCANState state = <_HDBSCANState?>clusterer._state cdef float* X_ptr = clusterer._raw_data.ptr - cdef int* labels_ptr = labels.ptr + cdef int64_t* labels_ptr = labels.ptr cdef float* points_to_predict_ptr = points_to_predict_m.ptr - cdef int* prediction_labels_ptr = prediction_labels.ptr + cdef int64_t* prediction_labels_ptr = prediction_labels.ptr cdef float* prediction_probs_ptr = prediction_probs.ptr cdef DistanceType metric = _metrics_mapping[clusterer.metric] cdef int min_samples = clusterer.min_samples or clusterer.min_cluster_size, @@ -1304,13 +1304,13 @@ def _extract_clusters( parents = input_to_cuml_array( condensed_tree['parent'], order='C', - convert_to_dtype=np.int32, + convert_to_dtype=np.int64, )[0] children = input_to_cuml_array( condensed_tree["child"], order='C', - convert_to_dtype=np.int32, + convert_to_dtype=np.int64, )[0] lambdas = input_to_cuml_array( @@ -1322,10 +1322,10 @@ def _extract_clusters( sizes = input_to_cuml_array( condensed_tree['child_size'], order='C', - convert_to_dtype=np.int32, + convert_to_dtype=np.int64, )[0] - labels = CumlArray.empty(n_leaves, dtype="int32") + labels = CumlArray.empty(n_leaves, dtype="int64") probabilities = CumlArray.empty(n_leaves, dtype="float32") cdef lib.CLUSTER_SELECTION_METHOD cluster_selection_method_val = { @@ -1341,15 +1341,15 @@ def _extract_clusters( handle_[0], n_leaves, n_edges, - (parents.ptr), - (children.ptr), + (parents.ptr), + (children.ptr), (lambdas.ptr), - (sizes.ptr), - (labels.ptr), + (sizes.ptr), + (labels.ptr), (probabilities.ptr), cluster_selection_method_val, allow_single_cluster, - max_cluster_size, + max_cluster_size, cluster_selection_epsilon, ) handle.sync() diff --git a/python/cuml/cuml/cluster/hdbscan/headers.pxd b/python/cuml/cuml/cluster/hdbscan/headers.pxd index e4bbb42d64..c40b0e1b82 100644 --- a/python/cuml/cuml/cluster/hdbscan/headers.pxd +++ b/python/cuml/cuml/cluster/hdbscan/headers.pxd @@ -12,6 +12,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from libc.stdint cimport int64_t from libcpp cimport bool from pylibraft.common.handle cimport handle_t from rmm.librmm.device_uvector cimport device_uvector @@ -44,24 +45,24 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML::HDBSCAN::Common" nogi value_idx get_n_leaves() except + int get_n_clusters() except + - cdef cppclass hdbscan_output[int, float]: + cdef cppclass hdbscan_output[int64_t, float]: hdbscan_output(const handle_t &handle, int n_leaves, - int *labels, + int64_t *labels, float *probabilities, - int *children, - int *sizes, + int64_t *children, + int64_t *sizes, float *deltas, - int *mst_src, - int *mst_dst, + int64_t *mst_src, + int64_t *mst_dst, float *mst_weights) except + - int get_n_leaves() except + + int64_t get_n_leaves() except + int get_n_clusters() except + float *get_stabilities() except + - int *get_labels() except + - int *get_inverse_label_map() except + + int64_t *get_labels() except + + int64_t *get_inverse_label_map() except + float *get_core_dists() except + - CondensedHierarchy[int, float] &get_condensed_tree() except + + CondensedHierarchy[int64_t, float] &get_condensed_tree() except + cdef cppclass HDBSCANParams: int min_samples @@ -74,21 +75,21 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML::HDBSCAN::Common" nogi bool allow_single_cluster, CLUSTER_SELECTION_METHOD cluster_selection_method, - cdef cppclass PredictionData[int, float]: + cdef cppclass PredictionData[int64_t, float]: PredictionData(const handle_t &handle, - int m, - int n, + int64_t m, + int64_t n, float *core_dists) except + size_t n_rows size_t n_cols void generate_prediction_data(const handle_t& handle, - CondensedHierarchy[int, float]& condensed_tree, - int* labels, - int* inverse_label_map, + CondensedHierarchy[int64_t, float]& condensed_tree, + int64_t* labels, + int64_t* inverse_label_map, int n_selected_clusters, - PredictionData[int, float]& prediction_data) except + + PredictionData[int64_t, float]& prediction_data) except + cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML" nogil: @@ -102,25 +103,25 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML" nogil: void build_condensed_hierarchy( const handle_t &handle, - const int *children, + const int64_t *children, const float *delta, - const int *sizes, + const int64_t *sizes, int min_cluster_size, int n_leaves, - CondensedHierarchy[int, float] &condensed_tree) except + + CondensedHierarchy[int64_t, float] &condensed_tree) except + void _extract_clusters(const handle_t &handle, size_t n_leaves, - int _n_edges, int *parents, int *children, - float *lambdas, int *sizes, int *labels, + int _n_edges, int64_t *parents, int64_t *children, + float *lambdas, int64_t *sizes, int64_t *labels, float *probabilities, CLUSTER_SELECTION_METHOD cluster_selection_method, - bool allow_single_cluster, int max_cluster_size, + bool allow_single_cluster, int64_t max_cluster_size, float cluster_selection_epsilon) except + void compute_all_points_membership_vectors( const handle_t &handle, - CondensedHierarchy[int, float] &condensed_tree, - PredictionData[int, float] &prediction_data_, + CondensedHierarchy[int64_t, float] &condensed_tree, + PredictionData[int64_t, float] &prediction_data_, float* X, DistanceType metric, float* membership_vec, @@ -128,8 +129,8 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML" nogil: void compute_membership_vector( const handle_t& handle, - CondensedHierarchy[int, float] &condensed_tree, - PredictionData[int, float] &prediction_data, + CondensedHierarchy[int64_t, float] &condensed_tree, + PredictionData[int64_t, float] &prediction_data, float* X, float* points_to_predict, size_t n_prediction_points, @@ -139,15 +140,15 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML" nogil: size_t batch_size) except + void out_of_sample_predict(const handle_t &handle, - CondensedHierarchy[int, float] &condensed_tree, - PredictionData[int, float] &prediction_data, + CondensedHierarchy[int64_t, float] &condensed_tree, + PredictionData[int64_t, float] &prediction_data, float* X, - int* labels, + int64_t* labels, float* points_to_predict, size_t n_prediction_points, DistanceType metric, int min_samples, - int* out_labels, + int64_t* out_labels, float* out_probabilities) except + cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML::HDBSCAN::HELPER" nogil: @@ -161,12 +162,12 @@ cdef extern from "cuml/cluster/hdbscan.hpp" namespace "ML::HDBSCAN::HELPER" nogi int min_samples) except + void compute_inverse_label_map(const handle_t& handle, - CondensedHierarchy[int, float]& + CondensedHierarchy[int64_t, float]& condensed_tree, size_t n_leaves, CLUSTER_SELECTION_METHOD cluster_selection_method, - device_uvector[int]& inverse_label_map, + device_uvector[int64_t]& inverse_label_map, bool allow_single_cluster, - int max_cluster_size, + int64_t max_cluster_size, float cluster_selection_epsilon) except +