From 1c3170124c4c7e6f0ec221587097a48db2f1ce42 Mon Sep 17 00:00:00 2001 From: Rory Mitchell Date: Fri, 26 Jun 2026 18:53:29 +0200 Subject: [PATCH 1/5] Implement weighted bootstrap sampling for RF --- cpp/include/cuml/ensemble/randomforest.hpp | 14 +- .../batched-levelalgo/builder.cuh | 2 +- .../decisiontree/batched-levelalgo/dataset.h | 2 +- cpp/src/decisiontree/decisiontree.cuh | 2 +- cpp/src/randomforest/randomforest.cu | 62 +--- cpp/src/randomforest/randomforest.cuh | 184 +++++++---- cpp/tests/sg/rf_test.cu | 310 ++++++++++-------- 7 files changed, 317 insertions(+), 259 deletions(-) diff --git a/cpp/include/cuml/ensemble/randomforest.hpp b/cpp/include/cuml/ensemble/randomforest.hpp index 3fb2f0fdb4..179b3a8ece 100644 --- a/cpp/include/cuml/ensemble/randomforest.hpp +++ b/cpp/include/cuml/ensemble/randomforest.hpp @@ -151,8 +151,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, - bool* bootstrap_masks = nullptr, - const float* sample_weight = nullptr); + const double* sample_weight = nullptr); void fit(const raft::handle_t& user_handle, RandomForestClassifierD* forest, double* input, @@ -162,7 +161,6 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, - bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); template @@ -174,10 +172,9 @@ void fit_treelite(const raft::handle_t& user_handle, L* labels, int n_unique_labels, RF_params rf_params, - bool* bootstrap_masks, T* feature_importances, rapids_logger::level_enum verbosity, - const T* sample_weight = nullptr); + const double* sample_weight = nullptr); void predict(const raft::handle_t& user_handle, const RandomForestClassifierF* forest, @@ -235,8 +232,7 @@ void fit(const raft::handle_t& user_handle, float* labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, - bool* bootstrap_masks = nullptr, - const float* sample_weight = nullptr); + const double* sample_weight = nullptr); void fit(const raft::handle_t& user_handle, RandomForestRegressorD* forest, double* input, @@ -245,7 +241,6 @@ void fit(const raft::handle_t& user_handle, double* labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, - bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); template @@ -256,10 +251,9 @@ void fit_treelite(const raft::handle_t& user_handle, int n_cols, L* labels, RF_params rf_params, - bool* bootstrap_masks, T* feature_importances, rapids_logger::level_enum verbosity, - const T* sample_weight = nullptr); + const double* sample_weight = nullptr); void predict(const raft::handle_t& user_handle, const RandomForestRegressorF* forest, diff --git a/cpp/src/decisiontree/batched-levelalgo/builder.cuh b/cpp/src/decisiontree/batched-levelalgo/builder.cuh index 483f20fa3d..1f60b9e87d 100644 --- a/cpp/src/decisiontree/batched-levelalgo/builder.cuh +++ b/cpp/src/decisiontree/batched-levelalgo/builder.cuh @@ -199,7 +199,7 @@ struct Builder { const DecisionTreeParams& p, const DataT* data, const LabelT* labels, - const DataT* sample_weight, + const double* sample_weight, IdxT n_rows, IdxT n_cols, rmm::device_uvector* row_ids, diff --git a/cpp/src/decisiontree/batched-levelalgo/dataset.h b/cpp/src/decisiontree/batched-levelalgo/dataset.h index df1398abe1..8f12152ce6 100644 --- a/cpp/src/decisiontree/batched-levelalgo/dataset.h +++ b/cpp/src/decisiontree/batched-levelalgo/dataset.h @@ -15,7 +15,7 @@ struct Dataset { /** input labels */ const LabelT* labels; /** optional input sample weights */ - const DataT* sample_weight; + const double* sample_weight; /** total rows in dataset */ IdxT M; /** total cols in dataset */ diff --git a/cpp/src/decisiontree/decisiontree.cuh b/cpp/src/decisiontree/decisiontree.cuh index 7b368eec2e..20d3656658 100644 --- a/cpp/src/decisiontree/decisiontree.cuh +++ b/cpp/src/decisiontree/decisiontree.cuh @@ -244,7 +244,7 @@ class DecisionTree { uint64_t seed, const Quantiles& quantiles, int treeid, - const DataT* sample_weight = nullptr) + const double* sample_weight = nullptr) { if (params.split_criterion == CRITERION::CRITERION_END) { // Set default to GINI (classification) or MSE (regression) diff --git a/cpp/src/randomforest/randomforest.cu b/cpp/src/randomforest/randomforest.cu index 79d4e947c3..6284633855 100644 --- a/cpp/src/randomforest/randomforest.cu +++ b/cpp/src/randomforest/randomforest.cu @@ -357,8 +357,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity, - bool* bootstrap_masks, - const float* sample_weight) + const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); ML::default_logger().set_level(verbosity); @@ -368,15 +367,8 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_classifier = std::make_shared>(rf_params, RF_type::CLASSIFICATION); - rf_classifier->fit(user_handle, - input, - n_rows, - n_cols, - labels, - n_unique_labels, - forest, - bootstrap_masks, - sample_weight); + rf_classifier->fit( + user_handle, input, n_rows, n_cols, labels, n_unique_labels, forest, sample_weight); } void fit(const raft::handle_t& user_handle, @@ -388,7 +380,6 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity, - bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -399,15 +390,8 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_classifier = std::make_shared>(rf_params, RF_type::CLASSIFICATION); - rf_classifier->fit(user_handle, - input, - n_rows, - n_cols, - labels, - n_unique_labels, - forest, - bootstrap_masks, - sample_weight); + rf_classifier->fit( + user_handle, input, n_rows, n_cols, labels, n_unique_labels, forest, sample_weight); } template @@ -419,10 +403,9 @@ void fit_treelite(const raft::handle_t& user_handle, label_t* labels, int n_unique_labels, RF_params rf_params, - bool* bootstrap_masks, value_t* feature_importances, rapids_logger::level_enum verbosity, - const value_t* sample_weight) + const double* sample_weight) { RandomForestMetaData metadata; fit(user_handle, @@ -434,7 +417,6 @@ void fit_treelite(const raft::handle_t& user_handle, n_unique_labels, rf_params, verbosity, - bootstrap_masks, sample_weight); // Compute feature importances if requested @@ -604,8 +586,7 @@ void fit(const raft::handle_t& user_handle, float* labels, RF_params rf_params, rapids_logger::level_enum verbosity, - bool* bootstrap_masks, - const float* sample_weight) + const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); ML::default_logger().set_level(verbosity); @@ -615,8 +596,7 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_regressor = std::make_shared>(rf_params, RF_type::REGRESSION); - rf_regressor->fit( - user_handle, input, n_rows, n_cols, labels, 1, forest, bootstrap_masks, sample_weight); + rf_regressor->fit(user_handle, input, n_rows, n_cols, labels, 1, forest, sample_weight); } void fit(const raft::handle_t& user_handle, @@ -627,7 +607,6 @@ void fit(const raft::handle_t& user_handle, double* labels, RF_params rf_params, rapids_logger::level_enum verbosity, - bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -638,8 +617,7 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_regressor = std::make_shared>(rf_params, RF_type::REGRESSION); - rf_regressor->fit( - user_handle, input, n_rows, n_cols, labels, 1, forest, bootstrap_masks, sample_weight); + rf_regressor->fit(user_handle, input, n_rows, n_cols, labels, 1, forest, sample_weight); } template @@ -650,22 +628,12 @@ void fit_treelite(const raft::handle_t& user_handle, int n_cols, label_t* labels, RF_params rf_params, - bool* bootstrap_masks, value_t* feature_importances, rapids_logger::level_enum verbosity, - const value_t* sample_weight) + const double* sample_weight) { RandomForestMetaData metadata; - fit(user_handle, - &metadata, - input, - n_rows, - n_cols, - labels, - rf_params, - verbosity, - bootstrap_masks, - sample_weight); + fit(user_handle, &metadata, input, n_rows, n_cols, labels, rf_params, verbosity, sample_weight); // Compute feature importances if requested if (feature_importances != nullptr) { @@ -870,10 +838,9 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_ha int* labels, int n_unique_labels, RF_params rf_params, - bool* bootstrap_masks, float* feature_importances, rapids_logger::level_enum verbosity, - const float* sample_weight); + const double* sample_weight); template CUML_EXPORT void fit_treelite(const raft::handle_t& user_handle, TreeliteModelHandle* model, double* input, @@ -882,7 +849,6 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_h int* labels, int n_unique_labels, RF_params rf_params, - bool* bootstrap_masks, double* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); @@ -893,10 +859,9 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_ int n_cols, float* labels, RF_params rf_params, - bool* bootstrap_masks, float* feature_importances, rapids_logger::level_enum verbosity, - const float* sample_weight); + const double* sample_weight); template CUML_EXPORT void fit_treelite(const raft::handle_t& user_handle, TreeliteModelHandle* model, double* input, @@ -904,7 +869,6 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& use int n_cols, double* labels, RF_params rf_params, - bool* bootstrap_masks, double* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh index 9e138cb1e6..fb3df55e94 100644 --- a/cpp/src/randomforest/randomforest.cuh +++ b/cpp/src/randomforest/randomforest.cuh @@ -17,10 +17,13 @@ #include +#include #include #include #include #include +#include +#include #include #include @@ -34,6 +37,7 @@ #define omp_get_max_threads() 1 #endif +#include #include namespace ML { @@ -43,36 +47,123 @@ template struct InvalidSampleWeight { __device__ bool operator()(T weight) const { return weight < T(0) || !isfinite(weight); } }; -} // namespace detail -template -class RandomForest { - protected: - RF_params rf_params; // structure containing RF hyperparameters - int rf_type; // 0 for classification 1 for regression +class RowSampler { + public: + RowSampler(const raft::handle_t& handle, + const RF_params& rf_params, + int n_rows, + int n_sampled_rows, + int n_streams, + const double* sample_weight) + : bootstrap_(rf_params.bootstrap), + seed_(rf_params.seed), + n_rows_(n_rows), + n_sampled_rows_(n_sampled_rows), + sample_weight_(sample_weight), + sample_weight_sum_(validate_sample_weight(handle, sample_weight_, n_rows_)), + sample_weight_cdf_(0, handle.get_stream()) + { + if (use_weighted_bootstrap()) { + sample_weight_cdf_.resize(n_rows_, handle.get_stream()); + thrust::inclusive_scan(rmm::exec_policy(handle.get_stream()), + sample_weight_, + sample_weight_ + n_rows_, + sample_weight_cdf_.begin()); + handle.sync_stream(); + } + + // Use a deque instead of vector because device_uvector has a deleted copy constructor. + for (int i = 0; i < n_streams; i++) { + auto stream = handle.get_stream_from_stream_pool(i); + selected_rows_.emplace_back(n_sampled_rows_, stream); + if (use_weighted_bootstrap()) { + weighted_draw_scratch_.emplace_back(n_sampled_rows_, stream); + } + } + } + + RowSampler(const RowSampler&) = delete; + RowSampler& operator=(const RowSampler&) = delete; - void get_row_sample(int tree_id, - int n_rows, - rmm::device_uvector* selected_rows, - const cudaStream_t stream) + rmm::device_uvector& sample(int tree_id, int stream_id, cudaStream_t stream) { raft::common::nvtx::range fun_scope("bootstrapping row IDs @randomforest.cuh"); - // Hash these together so they are uncorrelated + auto& selected_rows = selected_rows_[stream_id]; + + // Hash these together so per-tree row samples are uncorrelated. auto rs = DT::fnv1a32_basis; - rs = DT::fnv1a32(rs, rf_params.seed); + rs = DT::fnv1a32(rs, seed_); rs = DT::fnv1a32(rs, tree_id); raft::random::Rng rng(rs, raft::random::GenPhilox); - if (rf_params.bootstrap) { - // Use bootstrapped sample set - rng.uniformInt(selected_rows->data(), selected_rows->size(), 0, n_rows, stream); + if (bootstrap_) { + if (use_weighted_bootstrap()) { + auto& weighted_draw_scratch = weighted_draw_scratch_[stream_id]; + rng.uniform(weighted_draw_scratch.data(), + weighted_draw_scratch.size(), + 0.0, + sample_weight_sum_, + stream); + thrust::upper_bound(rmm::exec_policy(stream), + sample_weight_cdf_.data(), + sample_weight_cdf_.data() + n_rows_, + weighted_draw_scratch.begin(), + weighted_draw_scratch.end(), + selected_rows.begin()); + } else { + rng.uniformInt(selected_rows.data(), selected_rows.size(), 0, n_rows_, stream); + } } else { - // Use all the samples from the dataset - thrust::sequence(rmm::exec_policy(stream), selected_rows->begin(), selected_rows->end()); + thrust::sequence(rmm::exec_policy(stream), selected_rows.begin(), selected_rows.end()); } + + return selected_rows; + } + + const double* tree_sample_weight() const { return bootstrap_ ? nullptr : sample_weight_; } + + private: + static double validate_sample_weight(const raft::handle_t& handle, + const double* sample_weight, + int n_rows) + { + ASSERT(sample_weight == nullptr || DT::is_dev_ptr(sample_weight), + "sample_weight must be a GPU pointer"); + if (sample_weight == nullptr) { return 0.0; } + + bool has_invalid = thrust::any_of(rmm::exec_policy(handle.get_stream()), + sample_weight, + sample_weight + n_rows, + InvalidSampleWeight{}); + ASSERT(!has_invalid, "sample_weight values must be finite and non-negative"); + double weight_sum = thrust::reduce( + rmm::exec_policy(handle.get_stream()), sample_weight, sample_weight + n_rows, 0.0); + ASSERT(weight_sum > 0.0, "sample_weight values must contain at least one positive value"); + return weight_sum; } + bool use_weighted_bootstrap() const { return bootstrap_ && sample_weight_ != nullptr; } + + bool bootstrap_; + uint64_t seed_; + int n_rows_; + int n_sampled_rows_; + const double* sample_weight_; + double sample_weight_sum_; + rmm::device_uvector sample_weight_cdf_; + std::deque> selected_rows_; + std::deque> weighted_draw_scratch_; +}; +} // namespace detail + +template +class RandomForest { + protected: + RF_params rf_params; // structure containing RF hyperparameters + int rf_type; // 0 for classification 1 for regression + void error_checking(const T* input, L* predictions, int n_rows, int n_cols, bool predict) const { if (predict) { @@ -91,21 +182,6 @@ class RandomForest { } } - void validate_sample_weight(const raft::handle_t& handle, - const T* sample_weight, - int n_rows) const - { - ASSERT(sample_weight == nullptr || DT::is_dev_ptr(sample_weight), - "sample_weight must be a GPU pointer"); - if (sample_weight == nullptr) { return; } - - bool has_invalid = thrust::any_of(rmm::exec_policy(handle.get_stream()), - sample_weight, - sample_weight + n_rows, - detail::InvalidSampleWeight{}); - ASSERT(!has_invalid, "sample_weight values must be finite and non-negative"); - } - public: /** * @brief Construct RandomForest object. @@ -130,10 +206,9 @@ class RandomForest { * @param[in] n_unique_labels: (meaningful only for classification) #unique label values (known during preprocessing) * @param[in] forest: CPU point to RandomForestMetaData struct. - * @param[out] bootstrap_masks: optional device pointer to store bootstrap masks - * (n_trees * n_rows), only populated if a non-null pointer is provided - * @param[in] sample_weight: optional device pointer to per-row sample weights. Counts remain - * sample counts; weights are used only for impurity/objective math. + * @param[in] sample_weight: optional device pointer to per-row sample weights. With bootstrap + * enabled, rows are sampled with probability proportional to these weights and the sampled + * counts drive tree training. Without bootstrap, weights are used for impurity/objective math. */ void fit(const raft::handle_t& user_handle, const T* input, @@ -142,14 +217,12 @@ class RandomForest { L* labels, int n_unique_labels, RandomForestMetaData* forest, - bool* bootstrap_masks = nullptr, - const T* sample_weight = nullptr) + const double* sample_weight = nullptr) { raft::common::nvtx::range fun_scope("RandomForest::fit @randomforest.cuh"); this->error_checking(input, labels, n_rows, n_cols, false); const raft::handle_t& handle = user_handle; - this->validate_sample_weight(handle, sample_weight, n_rows); - int n_sampled_rows = 0; + int n_sampled_rows = 0; if (this->rf_params.bootstrap) { n_sampled_rows = std::round(this->rf_params.max_samples * n_rows); } else { @@ -174,15 +247,8 @@ class RandomForest { // n_streams should not be less than n_trees if (this->rf_params.n_trees < n_streams) n_streams = this->rf_params.n_trees; - // Select n_sampled_rows (with replacement) numbers from [0, n_rows) per tree. - // selected_rows: randomly generated IDs for bootstrapped samples (w/ replacement); a device - // ptr. - // Use a deque instead of vector because it can be used on objects with a deleted copy - // constructor - std::deque> selected_rows; - for (int i = 0; i < n_streams; i++) { - selected_rows.emplace_back(n_sampled_rows, handle.get_stream_from_stream_pool(i)); - } + detail::RowSampler row_sampler( + handle, this->rf_params, n_rows, n_sampled_rows, n_streams, sample_weight); forest->n_features = n_cols; @@ -191,7 +257,7 @@ class RandomForest { int stream_id = omp_get_thread_num(); auto s = handle.get_stream_from_stream_pool(stream_id); - this->get_row_sample(i, n_rows, &selected_rows[stream_id], s); + auto& selected_rows = row_sampler.sample(i, stream_id, s); /* Build individual tree in the forest. - input is a pointer to orig data that have n_cols features and n_rows rows. @@ -208,27 +274,13 @@ class RandomForest { n_cols, n_rows, labels, - &selected_rows[stream_id], + &selected_rows, n_unique_labels, this->rf_params.tree_params, this->rf_params.seed, quantiles, i, - sample_weight); - - // Store bootstrap mask if device buffer is provided - if (bootstrap_masks != nullptr) { - // Calculate pointer offset for this tree's mask - bool* tree_mask = bootstrap_masks + (i * n_rows); - - // Use Thrust to create boolean mask: first fill with false, then mark selected rows - thrust::fill(rmm::exec_policy(s), tree_mask, tree_mask + n_rows, false); - thrust::scatter(rmm::exec_policy(s), - thrust::make_constant_iterator(true), - thrust::make_constant_iterator(true) + n_sampled_rows, - selected_rows[stream_id].data(), - tree_mask); - } + row_sampler.tree_sample_weight()); } // Cleanup handle.sync_stream_pool(); diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu index 24091a8d98..0ff3917bbe 100644 --- a/cpp/tests/sg/rf_test.cu +++ b/cpp/tests/sg/rf_test.cu @@ -285,18 +285,18 @@ auto nvForestPredictProba(const raft::handle_t& handle, return pred; } -template +template RF_metrics Score(const raft::handle_t& handle, RfTestParams params, const LabelT* y, const LabelT* pred, - const DataT* sample_weight) + const double* sample_weight) { thrust::host_vector h_y(params.n_rows); thrust::host_vector h_pred(params.n_rows); raft::update_host(h_y.data(), y, params.n_rows, handle.get_stream()); raft::update_host(h_pred.data(), pred, params.n_rows, handle.get_stream()); - thrust::host_vector h_sample_weight; + thrust::host_vector h_sample_weight; if (sample_weight != nullptr) { h_sample_weight.resize(params.n_rows); raft::update_host(h_sample_weight.data(), sample_weight, params.n_rows, handle.get_stream()); @@ -334,7 +334,7 @@ auto TrainScore(const raft::handle_t& handle, DataT* X, DataT* X_transpose, LabelT* y, - const DataT* sample_weight) + const double* sample_weight) { RF_params rf_params = set_rf_params(params.max_depth, params.max_leaves, @@ -363,7 +363,6 @@ auto TrainScore(const raft::handle_t& handle, params.n_labels, rf_params, rapids_logger::level_enum::info, - nullptr, sample_weight); } else { fit(handle, @@ -374,7 +373,6 @@ auto TrainScore(const raft::handle_t& handle, y, rf_params, rapids_logger::level_enum::info, - nullptr, sample_weight); } @@ -438,10 +436,10 @@ class RfSpecialisedTest { raft::linalg::transpose( handle, X.data().get(), X_transpose.data().get(), params.n_rows, params.n_cols, nullptr); if (params.sample_weight) { - thrust::host_vector h_sample_weight(params.n_rows); + thrust::host_vector h_sample_weight(params.n_rows); for (std::size_t i = 0; i < params.n_rows; ++i) { int bucket = (int(i) * 37 + params.seed * 13) % 17; - h_sample_weight[i] = DataT(0.25) + DataT(bucket) * DataT(0.125); + h_sample_weight[i] = 0.25 + double(bucket) * 0.125; } sample_weight = h_sample_weight; } @@ -682,7 +680,7 @@ class RfSpecialisedTest { EXPECT_NEAR(sum, 1.0, 1e-6); } - const DataT* SampleWeightPtr() const + const double* SampleWeightPtr() const { return params.sample_weight ? sample_weight.data().get() : nullptr; } @@ -702,7 +700,7 @@ class RfSpecialisedTest { thrust::device_vector X; thrust::device_vector X_transpose; thrust::device_vector y; - thrust::device_vector sample_weight; + thrust::device_vector sample_weight; RfTestParams params; std::shared_ptr> forest; std::shared_ptr> predictions; @@ -850,7 +848,7 @@ TEST(RfTests, InvalidSampleWeightThrows) raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); thrust::device_vector X(n_rows * n_cols); thrust::device_vector y(n_rows); - thrust::device_vector sample_weight(n_rows, 1.0f); + thrust::device_vector sample_weight(n_rows, 1.0); raft::random::Rng r(8); r.normal(X.data().get(), X.size(), 0.0f, 1.0f, handle.get_stream()); thrust::host_vector h_y(n_rows); @@ -862,9 +860,9 @@ TEST(RfTests, InvalidSampleWeightThrows) RF_params rf_params = set_rf_params(3, 100, 1.0, 8, 1, 2, 0.0, false, 1, 1.0, 0, CRITERION::GINI, 1, 128); - auto expect_invalid_weight_throws = [&](float invalid_weight) { + auto expect_invalid_weight_throws = [&](double invalid_weight) { thrust::fill( - thrust::cuda::par.on(handle.get_stream()), sample_weight.begin(), sample_weight.end(), 1.0f); + thrust::cuda::par.on(handle.get_stream()), sample_weight.begin(), sample_weight.end(), 1.0); sample_weight[0] = invalid_weight; auto forest = std::make_shared>(); auto forest_ptr = forest.get(); @@ -877,13 +875,80 @@ TEST(RfTests, InvalidSampleWeightThrows) 2, rf_params, rapids_logger::level_enum::info, - nullptr, sample_weight.data().get()), raft::exception); }; - expect_invalid_weight_throws(-1.0f); - expect_invalid_weight_throws(std::numeric_limits::quiet_NaN()); + expect_invalid_weight_throws(-1.0); + expect_invalid_weight_throws(std::numeric_limits::quiet_NaN()); + + thrust::fill( + thrust::cuda::par.on(handle.get_stream()), sample_weight.begin(), sample_weight.end(), 0.0); + auto forest = std::make_shared>(); + auto forest_ptr = forest.get(); + EXPECT_THROW(fit(handle, + forest_ptr, + X.data().get(), + n_rows, + n_cols, + y.data().get(), + 2, + rf_params, + rapids_logger::level_enum::info, + sample_weight.data().get()), + raft::exception); +} + +TEST(RfTests, WeightedBootstrapSamplesOnlyPositiveWeightRows) +{ + constexpr int n_rows = 32; + constexpr int n_cols = 2; + constexpr int n_trees = 8; + constexpr int n_zero_weight_rows = 16; + + auto stream_pool = std::make_shared(2); + raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); + thrust::device_vector X(n_rows * n_cols); + thrust::device_vector y(n_rows); + thrust::device_vector sample_weight(n_rows); + + raft::random::Rng r(8); + r.normal(X.data().get(), X.size(), 0.0f, 1.0f, handle.get_stream()); + + thrust::host_vector h_y(n_rows); + thrust::host_vector h_sample_weight(n_rows); + for (int i = 0; i < n_rows; ++i) { + h_y[i] = i < n_zero_weight_rows ? 0 : 1; + h_sample_weight[i] = i < n_zero_weight_rows ? 0.0 : 1.0; + } + y = h_y; + sample_weight = h_sample_weight; + + RF_params rf_params = + set_rf_params(0, 100, 1.0, 8, 1, 2, 0.0, true, n_trees, 1.0, 0, CRITERION::GINI, 2, 128); + + auto forest = std::make_shared>(); + auto forest_ptr = forest.get(); + fit(handle, + forest_ptr, + X.data().get(), + n_rows, + n_cols, + y.data().get(), + 2, + rf_params, + rapids_logger::level_enum::info, + sample_weight.data().get()); + + for (auto const& tree_ptr : forest->trees) { + const auto& tree = *tree_ptr; + ASSERT_EQ(tree.sparsetree.size(), 1); + EXPECT_TRUE(tree.sparsetree[0].IsLeaf()); + EXPECT_EQ(tree.sparsetree[0].InstanceCount(), n_rows); + ASSERT_EQ(tree.vector_leaf.size(), 2); + EXPECT_NEAR(tree.vector_leaf[0], 0.0f, 1e-6f); + EXPECT_NEAR(tree.vector_leaf[1], 1.0f, 1e-6f); + } } //------------------------------------------------------------------------------------------------------------------------------------- @@ -1318,13 +1383,13 @@ TEST(RfWeightedTest, ClassificationRootLeafUsesWeights) RF_params rf_params = set_rf_params(0, -1, 1.0, 4, 1, 2, 0.0, false, 1, 1.0, 0, GINI, 1, 128); auto forest = std::make_shared>(); - std::vector X_host = {0.0f, 1.0f, 2.0f}; - thrust::device_vector X = X_host; - std::vector y_host = {0, 1, 1}; - thrust::device_vector y = y_host; - std::vector weight_host = {100.0f, 1.0f, 1.0f}; - thrust::device_vector weights = weight_host; - auto stream_pool = std::make_shared(1); + std::vector X_host = {0.0f, 1.0f, 2.0f}; + thrust::device_vector X = X_host; + std::vector y_host = {0, 1, 1}; + thrust::device_vector y = y_host; + std::vector weight_host = {100.0f, 1.0f, 1.0f}; + thrust::device_vector weights = weight_host; + auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); fit(handle, @@ -1336,7 +1401,6 @@ TEST(RfWeightedTest, ClassificationRootLeafUsesWeights) 2, rf_params, rapids_logger::level_enum::info, - nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1354,13 +1418,13 @@ TEST(RfWeightedTest, RegressionRootLeafUsesWeights) RF_params rf_params = set_rf_params(0, -1, 1.0, 4, 1, 2, 0.0, false, 1, 1.0, 0, MSE, 1, 128); auto forest = std::make_shared>(); - std::vector X_host = {0.0f, 1.0f, 2.0f}; - thrust::device_vector X = X_host; - std::vector y_host = {0.0f, 10.0f, 10.0f}; - thrust::device_vector y = y_host; - std::vector weight_host = {1.0f, 0.0f, 3.0f}; - thrust::device_vector weights = weight_host; - auto stream_pool = std::make_shared(1); + std::vector X_host = {0.0f, 1.0f, 2.0f}; + thrust::device_vector X = X_host; + std::vector y_host = {0.0f, 10.0f, 10.0f}; + thrust::device_vector y = y_host; + std::vector weight_host = {1.0f, 0.0f, 3.0f}; + thrust::device_vector weights = weight_host; + auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); fit(handle, @@ -1371,7 +1435,6 @@ TEST(RfWeightedTest, RegressionRootLeafUsesWeights) y.data().get(), rf_params, rapids_logger::level_enum::info, - nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1388,13 +1451,13 @@ TEST(RfWeightedTest, MinSamplesLeafUsesCountsNotWeights) RF_params rf_params = set_rf_params(1, -1, 1.0, 4, 2, 2, 0.0, false, 1, 1.0, 0, GINI, 1, 128); auto forest = std::make_shared>(); - std::vector X_host = {0.0f, 1.0f, 2.0f, 3.0f}; - thrust::device_vector X = X_host; - std::vector y_host = {0, 0, 1, 1}; - thrust::device_vector y = y_host; - std::vector weight_host = {0.1f, 0.1f, 100.0f, 100.0f}; - thrust::device_vector weights = weight_host; - auto stream_pool = std::make_shared(1); + std::vector X_host = {0.0f, 1.0f, 2.0f, 3.0f}; + thrust::device_vector X = X_host; + std::vector y_host = {0, 0, 1, 1}; + thrust::device_vector y = y_host; + std::vector weight_host = {0.1f, 0.1f, 100.0f, 100.0f}; + thrust::device_vector weights = weight_host; + auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); fit(handle, @@ -1406,7 +1469,6 @@ TEST(RfWeightedTest, MinSamplesLeafUsesCountsNotWeights) 2, rf_params, rapids_logger::level_enum::info, - nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1428,13 +1490,13 @@ TEST(RfWeightedTest, ZeroWeightSamplesDoNotCreatePositiveWeightSplit) RF_params rf_params = set_rf_params(1, -1, 1.0, 4, 1, 2, 0.0, false, 1, 1.0, 0, GINI, 1, 128); auto forest = std::make_shared>(); - std::vector X_host = {0.0f, 1.0f, 2.0f, 3.0f}; - thrust::device_vector X = X_host; - std::vector y_host = {0, 0, 1, 1}; - thrust::device_vector y = y_host; - std::vector weight_host = {0.0f, 0.0f, 1.0f, 1.0f}; - thrust::device_vector weights = weight_host; - auto stream_pool = std::make_shared(1); + std::vector X_host = {0.0f, 1.0f, 2.0f, 3.0f}; + thrust::device_vector X = X_host; + std::vector y_host = {0, 0, 1, 1}; + thrust::device_vector y = y_host; + std::vector weight_host = {0.0f, 0.0f, 1.0f, 1.0f}; + thrust::device_vector weights = weight_host; + auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); fit(handle, @@ -1446,7 +1508,6 @@ TEST(RfWeightedTest, ZeroWeightSamplesDoNotCreatePositiveWeightSplit) 2, rf_params, rapids_logger::level_enum::info, - nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1461,21 +1522,27 @@ TEST(RfWeightedTest, ZeroWeightSamplesDoNotCreatePositiveWeightSplit) TEST(RfWeightedTest, BootstrapDuplicatesContributePerOccurrence) { - std::vector X_host = {0.0f, 1.0f, 2.0f}; - thrust::device_vector X = X_host; - std::vector y_host = {0.0f, 10.0f, 100.0f}; - thrust::device_vector y = y_host; - std::vector weight_host = {1.0f, 2.0f, 5.0f}; - thrust::device_vector weights = weight_host; - auto stream_pool = std::make_shared(1); + std::vector X_host = {0.0f, 1.0f, 2.0f}; + thrust::device_vector X = X_host; + std::vector y_host = {0.0f, 10.0f, 100.0f}; + thrust::device_vector y = y_host; + std::vector weight_host = {1.0f, 2.0f, 5.0f}; + thrust::device_vector weights = weight_host; + auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); - constexpr int n_rows = 3; + constexpr int n_rows = 3; + auto mean_with_counts = [&](int count_0, int count_1, int count_2) { + auto label_sum = y_host[0] * count_0 + y_host[1] * count_1 + y_host[2] * count_2; + auto count_sum = count_0 + count_1 + count_2; + return label_sum / count_sum; + }; + auto unique_mean = mean_with_counts(1, 1, 1); + bool found_duplicate = false; for (uint64_t seed = 0; seed < 64 && !found_duplicate; ++seed) { RF_params rf_params = set_rf_params(0, -1, 1.0, 3, 1, 2, 0.0, true, 1, 1.0, seed, MSE, 1, 128); auto forest = std::make_shared>(); - rmm::device_uvector bootstrap_masks(n_rows, handle.get_stream()); fit(handle, forest.get(), @@ -1485,48 +1552,29 @@ TEST(RfWeightedTest, BootstrapDuplicatesContributePerOccurrence) y.data().get(), rf_params, rapids_logger::level_enum::info, - bootstrap_masks.data(), weights.data().get()); handle.sync_stream(); - std::array mask{}; - raft::update_host(mask.data(), bootstrap_masks.data(), mask.size(), handle.get_stream()); - handle.sync_stream(); - - std::array included{}; - int included_count = 0; - for (int i = 0; i < n_rows; ++i) { - if (mask[i]) { - if (included_count < 2) { included[included_count] = i; } - ++included_count; - } - } - if (included_count != 2) { continue; } - - found_duplicate = true; const auto& tree = *forest->trees[0]; ASSERT_EQ(tree.sparsetree.size(), 1); EXPECT_TRUE(tree.sparsetree[0].IsLeaf()); EXPECT_EQ(tree.sparsetree[0].InstanceCount(), n_rows); ASSERT_EQ(tree.vector_leaf.size(), 1); - auto mean_with_counts = [&](int count_a, int count_b) { - auto a = included[0]; - auto b = included[1]; - auto weighted_sum = - y_host[a] * weight_host[a] * count_a + y_host[b] * weight_host[b] * count_b; - auto weight_sum = weight_host[a] * count_a + weight_host[b] * count_b; - return weighted_sum / weight_sum; - }; - - auto unique_mean = mean_with_counts(1, 1); - auto duplicate_a_mean = mean_with_counts(2, 1); - auto duplicate_b_mean = mean_with_counts(1, 2); - auto observed = tree.vector_leaf[0]; - - EXPECT_GT(std::abs(observed - unique_mean), 1e-5f); - EXPECT_TRUE(std::abs(observed - duplicate_a_mean) < 1e-5f || - std::abs(observed - duplicate_b_mean) < 1e-5f); + auto observed = tree.vector_leaf[0]; + if (std::abs(observed - unique_mean) < 1e-5f) { continue; } + + found_duplicate = true; + bool matched_duplicate_mean = false; + for (int count_0 = 0; count_0 <= n_rows; ++count_0) { + for (int count_1 = 0; count_1 <= n_rows - count_0; ++count_1) { + int count_2 = n_rows - count_0 - count_1; + if (count_0 == 1 && count_1 == 1 && count_2 == 1) { continue; } + auto duplicate_mean = mean_with_counts(count_0, count_1, count_2); + matched_duplicate_mean |= std::abs(observed - duplicate_mean) < 1e-5f; + } + } + EXPECT_TRUE(matched_duplicate_mean); } EXPECT_TRUE(found_duplicate); @@ -1597,9 +1645,9 @@ class ObjectiveTest : public ::testing::TestWithParam { auto GenSampleWeights() { - std::vector sample_weights(params.n_rows, DataT(1)); + std::vector sample_weights(params.n_rows, 1.0); if constexpr (is_weighted) { - std::uniform_real_distribution weight_dist(DataT(0.2), DataT(3.0)); + std::uniform_real_distribution weight_dist(0.2, 3.0); for (auto& w : sample_weights) { w = weight_dist(rng); } @@ -1607,7 +1655,7 @@ class ObjectiveTest : public ::testing::TestWithParam { return sample_weights; } - auto GenHist(std::vector const& data, std::vector const& sample_weights) + auto GenHist(std::vector const& data, std::vector const& sample_weights) { std::vector cdf_hist, pdf_hist; auto bin_width = static_cast(raft::ceildiv(params.n_rows, params.max_n_bins)); @@ -1670,7 +1718,7 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto MSE(std::vector const& data, - std::vector const& + std::vector const& sample_weights) // 1/w * 1/2 * sum(w_i * (y - y_pred) * (y - y_pred)) { DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); @@ -1691,16 +1739,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto MSEGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto [parent_mse, label_sum, n, weight_sum] = MSE(data, sample_weights); auto [left_mse, label_sum_left, n_left, weight_sum_left] = MSE(left_data, left_sample_weights); @@ -1721,8 +1769,8 @@ class ObjectiveTest : public ::testing::TestWithParam { auto InverseGaussianHalfDeviance( std::vector const& data, - std::vector const& sample_weights) // 1/w * 2 * sum(w_i * (y - y_pred) * (y - - // y_pred)/(y * (y_pred) * (y_pred))) + std::vector const& sample_weights) // 1/w * 2 * sum(w_i * (y - y_pred) * (y - + // y_pred)/(y * (y_pred) * (y_pred))) { DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); DataT sum{0}; @@ -1742,16 +1790,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto InverseGaussianGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto [parent_ighd, label_sum, n, weight_sum] = InverseGaussianHalfDeviance(data, sample_weights); @@ -1773,7 +1821,7 @@ class ObjectiveTest : public ::testing::TestWithParam { return gain; } - auto GammaHalfDeviance(std::vector const& data, std::vector const& sample_weights) + auto GammaHalfDeviance(std::vector const& data, std::vector const& sample_weights) // 1/w * 2 * sum(w_i * (log(y_pred/y_true) + y_true/y_pred - 1)) { DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); @@ -1795,16 +1843,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto GammaGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto [parent_ghd, label_sum, n, weight_sum] = GammaHalfDeviance(data, sample_weights); auto [left_ghd, label_sum_left, n_left, weight_sum_left] = @@ -1827,7 +1875,7 @@ class ObjectiveTest : public ::testing::TestWithParam { auto PoissonHalfDeviance( std::vector const& data, - std::vector const& + std::vector const& sample_weights) // 1/w * sum(w_i * (y_true * log(y_true/y_pred) + y_pred - y_true)) { DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); @@ -1849,16 +1897,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto PoissonGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto [parent_phd, label_sum, n, weight_sum] = PoissonHalfDeviance(data, sample_weights); auto [left_phd, label_sum_left, n_left, weight_sum_left] = @@ -1878,7 +1926,7 @@ class ObjectiveTest : public ::testing::TestWithParam { return gain; } - auto Entropy(std::vector const& data, std::vector const& sample_weights) + auto Entropy(std::vector const& data, std::vector const& sample_weights) { // sum((n_c/n_total)*(log(n_c/n_total))) DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); DataT entropy(0); @@ -1895,16 +1943,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto EntropyGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t const split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto parent_entropy = Entropy(data, sample_weights); auto left_entropy = Entropy(left_data, left_sample_weights); @@ -1926,7 +1974,7 @@ class ObjectiveTest : public ::testing::TestWithParam { } } - auto GiniImpurity(std::vector const& data, std::vector const& sample_weights) + auto GiniImpurity(std::vector const& data, std::vector const& sample_weights) { // sum((n_c/n_total)(1-(n_c/n_total))) DataT weight_sum = std::accumulate(sample_weights.begin(), sample_weights.end(), DataT(0)); DataT gini(0); @@ -1942,16 +1990,16 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto GiniGroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t const split_bin_index) { auto split_offset = SplitOffset(split_bin_index); std::vector left_data(data.begin(), data.begin() + split_offset); std::vector right_data(data.begin() + split_offset, data.end()); - std::vector left_sample_weights(sample_weights.begin(), - sample_weights.begin() + split_offset); - std::vector right_sample_weights(sample_weights.begin() + split_offset, - sample_weights.end()); + std::vector left_sample_weights(sample_weights.begin(), + sample_weights.begin() + split_offset); + std::vector right_sample_weights(sample_weights.begin() + split_offset, + sample_weights.end()); auto parent_gini = GiniImpurity(data, sample_weights); auto left_gini = GiniImpurity(left_data, left_sample_weights); @@ -1974,7 +2022,7 @@ class ObjectiveTest : public ::testing::TestWithParam { } auto GroundTruthGain(std::vector const& data, - std::vector const& sample_weights, + std::vector const& sample_weights, std::size_t const split_bin_index) { if constexpr (ObjectiveConfig::splitCriteria == CRITERION::MSE) { From 836fbbcf2c6115ccf10342105746e293d82f7f7a Mon Sep 17 00:00:00 2001 From: Rory Mitchell Date: Fri, 26 Jun 2026 20:10:30 +0200 Subject: [PATCH 2/5] Restore RF bootstrap masks in row sampler --- cpp/include/cuml/ensemble/randomforest.hpp | 6 +++ cpp/src/randomforest/randomforest.cu | 50 +++++++++++++++++++--- cpp/src/randomforest/randomforest.cuh | 27 +++++++++++- cpp/tests/sg/rf_test.cu | 10 +++++ 4 files changed, 85 insertions(+), 8 deletions(-) diff --git a/cpp/include/cuml/ensemble/randomforest.hpp b/cpp/include/cuml/ensemble/randomforest.hpp index 179b3a8ece..1502adc9b4 100644 --- a/cpp/include/cuml/ensemble/randomforest.hpp +++ b/cpp/include/cuml/ensemble/randomforest.hpp @@ -151,6 +151,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, + bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); void fit(const raft::handle_t& user_handle, RandomForestClassifierD* forest, @@ -161,6 +162,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, + bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); template @@ -172,6 +174,7 @@ void fit_treelite(const raft::handle_t& user_handle, L* labels, int n_unique_labels, RF_params rf_params, + bool* bootstrap_masks, T* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight = nullptr); @@ -232,6 +235,7 @@ void fit(const raft::handle_t& user_handle, float* labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, + bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); void fit(const raft::handle_t& user_handle, RandomForestRegressorD* forest, @@ -241,6 +245,7 @@ void fit(const raft::handle_t& user_handle, double* labels, RF_params rf_params, rapids_logger::level_enum verbosity = rapids_logger::level_enum::info, + bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr); template @@ -251,6 +256,7 @@ void fit_treelite(const raft::handle_t& user_handle, int n_cols, L* labels, RF_params rf_params, + bool* bootstrap_masks, T* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight = nullptr); diff --git a/cpp/src/randomforest/randomforest.cu b/cpp/src/randomforest/randomforest.cu index 6284633855..77160e1df9 100644 --- a/cpp/src/randomforest/randomforest.cu +++ b/cpp/src/randomforest/randomforest.cu @@ -357,6 +357,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity, + bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -367,8 +368,15 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_classifier = std::make_shared>(rf_params, RF_type::CLASSIFICATION); - rf_classifier->fit( - user_handle, input, n_rows, n_cols, labels, n_unique_labels, forest, sample_weight); + rf_classifier->fit(user_handle, + input, + n_rows, + n_cols, + labels, + n_unique_labels, + forest, + bootstrap_masks, + sample_weight); } void fit(const raft::handle_t& user_handle, @@ -380,6 +388,7 @@ void fit(const raft::handle_t& user_handle, int n_unique_labels, RF_params rf_params, rapids_logger::level_enum verbosity, + bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -390,8 +399,15 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_classifier = std::make_shared>(rf_params, RF_type::CLASSIFICATION); - rf_classifier->fit( - user_handle, input, n_rows, n_cols, labels, n_unique_labels, forest, sample_weight); + rf_classifier->fit(user_handle, + input, + n_rows, + n_cols, + labels, + n_unique_labels, + forest, + bootstrap_masks, + sample_weight); } template @@ -403,6 +419,7 @@ void fit_treelite(const raft::handle_t& user_handle, label_t* labels, int n_unique_labels, RF_params rf_params, + bool* bootstrap_masks, value_t* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight) @@ -417,6 +434,7 @@ void fit_treelite(const raft::handle_t& user_handle, n_unique_labels, rf_params, verbosity, + bootstrap_masks, sample_weight); // Compute feature importances if requested @@ -586,6 +604,7 @@ void fit(const raft::handle_t& user_handle, float* labels, RF_params rf_params, rapids_logger::level_enum verbosity, + bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -596,7 +615,8 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_regressor = std::make_shared>(rf_params, RF_type::REGRESSION); - rf_regressor->fit(user_handle, input, n_rows, n_cols, labels, 1, forest, sample_weight); + rf_regressor->fit( + user_handle, input, n_rows, n_cols, labels, 1, forest, bootstrap_masks, sample_weight); } void fit(const raft::handle_t& user_handle, @@ -607,6 +627,7 @@ void fit(const raft::handle_t& user_handle, double* labels, RF_params rf_params, rapids_logger::level_enum verbosity, + bool* bootstrap_masks, const double* sample_weight) { raft::common::nvtx::range fun_scope("RF::fit @randomforest.cu"); @@ -617,7 +638,8 @@ void fit(const raft::handle_t& user_handle, std::shared_ptr> rf_regressor = std::make_shared>(rf_params, RF_type::REGRESSION); - rf_regressor->fit(user_handle, input, n_rows, n_cols, labels, 1, forest, sample_weight); + rf_regressor->fit( + user_handle, input, n_rows, n_cols, labels, 1, forest, bootstrap_masks, sample_weight); } template @@ -628,12 +650,22 @@ void fit_treelite(const raft::handle_t& user_handle, int n_cols, label_t* labels, RF_params rf_params, + bool* bootstrap_masks, value_t* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight) { RandomForestMetaData metadata; - fit(user_handle, &metadata, input, n_rows, n_cols, labels, rf_params, verbosity, sample_weight); + fit(user_handle, + &metadata, + input, + n_rows, + n_cols, + labels, + rf_params, + verbosity, + bootstrap_masks, + sample_weight); // Compute feature importances if requested if (feature_importances != nullptr) { @@ -838,6 +870,7 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_ha int* labels, int n_unique_labels, RF_params rf_params, + bool* bootstrap_masks, float* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); @@ -849,6 +882,7 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_h int* labels, int n_unique_labels, RF_params rf_params, + bool* bootstrap_masks, double* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); @@ -859,6 +893,7 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& user_ int n_cols, float* labels, RF_params rf_params, + bool* bootstrap_masks, float* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); @@ -869,6 +904,7 @@ template CUML_EXPORT void fit_treelite(const raft::handle_t& use int n_cols, double* labels, RF_params rf_params, + bool* bootstrap_masks, double* feature_importances, rapids_logger::level_enum verbosity, const double* sample_weight); diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh index fb3df55e94..f91fffae90 100644 --- a/cpp/src/randomforest/randomforest.cuh +++ b/cpp/src/randomforest/randomforest.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -55,11 +56,13 @@ class RowSampler { int n_rows, int n_sampled_rows, int n_streams, + bool* bootstrap_masks, const double* sample_weight) : bootstrap_(rf_params.bootstrap), seed_(rf_params.seed), n_rows_(n_rows), n_sampled_rows_(n_sampled_rows), + bootstrap_masks_(bootstrap_masks), sample_weight_(sample_weight), sample_weight_sum_(validate_sample_weight(handle, sample_weight_, n_rows_)), sample_weight_cdf_(0, handle.get_stream()) @@ -70,6 +73,8 @@ class RowSampler { sample_weight_, sample_weight_ + n_rows_, sample_weight_cdf_.begin()); + raft::update_host( + &sample_weight_sum_, sample_weight_cdf_.data() + (n_rows_ - 1), 1, handle.get_stream()); handle.sync_stream(); } @@ -119,12 +124,28 @@ class RowSampler { thrust::sequence(rmm::exec_policy(stream), selected_rows.begin(), selected_rows.end()); } + store_bootstrap_mask(tree_id, selected_rows, stream); return selected_rows; } const double* tree_sample_weight() const { return bootstrap_ ? nullptr : sample_weight_; } private: + void store_bootstrap_mask(int tree_id, + rmm::device_uvector& selected_rows, + cudaStream_t stream) + { + if (bootstrap_masks_ == nullptr) { return; } + + bool* tree_mask = bootstrap_masks_ + (std::size_t(tree_id) * n_rows_); + thrust::fill(rmm::exec_policy(stream), tree_mask, tree_mask + n_rows_, false); + thrust::scatter(rmm::exec_policy(stream), + thrust::make_constant_iterator(true), + thrust::make_constant_iterator(true) + n_sampled_rows_, + selected_rows.data(), + tree_mask); + } + static double validate_sample_weight(const raft::handle_t& handle, const double* sample_weight, int n_rows) @@ -150,6 +171,7 @@ class RowSampler { uint64_t seed_; int n_rows_; int n_sampled_rows_; + bool* bootstrap_masks_; const double* sample_weight_; double sample_weight_sum_; rmm::device_uvector sample_weight_cdf_; @@ -206,6 +228,8 @@ class RandomForest { * @param[in] n_unique_labels: (meaningful only for classification) #unique label values (known during preprocessing) * @param[in] forest: CPU point to RandomForestMetaData struct. + * @param[out] bootstrap_masks: optional device pointer to store bootstrap masks + * (n_trees * n_rows), only populated if a non-null pointer is provided. * @param[in] sample_weight: optional device pointer to per-row sample weights. With bootstrap * enabled, rows are sampled with probability proportional to these weights and the sampled * counts drive tree training. Without bootstrap, weights are used for impurity/objective math. @@ -217,6 +241,7 @@ class RandomForest { L* labels, int n_unique_labels, RandomForestMetaData* forest, + bool* bootstrap_masks = nullptr, const double* sample_weight = nullptr) { raft::common::nvtx::range fun_scope("RandomForest::fit @randomforest.cuh"); @@ -248,7 +273,7 @@ class RandomForest { if (this->rf_params.n_trees < n_streams) n_streams = this->rf_params.n_trees; detail::RowSampler row_sampler( - handle, this->rf_params, n_rows, n_sampled_rows, n_streams, sample_weight); + handle, this->rf_params, n_rows, n_sampled_rows, n_streams, bootstrap_masks, sample_weight); forest->n_features = n_cols; diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu index 0ff3917bbe..a1ee20d675 100644 --- a/cpp/tests/sg/rf_test.cu +++ b/cpp/tests/sg/rf_test.cu @@ -363,6 +363,7 @@ auto TrainScore(const raft::handle_t& handle, params.n_labels, rf_params, rapids_logger::level_enum::info, + nullptr, sample_weight); } else { fit(handle, @@ -373,6 +374,7 @@ auto TrainScore(const raft::handle_t& handle, y, rf_params, rapids_logger::level_enum::info, + nullptr, sample_weight); } @@ -875,6 +877,7 @@ TEST(RfTests, InvalidSampleWeightThrows) 2, rf_params, rapids_logger::level_enum::info, + nullptr, sample_weight.data().get()), raft::exception); }; @@ -895,6 +898,7 @@ TEST(RfTests, InvalidSampleWeightThrows) 2, rf_params, rapids_logger::level_enum::info, + nullptr, sample_weight.data().get()), raft::exception); } @@ -938,6 +942,7 @@ TEST(RfTests, WeightedBootstrapSamplesOnlyPositiveWeightRows) 2, rf_params, rapids_logger::level_enum::info, + nullptr, sample_weight.data().get()); for (auto const& tree_ptr : forest->trees) { @@ -1401,6 +1406,7 @@ TEST(RfWeightedTest, ClassificationRootLeafUsesWeights) 2, rf_params, rapids_logger::level_enum::info, + nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1435,6 +1441,7 @@ TEST(RfWeightedTest, RegressionRootLeafUsesWeights) y.data().get(), rf_params, rapids_logger::level_enum::info, + nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1469,6 +1476,7 @@ TEST(RfWeightedTest, MinSamplesLeafUsesCountsNotWeights) 2, rf_params, rapids_logger::level_enum::info, + nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1508,6 +1516,7 @@ TEST(RfWeightedTest, ZeroWeightSamplesDoNotCreatePositiveWeightSplit) 2, rf_params, rapids_logger::level_enum::info, + nullptr, weights.data().get()); ASSERT_EQ(forest->trees.size(), 1); @@ -1552,6 +1561,7 @@ TEST(RfWeightedTest, BootstrapDuplicatesContributePerOccurrence) y.data().get(), rf_params, rapids_logger::level_enum::info, + nullptr, weights.data().get()); handle.sync_stream(); From fe8a3e07833f8f794e055e54417db66d7d16c3d1 Mon Sep 17 00:00:00 2001 From: Rory Mitchell Date: Fri, 26 Jun 2026 20:17:51 +0200 Subject: [PATCH 3/5] Restore mask-based RF bootstrap tests --- cpp/tests/sg/rf_test.cu | 72 +++++++++++++++++++++++++++-------------- 1 file changed, 48 insertions(+), 24 deletions(-) diff --git a/cpp/tests/sg/rf_test.cu b/cpp/tests/sg/rf_test.cu index a1ee20d675..38bf9c6ab0 100644 --- a/cpp/tests/sg/rf_test.cu +++ b/cpp/tests/sg/rf_test.cu @@ -933,6 +933,7 @@ TEST(RfTests, WeightedBootstrapSamplesOnlyPositiveWeightRows) auto forest = std::make_shared>(); auto forest_ptr = forest.get(); + rmm::device_uvector bootstrap_masks(std::size_t(n_trees) * n_rows, handle.get_stream()); fit(handle, forest_ptr, X.data().get(), @@ -942,8 +943,20 @@ TEST(RfTests, WeightedBootstrapSamplesOnlyPositiveWeightRows) 2, rf_params, rapids_logger::level_enum::info, - nullptr, + bootstrap_masks.data(), sample_weight.data().get()); + handle.sync_stream(); + + thrust::host_vector h_bootstrap_masks(bootstrap_masks.size()); + raft::update_host( + h_bootstrap_masks.data(), bootstrap_masks.data(), bootstrap_masks.size(), handle.get_stream()); + handle.sync_stream(); + + for (int tree_id = 0; tree_id < n_trees; ++tree_id) { + for (int row_id = 0; row_id < n_zero_weight_rows; ++row_id) { + EXPECT_FALSE(h_bootstrap_masks[tree_id * n_rows + row_id]); + } + } for (auto const& tree_ptr : forest->trees) { const auto& tree = *tree_ptr; @@ -1540,18 +1553,12 @@ TEST(RfWeightedTest, BootstrapDuplicatesContributePerOccurrence) auto stream_pool = std::make_shared(1); raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); - constexpr int n_rows = 3; - auto mean_with_counts = [&](int count_0, int count_1, int count_2) { - auto label_sum = y_host[0] * count_0 + y_host[1] * count_1 + y_host[2] * count_2; - auto count_sum = count_0 + count_1 + count_2; - return label_sum / count_sum; - }; - auto unique_mean = mean_with_counts(1, 1, 1); - + constexpr int n_rows = 3; bool found_duplicate = false; for (uint64_t seed = 0; seed < 64 && !found_duplicate; ++seed) { RF_params rf_params = set_rf_params(0, -1, 1.0, 3, 1, 2, 0.0, true, 1, 1.0, seed, MSE, 1, 128); auto forest = std::make_shared>(); + rmm::device_uvector bootstrap_masks(n_rows, handle.get_stream()); fit(handle, forest.get(), @@ -1561,30 +1568,47 @@ TEST(RfWeightedTest, BootstrapDuplicatesContributePerOccurrence) y.data().get(), rf_params, rapids_logger::level_enum::info, - nullptr, + bootstrap_masks.data(), weights.data().get()); handle.sync_stream(); + std::array mask{}; + raft::update_host(mask.data(), bootstrap_masks.data(), mask.size(), handle.get_stream()); + handle.sync_stream(); + + std::array included{}; + int included_count = 0; + for (int i = 0; i < n_rows; ++i) { + if (mask[i]) { + if (included_count < 2) { included[included_count] = i; } + ++included_count; + } + } + if (included_count != 2) { continue; } + + found_duplicate = true; const auto& tree = *forest->trees[0]; ASSERT_EQ(tree.sparsetree.size(), 1); EXPECT_TRUE(tree.sparsetree[0].IsLeaf()); EXPECT_EQ(tree.sparsetree[0].InstanceCount(), n_rows); ASSERT_EQ(tree.vector_leaf.size(), 1); - auto observed = tree.vector_leaf[0]; - if (std::abs(observed - unique_mean) < 1e-5f) { continue; } - - found_duplicate = true; - bool matched_duplicate_mean = false; - for (int count_0 = 0; count_0 <= n_rows; ++count_0) { - for (int count_1 = 0; count_1 <= n_rows - count_0; ++count_1) { - int count_2 = n_rows - count_0 - count_1; - if (count_0 == 1 && count_1 == 1 && count_2 == 1) { continue; } - auto duplicate_mean = mean_with_counts(count_0, count_1, count_2); - matched_duplicate_mean |= std::abs(observed - duplicate_mean) < 1e-5f; - } - } - EXPECT_TRUE(matched_duplicate_mean); + auto mean_with_counts = [&](int count_a, int count_b) { + auto a = included[0]; + auto b = included[1]; + auto label_sum = y_host[a] * count_a + y_host[b] * count_b; + auto count_sum = count_a + count_b; + return label_sum / count_sum; + }; + + auto unique_mean = mean_with_counts(1, 1); + auto duplicate_a_mean = mean_with_counts(2, 1); + auto duplicate_b_mean = mean_with_counts(1, 2); + auto observed = tree.vector_leaf[0]; + + EXPECT_GT(std::abs(observed - unique_mean), 1e-5f); + EXPECT_TRUE(std::abs(observed - duplicate_a_mean) < 1e-5f || + std::abs(observed - duplicate_b_mean) < 1e-5f); } EXPECT_TRUE(found_duplicate); From c0cc2cd2022ca833404edae512db0148fdcafa29 Mon Sep 17 00:00:00 2001 From: Rory Mitchell Date: Fri, 26 Jun 2026 20:25:04 +0200 Subject: [PATCH 4/5] Harden RF bootstrap mask writes --- cpp/src/randomforest/randomforest.cuh | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh index f91fffae90..4c6d657dd2 100644 --- a/cpp/src/randomforest/randomforest.cuh +++ b/cpp/src/randomforest/randomforest.cuh @@ -5,6 +5,7 @@ #pragma once +#include #include #include @@ -67,6 +68,8 @@ class RowSampler { sample_weight_sum_(validate_sample_weight(handle, sample_weight_, n_rows_)), sample_weight_cdf_(0, handle.get_stream()) { + ASSERT(bootstrap_masks_ == nullptr || DT::is_dev_ptr(bootstrap_masks_), + "bootstrap_masks must be a GPU pointer"); if (use_weighted_bootstrap()) { sample_weight_cdf_.resize(n_rows_, handle.get_stream()); thrust::inclusive_scan(rmm::exec_policy(handle.get_stream()), @@ -137,7 +140,7 @@ class RowSampler { { if (bootstrap_masks_ == nullptr) { return; } - bool* tree_mask = bootstrap_masks_ + (std::size_t(tree_id) * n_rows_); + bool* tree_mask = bootstrap_masks_ + (ML::checked_mul(tree_id, n_rows_)); thrust::fill(rmm::exec_policy(stream), tree_mask, tree_mask + n_rows_, false); thrust::scatter(rmm::exec_policy(stream), thrust::make_constant_iterator(true), From e1883ec690fb08ec98157f3af5d207ab02048cf9 Mon Sep 17 00:00:00 2001 From: Rory Mitchell Date: Mon, 29 Jun 2026 10:36:52 +0200 Subject: [PATCH 5/5] Address RF row sampler review comments --- cpp/src/randomforest/randomforest.cuh | 59 ++++++++++++++++++--------- 1 file changed, 40 insertions(+), 19 deletions(-) diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh index 4c6d657dd2..b03fbed08e 100644 --- a/cpp/src/randomforest/randomforest.cuh +++ b/cpp/src/randomforest/randomforest.cuh @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -50,6 +51,8 @@ struct InvalidSampleWeight { __device__ bool operator()(T weight) const { return weight < T(0) || !isfinite(weight); } }; +// Matches estimator behavior: when bootstrapping is enabled and sample weights exist, +// those weights are materialized by drawing bootstrap rows according to them. class RowSampler { public: RowSampler(const raft::handle_t& handle, @@ -65,22 +68,25 @@ class RowSampler { n_sampled_rows_(n_sampled_rows), bootstrap_masks_(bootstrap_masks), sample_weight_(sample_weight), - sample_weight_sum_(validate_sample_weight(handle, sample_weight_, n_rows_)), + sample_weight_sum_(0.0), sample_weight_cdf_(0, handle.get_stream()) { ASSERT(bootstrap_masks_ == nullptr || DT::is_dev_ptr(bootstrap_masks_), "bootstrap_masks must be a GPU pointer"); + validate_sample_weight(handle, sample_weight_, n_rows_); if (use_weighted_bootstrap()) { sample_weight_cdf_.resize(n_rows_, handle.get_stream()); thrust::inclusive_scan(rmm::exec_policy(handle.get_stream()), sample_weight_, sample_weight_ + n_rows_, sample_weight_cdf_.begin()); - raft::update_host( - &sample_weight_sum_, sample_weight_cdf_.data() + (n_rows_ - 1), 1, handle.get_stream()); - handle.sync_stream(); } + if (sample_weight_ != nullptr) { + sample_weight_sum_ = compute_sample_weight_sum(handle); + ASSERT(sample_weight_sum_ > 0.0, + "sample_weight values must contain at least one positive value"); + } // Use a deque instead of vector because device_uvector has a deleted copy constructor. for (int i = 0; i < n_streams; i++) { auto stream = handle.get_stream_from_stream_pool(i); @@ -104,16 +110,19 @@ class RowSampler { auto rs = DT::fnv1a32_basis; rs = DT::fnv1a32(rs, seed_); rs = DT::fnv1a32(rs, tree_id); - raft::random::Rng rng(rs, raft::random::GenPhilox); + raft::random::RngState rng_state(rs, raft::random::GenPhilox); if (bootstrap_) { + raft::resources stream_resources; + raft::resource::set_cuda_stream(stream_resources, stream); if (use_weighted_bootstrap()) { auto& weighted_draw_scratch = weighted_draw_scratch_[stream_id]; - rng.uniform(weighted_draw_scratch.data(), - weighted_draw_scratch.size(), - 0.0, - sample_weight_sum_, - stream); + raft::random::uniform(stream_resources, + rng_state, + weighted_draw_scratch.data(), + weighted_draw_scratch.size(), + 0.0, + sample_weight_sum_); thrust::upper_bound(rmm::exec_policy(stream), sample_weight_cdf_.data(), sample_weight_cdf_.data() + n_rows_, @@ -121,7 +130,8 @@ class RowSampler { weighted_draw_scratch.end(), selected_rows.begin()); } else { - rng.uniformInt(selected_rows.data(), selected_rows.size(), 0, n_rows_, stream); + raft::random::uniformInt( + stream_resources, rng_state, selected_rows.data(), selected_rows.size(), 0, n_rows_); } } else { thrust::sequence(rmm::exec_policy(stream), selected_rows.begin(), selected_rows.end()); @@ -131,6 +141,7 @@ class RowSampler { return selected_rows; } + // Use sample weights in impurity / objective calculation only when bootstrapping is not enabled. const double* tree_sample_weight() const { return bootstrap_ ? nullptr : sample_weight_; } private: @@ -149,23 +160,33 @@ class RowSampler { tree_mask); } - static double validate_sample_weight(const raft::handle_t& handle, - const double* sample_weight, - int n_rows) + double compute_sample_weight_sum(const raft::handle_t& handle) const + { + if (use_weighted_bootstrap()) { + double weight_sum = 0.0; + raft::update_host( + &weight_sum, sample_weight_cdf_.data() + n_rows_ - 1, 1, handle.get_stream()); + handle.sync_stream(); + return weight_sum; + } + + return thrust::reduce( + rmm::exec_policy(handle.get_stream()), sample_weight_, sample_weight_ + n_rows_, 0.0); + } + + static void validate_sample_weight(const raft::handle_t& handle, + const double* sample_weight, + int n_rows) { ASSERT(sample_weight == nullptr || DT::is_dev_ptr(sample_weight), "sample_weight must be a GPU pointer"); - if (sample_weight == nullptr) { return 0.0; } + if (sample_weight == nullptr) { return; } bool has_invalid = thrust::any_of(rmm::exec_policy(handle.get_stream()), sample_weight, sample_weight + n_rows, InvalidSampleWeight{}); ASSERT(!has_invalid, "sample_weight values must be finite and non-negative"); - double weight_sum = thrust::reduce( - rmm::exec_policy(handle.get_stream()), sample_weight, sample_weight + n_rows, 0.0); - ASSERT(weight_sum > 0.0, "sample_weight values must contain at least one positive value"); - return weight_sum; } bool use_weighted_bootstrap() const { return bootstrap_ && sample_weight_ != nullptr; }