From 4957dc88818a0b7b4172c988f7fe43b98540af4e Mon Sep 17 00:00:00 2001 From: Dante Gama Dessavre Date: Sat, 28 Aug 2021 12:48:29 -0500 Subject: [PATCH 1/3] FIX Pin RMM back to branch-21.10 --- cpp/cmake/thirdparty/get_rmm.cmake | 6 +++--- 1 file changed, 3 insertions(+), 3 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}) From e98bf0ce951fc78689c076c408a76c821e8ab7a6 Mon Sep 17 00:00:00 2001 From: Dante Gama Dessavre Date: Sat, 28 Aug 2021 12:54:26 -0500 Subject: [PATCH 2/3] FIX Remove warnings from kmeans.hpp that only appear in 11.4 --- cpp/include/raft/spectral/kmeans.hpp | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/cpp/include/raft/spectral/kmeans.hpp b/cpp/include/raft/spectral/kmeans.hpp index b6f0105487..94e3135378 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,7 +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, + if (chooseNewCentroid(handle, n, d, uniformDist(rng), obs, dists, centroids)) WARNING("error in k-means++ (could not pick centroid)"); @@ -491,7 +489,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 +507,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 +550,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 +849,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, From 8201e226cabc7d81054b591e68fca01a860d19a2 Mon Sep 17 00:00:00 2001 From: Dante Gama Dessavre Date: Sat, 28 Aug 2021 13:06:20 -0500 Subject: [PATCH 3/3] FIX clang-format --- cpp/include/raft/spectral/kmeans.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/include/raft/spectral/kmeans.hpp b/cpp/include/raft/spectral/kmeans.hpp index 94e3135378..d089b85518 100644 --- a/cpp/include/raft/spectral/kmeans.hpp +++ b/cpp/include/raft/spectral/kmeans.hpp @@ -476,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, 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