[FEA] membership_vector for HDBSCAN #5247

tarang-jain · 2023-02-22T19:58:03Z

…n/cuml into fea-membership-vector

tarang-jain · 2023-02-28T00:36:10Z

cpp/src/hdbscan/detail/utils.h

+ * @param[out] indptr CSR indptr of parents array after sort
+ */
+template <typename value_idx, typename value_t>
+void softmax(const raft::handle_t& handle, value_t* data, value_idx n, size_t m)


We have to create a copy of the data in order to avoid inconsistency during normalization (for numerical stability)

That doesn't sound right- we shouldn't have to make a copy here. If these operations are being scheduled wrt the same cuda stream, it should be guaranteed that any prior operations scheduled on compute resources are complete before new operations are scheduled.

One thing we can do to test the theory that there's a data race somewhere is to use the input directly (instead of copying) and call handle.sync_stream() before and after computing the argmax. Allocating more memory will also cause a device synchronization, which would have the same effect as doing a stream sync

What I meant was that if this operation:
data[idx] = exp(data_copy[idx] - data_copy[n*(idx / n) + membership_argmax[idx / n]])
was replaced by this:
data[idx] = exp(data[idx] - data[n*(idx / n) + membership_argmax[idx / n]]), then it is possible that data[n*(idx / n) + membership_argmax[idx / n]] got modified while/before data[idx] was being computed since they are both being modified in parallel in the same thrust::for_each invocation.

I don't think you need to copy the whole dataset to do this, though. It looks like what you want is to subtract the value in a specific column for each row from each value in data, right? Instead of copying all m*n entries from the input, why not just create an array of size of m that contains all the max values for each row? That way you've reduced your problme down to the number of rows. I also think it would make your arithmetic easier to follow.

Also, I would highly suggest using raft::matrix::map_offset primitive in raft for this. The benefit to using the raft primitives (even if some of them end up just wrapping thrust calls) is that it provides a facade layer for future optimizations.

cjnolet · 2023-02-28T00:37:25Z

cpp/src/hdbscan/detail/utils.h

+
+  rmm::device_uvector<value_idx> membership_argmax(m, stream);
+
+    raft::matrix::argmax(


We should be calling the mdspan variants for new code. The raw pointer versions are deprecated.

cjnolet · 2023-02-28T20:34:56Z

cpp/src/hdbscan/detail/utils.h

+                       data_copy = data_copy.data(),
+                       membership_argmax = membership_argmax.data(),
+                       n] __device__(auto idx) {
+                        data[idx] = exp(data_copy[idx] - data_copy[n*(idx / n) + membership_argmax[idx / n]]);


We should probably consolidate this divison so we only have to do it once. We should also try and use the fast int division utility in raft.

cjnolet · 2023-03-07T20:59:30Z

cpp/src/hdbscan/detail/utils.h

+  auto data_const_view = raft::make_device_matrix_view<const value_t, value_idx, raft::row_major>(data, (int)m, n);
+  auto argmax_view = raft::make_device_vector_view<value_idx, value_idx>(argmax.data(), (int)m);
+
+  raft::matrix::argmax(


Do you really need the argmax here? It looks like you are computing the argmax just to then take the actual max. Could you jsut do a max instead?

If I'm reading the math correctly here, it looks like what you want is an Linf normalization, right (dividing each row by the max of the row)? If this is the case, there are other prims in raft::linalg that could make this easier as well. I guess this assumes log space so it's actually a subtraction, but same effect computationally.

I was looking at the raft docs and did not find raft::matrix::max, so I did an argmax. Yes, we are dividing each row by the max of the row. The only issue is that we don't want to row to be computed beforehand (i.e., we don't want the exponential to be computed) due to numerical overflow. So while exp(vector) followed by Linf normalization would mathematically be the same, it can still lead to numerical overflow.

I think what you want here is a row-level reduce that uses a max instead of addition. There are primitives to do this in raft::linalg that should take a custom lambda (and we have a max lambda that can be used).

…ership-vector

cjnolet

Looking good. Mostly minor things. I want to stress that we need to use raft calls wherever possible (mostly map_offset in this case) so that can centralized operations that we find we're using quite frequently. This helps optimizations to these functions propagate across the board.

cjnolet · 2023-03-27T20:17:42Z

cpp/src/hdbscan/detail/kernels/soft_clustering.cuh

@@ -28,6 +28,7 @@ __global__ void merge_height_kernel(value_t* heights,
                                    value_idx* parents,
                                    size_t m,
                                    value_idx n_selected_clusters,
+                                    MLCommon::FastIntDiv n,


Can you use the version of this in raft::util please? This should really be removed from cuml altogether now that we have a version in raft.

cjnolet · 2023-03-27T20:18:53Z

cpp/src/hdbscan/detail/soft_clustering.cuh

 {
  auto stream      = handle.get_stream();
  auto exec_policy = handle.get_thrust_policy();

  auto counting = thrust::make_counting_iterator<value_idx>(0);

-  rmm::device_uvector<value_t> exemplars_dense(n_exemplars * n, stream);
+  rmm::device_uvector<value_t> exemplars_dense(n_exemplars * 1000, stream);


Should this be hardcoded to 1k?

cjnolet · 2023-03-27T20:20:46Z

cpp/src/hdbscan/detail/soft_clustering.cuh

+  // hierarchy
+  rmm::device_uvector<value_t> nearest_cluster_max_lambda(n_prediction_points, stream);
+
+  thrust::for_each(exec_policy,


Please use raft::matrix::map_offset for this operation.

cjnolet · 2023-03-27T20:23:14Z

cpp/src/hdbscan/detail/kernels/soft_clustering.cuh

    prob_in_some_cluster[idx] =
-      heights[idx * n_selected_clusters + (int)height_argmax[idx]] / max_lambda;
+      heights[idx * n_selected_clusters + height_argmax[idx]] / max_lambda;


Please compute height_argmax[idx] once and store it off to reuse instead of having to do this load multiple times.

cjnolet · 2023-03-27T20:24:14Z

cpp/src/hdbscan/detail/kernels/soft_clustering.cuh

+                                            value_idx n_leaves,
+                                            size_t n_prediction_points)
+{
+  value_idx idx = blockDim.x * blockIdx.x + threadIdx.x;


This looks like another candidate for map_offset, right?

cjnolet · 2023-03-27T20:26:29Z

cpp/src/hdbscan/detail/soft_clustering.cuh

+    return;
+  };
+
+  thrust::for_each(


raft::matrix::map_offset please.

cjnolet · 2023-03-27T20:27:04Z

cpp/src/hdbscan/detail/utils.h

+
+  raft::linalg::norm(handle, data_const_view, linf_norm_view, raft::linalg::LinfNorm, raft::linalg::Apply::ALONG_ROWS);
+
+  raft::linalg::matrix_vector_op(handle, data_const_view, linf_norm_const_view, data_view, raft::linalg::Apply::ALONG_COLUMNS, [] __device__(value_t mat_in, value_t vec_in) {


much better!

cjnolet · 2023-03-27T20:28:02Z

cpp/test/CMakeLists.txt

@@ -122,7 +122,7 @@ if(BUILD_CUML_TESTS)
    ConfigureTest(PREFIX SG NAME GENETIC_PARAM_TEST PATH sg/genetic/param_test.cu OPTIONAL ML_INCLUDE)
  endif()

-  if("${CMAKE_CUDA_COMPILER_VERSION}" VERSION_LESS_EQUAL "11.2")
+  if("${CMAKE_CUDA_COMPILER_VERSION}" VERSION_GREATER_EQUAL "11.2")


Can you check if we can remove this conditional altogether? I think we want to strive to have these tests pass in all cuda versions.

cjnolet · 2023-03-27T20:34:54Z

cpp/src/hdbscan/detail/utils.h


 #include <algorithm>

 #include "../condensed_hierarchy.cu"
+#include <common/fast_int_div.cuh>

 #include <thrust/copy.h>
 #include <thrust/execution_policy.h>


We should eventually move Utils::normalize into RAFT. These are very useful functions. We have functions for norms, but normalization is just as important and pops up just as frequently.

cjnolet · 2023-03-27T20:35:20Z

cpp/src/hdbscan/detail/utils.h

+ * @param[out] m number of rows
+ */
+template <typename value_idx, typename value_t>
+void softmax(const raft::handle_t& handle, value_t* data, value_idx n, size_t m)


This could be moved into raft at some point too, since it's just a special case of normalizating into a multinomial distribution.

…n/cuml into fea-membership-vector

…ership-vector

cjnolet

LGTM! Can you create RAFT issues for the few comments I made about primitives we could move over to RAFT? It would also help if you could reference the issues in the code here with a small note that we'd like to eventually move those computations over. It just helps to keep it on our radar (and for future eyes to know our plans).

cjnolet · 2023-03-31T00:28:36Z

/merge

tarang-jain added 2 commits February 17, 2023 17:18

membership_vector initial commit

e49d06a

Further updates to membership_vector

436b180

github-actions bot added the CUDA/C++ label Feb 22, 2023

tarang-jain and others added 4 commits February 22, 2023 11:58

Merge branch 'branch-23.04' into fea-membership-vector

48030b8

Initial testing membership_vector

7912dba

Debug statements

4b41edb

Merge branch 'fea-membership-vector' of https://github.com/tarang-jai…

fe0fd34

…n/cuml into fea-membership-vector

github-actions bot added the Cython / Python Cython or Python issue label Feb 23, 2023

tarang-jain added 2 commits February 24, 2023 13:49

debugging membership_vector

9d5badc

membership_vector first working impl

19f9dd8

tarang-jain commented Feb 28, 2023

View reviewed changes

cjnolet reviewed Feb 28, 2023

View reviewed changes

tarang-jain added 2 commits February 27, 2023 17:13

GoogleTest intermediate commit

a4b565c

GTest working

1f4bf78

github-actions bot added the CMake label Feb 28, 2023

cjnolet reviewed Feb 28, 2023

View reviewed changes

tarang-jain added 2 commits February 28, 2023 15:49

working tests and styling changes

fdf100b

replace with raft mdspan primitives and add FastIntDiv

e18096a

tarang-jain marked this pull request as ready for review March 1, 2023 23:19

tarang-jain requested review from a team as code owners March 1, 2023 23:19

tarang-jain and others added 5 commits March 1, 2023 15:19

Merge branch 'branch-23.04' into fea-membership-vector

c2aa77e

cpu support

182ba31

Fix failing pytest

366ef26

Merge branch 'branch-23.04' into fea-membership-vector

b60d869

modification after merge

6bfaae2

cjnolet reviewed Mar 7, 2023

View reviewed changes

tarang-jain added 2 commits March 7, 2023 16:05

Update softmax with raft::linalg reduction

c4e0bf1

Remove sync stream

fb634e4

tarang-jain and others added 8 commits March 10, 2023 17:07

memory study commit (to be reversed)

a49ba87

Merge branch 'branch-23.04' of github.com:rapidsai/cuml into fea-memb…

4ed9fd7

…ership-vector

Style fix

fa7b44e

Merge branch 'branch-23.04' of github.com:rapidsai/cuml into fea-memb…

45f8ca4

…ership-vector

Remove print debug statements

367de04

Merge branch 'branch-23.04' of github.com:rapidsai/cuml into fea-memb…

980b1f7

…ership-vector

Merge branch 'branch-23.04' into fea-membership-vector

98aa237

Merge branch 'branch-23.04' into fea-membership-vector

d387026

cjnolet requested changes Mar 27, 2023

View reviewed changes

tarang-jain added 4 commits March 28, 2023 11:42

Updates after PR reviews

ed40e22

Merge branch 'fea-membership-vector' of https://github.com/tarang-jai…

387cde8

…n/cuml into fea-membership-vector

Merge branch 'branch-23.04' of github.com:rapidsai/cuml into fea-memb…

092b3f8

…ership-vector

Update height_argmax

ef85fd3

kuchenrolle mentioned this pull request Mar 29, 2023

.transform() does not generate probability distribution despite calculate_probabilites=True MaartenGr/BERTopic#1132

Closed

dantegd added feature request New feature or request non-breaking Non-breaking change labels Mar 29, 2023

cjnolet approved these changes Mar 30, 2023

View reviewed changes

tarang-jain added 2 commits March 30, 2023 09:24

Merge branch 'branch-23.04' into fea-membership-vector

17de9ec

Merge branch 'branch-23.04' into fea-membership-vector

38208ec

rapids-bot bot merged commit 79bfc47 into rapidsai:branch-23.04 Mar 31, 2023

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[FEA] membership_vector for HDBSCAN #5247

[FEA] membership_vector for HDBSCAN #5247

tarang-jain commented Feb 22, 2023 •

edited

tarang-jain Feb 28, 2023

cjnolet Feb 28, 2023

tarang-jain Feb 28, 2023

cjnolet Mar 7, 2023 •

edited

cjnolet Mar 7, 2023

cjnolet Feb 28, 2023

cjnolet Feb 28, 2023

cjnolet Mar 7, 2023 •

edited

cjnolet Mar 7, 2023 •

edited

tarang-jain Mar 7, 2023 •

edited

cjnolet Mar 7, 2023

cjnolet left a comment

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet Mar 27, 2023

cjnolet left a comment

cjnolet commented Mar 31, 2023


		rmm::device_uvector<value_idx> membership_argmax(m, stream);

		raft::matrix::argmax(


		raft::linalg::norm(handle, data_const_view, linf_norm_view, raft::linalg::LinfNorm, raft::linalg::Apply::ALONG_ROWS);

		raft::linalg::matrix_vector_op(handle, data_const_view, linf_norm_const_view, data_view, raft::linalg::Apply::ALONG_COLUMNS, [] __device__(value_t mat_in, value_t vec_in) {

[FEA] membership_vector for HDBSCAN #5247

[FEA] membership_vector for HDBSCAN #5247

Conversation

tarang-jain commented Feb 22, 2023 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

cjnolet Mar 7, 2023 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

cjnolet Mar 7, 2023 • edited

Choose a reason for hiding this comment

cjnolet Mar 7, 2023 • edited

Choose a reason for hiding this comment

tarang-jain Mar 7, 2023 • edited

Choose a reason for hiding this comment

Choose a reason for hiding this comment

cjnolet left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

cjnolet left a comment

Choose a reason for hiding this comment

cjnolet commented Mar 31, 2023

tarang-jain commented Feb 22, 2023 •

edited

cjnolet Mar 7, 2023 •

edited

cjnolet Mar 7, 2023 •

edited

cjnolet Mar 7, 2023 •

edited

tarang-jain Mar 7, 2023 •

edited