From 820e14d621f4c25a1e0dd6b879e01431dee2a300 Mon Sep 17 00:00:00 2001 From: Dante Gama Dessavre Date: Mon, 30 Aug 2021 17:08:41 -0500 Subject: [PATCH] Pin rmm to branch-21.10 and remove warnings from kmeans.hpp (#322) Authors: - Dante Gama Dessavre (https://github.com/dantegd) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Corey J. Nolet (https://github.com/cjnolet) - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/raft/pull/322 --- cpp/cmake/thirdparty/get_rmm.cmake | 6 +++--- cpp/include/raft/spectral/kmeans.hpp | 16 ++++++---------- 2 files changed, 9 insertions(+), 13 deletions(-) diff --git a/cpp/cmake/thirdparty/get_rmm.cmake b/cpp/cmake/thirdparty/get_rmm.cmake index e990ab1367..51f959a8d9 100644 --- a/cpp/cmake/thirdparty/get_rmm.cmake +++ b/cpp/cmake/thirdparty/get_rmm.cmake @@ -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}" @@ -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}) \ No newline at end of file +find_and_configure_rmm(${RAFT_MIN_VERSION_rmm}) diff --git a/cpp/include/raft/spectral/kmeans.hpp b/cpp/include/raft/spectral/kmeans.hpp index b6f0105487..d089b85518 100644 --- a/cpp/include/raft/spectral/kmeans.hpp +++ b/cpp/include/raft/spectral/kmeans.hpp @@ -256,7 +256,7 @@ static __global__ void minDistances2(index_type_t n, */ template 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) { @@ -341,7 +341,7 @@ static __global__ void divideCentroids( */ template 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) { @@ -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 @@ -450,7 +449,6 @@ static int initializeCentroids( thrust::uniform_real_distribution 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}; @@ -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 @@ -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)"); @@ -509,7 +506,7 @@ static int initializeCentroids( // Compute cluster sizes CUDA_TRY(cudaMemsetAsync(clusterSizes, 0, k * sizeof(index_type_t), stream)); - computeClusterSizes<<>>(n, k, codes, + computeClusterSizes<<>>(n, codes, clusterSizes); CHECK_CUDA(stream); @@ -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 @@ -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,