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

Pin rmm to branch-21.10 and remove warnings from kmeans.hpp #322

Merged
merged 3 commits into from
Aug 30, 2021
Merged

Pin rmm to branch-21.10 and remove warnings from kmeans.hpp #322

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
4957dc8
FIX Pin RMM back to branch-21.10
dantegd Aug 28, 2021
e98bf0c
FIX Remove warnings from kmeans.hpp that only appear in 11.4
dantegd Aug 28, 2021
8201e22
FIX clang-format
dantegd Aug 28, 2021
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
6 changes: 3 additions & 3 deletions cpp/cmake/thirdparty/get_rmm.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,8 +32,8 @@ function(find_and_configure_rmm VERSION)
INSTALL_EXPORT_SET raft-exports
CPM_ARGS
GIT_REPOSITORY https://github.com/rapidsai/rmm.git
GIT_TAG 23bbe745af1d988224b5498f7b8e3fe3720532d4
GIT_SHALLOW FALSE
GIT_TAG branch-${MAJOR_AND_MINOR}
GIT_SHALLOW TRUE
OPTIONS "BUILD_TESTS OFF"
"BUILD_BENCHMARKS OFF"
"CUDA_STATIC_RUNTIME ${CUDA_STATIC_RUNTIME}"
Expand All @@ -44,4 +44,4 @@ endfunction()

set(RAFT_MIN_VERSION_rmm "${RAFT_VERSION_MAJOR}.${RAFT_VERSION_MINOR}.00")

find_and_configure_rmm(${RAFT_MIN_VERSION_rmm})
find_and_configure_rmm(${RAFT_MIN_VERSION_rmm})
16 changes: 6 additions & 10 deletions cpp/include/raft/spectral/kmeans.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -256,7 +256,7 @@ static __global__ void minDistances2(index_type_t n,
*/
template <typename index_type_t>
static __global__ void computeClusterSizes(
index_type_t n, index_type_t k, const index_type_t* __restrict__ codes,
index_type_t n, const index_type_t* __restrict__ codes,
index_type_t* __restrict__ clusterSizes) {
index_type_t i = threadIdx.x + blockIdx.x * blockDim.x;
while (i < n) {
Expand Down Expand Up @@ -341,7 +341,7 @@ static __global__ void divideCentroids(
*/
template <typename index_type_t, typename value_type_t>
static int chooseNewCentroid(handle_t const& handle, index_type_t n,
index_type_t d, index_type_t k, value_type_t rand,
index_type_t d, value_type_t rand,
const value_type_t* __restrict__ obs,
value_type_t* __restrict__ dists,
value_type_t* __restrict__ centroid) {
Expand All @@ -353,7 +353,6 @@ static int chooseNewCentroid(handle_t const& handle, index_type_t n,
index_type_t obsIndex;

auto stream = handle.get_stream();
auto cublas_h = handle.get_cublas_handle();
auto thrust_exec_policy = handle.get_thrust_policy();

// Compute cumulative sum of distances
Expand Down Expand Up @@ -450,7 +449,6 @@ static int initializeCentroids(
thrust::uniform_real_distribution<value_type_t> uniformDist(0, 1);

auto stream = handle.get_stream();
auto cublas_h = handle.get_cublas_handle();
auto thrust_exec_policy = handle.get_thrust_policy();

constexpr index_type_t grid_lower_bound{65535};
Expand Down Expand Up @@ -478,8 +476,7 @@ static int initializeCentroids(
thrust::fill(thrust_exec_policy, thrust::device_pointer_cast(dists),
thrust::device_pointer_cast(dists + n), 1);
CHECK_CUDA(stream);
if (chooseNewCentroid(handle, n, d, k, uniformDist(rng), obs, dists,
centroids))
if (chooseNewCentroid(handle, n, d, uniformDist(rng), obs, dists, centroids))
WARNING("error in k-means++ (could not pick centroid)");

// Compute distances from first centroid
Expand All @@ -491,7 +488,7 @@ static int initializeCentroids(
// Choose remaining centroids
for (i = 1; i < k; ++i) {
// Choose ith centroid
if (chooseNewCentroid(handle, n, d, k, uniformDist(rng), obs, dists,
if (chooseNewCentroid(handle, n, d, uniformDist(rng), obs, dists,
centroids + IDX(0, i, d)))
WARNING("error in k-means++ (could not pick centroid)");

Expand All @@ -509,7 +506,7 @@ static int initializeCentroids(

// Compute cluster sizes
CUDA_TRY(cudaMemsetAsync(clusterSizes, 0, k * sizeof(index_type_t), stream));
computeClusterSizes<<<gridDim_block, BLOCK_SIZE, 0, stream>>>(n, k, codes,
computeClusterSizes<<<gridDim_block, BLOCK_SIZE, 0, stream>>>(n, codes,
clusterSizes);
CHECK_CUDA(stream);

Expand Down Expand Up @@ -552,7 +549,6 @@ static int assignCentroids(handle_t const& handle, index_type_t n,
index_type_t* __restrict__ clusterSizes,
value_type_t* residual_host) {
auto stream = handle.get_stream();
auto cublas_h = handle.get_cublas_handle();
auto thrust_exec_policy = handle.get_thrust_policy();

// Compute distance between centroids and observation vectors
Expand Down Expand Up @@ -852,7 +848,7 @@ int kmeans(handle_t const& handle, index_type_t n, index_type_t d,
// conditions, such as if obs is corrupt (as seen as a result of a
// DataFrame column of NULL edge vals used to create the Graph)
while (emptyCentroid < k) {
if (chooseNewCentroid(handle, n, d, k, uniformDist(rng), obs, work,
if (chooseNewCentroid(handle, n, d, uniformDist(rng), obs, work,
centroids + IDX(0, emptyCentroid, d)))
WARNING("could not replace empty centroid");
if (assignCentroids(handle, n, d, k, obs, centroids, work, codes,
Expand Down