From 6a818ca209db9162ee58545f479767045db04564 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 10 Feb 2023 21:10:00 -0500 Subject: [PATCH 1/3] Some fixes to kmeasn --- cpp/include/raft/cluster/detail/kmeans.cuh | 9 +++++++-- cpp/include/raft/random/detail/make_blobs.cuh | 4 ++-- 2 files changed, 9 insertions(+), 4 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index 9632fedb9d..43b703c5c9 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -869,10 +869,15 @@ void kmeans_fit(raft::device_resources const& handle, auto weight = raft::make_device_vector(handle, n_samples); if (sample_weight.has_value()) raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); - else + else { + printf("About to call fill\n"); thrust::fill( handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + weight.size(), 1); + handle.sync_stream(); + printf("Done.\n"); + } + // check if weights sum up to n_samples checkWeight(handle, weight.view(), workspace); @@ -1209,7 +1214,7 @@ void kmeans_transform(raft::device_resources const& handle, // n_clusters] for (IndexT dIdx = 0; dIdx < (IndexT)n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min(dataBatchSize, n_samples - dIdx); + auto ns = std::min((IndexT)dataBatchSize, (IndexT)n_samples - dIdx); // datasetView [ns x n_features] - view representing the current batch of // input dataset diff --git a/cpp/include/raft/random/detail/make_blobs.cuh b/cpp/include/raft/random/detail/make_blobs.cuh index b37dabb366..a27e1e9835 100644 --- a/cpp/include/raft/random/detail/make_blobs.cuh +++ b/cpp/include/raft/random/detail/make_blobs.cuh @@ -158,9 +158,9 @@ void generate_data(DataT* out, raft::random::RngState& rng_state) { constexpr IdxT block_size = 128; - IdxT items = n_rows * n_cols; + int64_t items = (int64_t)n_rows * n_cols; // Choose a grid size so that each thread can write two output values. - IdxT nBlocks = ceildiv(items, 2 * block_size); + int64_t nBlocks = ceildiv(items, 2 * block_size); // parentheses needed here for kernel, otherwise macro interprets the arguments // of triple chevron notation as macro arguments RAFT_CALL_RNG_FUNC(rng_state, From 38dc25f0b69e1a7ba5c3f878f595453eca1965c6 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 10 Feb 2023 21:15:49 -0500 Subject: [PATCH 2/3] A couple additional changes --- cpp/include/raft/cluster/detail/kmeans.cuh | 4 +--- cpp/include/raft/random/detail/make_blobs.cuh | 2 +- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index 43b703c5c9..0a4960e239 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -870,12 +870,10 @@ void kmeans_fit(raft::device_resources const& handle, if (sample_weight.has_value()) raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); else { - printf("About to call fill\n"); thrust::fill( handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + weight.size(), 1); handle.sync_stream(); - printf("Done.\n"); } // check if weights sum up to n_samples @@ -1214,7 +1212,7 @@ void kmeans_transform(raft::device_resources const& handle, // n_clusters] for (IndexT dIdx = 0; dIdx < (IndexT)n_samples; dIdx += dataBatchSize) { // # of samples for the current batch - auto ns = std::min((IndexT)dataBatchSize, (IndexT)n_samples - dIdx); + auto ns = std::min(static_cast(dataBatchSize), static_cast(n_samples - dIdx)); // datasetView [ns x n_features] - view representing the current batch of // input dataset diff --git a/cpp/include/raft/random/detail/make_blobs.cuh b/cpp/include/raft/random/detail/make_blobs.cuh index a27e1e9835..4e3da63ef6 100644 --- a/cpp/include/raft/random/detail/make_blobs.cuh +++ b/cpp/include/raft/random/detail/make_blobs.cuh @@ -158,7 +158,7 @@ void generate_data(DataT* out, raft::random::RngState& rng_state) { constexpr IdxT block_size = 128; - int64_t items = (int64_t)n_rows * n_cols; + int64_t items = static_cast(n_rows) * n_cols; // Choose a grid size so that each thread can write two output values. int64_t nBlocks = ceildiv(items, 2 * block_size); // parentheses needed here for kernel, otherwise macro interprets the arguments From 0ad75fecaa5a4fd42677701a06de6e40445f360d Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Fri, 10 Feb 2023 21:16:30 -0500 Subject: [PATCH 3/3] More changes --- cpp/include/raft/cluster/detail/kmeans.cuh | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans.cuh b/cpp/include/raft/cluster/detail/kmeans.cuh index 0a4960e239..40e6728dbe 100644 --- a/cpp/include/raft/cluster/detail/kmeans.cuh +++ b/cpp/include/raft/cluster/detail/kmeans.cuh @@ -869,13 +869,10 @@ void kmeans_fit(raft::device_resources const& handle, auto weight = raft::make_device_vector(handle, n_samples); if (sample_weight.has_value()) raft::copy(weight.data_handle(), sample_weight.value().data_handle(), n_samples, stream); - else { + else thrust::fill( handle.get_thrust_policy(), weight.data_handle(), weight.data_handle() + weight.size(), 1); - handle.sync_stream(); - } - // check if weights sum up to n_samples checkWeight(handle, weight.view(), workspace);