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.

Sign up for GitHub

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

Update to CCCL 2.2.0. #5702

Merged
merged 6 commits into from
Dec 19, 2023
Merged

Update to CCCL 2.2.0. #5702

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
eb02720
Use cuCollections built with CCCL 2.2.0.
bdice Dec 6, 2023
5aab380
Apply changes needed for CUB 2.
bdice Dec 11, 2023
81f5bf7
Use rapids_cpm_cccl.
bdice Dec 11, 2023
d830f0a
Use forked raft.
bdice Dec 19, 2023
c3d7a3c
Use upstream artifacts.
bdice Dec 19, 2023
966f742
Revert testing changes.
bdice Dec 19, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 2 additions & 3 deletions cpp/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -220,9 +220,8 @@ if(BUILD_CUML_TESTS OR BUILD_PRIMS_TESTS)
find_package(Threads)
endif()

# thrust before rmm, rmm before raft so we get the right version of thrust/rmm
include(cmake/thirdparty/get_thrust.cmake)
include(cmake/thirdparty/get_libcudacxx.cmake)
# CCCL before RMM, and RMM before RAFT
include(cmake/thirdparty/get_cccl.cmake)
vyasr marked this conversation as resolved.
Show resolved Hide resolved
include(cmake/thirdparty/get_rmm.cmake)
include(cmake/thirdparty/get_raft.cmake)

Expand Down
13 changes: 6 additions & 7 deletions cpp/cmake/thirdparty/get_thrust.cmake → cpp/cmake/thirdparty/get_cccl.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,12 +12,11 @@
# the License.
# =============================================================================

# Use CPM to find or clone thrust
function(find_and_configure_thrust)
include(${rapids-cmake-dir}/cpm/thrust.cmake)
rapids_cpm_thrust(NAMESPACE cuml
BUILD_EXPORT_SET cuml-exports
INSTALL_EXPORT_SET cuml-exports)
# Use CPM to find or clone CCCL
function(find_and_configure_cccl)
include(${rapids-cmake-dir}/cpm/cccl.cmake)
rapids_cpm_cccl(BUILD_EXPORT_SET cuml-exports
INSTALL_EXPORT_SET cuml-exports)
endfunction()

find_and_configure_thrust()
find_and_configure_cccl()
35 changes: 0 additions & 35 deletions cpp/cmake/thirdparty/get_libcudacxx.cmake
View file
Open in desktop

This file was deleted.

19 changes: 12 additions & 7 deletions cpp/src/hdbscan/detail/membership.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2021-2022, NVIDIA CORPORATION.
* Copyright (c) 2021-2023, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -76,13 +76,18 @@ void get_probabilities(const raft::handle_t& handle,
rmm::device_uvector<value_t> deaths(n_clusters, stream);
thrust::fill(exec_policy, deaths.begin(), deaths.end(), 0.0f);

cudaError_t (*reduce_func)(void*,
Copy link
Member

Choose a reason for hiding this comment

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

Why was the change to a function pointer needed?

Copy link
Contributor Author

@bdice bdice Dec 19, 2023 •
edited
Loading

Choose a reason for hiding this comment

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

I'm not sure. CUB was failing to compile the previous code because of some API change. This code is borrowed from #5623. @hcho3 Do you have anything to add here?

@divyegala We may have to merge this as-is and revisit the question, or else CI will be blocked once the rest of RAPIDS is updated to CCCL 2.2.0.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The change was made in b1302bd (#5623) but the commit message only says "Fix build."

Copy link
Member

Choose a reason for hiding this comment

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

Okay, it's not egregious so I will approve for now.

Copy link
Contributor

Choose a reason for hiding this comment

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

I kept getting the build error like

error: no instance of function template "detail::Utils::cub_segmented_reduce" matches the argument list
            argument types are: (float *, float *, int, int *, cudaStream_t, <unknown-type>)

because the compiler failed to build the type of cub::DeviceSegmentedReduce::Max. I had to explicitly declare the type of cub::DeviceSegmentedReduce::Max by declaring a function pointer.

size_t&,
const value_t*,
value_t*,
int,
const value_idx*,
const value_idx*,
cudaStream_t,
bool) =
cub::DeviceSegmentedReduce::Max<const value_t*, value_t*, const value_idx*, const value_idx*>;
Utils::cub_segmented_reduce(
lambdas,
deaths.data(),
n_clusters,
sorted_parents_offsets.data(),
stream,
cub::DeviceSegmentedReduce::Max<const value_t*, value_t*, const value_idx*, const value_idx*>);
lambdas, deaths.data(), n_clusters, sorted_parents_offsets.data(), stream, reduce_func);

// Calculate probability per point
thrust::fill(exec_policy, probabilities, probabilities + n_leaves, 0.0f);
Expand Down
23 changes: 16 additions & 7 deletions cpp/src/hdbscan/detail/stabilities.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -100,13 +100,22 @@ void compute_stabilities(const raft::handle_t& handle,
thrust::make_counting_iterator(n_edges),
births_init_op);

Utils::cub_segmented_reduce(
lambdas,
births_parent_min.data() + 1,
n_clusters - 1,
sorted_parents_offsets.data() + 1,
stream,
cub::DeviceSegmentedReduce::Min<const value_t*, value_t*, const value_idx*, const value_idx*>);
cudaError_t (*reduce_func)(void*,
size_t&,
const value_t*,
value_t*,
int,
const value_idx*,
const value_idx*,
cudaStream_t,
bool) =
cub::DeviceSegmentedReduce::Min<const value_t*, value_t*, const value_idx*, const value_idx*>;
Utils::cub_segmented_reduce(lambdas,
births_parent_min.data() + 1,
n_clusters - 1,
sorted_parents_offsets.data() + 1,
stream,
reduce_func);
// finally, we find minimum between initialized births where parent=child
// and births of parents for their children
auto births_zip =
Expand Down
16 changes: 9 additions & 7 deletions cpp/src/hdbscan/prediction_data.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -130,13 +130,15 @@ void generate_prediction_data(const raft::handle_t& handle,
prediction_data.set_n_clusters(handle, n_clusters);

// this is to find maximum lambdas of all children under a parent
detail::Utils::cub_segmented_reduce(
lambdas,
prediction_data.get_deaths(),
n_clusters,
sorted_parents_offsets.data(),
stream,
cub::DeviceSegmentedReduce::Max<const float*, float*, const int*, const int*>);
cudaError_t (*reduce_func)(
void*, size_t&, const float*, float*, int, const int*, const int*, cudaStream_t, bool) =
cub::DeviceSegmentedReduce::Max<const float*, float*, const int*, const int*>;
detail::Utils::cub_segmented_reduce(lambdas,
prediction_data.get_deaths(),
n_clusters,
sorted_parents_offsets.data(),
stream,
reduce_func);

rmm::device_uvector<int> is_leaf_cluster(n_clusters, stream);
thrust::fill(exec_policy, is_leaf_cluster.begin(), is_leaf_cluster.end(), 1);
Expand Down
4 changes: 3 additions & 1 deletion fetch_rapids.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
# =============================================================================
# Copyright (c) 2022, NVIDIA CORPORATION.
# Copyright (c) 2022-2023, NVIDIA CORPORATION.
#
# Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except
# in compliance with the License. You may obtain a copy of the License at
Expand All @@ -11,6 +11,8 @@
# or implied. See the License for the specific language governing permissions and limitations under
# the License.
# =============================================================================
set(rapids-cmake-repo bdice/rapids-cmake)
set(rapids-cmake-branch cuco-cccl-2.2.0)
bdice marked this conversation as resolved.
Show resolved Hide resolved
if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUML_RAPIDS.cmake)
file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-24.02/RAPIDS.cmake
${CMAKE_CURRENT_BINARY_DIR}/CUML_RAPIDS.cmake
Expand Down
Loading