From 1b06da165f06003875bf2999f379299f4145bcca Mon Sep 17 00:00:00 2001 From: Dmitry Razdoburdin Date: Fri, 18 Oct 2024 23:52:06 +0200 Subject: [PATCH] SYCL. Refactor on-device data structures (#10898) --- plugin/sycl/common/hist_util.cc | 4 +- plugin/sycl/common/partition_builder.h | 2 +- plugin/sycl/data.h | 63 --------- plugin/sycl/data/gradient_index.cc | 123 +++++++++--------- plugin/sycl/data/gradient_index.h | 79 ++--------- plugin/sycl/predictor/predictor.cc | 9 ++ plugin/sycl/tree/hist_updater.cc | 5 +- plugin/sycl/tree/updater_quantile_hist.cc | 6 +- tests/cpp/plugin/test_sycl_ghist_builder.cc | 5 +- tests/cpp/plugin/test_sycl_gradient_index.cc | 30 +---- tests/cpp/plugin/test_sycl_hist_updater.cc | 35 ++--- .../cpp/plugin/test_sycl_partition_builder.cc | 5 +- 12 files changed, 99 insertions(+), 267 deletions(-) diff --git a/plugin/sycl/common/hist_util.cc b/plugin/sycl/common/hist_util.cc index 9f35429678bc..136036253592 100644 --- a/plugin/sycl/common/hist_util.cc +++ b/plugin/sycl/common/hist_util.cc @@ -130,7 +130,7 @@ ::sycl::event BuildHistKernel(::sycl::queue* qu, const size_t n_columns = isDense ? gmat.nfeatures : gmat.row_stride; const auto* pgh = gpair_device.DataConst(); const BinIdxType* gradient_index = gmat.index.data(); - const uint32_t* offsets = gmat.index.Offset(); + const uint32_t* offsets = gmat.cut.cut_ptrs_.ConstDevicePointer(); const size_t nbins = gmat.nbins; const size_t max_work_group_size = @@ -210,7 +210,7 @@ ::sycl::event BuildHistKernel(::sycl::queue* qu, const GradientPair::ValueT* pgh = reinterpret_cast(gpair_device.DataConst()); const BinIdxType* gradient_index = gmat.index.data(); - const uint32_t* offsets = gmat.index.Offset(); + const uint32_t* offsets = gmat.cut.cut_ptrs_.ConstDevicePointer(); FPType* hist_data = reinterpret_cast(hist->Data()); const size_t nbins = gmat.nbins; diff --git a/plugin/sycl/common/partition_builder.h b/plugin/sycl/common/partition_builder.h index c520ff31fb8e..f9b968f9c0a6 100644 --- a/plugin/sycl/common/partition_builder.h +++ b/plugin/sycl/common/partition_builder.h @@ -85,7 +85,7 @@ inline ::sycl::event PartitionSparseKernel(::sycl::queue* qu, const BinIdxType* gradient_index = gmat.index.data(); const size_t* rid = rid_span.begin; const size_t range_size = rid_span.Size(); - const uint32_t* cut_ptrs = gmat.cut_device.Ptrs().DataConst(); + const uint32_t* cut_ptrs = gmat.cut.cut_ptrs_.ConstDevicePointer(); size_t* p_rid_buf = rid_buf->data(); return qu->submit([&](::sycl::handler& cgh) { diff --git a/plugin/sycl/data.h b/plugin/sycl/data.h index d5311a6d4383..57f5127fa510 100644 --- a/plugin/sycl/data.h +++ b/plugin/sycl/data.h @@ -224,69 +224,6 @@ class USMVector { std::shared_ptr data_; }; -/* Wrapper for DMatrix which stores all batches in a single USM buffer */ -struct DeviceMatrix { - DMatrix* p_mat; // Pointer to the original matrix on the host - ::sycl::queue* qu_; - USMVector row_ptr; - USMVector data; - size_t total_offset; - - DeviceMatrix() = default; - - void Init(::sycl::queue* qu, DMatrix* dmat) { - qu_ = qu; - p_mat = dmat; - - size_t num_row = 0; - size_t num_nonzero = 0; - for (auto &batch : dmat->GetBatches()) { - num_nonzero += batch.data.Size(); - num_row += batch.Size(); - } - - row_ptr.Resize(qu_, num_row + 1); - size_t* rows = row_ptr.Data(); - data.Resize(qu_, num_nonzero); - - size_t data_offset = 0; - ::sycl::event event; - for (auto &batch : dmat->GetBatches()) { - const auto& data_vec = batch.data.ConstHostVector(); - const auto& offset_vec = batch.offset.ConstHostVector(); - size_t batch_size = batch.Size(); - if (batch_size > 0) { - const auto base_rowid = batch.base_rowid; - event = qu->memcpy(row_ptr.Data() + base_rowid, offset_vec.data(), - sizeof(size_t) * batch_size, event); - if (base_rowid > 0) { - qu->submit([&](::sycl::handler& cgh) { - cgh.depends_on(event); - cgh.parallel_for<>(::sycl::range<1>(batch_size), [=](::sycl::id<1> pid) { - int row_id = pid[0]; - rows[row_id] += base_rowid; - }); - }); - } - event = qu->memcpy(data.Data() + data_offset, data_vec.data(), - sizeof(Entry) * offset_vec[batch_size], event); - data_offset += offset_vec[batch_size]; - qu->wait(); - } - } - qu_->submit([&](::sycl::handler& cgh) { - cgh.depends_on(event); - cgh.single_task<>([=] { - rows[num_row] = data_offset; - }); - }); - qu_->wait(); - total_offset = data_offset; - } - - ~DeviceMatrix() { - } -}; } // namespace sycl } // namespace xgboost diff --git a/plugin/sycl/data/gradient_index.cc b/plugin/sycl/data/gradient_index.cc index ad1fe5fe24ca..0123fae149d4 100644 --- a/plugin/sycl/data/gradient_index.cc +++ b/plugin/sycl/data/gradient_index.cc @@ -48,51 +48,53 @@ void mergeSort(BinIdxType* begin, BinIdxType* end, BinIdxType* buf) { } } -template +template void GHistIndexMatrix::SetIndexData(::sycl::queue* qu, BinIdxType* index_data, - const DeviceMatrix &dmat, + DMatrix *dmat, size_t nbins, - size_t row_stride, - uint32_t* offsets) { + size_t row_stride) { if (nbins == 0) return; - const xgboost::Entry *data_ptr = dmat.data.DataConst(); - const bst_idx_t *offset_vec = dmat.row_ptr.DataConst(); - const size_t num_rows = dmat.row_ptr.Size() - 1; - const bst_float* cut_values = cut_device.Values().DataConst(); - const uint32_t* cut_ptrs = cut_device.Ptrs().DataConst(); - size_t* hit_count_ptr = hit_count_buff.Data(); - - // Sparse case only - if (!offsets) { - // sort_buff has type uint8_t - sort_buff.Resize(qu, num_rows * row_stride * sizeof(BinIdxType)); - } + const bst_float* cut_values = cut.cut_values_.ConstDevicePointer(); + const uint32_t* cut_ptrs = cut.cut_ptrs_.ConstDevicePointer(); + size_t* hit_count_ptr = hit_count.DevicePointer(); + BinIdxType* sort_data = reinterpret_cast(sort_buff.Data()); - auto event = qu->submit([&](::sycl::handler& cgh) { - cgh.parallel_for<>(::sycl::range<1>(num_rows), [=](::sycl::item<1> pid) { - const size_t i = pid.get_id(0); - const size_t ibegin = offset_vec[i]; - const size_t iend = offset_vec[i + 1]; - const size_t size = iend - ibegin; - const size_t start = i * row_stride; - for (bst_uint j = 0; j < size; ++j) { - uint32_t idx = SearchBin(cut_values, cut_ptrs, data_ptr[ibegin + j]); - index_data[start + j] = offsets ? idx - offsets[j] : idx; - AtomicRef hit_count_ref(hit_count_ptr[idx]); - hit_count_ref.fetch_add(1); + ::sycl::event event; + for (auto &batch : dmat->GetBatches()) { + for (auto &batch : dmat->GetBatches()) { + const xgboost::Entry *data_ptr = batch.data.ConstDevicePointer(); + const bst_idx_t *offset_vec = batch.offset.ConstDevicePointer(); + size_t batch_size = batch.Size(); + if (batch_size > 0) { + const auto base_rowid = batch.base_rowid; + event = qu->submit([&](::sycl::handler& cgh) { + cgh.depends_on(event); + cgh.parallel_for<>(::sycl::range<1>(batch_size), [=](::sycl::item<1> pid) { + const size_t i = pid.get_id(0); + const size_t ibegin = offset_vec[i]; + const size_t iend = offset_vec[i + 1]; + const size_t size = iend - ibegin; + const size_t start = (i + base_rowid) * row_stride; + for (bst_uint j = 0; j < size; ++j) { + uint32_t idx = SearchBin(cut_values, cut_ptrs, data_ptr[ibegin + j]); + index_data[start + j] = isDense ? idx - cut_ptrs[j] : idx; + AtomicRef hit_count_ref(hit_count_ptr[idx]); + hit_count_ref.fetch_add(1); + } + if constexpr (!isDense) { + // Sparse case only + mergeSort(index_data + start, index_data + start + size, sort_data + start); + for (bst_uint j = size; j < row_stride; ++j) { + index_data[start + j] = nbins; + } + } + }); + }); } - if (!offsets) { - // Sparse case only - mergeSort(index_data + start, index_data + start + size, sort_data + start); - for (bst_uint j = size; j < row_stride; ++j) { - index_data[start + j] = nbins; - } - } - }); - }); - qu->memcpy(hit_count.data(), hit_count_ptr, nbins * sizeof(size_t), event); + } + } qu->wait(); } @@ -112,63 +114,60 @@ void GHistIndexMatrix::ResizeIndex(size_t n_index, bool isDense) { void GHistIndexMatrix::Init(::sycl::queue* qu, Context const * ctx, - const DeviceMatrix& p_fmat_device, + DMatrix *dmat, int max_bins) { - nfeatures = p_fmat_device.p_mat->Info().num_col_; + nfeatures = dmat->Info().num_col_; - cut = xgboost::common::SketchOnDMatrix(ctx, p_fmat_device.p_mat, max_bins); - cut_device.Init(qu, cut); + cut = xgboost::common::SketchOnDMatrix(ctx, dmat, max_bins); + cut.SetDevice(ctx->Device()); max_num_bins = max_bins; const uint32_t nbins = cut.Ptrs().back(); this->nbins = nbins; - hit_count.resize(nbins, 0); - hit_count_buff.Resize(qu, nbins, 0); - this->p_fmat = p_fmat_device.p_mat; - const bool isDense = p_fmat_device.p_mat->IsDense(); + hit_count.SetDevice(ctx->Device()); + hit_count.Resize(nbins, 0); + + this->p_fmat = dmat; + const bool isDense = dmat->IsDense(); this->isDense_ = isDense; index.setQueue(qu); row_stride = 0; - for (const auto& batch : p_fmat_device.p_mat->GetBatches()) { + size_t n_rows = 0; + for (const auto& batch : dmat->GetBatches()) { const auto& row_offset = batch.offset.ConstHostVector(); + batch.data.SetDevice(ctx->Device()); + batch.offset.SetDevice(ctx->Device()); + n_rows += batch.Size(); for (auto i = 1ull; i < row_offset.size(); i++) { row_stride = std::max(row_stride, static_cast(row_offset[i] - row_offset[i - 1])); } } - const size_t n_offsets = cut_device.Ptrs().Size() - 1; - const size_t n_rows = p_fmat_device.row_ptr.Size() - 1; + const size_t n_offsets = cut.cut_ptrs_.Size() - 1; const size_t n_index = n_rows * row_stride; ResizeIndex(n_index, isDense); - CHECK_GT(cut_device.Values().Size(), 0U); - - uint32_t* offsets = nullptr; - if (isDense) { - index.ResizeOffset(n_offsets); - offsets = index.Offset(); - qu->memcpy(offsets, cut_device.Ptrs().DataConst(), - sizeof(uint32_t) * n_offsets).wait_and_throw(); - } + CHECK_GT(cut.cut_values_.Size(), 0U); if (isDense) { BinTypeSize curent_bin_size = index.GetBinTypeSize(); if (curent_bin_size == BinTypeSize::kUint8BinsTypeSize) { - SetIndexData(qu, index.data(), p_fmat_device, nbins, row_stride, offsets); + SetIndexData(qu, index.data(), dmat, nbins, row_stride); } else if (curent_bin_size == BinTypeSize::kUint16BinsTypeSize) { - SetIndexData(qu, index.data(), p_fmat_device, nbins, row_stride, offsets); + SetIndexData(qu, index.data(), dmat, nbins, row_stride); } else { CHECK_EQ(curent_bin_size, BinTypeSize::kUint32BinsTypeSize); - SetIndexData(qu, index.data(), p_fmat_device, nbins, row_stride, offsets); + SetIndexData(qu, index.data(), dmat, nbins, row_stride); } /* For sparse DMatrix we have to store index of feature for each bin in index field to chose right offset. So offset is nullptr and index is not reduced */ } else { - SetIndexData(qu, index.data(), p_fmat_device, nbins, row_stride, offsets); + sort_buff.Resize(qu, n_rows * row_stride * sizeof(uint32_t)); + SetIndexData(qu, index.data(), dmat, nbins, row_stride); } } diff --git a/plugin/sycl/data/gradient_index.h b/plugin/sycl/data/gradient_index.h index 9183baf1ff08..7f9cf0bb4a44 100644 --- a/plugin/sycl/data/gradient_index.h +++ b/plugin/sycl/data/gradient_index.h @@ -16,40 +16,6 @@ namespace xgboost { namespace sycl { namespace common { -/*! - * \brief SYCL implementation of HistogramCuts stored in USM buffers to provide access from device kernels - */ -class HistogramCuts { - protected: - using BinIdx = uint32_t; - - public: - HistogramCuts() {} - - explicit HistogramCuts(::sycl::queue* qu) {} - - ~HistogramCuts() { - } - - void Init(::sycl::queue* qu, xgboost::common::HistogramCuts const& cuts) { - qu_ = qu; - cut_values_.Init(qu_, cuts.cut_values_.HostVector()); - cut_ptrs_.Init(qu_, cuts.cut_ptrs_.HostVector()); - min_vals_.Init(qu_, cuts.min_vals_.HostVector()); - } - - // Getters for USM buffers to pass pointers into device kernels - const USMVector& Ptrs() const { return cut_ptrs_; } - const USMVector& Values() const { return cut_values_; } - const USMVector& MinValues() const { return min_vals_; } - - private: - USMVector cut_values_; - USMVector cut_ptrs_; - USMVector min_vals_; - ::sycl::queue* qu_; -}; - using BinTypeSize = ::xgboost::common::BinTypeSize; /*! @@ -63,13 +29,6 @@ struct Index { Index& operator=(Index i) = delete; Index(Index&& i) = delete; Index& operator=(Index&& i) = delete; - uint32_t operator[](size_t i) const { - if (!offset_.Empty()) { - return func_(data_.DataConst(), i) + offset_[i%p_]; - } else { - return func_(data_.DataConst(), i); - } - } void SetBinTypeSize(BinTypeSize binTypeSize) { binTypeSize_ = binTypeSize; switch (binTypeSize) { @@ -102,14 +61,6 @@ struct Index { return reinterpret_cast(data_.DataConst()); } - uint32_t* Offset() { - return offset_.Data(); - } - - const uint32_t* Offset() const { - return offset_.DataConst(); - } - size_t Size() const { return data_.Size() / (binTypeSize_); } @@ -118,11 +69,6 @@ struct Index { data_.Resize(qu_, nBytesData); } - void ResizeOffset(const size_t nDisps) { - offset_.Resize(qu_, nDisps); - p_ = nDisps; - } - uint8_t* begin() const { return data_.Begin(); } @@ -149,10 +95,7 @@ struct Index { using Func = uint32_t (*)(const uint8_t*, size_t); USMVector data_; - // size of this field is equal to number of features - USMVector offset_; BinTypeSize binTypeSize_ {BinTypeSize::kUint8BinsTypeSize}; - size_t p_ {1}; Func func_; ::sycl::queue* qu_; @@ -168,13 +111,11 @@ struct GHistIndexMatrix { /*! \brief The index data */ Index index; /*! \brief hit count of each index */ - std::vector hit_count; - /*! \brief buffers for calculations */ - USMVector hit_count_buff; + HostDeviceVector hit_count; + USMVector sort_buff; /*! \brief The corresponding cuts */ xgboost::common::HistogramCuts cut; - HistogramCuts cut_device; DMatrix* p_fmat; size_t max_num_bins; size_t nbins; @@ -183,22 +124,22 @@ struct GHistIndexMatrix { // Create a global histogram matrix based on a given DMatrix device wrapper void Init(::sycl::queue* qu, Context const * ctx, - const sycl::DeviceMatrix& p_fmat_device, int max_num_bins); + DMatrix *dmat, int max_num_bins); - template + template void SetIndexData(::sycl::queue* qu, BinIdxType* index_data, - const sycl::DeviceMatrix &dmat_device, - size_t nbins, size_t row_stride, uint32_t* offsets); + DMatrix *dmat, + size_t nbins, size_t row_stride); void ResizeIndex(size_t n_index, bool isDense); inline void GetFeatureCounts(size_t* counts) const { - auto nfeature = cut_device.Ptrs().Size() - 1; + auto nfeature = cut.cut_ptrs_.Size() - 1; for (unsigned fid = 0; fid < nfeature; ++fid) { - auto ibegin = cut_device.Ptrs()[fid]; - auto iend = cut_device.Ptrs()[fid + 1]; + auto ibegin = cut.cut_ptrs_.ConstHostVector()[fid]; + auto iend = cut.cut_ptrs_.ConstHostVector()[fid + 1]; for (auto i = ibegin; i < iend; ++i) { - *(counts + fid) += hit_count[i]; + *(counts + fid) += hit_count.ConstHostVector()[i]; } } } diff --git a/plugin/sycl/predictor/predictor.cc b/plugin/sycl/predictor/predictor.cc index 32519f87dfb9..1ef12d1f707f 100755 --- a/plugin/sycl/predictor/predictor.cc +++ b/plugin/sycl/predictor/predictor.cc @@ -92,6 +92,11 @@ class DeviceModel { HostDeviceVector first_node_position; HostDeviceVector tree_group; + void SetDevice(DeviceOrd device) { + first_node_position.SetDevice(device); + tree_group.SetDevice(device); + } + void Init(::sycl::queue* qu, const gbm::GBTreeModel& model, size_t tree_begin, size_t tree_end) { int n_nodes = 0; first_node_position.Resize((tree_end - tree_begin) + 1); @@ -159,6 +164,7 @@ class Predictor : public xgboost::Predictor { void InitOutPredictions(const MetaInfo& info, HostDeviceVector* out_preds, const gbm::GBTreeModel& model) const override { + device_model.SetDevice(ctx_->Device()); CHECK_NE(model.learner_model_param->num_output_group, 0); size_t n = model.learner_model_param->num_output_group * info.num_row_; size_t base_margin_size = info.base_margin_.Data()->Size(); @@ -198,6 +204,7 @@ class Predictor : public xgboost::Predictor { const gbm::GBTreeModel &model, uint32_t tree_begin, uint32_t tree_end = 0) const override { auto* out_preds = &predts->predictions; + out_preds->SetDevice(ctx_->Device()); if (tree_end == 0) { tree_end = model.trees.size(); } @@ -336,6 +343,8 @@ class Predictor : public xgboost::Predictor { float* out_predictions = out_preds->DevicePointer(); ::sycl::event event; for (auto &batch : dmat->GetBatches()) { + batch.data.SetDevice(ctx_->Device()); + batch.offset.SetDevice(ctx_->Device()); const Entry* data = batch.data.ConstDevicePointer(); const size_t* row_ptr = batch.offset.ConstDevicePointer(); size_t batch_size = batch.Size(); diff --git a/plugin/sycl/tree/hist_updater.cc b/plugin/sycl/tree/hist_updater.cc index 506e05499cf0..0a6354e76c58 100644 --- a/plugin/sycl/tree/hist_updater.cc +++ b/plugin/sycl/tree/hist_updater.cc @@ -793,9 +793,8 @@ void HistUpdater::EvaluateSplits( auto evaluator = tree_evaluator_.GetEvaluator(); SplitQuery* split_queries_device = split_queries_device_.Data(); - const uint32_t* cut_ptr = gmat.cut_device.Ptrs().DataConst(); - const bst_float* cut_val = gmat.cut_device.Values().DataConst(); - const bst_float* cut_minval = gmat.cut_device.MinValues().DataConst(); + const uint32_t* cut_ptr = gmat.cut.cut_ptrs_.ConstDevicePointer(); + const bst_float* cut_val = gmat.cut.cut_values_.ConstDevicePointer(); snode_device_.ResizeNoCopy(qu_, snode_host_.size()); event = qu_->memcpy(snode_device_.Data(), snode_host_.data(), diff --git a/plugin/sycl/tree/updater_quantile_hist.cc b/plugin/sycl/tree/updater_quantile_hist.cc index 7d92c5778190..cde5fefba19d 100644 --- a/plugin/sycl/tree/updater_quantile_hist.cc +++ b/plugin/sycl/tree/updater_quantile_hist.cc @@ -83,12 +83,8 @@ void QuantileHistMaker::Update(xgboost::tree::TrainParam const *param, xgboost::common::Span> out_position, const std::vector &trees) { if (dmat != p_last_dmat_ || is_gmat_initialized_ == false) { - updater_monitor_.Start("DeviceMatrixInitialization"); - sycl::DeviceMatrix dmat_device; - dmat_device.Init(qu_, dmat); - updater_monitor_.Stop("DeviceMatrixInitialization"); updater_monitor_.Start("GmatInitialization"); - gmat_.Init(qu_, ctx_, dmat_device, static_cast(param_.max_bin)); + gmat_.Init(qu_, ctx_, dmat, static_cast(param_.max_bin)); updater_monitor_.Stop("GmatInitialization"); is_gmat_initialized_ = true; } diff --git a/tests/cpp/plugin/test_sycl_ghist_builder.cc b/tests/cpp/plugin/test_sycl_ghist_builder.cc index 0b3d8a60bae2..2b62b5d1fd41 100644 --- a/tests/cpp/plugin/test_sycl_ghist_builder.cc +++ b/tests/cpp/plugin/test_sycl_ghist_builder.cc @@ -30,11 +30,8 @@ void GHistBuilderTest(float sparsity, bool force_atomic_use) { auto qu = device_manager.GetQueue(ctx.Device()); auto p_fmat = RandomDataGenerator{num_rows, num_columns, sparsity}.GenerateDMatrix(); - sycl::DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); - GHistIndexMatrix gmat_sycl; - gmat_sycl.Init(qu, &ctx, dmat, n_bins); + gmat_sycl.Init(qu, &ctx, p_fmat.get(), n_bins); xgboost::GHistIndexMatrix gmat{&ctx, p_fmat.get(), n_bins, 0.3, false}; diff --git a/tests/cpp/plugin/test_sycl_gradient_index.cc b/tests/cpp/plugin/test_sycl_gradient_index.cc index 4d605ce7aabe..cfa625080254 100644 --- a/tests/cpp/plugin/test_sycl_gradient_index.cc +++ b/tests/cpp/plugin/test_sycl_gradient_index.cc @@ -15,28 +15,6 @@ namespace xgboost::sycl::data { -TEST(SyclGradientIndex, HistogramCuts) { - size_t max_bins = 8; - - Context ctx; - ctx.UpdateAllowUnknown(Args{{"device", "sycl"}}); - - DeviceManager device_manager; - auto qu = device_manager.GetQueue(ctx.Device()); - - auto p_fmat = RandomDataGenerator{512, 16, 0.5}.GenerateDMatrix(true); - - xgboost::common::HistogramCuts cut = - xgboost::common::SketchOnDMatrix(&ctx, p_fmat.get(), max_bins); - - common::HistogramCuts cut_sycl; - cut_sycl.Init(qu, cut); - - VerifySyclVector(cut_sycl.Ptrs(), cut.cut_ptrs_.HostVector()); - VerifySyclVector(cut_sycl.Values(), cut.cut_values_.HostVector()); - VerifySyclVector(cut_sycl.MinValues(), cut.min_vals_.HostVector()); -} - TEST(SyclGradientIndex, Init) { size_t n_rows = 128; size_t n_columns = 7; @@ -48,13 +26,9 @@ TEST(SyclGradientIndex, Init) { auto qu = device_manager.GetQueue(ctx.Device()); auto p_fmat = RandomDataGenerator{n_rows, n_columns, 0.3}.GenerateDMatrix(); - - sycl::DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); - int max_bins = 256; common::GHistIndexMatrix gmat_sycl; - gmat_sycl.Init(qu, &ctx, dmat, max_bins); + gmat_sycl.Init(qu, &ctx, p_fmat.get(), max_bins); xgboost::GHistIndexMatrix gmat{&ctx, p_fmat.get(), max_bins, 0.3, false}; @@ -64,7 +38,7 @@ TEST(SyclGradientIndex, Init) { } { - VerifySyclVector(gmat_sycl.hit_count, gmat.hit_count); + VerifySyclVector(gmat_sycl.hit_count.ConstHostVector(), gmat.hit_count); } { diff --git a/tests/cpp/plugin/test_sycl_hist_updater.cc b/tests/cpp/plugin/test_sycl_hist_updater.cc index 8e5a1d9d9ad6..cead3572763d 100644 --- a/tests/cpp/plugin/test_sycl_hist_updater.cc +++ b/tests/cpp/plugin/test_sycl_hist_updater.cc @@ -171,10 +171,8 @@ void TestHistUpdaterInitData(const xgboost::tree::TrainParam& param, bool has_ne USMVector gpair(qu, num_rows); GenerateRandomGPairs(qu, gpair.Data(), num_rows, has_neg_hess); - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); RegTree tree; auto* row_set_collection = updater.TestInitData(gmat, gpair, *p_fmat, tree); @@ -228,10 +226,8 @@ void TestHistUpdaterBuildHistogramsLossGuide(const xgboost::tree::TrainParam& pa auto* gpair_ptr = gpair.Data(); GenerateRandomGPairs(qu, gpair_ptr, num_rows, false); - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); RegTree tree; tree.ExpandNode(0, 0, 0, false, 0, 0, 0, 0, 0, 0, 0); @@ -290,10 +286,8 @@ void TestHistUpdaterInitNewNode(const xgboost::tree::TrainParam& param, float sp auto* gpair_ptr = gpair.Data(); GenerateRandomGPairs(qu, gpair_ptr, num_rows, false); - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); RegTree tree; tree.ExpandNode(0, 0, 0, false, 0, 0, 0, 0, 0, 0, 0); @@ -348,10 +342,8 @@ void TestHistUpdaterEvaluateSplits(const xgboost::tree::TrainParam& param) { auto* gpair_ptr = gpair.Data(); GenerateRandomGPairs(qu, gpair_ptr, num_rows, false); - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); RegTree tree; tree.ExpandNode(0, 0, 0, false, 0, 0, 0, 0, 0, 0, 0); @@ -371,8 +363,8 @@ void TestHistUpdaterEvaluateSplits(const xgboost::tree::TrainParam& param) { // Check all splits manually. Save the best one and compare with the ans TreeEvaluator tree_evaluator(qu, param, num_columns); auto evaluator = tree_evaluator.GetEvaluator(); - const uint32_t* cut_ptr = gmat.cut_device.Ptrs().DataConst(); - const size_t size = gmat.cut_device.Ptrs().Size(); + const uint32_t* cut_ptr = gmat.cut.cut_ptrs_.ConstDevicePointer(); + const size_t size = gmat.cut.cut_ptrs_.Size(); int n_better_splits = 0; const auto* hist_ptr = (*hist)[0].DataConst(); std::vector best_loss_chg_des(1, -1); @@ -412,11 +404,8 @@ void TestHistUpdaterApplySplit(const xgboost::tree::TrainParam& param, float spa auto qu = device_manager.GetQueue(ctx.Device()); auto p_fmat = RandomDataGenerator{num_rows, num_columns, sparsity}.GenerateDMatrix(); - sycl::DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); - common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, max_bins); + gmat.Init(qu, &ctx, p_fmat.get(), max_bins); RegTree tree; tree.ExpandNode(0, 0, 0, false, 0, 0, 0, 0, 0, 0, 0); @@ -499,11 +488,8 @@ void TestHistUpdaterExpandWithLossGuide(const xgboost::tree::TrainParam& param) std::vector data = {7, 3, 15}; auto p_fmat = GetDMatrixFromData(data, num_rows, num_columns); - - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); std::vector gpair_host = {{1, 2}, {3, 1}, {1, 1}}; USMVector gpair(qu, gpair_host); @@ -547,11 +533,8 @@ void TestHistUpdaterExpandWithDepthWise(const xgboost::tree::TrainParam& param) std::vector data = {7, 3, 15}; auto p_fmat = GetDMatrixFromData(data, num_rows, num_columns); - - DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, n_bins); + gmat.Init(qu, &ctx, p_fmat.get(), n_bins); std::vector gpair_host = {{1, 2}, {3, 1}, {1, 1}}; USMVector gpair(qu, gpair_host); diff --git a/tests/cpp/plugin/test_sycl_partition_builder.cc b/tests/cpp/plugin/test_sycl_partition_builder.cc index 03db81c4f55a..5928988c6441 100644 --- a/tests/cpp/plugin/test_sycl_partition_builder.cc +++ b/tests/cpp/plugin/test_sycl_partition_builder.cc @@ -24,11 +24,8 @@ void TestPartitioning(float sparsity, int max_bins) { auto qu = device_manager.GetQueue(ctx.Device()); auto p_fmat = RandomDataGenerator{num_rows, num_columns, sparsity}.GenerateDMatrix(); - sycl::DeviceMatrix dmat; - dmat.Init(qu, p_fmat.get()); - common::GHistIndexMatrix gmat; - gmat.Init(qu, &ctx, dmat, max_bins); + gmat.Init(qu, &ctx, p_fmat.get(), max_bins); RowSetCollection row_set_collection; auto& row_indices = row_set_collection.Data();