From 43c9b93aebc541b99a1b755a37e65bca608146f8 Mon Sep 17 00:00:00 2001 From: divyegala Date: Tue, 26 Jul 2022 17:10:46 -0700 Subject: [PATCH 1/8] working through --- cpp/include/raft/cluster/detail/kmeans.cuh | 104 +++++----- .../raft/cluster/detail/kmeans_common.cuh | 186 +++++++++--------- cpp/include/raft/cluster/kmeans.cuh | 78 ++++---- cpp/include/raft/detail/mdarray.hpp | 2 +- cpp/test/cluster/kmeans.cu | 18 +- 5 files changed, 194 insertions(+), 194 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index c3ca60973a..0d39afeeb8 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -55,8 +55,8 @@ namespace detail { template void initRandom(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroids) + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroids) { cudaStream_t stream = handle.get_stream(); auto n_clusters = params.n_clusters; @@ -80,8 +80,8 @@ void initRandom(const raft::handle_t& handle, template void kmeansPlusPlus(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroidsRawData, + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroidsRawData, rmm::device_uvector& workspace) { cudaStream_t stream = handle.get_stream(); @@ -104,22 +104,22 @@ void kmeansPlusPlus(const raft::handle_t& handle, // temporary buffers std::vector h_wt(n_samples); - auto centroidCandidates = raft::make_device_matrix(n_trials, n_features, stream); - auto costPerCandidate = raft::make_device_vector(n_trials, stream); - auto minClusterDistance = raft::make_device_vector(n_samples, stream); - auto distBuffer = raft::make_device_matrix(n_trials, n_samples, stream); + auto centroidCandidates = raft::make_device_matrix(n_trials, n_features, stream); + auto costPerCandidate = raft::make_device_vector(n_trials, stream); + auto minClusterDistance = raft::make_device_vector(n_samples, stream); + auto distBuffer = raft::make_device_matrix(n_trials, n_samples, stream); rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); rmm::device_scalar clusterCost(stream); rmm::device_scalar> minClusterIndexAndDistance(stream); // L2 norm of X: ||c||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(n_samples, stream); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { raft::linalg::rowNorm( - L2NormX.data(), X.data(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + L2NormX.data_handle(), X.data_handle(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); } std::mt19937 gen(params.rng_state.seed); @@ -127,15 +127,15 @@ void kmeansPlusPlus(const raft::handle_t& handle, // <<< Step-1 >>>: C <-- sample a point uniformly at random from X auto initialCentroid = - raft::make_device_matrix_view(X.data() + dis(gen) * n_features, 1, n_features); + raft::make_device_matrix_view(X.data_handle() + dis(gen) * n_features, 1, n_features); int n_clusters_picked = 1; // store the chosen centroid in the buffer - raft::copy(centroidsRawData.data(), initialCentroid.data(), initialCentroid.size(), stream); + raft::copy(centroidsRawData.data_handle(), initialCentroid.data_handle(), initialCentroid.size(), stream); // C = initial set of centroids - auto centroids = raft::make_device_matrix_view( - centroidsRawData.data(), initialCentroid.extent(0), initialCentroid.extent(1)); + auto centroids = raft::make_device_matrix_view( + centroidsRawData.data_handle(), initialCentroid.extent(0), initialCentroid.extent(1)); // <<< End of Step-1 >>> // Calculate cluster distance, d^2(x, C), for all the points x in X to the nearest centroid @@ -155,7 +155,7 @@ void kmeansPlusPlus(const raft::handle_t& handle, // <<< Step-3 >>> : Sample x in X with probability p_x = d^2(x, C) / phi_X (C) // Choose 'n_trials' centroid candidates from X with probability proportional to the squared // distance to the nearest existing cluster - raft::copy(h_wt.data(), minClusterDistance.data(), minClusterDistance.size(), stream); + raft::copy(h_wt.data(), minClusterDistance.data_handle(), minClusterDistance.size(), stream); handle.sync_stream(stream); // Note - n_trials is relative small here, we don't need raft::gather call @@ -163,9 +163,9 @@ void kmeansPlusPlus(const raft::handle_t& handle, for (int cIdx = 0; cIdx < n_trials; ++cIdx) { auto rand_idx = d(gen); auto randCentroid = - raft::make_device_matrix_view(X.data() + n_features * rand_idx, 1, n_features); - raft::copy(centroidCandidates.data() + cIdx * n_features, - randCentroid.data(), + raft::make_device_matrix_view(X.data_handle() + n_features * rand_idx, 1, n_features); + raft::copy(centroidCandidates.data_handle() + cIdx * n_features, + randCentroid.data_handle(), randCentroid.size(), stream); } @@ -182,9 +182,9 @@ void kmeansPlusPlus(const raft::handle_t& handle, // minClusterDistance that includes candidate-i auto minDistBuf = distBuffer.view(); raft::linalg::matrixVectorOp( - minDistBuf.data(), - pwd.data(), - minClusterDistance.data(), + minDistBuf.data_handle(), + pwd.data_handle(), + minClusterDistance.data_handle(), pwd.extent(1), pwd.extent(0), true, @@ -194,8 +194,8 @@ void kmeansPlusPlus(const raft::handle_t& handle, // Calculate costPerCandidate[n_trials] where costPerCandidate[i] is the cluster cost when using // centroid candidate-i - raft::linalg::reduce(costPerCandidate.data(), - minDistBuf.data(), + raft::linalg::reduce(costPerCandidate.data_handle(), + minDistBuf.data_handle(), minDistBuf.extent(1), minDistBuf.extent(0), static_cast(0), @@ -210,7 +210,7 @@ void kmeansPlusPlus(const raft::handle_t& handle, size_t temp_storage_bytes = 0; cub::DeviceReduce::ArgMin(nullptr, temp_storage_bytes, - costPerCandidate.data(), + costPerCandidate.data_handle(), minClusterIndexAndDistance.data(), costPerCandidate.extent(0)); @@ -220,7 +220,7 @@ void kmeansPlusPlus(const raft::handle_t& handle, // Run argmin-reduction cub::DeviceReduce::ArgMin(workspace.data(), temp_storage_bytes, - costPerCandidate.data(), + costPerCandidate.data_handle(), minClusterIndexAndDistance.data(), costPerCandidate.extent(0)); @@ -230,13 +230,13 @@ void kmeansPlusPlus(const raft::handle_t& handle, /// <<< Step-4 >>>: C = C U {x} // Update minimum cluster distance corresponding to the chosen centroid candidate - raft::copy(minClusterDistance.data(), - minDistBuf.data() + bestCandidateIdx * n_samples, + raft::copy(minClusterDistance.data_handle(), + minDistBuf.data_handle() + bestCandidateIdx * n_samples, n_samples, stream); - raft::copy(centroidsRawData.data() + n_clusters_picked * n_features, - centroidCandidates.data() + bestCandidateIdx * n_features, + raft::copy(centroidsRawData.data_handle() + n_clusters_picked * n_features, + centroidCandidates.data_handle() + bestCandidateIdx * n_features, n_features, stream); @@ -252,9 +252,9 @@ void kmeansPlusPlus(const raft::handle_t& handle, template void kmeans_fit_main(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_vector_view& weight, - const raft::device_matrix_view& centroidsRawData, + const raft::device_matrix_view& X, + const raft::device_vector_view& weight, + const raft::device_matrix_view& centroidsRawData, const raft::host_scalar_view& inertia, const raft::host_scalar_view& n_iter, rmm::device_uvector& workspace) @@ -270,7 +270,7 @@ void kmeans_fit_main(const raft::handle_t& handle, // - key is the index of nearest cluster // - value is the distance to the nearest cluster auto minClusterAndDistance = - raft::make_device_vector>(n_samples, stream); + raft::make_device_vector, IndexT>(n_samples, stream); // temporary buffer to store L2 norm of centroids or distance matrix, // destructor releases the resource @@ -278,20 +278,20 @@ void kmeans_fit_main(const raft::handle_t& handle, // temporary buffer to store intermediate centroids, destructor releases the // resource - auto newCentroids = raft::make_device_matrix(n_clusters, n_features, stream); + auto newCentroids = raft::make_device_matrix(n_clusters, n_features, stream); // temporary buffer to store weights per cluster, destructor releases the // resource - auto wtInCluster = raft::make_device_vector(n_clusters, stream); + auto wtInCluster = raft::make_device_vector(n_clusters, stream); rmm::device_scalar> clusterCostD(stream); // L2 norm of X: ||x||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(n_samples, stream); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { raft::linalg::rowNorm( - L2NormX.data(), X.data(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + L2NormX.data_handle(), X.data_handle(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); } RAFT_LOG_DEBUG( @@ -306,7 +306,7 @@ void kmeans_fit_main(const raft::handle_t& handle, "cluster centers", n_iter[0]); - auto centroids = raft::make_device_matrix_view(centroidsRawData.data(), n_clusters, n_features); + auto centroids = raft::make_device_matrix_view(centroidsRawData.data_handle(), n_clusters, n_features); // computes minClusterAndDistance[0:n_samples) where // minClusterAndDistance[i] is a pair where @@ -329,27 +329,27 @@ void kmeans_fit_main(const raft::handle_t& handle, cub::TransformInputIterator, cub::KeyValuePair*> - itr(minClusterAndDistance.data(), conversion_op); + itr(minClusterAndDistance.data_handle(), conversion_op); workspace.resize(n_samples, stream); // Calculates weighted sum of all the samples assigned to cluster-i and store the // result in newCentroids[i] - raft::linalg::reduce_rows_by_key((DataT*)X.data(), + raft::linalg::reduce_rows_by_key((DataT*)X.data_handle(), X.extent(1), itr, - weight.data(), + weight.data_handle(), workspace.data(), X.extent(0), X.extent(1), n_clusters, - newCentroids.data(), + newCentroids.data_handle(), stream); // Reduce weights by key to compute weight in each cluster - raft::linalg::reduce_cols_by_key(weight.data(), + raft::linalg::reduce_cols_by_key(weight.data_handle(), itr, - wtInCluster.data(), + wtInCluster.data_handle(), (IndexT)1, (IndexT)weight.extent(0), (IndexT)n_clusters, @@ -361,9 +361,9 @@ void kmeans_fit_main(const raft::handle_t& handle, // of samples in cluster-i. // Note - when wtInCluster[i] is 0, newCentroid[i] is reset to 0 raft::linalg::matrixVectorOp( - newCentroids.data(), - newCentroids.data(), - wtInCluster.data(), + newCentroids.data_handle(), + newCentroids.data_handle(), + wtInCluster.data_handle(), newCentroids.extent(1), newCentroids.extent(0), true, @@ -377,15 +377,15 @@ void kmeans_fit_main(const raft::handle_t& handle, stream); // copy centroids[i] to newCentroids[i] when wtInCluster[i] is 0 - cub::ArgIndexInputIterator itr_wt(wtInCluster.data()); + cub::ArgIndexInputIterator itr_wt(wtInCluster.data_handle()); raft::matrix::gather_if( - centroids.data(), + centroids.data_handle(), centroids.extent(1), centroids.extent(0), itr_wt, itr_wt, wtInCluster.size(), - newCentroids.data(), + newCentroids.data_handle(), [=] __device__(cub::KeyValuePair map) { // predicate // copy when the # of samples in the cluster is 0 if (map.value == 0) @@ -413,9 +413,9 @@ void kmeans_fit_main(const raft::handle_t& handle, newCentroids.data()); DataT sqrdNormError = 0; - raft::copy(&sqrdNormError, sqrdNorm.data(), sqrdNorm.size(), stream); + raft::copy(&sqrdNormError, sqrdNorm.data_handle(), sqrdNorm.size(), stream); - raft::copy(centroidsRawData.data(), newCentroids.data(), newCentroids.size(), stream); + raft::copy(centroidsRawData.data_handle(), newCentroids.data_handle(), newCentroids.size(), stream); bool done = false; if (params.inertia_check) { diff --git a/cpp/include/raft/cluster/detail/kmeans_common.cuh b/cpp/include/raft/cluster/detail/kmeans_common.cuh index 0d46b532c4..13c2bc975a 100644 --- a/cpp/include/raft/cluster/detail/kmeans_common.cuh +++ b/cpp/include/raft/cluster/detail/kmeans_common.cuh @@ -150,9 +150,9 @@ void countLabels(const raft::handle_t& handle, stream)); } -template +template void checkWeight(const raft::handle_t& handle, - const raft::device_vector_view& weight, + const raft::device_vector_view& weight, rmm::device_uvector& workspace) { cudaStream_t stream = handle.get_stream(); @@ -161,14 +161,14 @@ void checkWeight(const raft::handle_t& handle, size_t temp_storage_bytes = 0; RAFT_CUDA_TRY(cub::DeviceReduce::Sum( - nullptr, temp_storage_bytes, weight.data(), wt_aggr.data(), n_samples, stream)); + nullptr, temp_storage_bytes, weight.data_handle(), wt_aggr.data_handle(), n_samples, stream)); workspace.resize(temp_storage_bytes, stream); RAFT_CUDA_TRY(cub::DeviceReduce::Sum( - workspace.data(), temp_storage_bytes, weight.data(), wt_aggr.data(), n_samples, stream)); + workspace.data(), temp_storage_bytes, weight.data_handle(), wt_aggr.data_handle(), n_samples, stream)); DataT wt_sum = 0; - raft::copy(&wt_sum, wt_aggr.data(), 1, stream); + raft::copy(&wt_sum, wt_aggr.data_handle(), 1, stream); handle.sync_stream(stream); if (wt_sum != n_samples) { @@ -179,8 +179,8 @@ void checkWeight(const raft::handle_t& handle, auto scale = static_cast(n_samples) / wt_sum; raft::linalg::unaryOp( - weight.data(), - weight.data(), + weight.data_handle(), + weight.data_handle(), n_samples, [=] __device__(const DataT& wt) { return wt * scale; }, stream); @@ -201,9 +201,9 @@ IndexT getCentroidsBatchSize(const KMeansParams& params, IndexT n_local_clusters return (minVal == 0) ? n_local_clusters : minVal; } -template +template void computeClusterCost(const raft::handle_t& handle, - const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& minClusterDistance, rmm::device_uvector& workspace, const raft::device_scalar_view& clusterCost, ReductionOpT reduction_op) @@ -212,8 +212,8 @@ void computeClusterCost(const raft::handle_t& handle, size_t temp_storage_bytes = 0; RAFT_CUDA_TRY(cub::DeviceReduce::Reduce(nullptr, temp_storage_bytes, - minClusterDistance.data(), - clusterCost.data(), + minClusterDistance.data_handle(), + clusterCost.data_handle(), minClusterDistance.size(), reduction_op, DataT(), @@ -223,8 +223,8 @@ void computeClusterCost(const raft::handle_t& handle, RAFT_CUDA_TRY(cub::DeviceReduce::Reduce(workspace.data(), temp_storage_bytes, - minClusterDistance.data(), - clusterCost.data(), + minClusterDistance.data_handle(), + clusterCost.data_handle(), minClusterDistance.size(), reduction_op, DataT(), @@ -233,9 +233,9 @@ void computeClusterCost(const raft::handle_t& handle, template void sampleCentroids(const raft::handle_t& handle, - const raft::device_matrix_view& X, - const raft::device_vector_view& minClusterDistance, - const raft::device_vector_view& isSampleCentroid, + const raft::device_matrix_view& X, + const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& isSampleCentroid, SamplingOp& select_op, rmm::device_uvector& inRankCp, rmm::device_uvector& workspace) @@ -247,13 +247,13 @@ void sampleCentroids(const raft::handle_t& handle, auto nSelected = raft::make_device_scalar(0, stream); cub::ArgIndexInputIterator ip_itr(minClusterDistance.data()); auto sampledMinClusterDistance = - raft::make_device_vector>(n_local_samples, stream); + raft::make_device_vector, IndexT>(n_local_samples, stream); size_t temp_storage_bytes = 0; RAFT_CUDA_TRY(cub::DeviceSelect::If(nullptr, temp_storage_bytes, ip_itr, - sampledMinClusterDistance.data(), - nSelected.data(), + sampledMinClusterDistance.data_handle(), + nSelected.data_handle(), n_local_samples, select_op, stream)); @@ -263,19 +263,19 @@ void sampleCentroids(const raft::handle_t& handle, RAFT_CUDA_TRY(cub::DeviceSelect::If(workspace.data(), temp_storage_bytes, ip_itr, - sampledMinClusterDistance.data(), - nSelected.data(), + sampledMinClusterDistance.data_handle(), + nSelected.data_handle(), n_local_samples, select_op, stream)); IndexT nPtsSampledInRank = 0; - raft::copy(&nPtsSampledInRank, nSelected.data(), 1, stream); + raft::copy(&nPtsSampledInRank, nSelected.data_handle(), 1, stream); handle.sync_stream(stream); - IndexT* rawPtr_isSampleCentroid = isSampleCentroid.data(); + IndexT* rawPtr_isSampleCentroid = isSampleCentroid.data_handle(); thrust::for_each_n(handle.get_thrust_policy(), - sampledMinClusterDistance.data(), + sampledMinClusterDistance.data_handle(), nPtsSampledInRank, [=] __device__(cub::KeyValuePair val) { rawPtr_isSampleCentroid[val.key] = 1; @@ -283,10 +283,10 @@ void sampleCentroids(const raft::handle_t& handle, inRankCp.resize(nPtsSampledInRank * n_features, stream); - raft::matrix::gather((DataT*)X.data(), + raft::matrix::gather((DataT*)X.data_handle(), X.extent(1), X.extent(0), - sampledMinClusterDistance.data(), + sampledMinClusterDistance.data_handle(), nPtsSampledInRank, inRankCp.data(), [=] __device__(cub::KeyValuePair val) { // MapTransformOp @@ -299,9 +299,9 @@ void sampleCentroids(const raft::handle_t& handle, // result will be stored in 'pairwiseDistance[n x k]' template void pairwise_distance_kmeans(const raft::handle_t& handle, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroids, - const raft::device_matrix_view& pairwiseDistance, + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroids, + const raft::device_matrix_view& pairwiseDistance, rmm::device_uvector& workspace, raft::distance::DistanceType metric) { @@ -312,10 +312,10 @@ void pairwise_distance_kmeans(const raft::handle_t& handle, ASSERT(X.extent(1) == centroids.extent(1), "# features in dataset and centroids are different (must be same)"); - raft::distance::pairwise_distance(handle, - X.data(), - centroids.data(), - pairwiseDistance.data(), + raft::distance::pairwise_distance(handle, + X.data_handle(), + centroids.data_handle(), + pairwiseDistance.data_handle(), n_samples, n_clusters, n_features, @@ -327,8 +327,8 @@ void pairwise_distance_kmeans(const raft::handle_t& handle, // in 'out' does not modify the input template void shuffleAndGather(const raft::handle_t& handle, - const raft::device_matrix_view& in, - const raft::device_matrix_view& out, + const raft::device_matrix_view& in, + const raft::device_matrix_view& out, uint32_t n_samples_to_gather, uint64_t seed, rmm::device_uvector* workspace = nullptr) @@ -337,12 +337,12 @@ void shuffleAndGather(const raft::handle_t& handle, auto n_samples = in.extent(0); auto n_features = in.extent(1); - auto indices = raft::make_device_vector(n_samples, stream); + auto indices = raft::make_device_vector(n_samples, stream); if (workspace) { // shuffle indices on device raft::random::permute( - indices.data(), nullptr, nullptr, (IndexT)in.extent(1), (IndexT)in.extent(0), true, stream); + indices.data_handle(), nullptr, nullptr, (IndexT)in.extent(1), (IndexT)in.extent(0), true, stream); } else { // shuffle indices on host and copy to device... std::vector ht_indices(n_samples); @@ -352,15 +352,15 @@ void shuffleAndGather(const raft::handle_t& handle, std::mt19937 gen(seed); std::shuffle(ht_indices.begin(), ht_indices.end(), gen); - raft::copy(indices.data(), ht_indices.data(), indices.size(), stream); + raft::copy(indices.data_handle(), ht_indices.data(), indices.size(), stream); } - raft::matrix::gather((DataT*)in.data(), + raft::matrix::gather((DataT*)in.data_handle(), in.extent(1), in.extent(0), - indices.data(), + indices.data_handle(), n_samples_to_gather, - out.data(), + out.data_handle(), stream); } @@ -371,10 +371,10 @@ template void minClusterAndDistanceCompute( const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view X, - const raft::device_matrix_view centroids, - const raft::device_vector_view>& minClusterAndDistance, - const raft::device_vector_view& L2NormX, + const raft::device_matrix_view X, + const raft::device_matrix_view centroids, + const raft::device_vector_view, IndexT>& minClusterAndDistance, + const raft::device_vector_view& L2NormX, rmm::device_uvector& L2NormBuf_OR_DistBuf, rmm::device_uvector& workspace) { @@ -390,7 +390,7 @@ void minClusterAndDistanceCompute( metric == raft::distance::DistanceType::L2SqrtExpanded) { L2NormBuf_OR_DistBuf.resize(n_clusters, stream); raft::linalg::rowNorm(L2NormBuf_OR_DistBuf.data(), - centroids.data(), + centroids.data_handle(), centroids.extent(1), centroids.extent(0), raft::linalg::L2Norm, @@ -402,16 +402,16 @@ void minClusterAndDistanceCompute( // Note - pairwiseDistance and centroidsNorm share the same buffer // centroidsNorm [n_clusters] - tensor wrapper around centroids L2 Norm - auto centroidsNorm = raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); + auto centroidsNorm = raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); // pairwiseDistance[ns x nc] - tensor wrapper around the distance buffer auto pairwiseDistance = - raft::make_device_matrix_view(L2NormBuf_OR_DistBuf.data(), dataBatchSize, centroidsBatchSize); + raft::make_device_matrix_view(L2NormBuf_OR_DistBuf.data(), dataBatchSize, centroidsBatchSize); cub::KeyValuePair initial_value(0, std::numeric_limits::max()); thrust::fill(handle.get_thrust_policy(), - minClusterAndDistance.data(), - minClusterAndDistance.data() + minClusterAndDistance.size(), + minClusterAndDistance.data_handle(), + minClusterAndDistance.data_handle() + minClusterAndDistance.size(), initial_value); // tile over the input dataset @@ -422,13 +422,13 @@ void minClusterAndDistanceCompute( // datasetView [ns x n_features] - view representing the current batch of // input dataset auto datasetView = - raft::make_device_matrix_view(X.data() + (dIdx * n_features), ns, n_features); + raft::make_device_matrix_view(X.data_handle() + (dIdx * n_features), ns, n_features); // minClusterAndDistanceView [ns x n_clusters] auto minClusterAndDistanceView = - raft::make_device_vector_view(minClusterAndDistance.data() + dIdx, ns); + raft::make_device_vector_view(minClusterAndDistance.data_handle() + dIdx, ns); - auto L2NormXView = raft::make_device_vector_view(L2NormX.data() + dIdx, ns); + auto L2NormXView = raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { @@ -438,22 +438,22 @@ void minClusterAndDistanceCompute( // centroidsView [nc x n_features] - view representing the current batch // of centroids auto centroidsView = - raft::make_device_matrix_view(centroids.data() + (cIdx * n_features), nc, n_features); + raft::make_device_matrix_view(centroids.data_handle() + (cIdx * n_features), nc, n_features); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - auto centroidsNormView = raft::make_device_vector_view(centroidsNorm.data() + cIdx, nc); + auto centroidsNormView = raft::make_device_vector_view(centroidsNorm.data_handle() + cIdx, nc); workspace.resize((sizeof(int)) * ns, stream); FusedL2NNReduceOp redOp(cIdx); raft::distance::KVPMinReduce pairRedOp; raft::distance::fusedL2NN, IndexT>( - minClusterAndDistanceView.data(), - datasetView.data(), - centroidsView.data(), - L2NormXView.data(), - centroidsNormView.data(), + minClusterAndDistanceView.data_handle(), + datasetView.data_handle(), + centroidsView.data_handle(), + L2NormXView.data_handle(), + centroidsNormView.data_handle(), ns, nc, n_features, @@ -466,7 +466,7 @@ void minClusterAndDistanceCompute( } else { // pairwiseDistanceView [ns x nc] - view representing the pairwise // distance for current batch - auto pairwiseDistanceView = raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); + auto pairwiseDistanceView = raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); // calculate pairwise distance between current tile of cluster centroids // and input dataset @@ -477,8 +477,8 @@ void minClusterAndDistanceCompute( // calculates the closest centroid and the distance to the closest // centroid raft::linalg::coalescedReduction( - minClusterAndDistanceView.data(), - pairwiseDistanceView.data(), + minClusterAndDistanceView.data_handle(), + pairwiseDistanceView.data_handle(), pairwiseDistanceView.extent(1), pairwiseDistanceView.extent(0), initial_value, @@ -502,10 +502,10 @@ void minClusterAndDistanceCompute( template void minClusterDistanceCompute(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroids, - const raft::device_vector_view& minClusterDistance, - const raft::device_vector_view& L2NormX, + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroids, + const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& L2NormX, rmm::device_uvector& L2NormBuf_OR_DistBuf, rmm::device_uvector& workspace) { @@ -522,7 +522,7 @@ void minClusterDistanceCompute(const raft::handle_t& handle, metric == raft::distance::DistanceType::L2SqrtExpanded) { L2NormBuf_OR_DistBuf.resize(n_clusters, stream); raft::linalg::rowNorm(L2NormBuf_OR_DistBuf.data(), - centroids.data(), + centroids.data_handle(), centroids.extent(1), centroids.extent(0), raft::linalg::L2Norm, @@ -535,14 +535,14 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // Note - pairwiseDistance and centroidsNorm share the same buffer // centroidsNorm [n_clusters] - tensor wrapper around centroids L2 Norm auto centroidsNorm = - raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); + raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); // pairwiseDistance[ns x nc] - tensor wrapper around the distance buffer - auto pairwiseDistance = raft::make_device_matrix_view( + auto pairwiseDistance = raft::make_device_matrix_view( L2NormBuf_OR_DistBuf.data(), dataBatchSize, centroidsBatchSize); thrust::fill(handle.get_thrust_policy(), - minClusterDistance.data(), - minClusterDistance.data() + minClusterDistance.size(), + minClusterDistance.data_handle(), + minClusterDistance.data_handle() + minClusterDistance.size(), std::numeric_limits::max()); // tile over the input data and calculate distance matrix [n_samples x @@ -554,13 +554,13 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // datasetView [ns x n_features] - view representing the current batch of // input dataset auto datasetView = - raft::make_device_matrix_view(X.data() + dIdx * n_features, ns, n_features); + raft::make_device_matrix_view(X.data_handle() + dIdx * n_features, ns, n_features); // minClusterDistanceView [ns x n_clusters] auto minClusterDistanceView = - raft::make_device_vector_view(minClusterDistance.data() + dIdx, ns); + raft::make_device_vector_view(minClusterDistance.data_handle() + dIdx, ns); - auto L2NormXView = raft::make_device_vector_view(L2NormX.data() + dIdx, ns); + auto L2NormXView = raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { @@ -570,22 +570,22 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // centroidsView [nc x n_features] - view representing the current batch // of centroids auto centroidsView = - raft::make_device_matrix_view(centroids.data() + cIdx * n_features, nc, n_features); + raft::make_device_matrix_view(centroids.data() + cIdx * n_features, nc, n_features); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { auto centroidsNormView = - raft::make_device_vector_view(centroidsNorm.data() + cIdx, nc); + raft::make_device_vector_view(centroidsNorm.data() + cIdx, nc); workspace.resize((sizeof(IndexT)) * ns, stream); FusedL2NNReduceOp redOp(cIdx); raft::distance::KVPMinReduce pairRedOp; raft::distance::fusedL2NN( - minClusterDistanceView.data(), - datasetView.data(), - centroidsView.data(), - L2NormXView.data(), - centroidsNormView.data(), + minClusterDistanceView.data_handle(), + datasetView.data_handle(), + centroidsView.data_handle(), + L2NormXView.data_handle(), + centroidsNormView.data_handle(), ns, nc, n_features, @@ -599,7 +599,7 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // pairwiseDistanceView [ns x nc] - view representing the pairwise // distance for current batch auto pairwiseDistanceView = - raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); + raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); // calculate pairwise distance between current tile of cluster centroids // and input dataset @@ -607,8 +607,8 @@ void minClusterDistanceCompute(const raft::handle_t& handle, handle, datasetView, centroidsView, pairwiseDistanceView, workspace, metric); raft::linalg::coalescedReduction( - minClusterDistanceView.data(), - pairwiseDistanceView.data(), + minClusterDistanceView.data_handle(), + pairwiseDistanceView.data_handle(), pairwiseDistanceView.extent(1), pairwiseDistanceView.extent(0), std::numeric_limits::max(), @@ -631,11 +631,11 @@ void minClusterDistanceCompute(const raft::handle_t& handle, template void countSamplesInCluster(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_vector_view& L2NormX, - const raft::device_matrix_view& centroids, + const raft::device_matrix_view& X, + const raft::device_vector_view& L2NormX, + const raft::device_matrix_view& centroids, rmm::device_uvector& workspace, - const raft::device_vector_view& sampleCountInCluster) + const raft::device_vector_view& sampleCountInCluster) { cudaStream_t stream = handle.get_stream(); auto n_samples = X.extent(0); @@ -646,7 +646,7 @@ void countSamplesInCluster(const raft::handle_t& handle, // - key is the index of nearest cluster // - value is the distance to the nearest cluster auto minClusterAndDistance = - raft::make_device_vector>(n_samples, stream); + raft::make_device_vector, IndexT>(n_samples, stream); // temporary buffer to store distance matrix, destructor releases the resource rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); @@ -659,7 +659,7 @@ void countSamplesInCluster(const raft::handle_t& handle, detail::minClusterAndDistanceCompute(handle, params, X, - (raft::device_matrix_view)centroids, + (raft::device_matrix_view)centroids, minClusterAndDistance.view(), L2NormX, L2NormBuf_OR_DistBuf, @@ -672,11 +672,11 @@ void countSamplesInCluster(const raft::handle_t& handle, cub::TransformInputIterator, cub::KeyValuePair*> - itr(minClusterAndDistance.data(), conversion_op); + itr(minClusterAndDistance.data_handle(), conversion_op); // count # of samples in each cluster countLabels( - handle, itr, sampleCountInCluster.data(), (IndexT)n_samples, (IndexT)n_clusters, workspace); + handle, itr, sampleCountInCluster.data_handle(), (IndexT)n_samples, (IndexT)n_clusters, workspace); } } // namespace detail } // namespace cluster diff --git a/cpp/include/raft/cluster/kmeans.cuh b/cpp/include/raft/cluster/kmeans.cuh index 3285a98083..d46f53d9c1 100644 --- a/cpp/include/raft/cluster/kmeans.cuh +++ b/cpp/include/raft/cluster/kmeans.cuh @@ -50,9 +50,9 @@ namespace cluster { template void kmeans_fit(handle_t const& handle, const KMeansParams& params, - raft::device_matrix_view X, - std::optional> sample_weight, - raft::device_matrix_view centroids, + raft::device_matrix_view X, + std::optional> sample_weight, + raft::device_matrix_view centroids, raft::host_scalar_view inertia, raft::host_scalar_view n_iter) { @@ -97,10 +97,10 @@ void kmeans_fit(handle_t const& handle, template void kmeans_predict(handle_t const& handle, const KMeansParams& params, - raft::device_matrix_view X, - std::optional> sample_weight, - raft::device_matrix_view centroids, - raft::device_vector_view labels, + raft::device_matrix_view X, + std::optional> sample_weight, + raft::device_matrix_view centroids, + raft::device_vector_view labels, bool normalize_weight, raft::host_scalar_view inertia) { @@ -162,10 +162,10 @@ void kmeans_predict(handle_t const& handle, template void kmeans_fit_predict(handle_t const& handle, const KMeansParams& params, - raft::device_matrix_view X, - std::optional> sample_weight, - std::optional> centroids, - raft::device_vector_view labels, + raft::device_matrix_view X, + std::optional> sample_weight, + std::optional> centroids, + raft::device_vector_view labels, raft::host_scalar_view inertia, raft::host_scalar_view n_iter) { @@ -207,9 +207,9 @@ void kmeans_fit_predict(handle_t const& handle, template void kmeans_transform(const raft::handle_t& handle, const KMeansParams& params, - raft::device_matrix_view X, - raft::device_matrix_view centroids, - raft::device_matrix_view X_new) + raft::device_matrix_view X, + raft::device_matrix_view centroids, + raft::device_matrix_view X_new) { detail::kmeans_transform(handle, params, X, centroids, X_new); } @@ -254,9 +254,9 @@ using KeyValueIndexOp = detail::KeyValueIndexOp; */ template void sampleCentroids(const raft::handle_t& handle, - const raft::device_matrix_view& X, - const raft::device_vector_view& minClusterDistance, - const raft::device_vector_view& isSampleCentroid, + const raft::device_matrix_view& X, + const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& isSampleCentroid, SamplingOp& select_op, rmm::device_uvector& inRankCp, rmm::device_uvector& workspace) @@ -279,14 +279,14 @@ void sampleCentroids(const raft::handle_t& handle, * @param[in] reduction_op The reduction operation used for the cost * */ -template +template void computeClusterCost(const raft::handle_t& handle, - const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& minClusterDistance, rmm::device_uvector& workspace, const raft::device_scalar_view& clusterCost, ReductionOpT reduction_op) { - detail::computeClusterCost( + detail::computeClusterCost( handle, minClusterDistance, workspace, clusterCost, reduction_op); } @@ -314,10 +314,10 @@ void computeClusterCost(const raft::handle_t& handle, template void minClusterDistanceCompute(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroids, - const raft::device_vector_view& minClusterDistance, - const raft::device_vector_view& L2NormX, + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroids, + const raft::device_vector_view& minClusterDistance, + const raft::device_vector_view& L2NormX, rmm::device_uvector& L2NormBuf_OR_DistBuf, rmm::device_uvector& workspace) { @@ -353,10 +353,10 @@ template void minClusterAndDistanceCompute( const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view X, - const raft::device_matrix_view centroids, - const raft::device_vector_view>& minClusterAndDistance, - const raft::device_vector_view& L2NormX, + const raft::device_matrix_view X, + const raft::device_matrix_view centroids, + const raft::device_vector_view, IndexT>& minClusterAndDistance, + const raft::device_vector_view& L2NormX, rmm::device_uvector& L2NormBuf_OR_DistBuf, rmm::device_uvector& workspace) { @@ -383,8 +383,8 @@ void minClusterAndDistanceCompute( */ template void shuffleAndGather(const raft::handle_t& handle, - const raft::device_matrix_view& in, - const raft::device_matrix_view& out, + const raft::device_matrix_view& in, + const raft::device_matrix_view& out, uint32_t n_samples_to_gather, uint64_t seed, rmm::device_uvector* workspace = nullptr) @@ -414,11 +414,11 @@ void shuffleAndGather(const raft::handle_t& handle, template void countSamplesInCluster(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_vector_view& L2NormX, - const raft::device_matrix_view& centroids, + const raft::device_matrix_view& X, + const raft::device_vector_view& L2NormX, + const raft::device_matrix_view& centroids, rmm::device_uvector& workspace, - const raft::device_vector_view& sampleCountInCluster) + const raft::device_vector_view& sampleCountInCluster) { detail::countSamplesInCluster( handle, params, X, L2NormX, centroids, workspace, sampleCountInCluster); @@ -445,8 +445,8 @@ void countSamplesInCluster(const raft::handle_t& handle, template void kmeansPlusPlus(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroidsRawData, + const raft::device_matrix_view& X, + const raft::device_matrix_view& centroidsRawData, rmm::device_uvector& workspace) { detail::kmeansPlusPlus(handle, params, X, centroidsRawData, workspace); @@ -478,9 +478,9 @@ void kmeansPlusPlus(const raft::handle_t& handle, template void kmeans_fit_main(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_vector_view& weight, - const raft::device_matrix_view& centroidsRawData, + const raft::device_matrix_view& X, + const raft::device_vector_view& weight, + const raft::device_matrix_view& centroidsRawData, const raft::host_scalar_view& inertia, const raft::host_scalar_view& n_iter, rmm::device_uvector& workspace) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index 069705f166..9d749ee47b 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -256,7 +256,7 @@ using vector_extent = stdex::extents; template using matrix_extent = stdex::extents; -template +template using scalar_extent = stdex::extents; template diff --git a/cpp/test/cluster/kmeans.cu b/cpp/test/cluster/kmeans.cu index f54484c9ba..1d5b623e8b 100644 --- a/cpp/test/cluster/kmeans.cu +++ b/cpp/test/cluster/kmeans.cu @@ -64,8 +64,8 @@ class KmeansTest : public ::testing::TestWithParam> { params.rng_state.seed = 1; params.oversampling_factor = 0; - auto X = raft::make_device_matrix(n_samples, n_features, stream); - auto labels = raft::make_device_vector(n_samples, stream); + auto X = raft::make_device_matrix(n_samples, n_features, stream); + auto labels = raft::make_device_vector(n_samples, stream); raft::random::make_blobs(X.data(), labels.data(), @@ -88,23 +88,23 @@ class KmeansTest : public ::testing::TestWithParam> { std::optional> d_sw = std::nullopt; auto d_centroids_view = - raft::make_device_matrix_view(d_centroids.data(), params.n_clusters, n_features); + raft::make_device_matrix_view(d_centroids.data(), params.n_clusters, n_features); if (testparams.weighted) { d_sample_weight.resize(n_samples, stream); d_sw = std::make_optional( - raft::make_device_vector_view(d_sample_weight.data(), n_samples)); + raft::make_device_vector_view(d_sample_weight.data(), n_samples)); thrust::fill(thrust::cuda::par.on(stream), d_sample_weight.data(), d_sample_weight.data() + n_samples, 1); } - raft::copy(d_labels_ref.data(), labels.data(), n_samples, stream); + raft::copy(d_labels_ref.data(), labels.data_handle(), n_samples, stream); handle.sync_stream(stream); T inertia = 0; int n_iter = 0; - auto X_view = (raft::device_matrix_view)X.view(); + auto X_view = (raft::device_matrix_view)X.view(); raft::cluster::kmeans_fit_predict( handle, @@ -112,9 +112,9 @@ class KmeansTest : public ::testing::TestWithParam> { X_view, d_sw, d_centroids_view, - raft::make_device_vector_view(d_labels.data(), n_samples), - raft::make_host_scalar_view(&inertia), - raft::make_host_scalar_view(&n_iter)); + raft::make_device_vector_view(d_labels.data(), n_samples), + raft::make_host_scalar_view(&inertia), + raft::make_host_scalar_view(&n_iter)); handle.sync_stream(stream); From 1d17d506bc486b55670641b183baf2d500473a64 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 26 Jul 2022 21:28:17 -0400 Subject: [PATCH 2/8] iFixing style --- cpp/include/raft/cluster/detail/kmeans.cuh | 37 ++++++--- .../raft/cluster/detail/kmeans_common.cuh | 79 ++++++++++++------- 2 files changed, 74 insertions(+), 42 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index 0d39afeeb8..bb0c177f1b 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -118,20 +118,26 @@ void kmeansPlusPlus(const raft::handle_t& handle, if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - raft::linalg::rowNorm( - L2NormX.data_handle(), X.data_handle(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + raft::linalg::rowNorm(L2NormX.data_handle(), + X.data_handle(), + X.extent(1), + X.extent(0), + raft::linalg::L2Norm, + true, + stream); } std::mt19937 gen(params.rng_state.seed); std::uniform_int_distribution<> dis(0, n_samples - 1); // <<< Step-1 >>>: C <-- sample a point uniformly at random from X - auto initialCentroid = - raft::make_device_matrix_view(X.data_handle() + dis(gen) * n_features, 1, n_features); + auto initialCentroid = raft::make_device_matrix_view( + X.data_handle() + dis(gen) * n_features, 1, n_features); int n_clusters_picked = 1; // store the chosen centroid in the buffer - raft::copy(centroidsRawData.data_handle(), initialCentroid.data_handle(), initialCentroid.size(), stream); + raft::copy( + centroidsRawData.data_handle(), initialCentroid.data_handle(), initialCentroid.size(), stream); // C = initial set of centroids auto centroids = raft::make_device_matrix_view( @@ -161,9 +167,9 @@ void kmeansPlusPlus(const raft::handle_t& handle, // Note - n_trials is relative small here, we don't need raft::gather call std::discrete_distribution<> d(h_wt.begin(), h_wt.end()); for (int cIdx = 0; cIdx < n_trials; ++cIdx) { - auto rand_idx = d(gen); - auto randCentroid = - raft::make_device_matrix_view(X.data_handle() + n_features * rand_idx, 1, n_features); + auto rand_idx = d(gen); + auto randCentroid = raft::make_device_matrix_view( + X.data_handle() + n_features * rand_idx, 1, n_features); raft::copy(centroidCandidates.data_handle() + cIdx * n_features, randCentroid.data_handle(), randCentroid.size(), @@ -290,8 +296,13 @@ void kmeans_fit_main(const raft::handle_t& handle, auto L2NormX = raft::make_device_vector(n_samples, stream); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - raft::linalg::rowNorm( - L2NormX.data_handle(), X.data_handle(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + raft::linalg::rowNorm(L2NormX.data_handle(), + X.data_handle(), + X.extent(1), + X.extent(0), + raft::linalg::L2Norm, + true, + stream); } RAFT_LOG_DEBUG( @@ -306,7 +317,8 @@ void kmeans_fit_main(const raft::handle_t& handle, "cluster centers", n_iter[0]); - auto centroids = raft::make_device_matrix_view(centroidsRawData.data_handle(), n_clusters, n_features); + auto centroids = raft::make_device_matrix_view( + centroidsRawData.data_handle(), n_clusters, n_features); // computes minClusterAndDistance[0:n_samples) where // minClusterAndDistance[i] is a pair where @@ -415,7 +427,8 @@ void kmeans_fit_main(const raft::handle_t& handle, DataT sqrdNormError = 0; raft::copy(&sqrdNormError, sqrdNorm.data_handle(), sqrdNorm.size(), stream); - raft::copy(centroidsRawData.data_handle(), newCentroids.data_handle(), newCentroids.size(), stream); + raft::copy( + centroidsRawData.data_handle(), newCentroids.data_handle(), newCentroids.size(), stream); bool done = false; if (params.inertia_check) { diff --git a/cpp/include/raft/cluster/detail/kmeans_common.cuh b/cpp/include/raft/cluster/detail/kmeans_common.cuh index 13c2bc975a..46264b103d 100644 --- a/cpp/include/raft/cluster/detail/kmeans_common.cuh +++ b/cpp/include/raft/cluster/detail/kmeans_common.cuh @@ -165,8 +165,12 @@ void checkWeight(const raft::handle_t& handle, workspace.resize(temp_storage_bytes, stream); - RAFT_CUDA_TRY(cub::DeviceReduce::Sum( - workspace.data(), temp_storage_bytes, weight.data_handle(), wt_aggr.data_handle(), n_samples, stream)); + RAFT_CUDA_TRY(cub::DeviceReduce::Sum(workspace.data(), + temp_storage_bytes, + weight.data_handle(), + wt_aggr.data_handle(), + n_samples, + stream)); DataT wt_sum = 0; raft::copy(&wt_sum, wt_aggr.data_handle(), 1, stream); handle.sync_stream(stream); @@ -312,15 +316,16 @@ void pairwise_distance_kmeans(const raft::handle_t& handle, ASSERT(X.extent(1) == centroids.extent(1), "# features in dataset and centroids are different (must be same)"); - raft::distance::pairwise_distance(handle, - X.data_handle(), - centroids.data_handle(), - pairwiseDistance.data_handle(), - n_samples, - n_clusters, - n_features, - workspace, - metric); + raft::distance::pairwise_distance( + handle, + X.data_handle(), + centroids.data_handle(), + pairwiseDistance.data_handle(), + n_samples, + n_clusters, + n_features, + workspace, + metric); } // shuffle and randomly select 'n_samples_to_gather' from input 'in' and stores @@ -341,8 +346,13 @@ void shuffleAndGather(const raft::handle_t& handle, if (workspace) { // shuffle indices on device - raft::random::permute( - indices.data_handle(), nullptr, nullptr, (IndexT)in.extent(1), (IndexT)in.extent(0), true, stream); + raft::random::permute(indices.data_handle(), + nullptr, + nullptr, + (IndexT)in.extent(1), + (IndexT)in.extent(0), + true, + stream); } else { // shuffle indices on host and copy to device... std::vector ht_indices(n_samples); @@ -402,10 +412,11 @@ void minClusterAndDistanceCompute( // Note - pairwiseDistance and centroidsNorm share the same buffer // centroidsNorm [n_clusters] - tensor wrapper around centroids L2 Norm - auto centroidsNorm = raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); + auto centroidsNorm = + raft::make_device_vector_view(L2NormBuf_OR_DistBuf.data(), n_clusters); // pairwiseDistance[ns x nc] - tensor wrapper around the distance buffer - auto pairwiseDistance = - raft::make_device_matrix_view(L2NormBuf_OR_DistBuf.data(), dataBatchSize, centroidsBatchSize); + auto pairwiseDistance = raft::make_device_matrix_view( + L2NormBuf_OR_DistBuf.data(), dataBatchSize, centroidsBatchSize); cub::KeyValuePair initial_value(0, std::numeric_limits::max()); @@ -421,14 +432,15 @@ void minClusterAndDistanceCompute( // datasetView [ns x n_features] - view representing the current batch of // input dataset - auto datasetView = - raft::make_device_matrix_view(X.data_handle() + (dIdx * n_features), ns, n_features); + auto datasetView = raft::make_device_matrix_view( + X.data_handle() + (dIdx * n_features), ns, n_features); // minClusterAndDistanceView [ns x n_clusters] auto minClusterAndDistanceView = raft::make_device_vector_view(minClusterAndDistance.data_handle() + dIdx, ns); - auto L2NormXView = raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); + auto L2NormXView = + raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { @@ -437,12 +449,13 @@ void minClusterAndDistanceCompute( // centroidsView [nc x n_features] - view representing the current batch // of centroids - auto centroidsView = - raft::make_device_matrix_view(centroids.data_handle() + (cIdx * n_features), nc, n_features); + auto centroidsView = raft::make_device_matrix_view( + centroids.data_handle() + (cIdx * n_features), nc, n_features); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - auto centroidsNormView = raft::make_device_vector_view(centroidsNorm.data_handle() + cIdx, nc); + auto centroidsNormView = + raft::make_device_vector_view(centroidsNorm.data_handle() + cIdx, nc); workspace.resize((sizeof(int)) * ns, stream); FusedL2NNReduceOp redOp(cIdx); @@ -466,7 +479,8 @@ void minClusterAndDistanceCompute( } else { // pairwiseDistanceView [ns x nc] - view representing the pairwise // distance for current batch - auto pairwiseDistanceView = raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); + auto pairwiseDistanceView = + raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); // calculate pairwise distance between current tile of cluster centroids // and input dataset @@ -553,14 +567,15 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // datasetView [ns x n_features] - view representing the current batch of // input dataset - auto datasetView = - raft::make_device_matrix_view(X.data_handle() + dIdx * n_features, ns, n_features); + auto datasetView = raft::make_device_matrix_view( + X.data_handle() + dIdx * n_features, ns, n_features); // minClusterDistanceView [ns x n_clusters] auto minClusterDistanceView = raft::make_device_vector_view(minClusterDistance.data_handle() + dIdx, ns); - auto L2NormXView = raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); + auto L2NormXView = + raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { @@ -569,8 +584,8 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // centroidsView [nc x n_features] - view representing the current batch // of centroids - auto centroidsView = - raft::make_device_matrix_view(centroids.data() + cIdx * n_features, nc, n_features); + auto centroidsView = raft::make_device_matrix_view( + centroids.data() + cIdx * n_features, nc, n_features); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { @@ -675,8 +690,12 @@ void countSamplesInCluster(const raft::handle_t& handle, itr(minClusterAndDistance.data_handle(), conversion_op); // count # of samples in each cluster - countLabels( - handle, itr, sampleCountInCluster.data_handle(), (IndexT)n_samples, (IndexT)n_clusters, workspace); + countLabels(handle, + itr, + sampleCountInCluster.data_handle(), + (IndexT)n_samples, + (IndexT)n_clusters, + workspace); } } // namespace detail } // namespace cluster From a18ad3aeb473fa1778d1721e39457848588c0969 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 26 Jul 2022 23:01:29 -0400 Subject: [PATCH 3/8] Fixing remaining kmeans build issues --- cpp/include/raft/cluster/detail/kmeans.cuh | 208 ++++++++++-------- .../raft/cluster/detail/kmeans_common.cuh | 66 +++--- cpp/include/raft/spectral/cluster_solvers.cuh | 7 +- cpp/test/cluster/kmeans.cu | 8 +- 4 files changed, 161 insertions(+), 128 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index bb0c177f1b..b7e1d73a4b 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -104,17 +104,17 @@ void kmeansPlusPlus(const raft::handle_t& handle, // temporary buffers std::vector h_wt(n_samples); - auto centroidCandidates = raft::make_device_matrix(n_trials, n_features, stream); - auto costPerCandidate = raft::make_device_vector(n_trials, stream); - auto minClusterDistance = raft::make_device_vector(n_samples, stream); - auto distBuffer = raft::make_device_matrix(n_trials, n_samples, stream); + auto centroidCandidates = raft::make_device_matrix(handle, n_trials, n_features); + auto costPerCandidate = raft::make_device_vector(handle, n_trials); + auto minClusterDistance = raft::make_device_vector(handle, n_samples); + auto distBuffer = raft::make_device_matrix(handle, n_trials, n_samples); rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); rmm::device_scalar clusterCost(stream); rmm::device_scalar> minClusterIndexAndDistance(stream); // L2 norm of X: ||c||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(handle, n_samples); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { @@ -276,7 +276,7 @@ void kmeans_fit_main(const raft::handle_t& handle, // - key is the index of nearest cluster // - value is the distance to the nearest cluster auto minClusterAndDistance = - raft::make_device_vector, IndexT>(n_samples, stream); + raft::make_device_vector, IndexT>(handle, n_samples); // temporary buffer to store L2 norm of centroids or distance matrix, // destructor releases the resource @@ -284,16 +284,16 @@ void kmeans_fit_main(const raft::handle_t& handle, // temporary buffer to store intermediate centroids, destructor releases the // resource - auto newCentroids = raft::make_device_matrix(n_clusters, n_features, stream); + auto newCentroids = raft::make_device_matrix(handle, n_clusters, n_features); // temporary buffer to store weights per cluster, destructor releases the // resource - auto wtInCluster = raft::make_device_vector(n_clusters, stream); + auto wtInCluster = raft::make_device_vector(handle, n_clusters); rmm::device_scalar> clusterCostD(stream); // L2 norm of X: ||x||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(handle, n_samples); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { raft::linalg::rowNorm(L2NormX.data_handle(), @@ -412,17 +412,17 @@ void kmeans_fit_main(const raft::handle_t& handle, // compute the squared norm between the newCentroids and the original // centroids, destructor releases the resource - auto sqrdNorm = raft::make_device_scalar(DataT(0), stream); + auto sqrdNorm = raft::make_device_scalar(handle, DataT(0)); raft::linalg::mapThenSumReduce( - sqrdNorm.data(), + sqrdNorm.data_handle(), newCentroids.size(), [=] __device__(const DataT a, const DataT b) { DataT diff = a - b; return diff * diff; }, stream, - centroids.data(), - newCentroids.data()); + centroids.data_handle(), + newCentroids.data_handle()); DataT sqrdNormError = 0; raft::copy(&sqrdNormError, sqrdNorm.data_handle(), sqrdNorm.size(), stream); @@ -469,7 +469,8 @@ void kmeans_fit_main(const raft::handle_t& handle, } } - auto centroids = raft::make_device_matrix_view(centroidsRawData.data(), n_clusters, n_features); + auto centroids = raft::make_device_matrix_view( + centroidsRawData.data_handle(), n_clusters, n_features); detail::minClusterAndDistanceCompute(handle, params, @@ -482,10 +483,10 @@ void kmeans_fit_main(const raft::handle_t& handle, // TODO: add different templates for InType of binaryOp to avoid thrust transform thrust::transform(handle.get_thrust_policy(), - minClusterAndDistance.data(), - minClusterAndDistance.data() + minClusterAndDistance.size(), - weight.data(), - minClusterAndDistance.data(), + minClusterAndDistance.data_handle(), + minClusterAndDistance.data_handle() + minClusterAndDistance.size(), + weight.data_handle(), + minClusterAndDistance.data_handle(), [=] __device__(const cub::KeyValuePair kvp, DataT wt) { cub::KeyValuePair res; res.value = kvp.value * wt; @@ -506,7 +507,7 @@ void kmeans_fit_main(const raft::handle_t& handle, return res; }); - raft::copy(inertia.data(), &(clusterCostD.data()->value), 1, stream); + raft::copy(inertia.data_handle(), &(clusterCostD.data()->value), 1, stream); RAFT_LOG_DEBUG("KMeans.fit: completed after %d iterations with %f inertia[0] ", n_iter[0] > params.max_iter ? n_iter[0] - 1 : n_iter[0], @@ -539,8 +540,8 @@ void kmeans_fit_main(const raft::handle_t& handle, template void initScalableKMeansPlusPlus(const raft::handle_t& handle, const KMeansParams& params, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroidsRawData, + const raft::device_matrix_view X, + const raft::device_matrix_view centroidsRawData, rmm::device_uvector& workspace) { cudaStream_t stream = handle.get_stream(); @@ -556,7 +557,8 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, std::uniform_int_distribution<> dis(0, n_samples - 1); auto cIdx = dis(gen); - auto initialCentroid = raft::make_device_matrix_view(X.data() + cIdx * n_features, 1, n_features); + auto initialCentroid = raft::make_device_matrix_view( + X.data_handle() + cIdx * n_features, 1, n_features); // flag the sample that is chosen as initial centroid std::vector h_isSampleCentroid(n_samples); @@ -564,16 +566,17 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, h_isSampleCentroid[cIdx] = 1; // device buffer to flag the sample that is chosen as initial centroid - auto isSampleCentroid = raft::make_device_vector(n_samples, stream); + auto isSampleCentroid = raft::make_device_vector(handle, n_samples); - raft::copy(isSampleCentroid.data(), h_isSampleCentroid.data(), isSampleCentroid.size(), stream); + raft::copy( + isSampleCentroid.data_handle(), h_isSampleCentroid.data(), isSampleCentroid.size(), stream); rmm::device_uvector centroidsBuf(initialCentroid.size(), stream); // reset buffer to store the chosen centroid - raft::copy(centroidsBuf.data(), initialCentroid.data(), initialCentroid.size(), stream); + raft::copy(centroidsBuf.data(), initialCentroid.data_handle(), initialCentroid.size(), stream); - auto potentialCentroids = raft::make_device_matrix_view( + auto potentialCentroids = raft::make_device_matrix_view( centroidsBuf.data(), initialCentroid.extent(0), initialCentroid.extent(1)); // <<< End of Step-1 >>> @@ -582,15 +585,20 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); // L2 norm of X: ||x||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(handle, n_samples); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - raft::linalg::rowNorm( - L2NormX.data(), X.data(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + raft::linalg::rowNorm(L2NormX.data_handle(), + X.data_handle(), + X.extent(1), + X.extent(0), + raft::linalg::L2Norm, + true, + stream); } - auto minClusterDistanceVec = raft::make_device_vector(n_samples, stream); - auto uniformRands = raft::make_device_vector(n_samples, stream); + auto minClusterDistanceVec = raft::make_device_vector(handle, n_samples); + auto uniformRands = raft::make_device_vector(handle, n_samples); rmm::device_scalar clusterCost(stream); // <<< Step-2 >>>: psi <- phi_X (C) @@ -637,7 +645,7 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, detail::computeClusterCost(handle, minClusterDistanceVec.view(), workspace, - raft::make_device_scalar_view(clusterCost.data()), + raft::make_device_scalar_view(clusterCost.data()), [] __device__(const DataT& a, const DataT& b) { return a + b; }); psi = clusterCost.value(stream); @@ -645,10 +653,13 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, // <<<< Step-4 >>> : Sample each point x in X independently and identify new // potentialCentroids raft::random::uniform( - handle, rng, uniformRands.data(), uniformRands.extent(0), (DataT)0, (DataT)1); + handle, rng, uniformRands.data_handle(), uniformRands.extent(0), (DataT)0, (DataT)1); - detail::SamplingOp select_op( - psi, params.oversampling_factor, n_clusters, uniformRands.data(), isSampleCentroid.data()); + detail::SamplingOp select_op(psi, + params.oversampling_factor, + n_clusters, + uniformRands.data_handle(), + isSampleCentroid.data_handle()); rmm::device_uvector CpRaw(0, stream); detail::sampleCentroids(handle, @@ -658,17 +669,19 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, select_op, CpRaw, workspace); - auto Cp = raft::make_device_matrix_view(CpRaw.data(), CpRaw.size() / n_features, n_features); + auto Cp = raft::make_device_matrix_view( + CpRaw.data(), CpRaw.size() / n_features, n_features); /// <<<< End of Step-4 >>>> /// <<<< Step-5 >>> : C = C U C' // append the data in Cp to the buffer holding the potentialCentroids centroidsBuf.resize(centroidsBuf.size() + Cp.size(), stream); - raft::copy(centroidsBuf.data() + centroidsBuf.size() - Cp.size(), Cp.data(), Cp.size(), stream); + raft::copy( + centroidsBuf.data() + centroidsBuf.size() - Cp.size(), Cp.data_handle(), Cp.size(), stream); IndexT tot_centroids = potentialCentroids.extent(0) + Cp.extent(0); potentialCentroids = - raft::make_device_matrix_view(centroidsBuf.data(), tot_centroids, n_features); + raft::make_device_matrix_view(centroidsBuf.data(), tot_centroids, n_features); /// <<<< End of Step-5 >>> } /// <<<< Step-6 >>> @@ -679,7 +692,7 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, // <<< Step-7 >>>: For x in C, set w_x to be the number of pts closest to X // temporary buffer to store the sample count per cluster, destructor // releases the resource - auto weight = raft::make_device_vector(potentialCentroids.extent(0), stream); + auto weight = raft::make_device_vector(handle, potentialCentroids.extent(0)); detail::countSamplesInCluster( handle, params, X, L2NormX.view(), potentialCentroids, workspace, weight.view()); @@ -723,14 +736,16 @@ void initScalableKMeansPlusPlus(const raft::handle_t& handle, initRandom(handle, rand_params, X, centroidsRawData); // copy centroids generated during kmeans|| iteration to the buffer - raft::copy(centroidsRawData.data() + n_random_clusters * n_features, - potentialCentroids.data(), + raft::copy(centroidsRawData.data_handle() + n_random_clusters * n_features, + potentialCentroids.data_handle(), potentialCentroids.size(), stream); } else { // found the required n_clusters - raft::copy( - centroidsRawData.data(), potentialCentroids.data(), potentialCentroids.size(), stream); + raft::copy(centroidsRawData.data_handle(), + potentialCentroids.data_handle(), + potentialCentroids.size(), + stream); } } @@ -789,14 +804,15 @@ void kmeans_fit(handle_t const& handle, rmm::device_uvector workspace(0, stream); auto weight = raft::make_device_vector(handle, n_samples); if (sample_weight.has_value()) - raft::copy(weight.data(), sample_weight.value().data(), n_samples, stream); + raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); else - thrust::fill(handle.get_thrust_policy(), weight.data(), weight.data() + weight.size(), 1); + thrust::fill( + handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + weight.size(), 1); // check if weights sum up to n_samples checkWeight(handle, weight.view(), workspace); - auto centroidsRawData = raft::make_device_matrix(n_clusters, n_features, stream); + auto centroidsRawData = raft::make_device_matrix(handle, n_clusters, n_features); auto n_init = params.n_init; if (params.init == KMeansParams::InitMethod::Array && n_init != 1) { @@ -845,7 +861,8 @@ void kmeans_fit(handle_t const& handle, "passed to init arguement.", seed_iter + 1, n_init); - raft::copy(centroidsRawData.data(), centroids.data(), n_clusters * n_features, stream); + raft::copy( + centroidsRawData.data_handle(), centroids.data_handle(), n_clusters * n_features, stream); } else { THROW("unknown initialization method to select initial centers"); } @@ -861,7 +878,8 @@ void kmeans_fit(handle_t const& handle, if (iter_inertia < inertia[0]) { inertia[0] = iter_inertia; n_iter[0] = n_current_iter; - raft::copy(centroids.data(), centroidsRawData.data(), n_clusters * n_features, stream); + raft::copy( + centroids.data_handle(), centroidsRawData.data_handle(), n_clusters * n_features, stream); } RAFT_LOG_DEBUG("KMeans.fit after iteration-%d/%d: inertia - %f, n_iter[0] - %d", seed_iter + 1, @@ -883,12 +901,14 @@ void kmeans_fit(handle_t const& handle, DataT& inertia, IndexT& n_iter) { - auto XView = raft::make_device_matrix_view(X, n_samples, n_features); - auto centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); + auto XView = raft::make_device_matrix_view(X, n_samples, n_features); + auto centroidsView = + raft::make_device_matrix_view(centroids, params.n_clusters, n_features); std::optional> sample_weightView = std::nullopt; - if (sample_weight) sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); - auto inertiaView = raft::make_host_scalar_view(&inertia); - auto n_iterView = raft::make_host_scalar_view(&n_iter); + if (sample_weight) + sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); + auto inertiaView = raft::make_host_scalar_view(&inertia); + auto n_iterView = raft::make_host_scalar_view(&n_iter); detail::kmeans_fit( handle, params, XView, sample_weightView, centroidsView, inertiaView, n_iterView); @@ -925,25 +945,31 @@ void kmeans_predict(handle_t const& handle, // Allocate memory // Device-accessible allocation of expandable storage used as temorary buffers rmm::device_uvector workspace(0, stream); - auto weight = raft::make_device_vector(handle, n_samples); + auto weight = raft::make_device_vector(handle, n_samples); if (sample_weight.has_value()) - raft::copy(weight.data(), sample_weight.value().data(), n_samples, stream); + raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); else - thrust::fill(handle.get_thrust_policy(), weight.data(), weight.data() + weight.size(), 1); + thrust::fill( + handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + weight.size(), 1); // check if weights sum up to n_samples if (normalize_weight) checkWeight(handle, weight.view(), workspace); auto minClusterAndDistance = - raft::make_device_vector>(n_samples, stream); + raft::make_device_vector, IndexT>(handle, n_samples); rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); // L2 norm of X: ||x||^2 - auto L2NormX = raft::make_device_vector(n_samples, stream); + auto L2NormX = raft::make_device_vector(handle, n_samples); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { - raft::linalg::rowNorm( - L2NormX.data(), X.data(), X.extent(1), X.extent(0), raft::linalg::L2Norm, true, stream); + raft::linalg::rowNorm(L2NormX.data_handle(), + X.data_handle(), + X.extent(1), + X.extent(0), + raft::linalg::L2Norm, + true, + stream); } // computes minClusterAndDistance[0:n_samples) where minClusterAndDistance[i] @@ -964,10 +990,10 @@ void kmeans_predict(handle_t const& handle, rmm::device_scalar> clusterCostD(stream); // TODO: add different templates for InType of binaryOp to avoid thrust transform thrust::transform(handle.get_thrust_policy(), - minClusterAndDistance.data(), - minClusterAndDistance.data() + minClusterAndDistance.size(), - weight.data(), - minClusterAndDistance.data(), + minClusterAndDistance.data_handle(), + minClusterAndDistance.data_handle() + minClusterAndDistance.size(), + weight.data_handle(), + minClusterAndDistance.data_handle(), [=] __device__(const cub::KeyValuePair kvp, DataT wt) { cub::KeyValuePair res; res.value = kvp.value * wt; @@ -987,12 +1013,12 @@ void kmeans_predict(handle_t const& handle, return res; }); - raft::copy(inertia.data(), &(clusterCostD.data()->value), 1, stream); + raft::copy(inertia.data_handle(), &(clusterCostD.data()->value), 1, stream); thrust::transform(handle.get_thrust_policy(), - minClusterAndDistance.data(), - minClusterAndDistance.data() + minClusterAndDistance.size(), - labels.data(), + minClusterAndDistance.data_handle(), + minClusterAndDistance.data_handle() + minClusterAndDistance.size(), + labels.data_handle(), [=] __device__(cub::KeyValuePair pair) { return pair.key; }); } @@ -1008,11 +1034,13 @@ void kmeans_predict(handle_t const& handle, bool normalize_weight, DataT& inertia) { - auto XView = raft::make_device_matrix_view(X, n_samples, n_features); - auto centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); - std::optional> sample_weightView = std::nullopt; - if (sample_weight) sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); - auto labelsView = raft::make_device_vector_view(labels, n_samples); + auto XView = raft::make_device_matrix_view w(X, n_samples, n_features); + auto centroidsView = + raft::make_device_matrix_view(centroids, params.n_clusters, n_features); + std::optional, IndexT> sample_weightView = std::nullopt; + if (sample_weight) + sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); + auto labelsView = raft::make_device_vector_view(labels, n_samples); auto inertiaView = raft::make_host_scalar_view(&inertia); detail::kmeans_predict(handle, @@ -1038,7 +1066,7 @@ void kmeans_fit_predict(handle_t const& handle, if (!centroids.has_value()) { auto n_features = X.extent(1); auto centroids_matrix = - raft::make_device_matrix(params.n_clusters, n_features, handle.get_stream()); + raft::make_device_matrix(handle, params.n_clusters, n_features); detail::kmeans_fit( handle, params, X, sample_weight, centroids_matrix.view(), inertia, n_iter); detail::kmeans_predict( @@ -1063,15 +1091,17 @@ void kmeans_fit_predict(handle_t const& handle, DataT& inertia, IndexT& n_iter) { - auto XView = raft::make_device_matrix_view(X, n_samples, n_features); - std::optional> sample_weightView = std::nullopt; - if (sample_weight) sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); - std::optional> centroidsView = std::nullopt; + auto XView = raft::make_device_matrix_view(X, n_samples, n_features); + std::optional, IndexT> sample_weightView = std::nullopt; + if (sample_weight) + sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); + std::optional> centroidsView = std::nullopt; if (centroids) - centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); - auto labelsView = raft::make_device_vector_view(labels, n_samples); - auto inertiaView = raft::make_host_scalar_view(&inertia); - auto n_iterView = raft::make_host_scalar_view(&n_iter); + centroidsView = + raft::make_device_matrix_view(centroids, params.n_clusters, n_features); + auto labelsView = raft::make_device_vector_view(labels, n_samples); + auto inertiaView = raft::make_host_scalar_view(&inertia); + auto n_iterView = raft::make_host_scalar_view(&n_iter); detail::kmeans_fit_predict( handle, params, XView, sample_weightView, centroidsView, labelsView, inertiaView, n_iterView); @@ -1114,11 +1144,12 @@ void kmeans_transform(const raft::handle_t& handle, // datasetView [ns x n_features] - view representing the current batch of // input dataset - auto datasetView = raft::make_device_matrix_view(X.data() + n_features * dIdx, ns, n_features); + auto datasetView = + raft::make_device_matrix_view(X.data() + n_features * dIdx, ns, n_features); // pairwiseDistanceView [ns x n_clusters] - auto pairwiseDistanceView = - raft::make_device_matrix_view(X_new.data() + n_clusters * dIdx, ns, n_clusters); + auto pairwiseDistanceView = raft::make_device_matrix_view( + X_new.data() + n_clusters * dIdx, ns, n_clusters); // calculate pairwise distance between cluster centroids and current batch // of input dataset @@ -1136,9 +1167,10 @@ void kmeans_transform(const raft::handle_t& handle, IndexT n_features, DataT* X_new) { - auto XView = raft::make_device_matrix_view(X, n_samples, n_features); - auto centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); - auto X_newView = raft::make_device_matrix_view(X_new, n_samples, n_features); + auto XView = raft::make_device_matrix_view(X, n_samples, n_features); + auto centroidsView = + raft::make_device_matrix_view(centroids, params.n_clusters, n_features); + auto X_newView = raft::make_device_matrix_view(X_new, n_samples, n_features); detail::kmeans_transform(handle, params, XView, centroidsView, X_newView); } diff --git a/cpp/include/raft/cluster/detail/kmeans_common.cuh b/cpp/include/raft/cluster/detail/kmeans_common.cuh index 46264b103d..943e2647c8 100644 --- a/cpp/include/raft/cluster/detail/kmeans_common.cuh +++ b/cpp/include/raft/cluster/detail/kmeans_common.cuh @@ -156,7 +156,7 @@ void checkWeight(const raft::handle_t& handle, rmm::device_uvector& workspace) { cudaStream_t stream = handle.get_stream(); - auto wt_aggr = raft::make_device_scalar(0, stream); + auto wt_aggr = raft::make_device_scalar(handle, 0); auto n_samples = weight.extent(0); size_t temp_storage_bytes = 0; @@ -248,10 +248,10 @@ void sampleCentroids(const raft::handle_t& handle, auto n_local_samples = X.extent(0); auto n_features = X.extent(1); - auto nSelected = raft::make_device_scalar(0, stream); - cub::ArgIndexInputIterator ip_itr(minClusterDistance.data()); + auto nSelected = raft::make_device_scalar(handle, 0); + cub::ArgIndexInputIterator ip_itr(minClusterDistance.data_handle()); auto sampledMinClusterDistance = - raft::make_device_vector, IndexT>(n_local_samples, stream); + raft::make_device_vector, IndexT>(handle, n_local_samples); size_t temp_storage_bytes = 0; RAFT_CUDA_TRY(cub::DeviceSelect::If(nullptr, temp_storage_bytes, @@ -303,9 +303,9 @@ void sampleCentroids(const raft::handle_t& handle, // result will be stored in 'pairwiseDistance[n x k]' template void pairwise_distance_kmeans(const raft::handle_t& handle, - const raft::device_matrix_view& X, - const raft::device_matrix_view& centroids, - const raft::device_matrix_view& pairwiseDistance, + const raft::device_matrix_view X, + const raft::device_matrix_view centroids, + const raft::device_matrix_view pairwiseDistance, rmm::device_uvector& workspace, raft::distance::DistanceType metric) { @@ -316,16 +316,15 @@ void pairwise_distance_kmeans(const raft::handle_t& handle, ASSERT(X.extent(1) == centroids.extent(1), "# features in dataset and centroids are different (must be same)"); - raft::distance::pairwise_distance( - handle, - X.data_handle(), - centroids.data_handle(), - pairwiseDistance.data_handle(), - n_samples, - n_clusters, - n_features, - workspace, - metric); + raft::distance::pairwise_distance(handle, + X.data_handle(), + centroids.data_handle(), + pairwiseDistance.data_handle(), + n_samples, + n_clusters, + n_features, + workspace, + metric); } // shuffle and randomly select 'n_samples_to_gather' from input 'in' and stores @@ -342,7 +341,7 @@ void shuffleAndGather(const raft::handle_t& handle, auto n_samples = in.extent(0); auto n_features = in.extent(1); - auto indices = raft::make_device_vector(n_samples, stream); + auto indices = raft::make_device_vector(handle, n_samples); if (workspace) { // shuffle indices on device @@ -383,8 +382,8 @@ void minClusterAndDistanceCompute( const KMeansParams& params, const raft::device_matrix_view X, const raft::device_matrix_view centroids, - const raft::device_vector_view, IndexT>& minClusterAndDistance, - const raft::device_vector_view& L2NormX, + const raft::device_vector_view, IndexT> minClusterAndDistance, + const raft::device_vector_view L2NormX, rmm::device_uvector& L2NormBuf_OR_DistBuf, rmm::device_uvector& workspace) { @@ -428,7 +427,7 @@ void minClusterAndDistanceCompute( // tile over the input dataset for (std::size_t dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min(dataBatchSize, n_samples - dIdx); + auto ns = std::min((std::size_t)dataBatchSize, n_samples - dIdx); // datasetView [ns x n_features] - view representing the current batch of // input dataset @@ -437,7 +436,8 @@ void minClusterAndDistanceCompute( // minClusterAndDistanceView [ns x n_clusters] auto minClusterAndDistanceView = - raft::make_device_vector_view(minClusterAndDistance.data_handle() + dIdx, ns); + raft::make_device_vector_view, IndexT>( + minClusterAndDistance.data_handle() + dIdx, ns); auto L2NormXView = raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); @@ -445,7 +445,7 @@ void minClusterAndDistanceCompute( // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { // # of centroids for the current batch - auto nc = std::min(centroidsBatchSize, n_clusters - cIdx); + auto nc = std::min((std::size_t)centroidsBatchSize, n_clusters - cIdx); // centroidsView [nc x n_features] - view representing the current batch // of centroids @@ -480,7 +480,7 @@ void minClusterAndDistanceCompute( // pairwiseDistanceView [ns x nc] - view representing the pairwise // distance for current batch auto pairwiseDistanceView = - raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); + raft::make_device_matrix_view(pairwiseDistance.data_handle(), ns, nc); // calculate pairwise distance between current tile of cluster centroids // and input dataset @@ -563,7 +563,7 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // n_clusters] for (std::size_t dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min(dataBatchSize, n_samples - dIdx); + auto ns = std::min((std::size_t)dataBatchSize, n_samples - dIdx); // datasetView [ns x n_features] - view representing the current batch of // input dataset @@ -580,17 +580,17 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // tile over the centroids for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { // # of centroids for the current batch - auto nc = std::min(centroidsBatchSize, n_clusters - cIdx); + auto nc = std::min((std::size_t)centroidsBatchSize, n_clusters - cIdx); // centroidsView [nc x n_features] - view representing the current batch // of centroids auto centroidsView = raft::make_device_matrix_view( - centroids.data() + cIdx * n_features, nc, n_features); + centroids.data_handle() + cIdx * n_features, nc, n_features); if (metric == raft::distance::DistanceType::L2Expanded || metric == raft::distance::DistanceType::L2SqrtExpanded) { auto centroidsNormView = - raft::make_device_vector_view(centroidsNorm.data() + cIdx, nc); + raft::make_device_vector_view(centroidsNorm.data_handle() + cIdx, nc); workspace.resize((sizeof(IndexT)) * ns, stream); FusedL2NNReduceOp redOp(cIdx); @@ -614,7 +614,7 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // pairwiseDistanceView [ns x nc] - view representing the pairwise // distance for current batch auto pairwiseDistanceView = - raft::make_device_matrix_view(pairwiseDistance.data(), ns, nc); + raft::make_device_matrix_view(pairwiseDistance.data_handle(), ns, nc); // calculate pairwise distance between current tile of cluster centroids // and input dataset @@ -647,10 +647,10 @@ template void countSamplesInCluster(const raft::handle_t& handle, const KMeansParams& params, const raft::device_matrix_view& X, - const raft::device_vector_view& L2NormX, - const raft::device_matrix_view& centroids, + const raft::device_vector_view L2NormX, + const raft::device_matrix_view centroids, rmm::device_uvector& workspace, - const raft::device_vector_view& sampleCountInCluster) + const raft::device_vector_view sampleCountInCluster) { cudaStream_t stream = handle.get_stream(); auto n_samples = X.extent(0); @@ -661,7 +661,7 @@ void countSamplesInCluster(const raft::handle_t& handle, // - key is the index of nearest cluster // - value is the distance to the nearest cluster auto minClusterAndDistance = - raft::make_device_vector, IndexT>(n_samples, stream); + raft::make_device_vector, IndexT>(handle, n_samples); // temporary buffer to store distance matrix, destructor releases the resource rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); diff --git a/cpp/include/raft/spectral/cluster_solvers.cuh b/cpp/include/raft/spectral/cluster_solvers.cuh index 8d56f3b172..6f9ebcd6af 100644 --- a/cpp/include/raft/spectral/cluster_solvers.cuh +++ b/cpp/include/raft/spectral/cluster_solvers.cuh @@ -66,9 +66,10 @@ struct kmeans_solver_t { auto X = raft::make_device_matrix_view(obs, n_obs_vecs, dim); auto labels = raft::make_device_vector_view(codes, n_obs_vecs); auto centroids = - raft::make_device_matrix(config_.n_clusters, dim, handle.get_stream()); - auto weight = raft::make_device_vector(n_obs_vecs, handle.get_stream()); - thrust::fill(handle.get_thrust_policy(), weight.data(), weight.data() + n_obs_vecs, 1); + raft::make_device_matrix(handle, config_.n_clusters, dim); + auto weight = raft::make_device_vector(handle, n_obs_vecs); + thrust::fill( + handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + n_obs_vecs, 1); auto sw = std::make_optional((raft::device_vector_view)weight.view()); raft::cluster::kmeans_fit_predict( diff --git a/cpp/test/cluster/kmeans.cu b/cpp/test/cluster/kmeans.cu index 1d5b623e8b..247386b550 100644 --- a/cpp/test/cluster/kmeans.cu +++ b/cpp/test/cluster/kmeans.cu @@ -64,11 +64,11 @@ class KmeansTest : public ::testing::TestWithParam> { params.rng_state.seed = 1; params.oversampling_factor = 0; - auto X = raft::make_device_matrix(n_samples, n_features, stream); - auto labels = raft::make_device_vector(n_samples, stream); + auto X = raft::make_device_matrix(handle, n_samples, n_features); + auto labels = raft::make_device_vector(handle, n_samples); - raft::random::make_blobs(X.data(), - labels.data(), + raft::random::make_blobs(X.data_handle(), + labels.data_handle(), n_samples, n_features, params.n_clusters, From 5dd730d140335ecbbe68d36ddf6dab9e2cdfd47b Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Tue, 26 Jul 2022 23:30:18 -0400 Subject: [PATCH 4/8] Fixing typo --- cpp/include/raft/cluster/detail/kmeans.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index b7e1d73a4b..53a70a2f16 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -1034,7 +1034,7 @@ void kmeans_predict(handle_t const& handle, bool normalize_weight, DataT& inertia) { - auto XView = raft::make_device_matrix_view w(X, n_samples, n_features); + auto XView = raft::make_device_matrix_view(X, n_samples, n_features); auto centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); std::optional, IndexT> sample_weightView = std::nullopt; From eb0885ff03b43f1cd3545295e547be0c38179ef6 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 27 Jul 2022 00:33:40 -0400 Subject: [PATCH 5/8] Final fixes --- cpp/include/raft/cluster/detail/kmeans.cuh | 16 +++++++++------- .../raft/cluster/detail/kmeans_common.cuh | 16 ++++++++-------- 2 files changed, 17 insertions(+), 15 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index 53a70a2f16..f34005fc29 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -1037,9 +1037,10 @@ void kmeans_predict(handle_t const& handle, auto XView = raft::make_device_matrix_view(X, n_samples, n_features); auto centroidsView = raft::make_device_matrix_view(centroids, params.n_clusters, n_features); - std::optional, IndexT> sample_weightView = std::nullopt; + std::optional> sample_weightView{std::nullopt}; if (sample_weight) - sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); + sample_weightView.emplace( + raft::make_device_vector_view(sample_weight, n_samples)); auto labelsView = raft::make_device_vector_view(labels, n_samples); auto inertiaView = raft::make_host_scalar_view(&inertia); @@ -1092,13 +1093,14 @@ void kmeans_fit_predict(handle_t const& handle, IndexT& n_iter) { auto XView = raft::make_device_matrix_view(X, n_samples, n_features); - std::optional, IndexT> sample_weightView = std::nullopt; + std::optional> sample_weightView{std::nullopt}; if (sample_weight) - sample_weightView = raft::make_device_vector_view(sample_weight, n_samples); - std::optional> centroidsView = std::nullopt; + sample_weightView.emplace( + raft::make_device_vector_view(sample_weight, n_samples)); + std::optional> centroidsView{std::nullopt}; if (centroids) - centroidsView = - raft::make_device_matrix_view(centroids, params.n_clusters, n_features); + centroidsView.emplace( + raft::make_device_matrix_view(centroids, params.n_clusters, n_features)); auto labelsView = raft::make_device_vector_view(labels, n_samples); auto inertiaView = raft::make_host_scalar_view(&inertia); auto n_iterView = raft::make_host_scalar_view(&n_iter); diff --git a/cpp/include/raft/cluster/detail/kmeans_common.cuh b/cpp/include/raft/cluster/detail/kmeans_common.cuh index 943e2647c8..358c8ce16e 100644 --- a/cpp/include/raft/cluster/detail/kmeans_common.cuh +++ b/cpp/include/raft/cluster/detail/kmeans_common.cuh @@ -425,9 +425,9 @@ void minClusterAndDistanceCompute( initial_value); // tile over the input dataset - for (std::size_t dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { + for (IndexT dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min((std::size_t)dataBatchSize, n_samples - dIdx); + auto ns = std::min((IndexT)dataBatchSize, n_samples - dIdx); // datasetView [ns x n_features] - view representing the current batch of // input dataset @@ -443,9 +443,9 @@ void minClusterAndDistanceCompute( raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids - for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { + for (IndexT cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { // # of centroids for the current batch - auto nc = std::min((std::size_t)centroidsBatchSize, n_clusters - cIdx); + auto nc = std::min((IndexT)centroidsBatchSize, n_clusters - cIdx); // centroidsView [nc x n_features] - view representing the current batch // of centroids @@ -561,9 +561,9 @@ void minClusterDistanceCompute(const raft::handle_t& handle, // tile over the input data and calculate distance matrix [n_samples x // n_clusters] - for (std::size_t dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { + for (IndexT dIdx = 0; dIdx < n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min((std::size_t)dataBatchSize, n_samples - dIdx); + auto ns = std::min((IndexT)dataBatchSize, n_samples - dIdx); // datasetView [ns x n_features] - view representing the current batch of // input dataset @@ -578,9 +578,9 @@ void minClusterDistanceCompute(const raft::handle_t& handle, raft::make_device_vector_view(L2NormX.data_handle() + dIdx, ns); // tile over the centroids - for (std::size_t cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { + for (IndexT cIdx = 0; cIdx < n_clusters; cIdx += centroidsBatchSize) { // # of centroids for the current batch - auto nc = std::min((std::size_t)centroidsBatchSize, n_clusters - cIdx); + auto nc = std::min((IndexT)centroidsBatchSize, n_clusters - cIdx); // centroidsView [nc x n_features] - view representing the current batch // of centroids From d84351ed169a7c97e59ae744c00fe20a0a4631c8 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 27 Jul 2022 14:02:57 -0400 Subject: [PATCH 6/8] iAdding distance/nn specializations for cluster and spectral --- cpp/include/raft/cluster/specializations.cuh | 24 +++++++++++++++++++ cpp/include/raft/spectral/specializations.cuh | 24 +++++++++++++++++++ cpp/test/cluster/kmeans.cu | 4 ++++ cpp/test/cluster_solvers.cu | 5 ++++ 4 files changed, 57 insertions(+) create mode 100644 cpp/include/raft/cluster/specializations.cuh create mode 100644 cpp/include/raft/spectral/specializations.cuh diff --git a/cpp/include/raft/cluster/specializations.cuh b/cpp/include/raft/cluster/specializations.cuh new file mode 100644 index 0000000000..3bb5a26ace --- /dev/null +++ b/cpp/include/raft/cluster/specializations.cuh @@ -0,0 +1,24 @@ +/* + * Copyright (c) 2022, 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#ifndef __CLUSTER_SPECIALIZATIONS_H +#define __CLUSTER_SPECIALIZATIONS_H + +#pragma once + +#include +#include + +#endif \ No newline at end of file diff --git a/cpp/include/raft/spectral/specializations.cuh b/cpp/include/raft/spectral/specializations.cuh new file mode 100644 index 0000000000..2303b426fd --- /dev/null +++ b/cpp/include/raft/spectral/specializations.cuh @@ -0,0 +1,24 @@ +/* + * Copyright (c) 2022, 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#ifndef __SPECTRAL_SPECIALIZATIONS_H +#define __SPECTRAL_SPECIALIZATIONS_H + +#pragma once + +#include +#include + +#endif \ No newline at end of file diff --git a/cpp/test/cluster/kmeans.cu b/cpp/test/cluster/kmeans.cu index 247386b550..4bcad14f26 100644 --- a/cpp/test/cluster/kmeans.cu +++ b/cpp/test/cluster/kmeans.cu @@ -29,6 +29,10 @@ #include #include +#if defined RAFT_DISTANCE_COMPILED && RAFT_NN_COMPILED +#include +#endif + namespace raft { template diff --git a/cpp/test/cluster_solvers.cu b/cpp/test/cluster_solvers.cu index d475fd2a69..e334ad3fa4 100644 --- a/cpp/test/cluster_solvers.cu +++ b/cpp/test/cluster_solvers.cu @@ -19,6 +19,11 @@ #include #include +#if defined RAFT_DISTANCE_COMPILED && RAFT_NN_COMPILED +#include +#endif + + #include #include From 748a96c54fc56b49b09f5d7411c357471c94281c Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 27 Jul 2022 14:08:01 -0400 Subject: [PATCH 7/8] Fixing style --- cpp/test/cluster_solvers.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/test/cluster_solvers.cu b/cpp/test/cluster_solvers.cu index e334ad3fa4..1e07a87280 100644 --- a/cpp/test/cluster_solvers.cu +++ b/cpp/test/cluster_solvers.cu @@ -23,7 +23,6 @@ #include #endif - #include #include From 06cb8533cd32ccf680a6006b03449d8e7cb44070 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Wed, 27 Jul 2022 14:16:50 -0400 Subject: [PATCH 8/8] Fixing macros --- cpp/test/cluster/kmeans.cu | 2 +- cpp/test/cluster_solvers.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/test/cluster/kmeans.cu b/cpp/test/cluster/kmeans.cu index 4bcad14f26..24fe2c03cd 100644 --- a/cpp/test/cluster/kmeans.cu +++ b/cpp/test/cluster/kmeans.cu @@ -29,7 +29,7 @@ #include #include -#if defined RAFT_DISTANCE_COMPILED && RAFT_NN_COMPILED +#if defined RAFT_DISTANCE_COMPILED && defined RAFT_NN_COMPILED #include #endif diff --git a/cpp/test/cluster_solvers.cu b/cpp/test/cluster_solvers.cu index 1e07a87280..0c74b81e99 100644 --- a/cpp/test/cluster_solvers.cu +++ b/cpp/test/cluster_solvers.cu @@ -19,7 +19,7 @@ #include #include -#if defined RAFT_DISTANCE_COMPILED && RAFT_NN_COMPILED +#if defined RAFT_DISTANCE_COMPILED && defined RAFT_NN_COMPILED #include #endif