Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Build a Dendrogram class, adapt Louvain/Leiden/ECG to use it #1359

Merged
merged 21 commits into from
Feb 5, 2021

Conversation

ChuckHastings
Copy link
Collaborator

Preparing for MNMG Leiden and ECG identified an area for code cleanup.

The original cuGraph implementation of Louvain would flatten the hierarchical clustering as it was computed, filling (and returning) the final clustering. This adds an awkward step in the middle of the Louvain computation. Additionally, since Louvain (and Leiden and ECG which derive from it) is actually a hierarchical clustering algorithm it would be nice to generate the actual Dendrogram.

This PR implements a Dendrogram class, a function for flattening the Dendrogram, and modifies Louvain, Leiden and ECG to use the Dendrogram class.

It was suggested that the Dendrogram class could be moved to raft, decided to defer that until later, it's easy enough to move.

@ChuckHastings ChuckHastings requested review from a team as code owners January 27, 2021 18:48
@codecov-io
Copy link

codecov-io commented Jan 27, 2021

Codecov Report

Merging #1359 (9532b13) into branch-0.18 (2fb0725) will increase coverage by 0.25%.
The diff coverage is 52.38%.

Impacted file tree graph

@@               Coverage Diff               @@
##           branch-0.18    #1359      +/-   ##
===============================================
+ Coverage        60.38%   60.64%   +0.25%     
===============================================
  Files               67       69       +2     
  Lines             3029     3120      +91     
===============================================
+ Hits              1829     1892      +63     
- Misses            1200     1228      +28     
Impacted Files Coverage Δ
python/cugraph/centrality/__init__.py 100.00% <ø> (ø)
python/cugraph/dask/structure/renumber.py 0.00% <0.00%> (ø)
python/cugraph/link_analysis/pagerank.py 100.00% <ø> (ø)
python/cugraph/comms/comms.py 34.52% <25.00%> (ø)
python/cugraph/dask/common/input_utils.py 23.07% <28.57%> (+1.14%) ⬆️
python/cugraph/dask/common/mg_utils.py 37.50% <38.09%> (-2.50%) ⬇️
python/cugraph/community/spectral_clustering.py 72.54% <38.46%> (-11.67%) ⬇️
python/cugraph/structure/number_map.py 59.20% <50.00%> (+3.24%) ⬆️
python/cugraph/structure/graph.py 66.99% <76.47%> (+0.19%) ⬆️
python/cugraph/utilities/utils.py 72.44% <85.71%> (+0.88%) ⬆️
... and 16 more

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 7266cdb...9532b13. Read the comment docs.

@ChuckHastings
Copy link
Collaborator Author

rerun tests

@BradReesWork BradReesWork added improvement Improvement / enhancement to an existing function non-breaking Non-breaking change labels Jan 29, 2021
@BradReesWork BradReesWork added this to the 0.18 milestone Jan 29, 2021

#include <memory>
#include <rmm/device_buffer.hpp>
#include <vector>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we have any guidance on ordering include statements?

What I do is including more local headers first (rmm headers are more local than std C++ headers) like the following.

#include <rmm/device_buffer.hpp>

#include <memory>
#include <vector>

AFAIK, this is quite widely used in RAPIDS, and I started to follow this practice, and we may better stick to this unless we have other guidelines?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I had forgotten to leave the blank line and clang-format reordered them. This will be updated in my next push

template <typename vertex_t>
class Dendrogram {
public:
Dendrogram() : level_size_(), level_ptr_() {}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we need this? This does not do anything special, so can't we just rely on the default constructor (automatically generated by the compiler).

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Residual. An earlier version (pre-PR) had something that was actually initialized here. When I separated the code for flattening the Dendrogram out of the class (in preparation for moving to RAFT), the constructor became simpler.

Removed in next push.

namespace cugraph {

template <typename vertex_t>
class Dendrogram {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What's our guideline on class naming? IIRC, we're asked to use dendrogram_t instead of Dendrogram.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I like class names to start with a capital letter. dendrogram_t should be a type (which technically could be a class )

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We have a mix in our code base. I prefer the capital letter starting a class name, and using the _t suffix on type parameters in the templates.

We should codify our coding standards better.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I thought the same: naming convention is dendogram_t. If not, then what is the rule of naming types (a class is a type...)?


void add_level(vertex_t num_verts)
{
cudaStream_t stream{0};
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is not stream ready. Better follow rmm's approach?

https://github.com/rapidsai/rmm/blob/branch-0.18/include/rmm/device_buffer.hpp#L289

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That might be more than we need, especially since RAFT (where this class will be moved in 0.19) doesn't use RMM the same way. I'll push a change to partially address this. Let me know if you think I should go further than the next push or if that's sufficient for now (since we'll have to refactor this method for RAFT integration).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We at least need a mechanism to specify the stream to execute this. The RMM way is one example (and I think following the RMM style is good to enforce some consistency across RAPIDS projects) but if this is too much and there is a simpler way, I am OK as well. But we at least need a mechanism to specify the stream to execute this.


size_t num_levels() const { return level_size_.size(); }

vertex_t *get_level_ptr_unsafe(size_t level) const
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just FYI,

I have used nocheck instead of unsafe, but I think unsafe is a better name.

If no one opposes, I will submit a PR replacing nocheck with unsafe to be more consistent.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Little more on this.

I just realized that I used nocheck following cudf.

https://github.com/rapidsai/cudf/blob/branch-0.18/cpp/include/cudf/column/column_device_view.cuh#L167

cudf is using unsafe as well (e.g. set_bit_unsafe) but nocheck is used in cudf::column_device_view which is more widely used than bit manipulation utilities.

This being said, any preference between nocheck vs unsafe?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I should probably switch to nocheck. I recalled using the unsafe suffix in some previous work and it seemed to be appropriate here. However, after your comment I looked back at some of that. Generally what I have done in the past is:

  • nocheck suffix indicating don't check bounds
  • unsafe suffix indicating the method is not thread-safe

That seems like a better and more consistent naming approach.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Changed to nocheck in next push

#include <converters/permute_graph.cuh>
#include <ctime>
#include <utilities/error.hpp>
#include "utilities/graph_utils.cuh"
#include <utilities/graph_utils.cuh>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Better order include statements?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I can do that in the next push.

IdxT *parts,
ValT *weights)
__global__ void match_check_kernel(
IdxT size, IdxT num_verts, IdxT *offsets, IdxT *indices, IdxT *parts, ValT *weights)
Copy link
Contributor

@seunghwak seunghwak Feb 2, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we better use thrust?

thrust::transform(
...
[...]__device__(auto ...) {
  auto source = thrust::upper_bound(thrust::seq, ...);
  ...
}
);

What do we get by writing a custom kernel?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Don't know. This custom kernel was already there. My clang-format restructured this line.

I'd rather address this issue in the MNMG ECG implementation when we refactor this than to do this now. I'm adding a FIXME to remind me to do that.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree w/ @seunghwak . The long-held policy has been to avoid custom-written kernels as much as possible. And use thrust / CUB instead.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let me restate my response, perhaps I was too terse.

I agree, this seems like a place where we should just use thrust/cub. There's no reason that I can think of that this custom kernel would be better than a thrust/cub implementation. There's nothing in this kernel that takes advantage of features that would make a custom kernel perform better, and generally it increases the maintenance cost of the code.

I modified this file (ecg.cu) solely for the purpose of using the Dendrogram class, since the ECG algorithm calls Louvain and I modified Louvain. I'm reluctant, in general, to rewrite things in the code that I tangentially touch. It tends to make the work continue without getting anything finished and merged into the baseline.

The only reason this was even in the list of changes is because I ran clang-format and it changed the format of these lines of code. I made no changes relevant to this kernel or the calls to it.

I have a branch where I have started the MNMG ECG work. In that branch I will completely refactor how the ECG code works. My intention is to remove this kernel in that branch which hopefully will be merged sometime during release 0.19.

If we think it's critical that I do this for release 0.18, I can certainly take a few hours and refactor this once I'm done addressing the device_uvector request.

#include <rmm/thrust_rmm_allocator.h>
#include <community/dendrogram.cuh>
#include <experimental/graph_functions.hpp>
#include <raft/handle.hpp>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Better order include statements?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in next push.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@seunghwak, include statements order is dictated by clang-format

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If you put blank lines in between blocks of statements, clang-format will honor the blocks.


return runner(clustering, max_level, resolution);
thrust::device_vector<vertex_t> vertex_ids_v(graph.number_of_vertices);
Copy link
Contributor

@seunghwak seunghwak Feb 2, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If you're not dealing with non-arithmetic types, rmm::device_uvector will be more efficient and more stream ready.

Note that rmm::device_uvector does not invoke default constructor to initialize vector elements, so if your code expects initialization, you need to call thrust::fill().

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Changed in the next push.

@BradReesWork
Copy link
Member

@ChuckHastings any thought on moving Dendrogram into RAFT? I believe @cjnolet has a need

@ChuckHastings
Copy link
Collaborator Author

@ChuckHastings any thought on moving Dendrogram into RAFT? I believe @cjnolet has a need

My plan was to get everything working here first and move to raft in 0.19. Fewer moving parts that way. I think @cjnolet was not in a big hurry.

@cjnolet
Copy link
Member

cjnolet commented Feb 2, 2021

My plan was to get everything working here first and move to raft in 0.19. Fewer moving parts that way. I think @cjnolet was not in a big hurry.

Sounds good. I'll be moving a bunch of sparse prims into RAFT in 0.19 as well, including the single-linkage clustering.


size_t num_levels() const { return level_size_.size(); }

vertex_t *get_level_ptr_unsafe(size_t level) const
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is dangerous to have a const getter return a non-const pointer to a member. Something like this, for example, can happen:

template<typename T>
void foo(Dendrogram<T> const& d)
{
  d.get_level_ptr_unsafe(0)[0] = T{3};
}

int main(void)
{
  Dendrogram<int> d;

  foo(d);

  return 0;
}

foo() should not be permitted to modify the const& argument, d. But it can. This breaks the contract that the getter is supposed to fulfill (namely being an immutable getter).
(to compile this simple example, I changed device_buffer to std::vector<vertex_t>, but the idea is the same).

At the very least, get_level_ptr_unsafe() should not be a const member.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, that was careless. Next push will address this.

There should be a const version that returns a const pointer and a non-const version that returns a non-const pointer as it is used in both contexts.


vertex_t get_level_size_unsafe(size_t level) const { return level_size_[level]; }

vertex_t *current_level_begin() const { return get_level_ptr_unsafe(current_level()); }
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same comment about making this a non-const getter.


vertex_t *current_level_begin() const { return get_level_ptr_unsafe(current_level()); }

vertex_t *current_level_end() const { return current_level_begin() + current_level_size(); }
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same comment about making this a non-const getter.

namespace cugraph {

template <typename vertex_t>
class Dendrogram {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I thought the same: naming convention is dendogram_t. If not, then what is the rule of naming types (a class is a type...)?

IdxT *parts,
ValT *weights)
__global__ void match_check_kernel(
IdxT size, IdxT num_verts, IdxT *offsets, IdxT *indices, IdxT *parts, ValT *weights)
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree w/ @seunghwak . The long-held policy has been to avoid custom-written kernels as much as possible. And use thrust / CUB instead.

@@ -114,37 +135,51 @@ void ecg(raft::handle_t const &handle,
vertex_t ensemble_size,
vertex_t *clustering)
{
using graph_type = GraphCSRView<vertex_t, edge_t, weight_t>;

CUGRAPH_EXPECTS(graph.edge_data != nullptr,
"Invalid input argument: louvain expects a weighted graph");
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why public inheritance? If no method is overriden (virtual) then public inheritance is too much of an exposure (hence creating unnecessary coupling). Why not use private inheritance?

Moreover, if Louvain class is meant to be publicly derived (and it seems it is since it exposes at least one virtual method) then it should have a virtual destructor.

@@ -114,37 +135,51 @@ void ecg(raft::handle_t const &handle,
vertex_t ensemble_size,
vertex_t *clustering)
{
using graph_type = GraphCSRView<vertex_t, edge_t, weight_t>;

CUGRAPH_EXPECTS(graph.edge_data != nullptr,
"Invalid input argument: louvain expects a weighted graph");
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Again, confusion about class naming convention, shouldn't be <small_caps>_t? I understand that some are older (legacy) classes, but perhaps new classed might follow the rule?

#include <rmm/thrust_rmm_allocator.h>
#include <community/dendrogram.cuh>
#include <experimental/graph_functions.hpp>
#include <raft/handle.hpp>
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@seunghwak, include statements order is dictated by clang-format

Copy link
Member

@afender afender left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good.

@@ -114,37 +138,34 @@ void ecg(raft::handle_t const &handle,
vertex_t ensemble_size,
vertex_t *clustering)
{
using graph_type = GraphCSRView<vertex_t, edge_t, weight_t>;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should try and get rid of legacy classes in SG Louvain and ECG to use graph_t. Or the MG path could be used to support the 1 GPU case.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There's work currently scheduled for 0.19 that will adapt all of these to use the graph primitives. We should be able to consider getting rid of the legacy versions at that point.

@ChuckHastings
Copy link
Collaborator Author

rerun tests

1 similar comment
@BradReesWork
Copy link
Member

rerun tests

@ChuckHastings
Copy link
Collaborator Author

rerun tests

This is now failing on Pascal (using a method that only works on > Pascal). Working on an update to the python tests to disable these tests on a Pascal system.

@ChuckHastings ChuckHastings requested a review from a team as a code owner February 4, 2021 17:30
@ChuckHastings ChuckHastings requested review from a team as code owners February 5, 2021 00:20
@BradReesWork
Copy link
Member

@gpucibot merge

Copy link
Contributor

@rlratzel rlratzel left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, also happy to see the new notebook test skipping mechanism. One suggestion which need not hold up my approval (and I can file an issue if you agree and want to defer this):

I'm thinking it might be better to generalize the keywords a bit more with the side effect of making them more self documenting to NB users who run across them outside the context of our test infra. Maybe for example:

# AUTOMATED TESTING: skip

skips that NB when run from our test scripts

# AUTOMATED TESTING: skip on Pascal

skips that NB when run from our test scripts on a Pascal system

I'm just thinking the AUTOMATED TESTING tag makes it obvious this isn't a comment an unknowing user can just reword for clarity at some point later.

@rapids-bot rapids-bot bot merged commit 039b857 into rapidsai:branch-0.18 Feb 5, 2021
@ChuckHastings ChuckHastings deleted the fea_louvain_dendogram branch February 10, 2021 16:09
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
improvement Improvement / enhancement to an existing function non-breaking Non-breaking change
Projects
None yet
Development

Successfully merging this pull request may close these issues.

9 participants