/*! * Copyright 2017-2024 by Contributors * \file hist_updater.cc */ #include "hist_updater.h" #include namespace xgboost { namespace sycl { namespace tree { template void HistUpdater::InitSampling( const USMVector &gpair, USMVector* row_indices) { const size_t num_rows = row_indices->Size(); auto* row_idx = row_indices->Data(); const auto* gpair_ptr = gpair.DataConst(); uint64_t num_samples = 0; const auto subsample = param_.subsample; ::sycl::event event; { ::sycl::buffer flag_buf(&num_samples, 1); uint64_t seed = seed_; seed_ += num_rows; event = qu_.submit([&](::sycl::handler& cgh) { auto flag_buf_acc = flag_buf.get_access<::sycl::access::mode::read_write>(cgh); cgh.parallel_for<>(::sycl::range<1>(::sycl::range<1>(num_rows)), [=](::sycl::item<1> pid) { uint64_t i = pid.get_id(0); // Create minstd_rand engine oneapi::dpl::minstd_rand engine(seed, i); oneapi::dpl::bernoulli_distribution coin_flip(subsample); auto rnd = coin_flip(engine); if (gpair_ptr[i].GetHess() >= 0.0f && rnd) { AtomicRef num_samples_ref(flag_buf_acc[0]); row_idx[num_samples_ref++] = i; } }); }); /* After calling a destructor for flag_buf, content will be copyed to num_samples */ } row_indices->Resize(&qu_, num_samples, 0, &event); qu_.wait(); } template void HistUpdater::InitData( Context const * ctx, const common::GHistIndexMatrix& gmat, const USMVector &gpair, const DMatrix& fmat, const RegTree& tree) { CHECK((param_.max_depth > 0 || param_.max_leaves > 0)) << "max_depth or max_leaves cannot be both 0 (unlimited); " << "at least one should be a positive quantity."; if (param_.grow_policy == xgboost::tree::TrainParam::kDepthWise) { CHECK(param_.max_depth > 0) << "max_depth cannot be 0 (unlimited) " << "when grow_policy is depthwise."; } builder_monitor_.Start("InitData"); const auto& info = fmat.Info(); // initialize the row set { row_set_collection_.Clear(); USMVector* row_indices = &(row_set_collection_.Data()); row_indices->Resize(&qu_, info.num_row_); size_t* p_row_indices = row_indices->Data(); // mark subsample and build list of member rows if (param_.subsample < 1.0f) { CHECK_EQ(param_.sampling_method, xgboost::tree::TrainParam::kUniform) << "Only uniform sampling is supported, " << "gradient-based sampling is only support by GPU Hist."; InitSampling(gpair, row_indices); } else { int has_neg_hess = 0; const GradientPair* gpair_ptr = gpair.DataConst(); ::sycl::event event; { ::sycl::buffer flag_buf(&has_neg_hess, 1); event = qu_.submit([&](::sycl::handler& cgh) { auto flag_buf_acc = flag_buf.get_access<::sycl::access::mode::read_write>(cgh); cgh.parallel_for<>(::sycl::range<1>(::sycl::range<1>(info.num_row_)), [=](::sycl::item<1> pid) { const size_t idx = pid.get_id(0); p_row_indices[idx] = idx; if (gpair_ptr[idx].GetHess() < 0.0f) { AtomicRef has_neg_hess_ref(flag_buf_acc[0]); has_neg_hess_ref.fetch_max(1); } }); }); } if (has_neg_hess) { size_t max_idx = 0; { ::sycl::buffer flag_buf(&max_idx, 1); event = qu_.submit([&](::sycl::handler& cgh) { cgh.depends_on(event); auto flag_buf_acc = flag_buf.get_access<::sycl::access::mode::read_write>(cgh); cgh.parallel_for<>(::sycl::range<1>(::sycl::range<1>(info.num_row_)), [=](::sycl::item<1> pid) { const size_t idx = pid.get_id(0); if (gpair_ptr[idx].GetHess() >= 0.0f) { AtomicRef max_idx_ref(flag_buf_acc[0]); p_row_indices[max_idx_ref++] = idx; } }); }); } row_indices->Resize(&qu_, max_idx, 0, &event); } qu_.wait_and_throw(); } } row_set_collection_.Init(); } template class HistUpdater; template class HistUpdater; } // namespace tree } // namespace sycl } // namespace xgboost