Skip to content

Commit

Permalink
introduce host_device_vector; refactor device_manager
Browse files Browse the repository at this point in the history
  • Loading branch information
Dmitry Razdoburdin committed Sep 20, 2024
1 parent de00e07 commit 1a7b341
Show file tree
Hide file tree
Showing 24 changed files with 902 additions and 266 deletions.
40 changes: 20 additions & 20 deletions plugin/sycl/common/hist_util.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,23 +19,23 @@ namespace common {
* \brief Fill histogram with zeroes
*/
template<typename GradientSumT>
void InitHist(::sycl::queue qu, GHistRow<GradientSumT, MemoryType::on_device>* hist,
void InitHist(::sycl::queue* qu, GHistRow<GradientSumT, MemoryType::on_device>* hist,
size_t size, ::sycl::event* event) {
*event = qu.fill(hist->Begin(),
*event = qu->fill(hist->Begin(),
xgboost::detail::GradientPairInternal<GradientSumT>(), size, *event);
}
template void InitHist(::sycl::queue qu,
template void InitHist(::sycl::queue* qu,
GHistRow<float, MemoryType::on_device>* hist,
size_t size, ::sycl::event* event);
template void InitHist(::sycl::queue qu,
template void InitHist(::sycl::queue* qu,
GHistRow<double, MemoryType::on_device>* hist,
size_t size, ::sycl::event* event);

/*!
* \brief Compute Subtraction: dst = src1 - src2
*/
template<typename GradientSumT>
::sycl::event SubtractionHist(::sycl::queue qu,
::sycl::event SubtractionHist(::sycl::queue* qu,
GHistRow<GradientSumT, MemoryType::on_device>* dst,
const GHistRow<GradientSumT, MemoryType::on_device>& src1,
const GHistRow<GradientSumT, MemoryType::on_device>& src2,
Expand All @@ -44,7 +44,7 @@ ::sycl::event SubtractionHist(::sycl::queue qu,
const GradientSumT* psrc1 = reinterpret_cast<const GradientSumT*>(src1.DataConst());
const GradientSumT* psrc2 = reinterpret_cast<const GradientSumT*>(src2.DataConst());

auto event_final = qu.submit([&](::sycl::handler& cgh) {
auto event_final = qu->submit([&](::sycl::handler& cgh) {
cgh.depends_on(event_priv);
cgh.parallel_for<>(::sycl::range<1>(2 * size), [pdst, psrc1, psrc2](::sycl::item<1> pid) {
const size_t i = pid.get_id(0);
Expand All @@ -53,25 +53,25 @@ ::sycl::event SubtractionHist(::sycl::queue qu,
});
return event_final;
}
template ::sycl::event SubtractionHist(::sycl::queue qu,
template ::sycl::event SubtractionHist(::sycl::queue* qu,
GHistRow<float, MemoryType::on_device>* dst,
const GHistRow<float, MemoryType::on_device>& src1,
const GHistRow<float, MemoryType::on_device>& src2,
size_t size, ::sycl::event event_priv);
template ::sycl::event SubtractionHist(::sycl::queue qu,
template ::sycl::event SubtractionHist(::sycl::queue* qu,
GHistRow<double, MemoryType::on_device>* dst,
const GHistRow<double, MemoryType::on_device>& src1,
const GHistRow<double, MemoryType::on_device>& src2,
size_t size, ::sycl::event event_priv);

inline auto GetBlocksParameters(const ::sycl::queue& qu, size_t size, size_t max_nblocks) {
inline auto GetBlocksParameters(::sycl::queue* qu, size_t size, size_t max_nblocks) {
struct _ {
size_t block_size, nblocks;
};

const size_t min_block_size = 32;
const size_t max_compute_units =
qu.get_device().get_info<::sycl::info::device::max_compute_units>();
qu->get_device().get_info<::sycl::info::device::max_compute_units>();

size_t nblocks = max_compute_units;

Expand All @@ -90,7 +90,7 @@ inline auto GetBlocksParameters(const ::sycl::queue& qu, size_t size, size_t max

// Kernel with buffer using
template<typename FPType, typename BinIdxType, bool isDense>
::sycl::event BuildHistKernel(::sycl::queue qu,
::sycl::event BuildHistKernel(::sycl::queue* qu,
const USMVector<GradientPair, MemoryType::on_device>& gpair_device,
const RowSetCollection::Elem& row_indices,
const GHistIndexMatrix& gmat,
Expand All @@ -107,7 +107,7 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,
const size_t nbins = gmat.nbins;

const size_t max_work_group_size =
qu.get_device().get_info<::sycl::info::device::max_work_group_size>();
qu->get_device().get_info<::sycl::info::device::max_work_group_size>();
const size_t work_group_size = n_columns < max_work_group_size ? n_columns : max_work_group_size;

// Captured structured bindings are a C++20 extension
Expand All @@ -116,8 +116,8 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,
const size_t nblocks = block_params.nblocks;

GradientPairT* hist_buffer_data = hist_buffer->Data();
auto event_fill = qu.fill(hist_buffer_data, GradientPairT(0, 0), nblocks * nbins * 2, event_priv);
auto event_main = qu.submit([&](::sycl::handler& cgh) {
auto event_fill = qu->fill(hist_buffer_data, GradientPairT(0, 0), nblocks * nbins * 2, event_priv);
auto event_main = qu->submit([&](::sycl::handler& cgh) {
cgh.depends_on(event_fill);
cgh.parallel_for<>(::sycl::nd_range<2>(::sycl::range<2>(nblocks, work_group_size),
::sycl::range<2>(1, work_group_size)),
Expand Down Expand Up @@ -151,7 +151,7 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,
});

GradientPairT* hist_data = hist->Data();
auto event_save = qu.submit([&](::sycl::handler& cgh) {
auto event_save = qu->submit([&](::sycl::handler& cgh) {
cgh.depends_on(event_main);
cgh.parallel_for<>(::sycl::range<1>(nbins), [=](::sycl::item<1> pid) {
size_t idx_bin = pid.get_id(0);
Expand All @@ -170,7 +170,7 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,

// Kernel with atomic using
template<typename FPType, typename BinIdxType, bool isDense>
::sycl::event BuildHistKernel(::sycl::queue qu,
::sycl::event BuildHistKernel(::sycl::queue* qu,
const USMVector<GradientPair, MemoryType::on_device>& gpair_device,
const RowSetCollection::Elem& row_indices,
const GHistIndexMatrix& gmat,
Expand All @@ -189,8 +189,8 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,
constexpr size_t work_group_size = 32;
const size_t n_work_groups = n_columns / work_group_size + (n_columns % work_group_size > 0);

auto event_fill = qu.fill(hist_data, FPType(0), nbins * 2, event_priv);
auto event_main = qu.submit([&](::sycl::handler& cgh) {
auto event_fill = qu->fill(hist_data, FPType(0), nbins * 2, event_priv);
auto event_main = qu->submit([&](::sycl::handler& cgh) {
cgh.depends_on(event_fill);
cgh.parallel_for<>(::sycl::nd_range<2>(::sycl::range<2>(size, n_work_groups * work_group_size),
::sycl::range<2>(1, work_group_size)),
Expand Down Expand Up @@ -225,7 +225,7 @@ ::sycl::event BuildHistKernel(::sycl::queue qu,

template<typename FPType, typename BinIdxType>
::sycl::event BuildHistDispatchKernel(
::sycl::queue qu,
::sycl::queue* qu,
const USMVector<GradientPair, MemoryType::on_device>& gpair_device,
const RowSetCollection::Elem& row_indices,
const GHistIndexMatrix& gmat,
Expand Down Expand Up @@ -265,7 +265,7 @@ ::sycl::event BuildHistDispatchKernel(
}

template<typename FPType>
::sycl::event BuildHistKernel(::sycl::queue qu,
::sycl::event BuildHistKernel(::sycl::queue* qu,
const USMVector<GradientPair, MemoryType::on_device>& gpair_device,
const RowSetCollection::Elem& row_indices,
const GHistIndexMatrix& gmat, const bool isDense,
Expand Down
22 changes: 11 additions & 11 deletions plugin/sycl/common/hist_util.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -32,15 +32,15 @@ class ColumnMatrix;
* \brief Fill histogram with zeroes
*/
template<typename GradientSumT>
void InitHist(::sycl::queue qu,
void InitHist(::sycl::queue* qu,
GHistRow<GradientSumT, MemoryType::on_device>* hist,
size_t size, ::sycl::event* event);

/*!
* \brief Compute subtraction: dst = src1 - src2
*/
template<typename GradientSumT>
::sycl::event SubtractionHist(::sycl::queue qu,
::sycl::event SubtractionHist(::sycl::queue* qu,
GHistRow<GradientSumT, MemoryType::on_device>* dst,
const GHistRow<GradientSumT, MemoryType::on_device>& src1,
const GHistRow<GradientSumT, MemoryType::on_device>& src2,
Expand All @@ -64,7 +64,7 @@ class HistCollection {
}

// Initialize histogram collection
void Init(::sycl::queue qu, uint32_t nbins) {
void Init(::sycl::queue* qu, uint32_t nbins) {
qu_ = qu;
if (nbins_ != nbins) {
nbins_ = nbins;
Expand All @@ -77,11 +77,11 @@ class HistCollection {
::sycl::event event;
if (data_.count(nid) == 0) {
data_[nid] =
std::make_shared<GHistRowT>(&qu_, nbins_,
std::make_shared<GHistRowT>(qu_, nbins_,
xgboost::detail::GradientPairInternal<GradientSumT>(0, 0),
&event);
} else {
data_[nid]->Resize(&qu_, nbins_,
data_[nid]->Resize(qu_, nbins_,
xgboost::detail::GradientPairInternal<GradientSumT>(0, 0),
&event);
}
Expand All @@ -94,7 +94,7 @@ class HistCollection {

std::unordered_map<uint32_t, std::shared_ptr<GHistRowT>> data_;

::sycl::queue qu_;
::sycl::queue* qu_;
};

/*!
Expand All @@ -105,7 +105,7 @@ class ParallelGHistBuilder {
public:
using GHistRowT = GHistRow<GradientSumT, MemoryType::on_device>;

void Init(::sycl::queue qu, size_t nbins) {
void Init(::sycl::queue* qu, size_t nbins) {
qu_ = qu;
if (nbins != nbins_) {
hist_buffer_.Init(qu_, nbins);
Expand All @@ -114,7 +114,7 @@ class ParallelGHistBuilder {
}

void Reset(size_t nblocks) {
hist_device_buffer_.Resize(&qu_, nblocks * nbins_ * 2);
hist_device_buffer_.Resize(qu_, nblocks * nbins_ * 2);
}

GHistRowT& GetDeviceBuffer() {
Expand All @@ -130,7 +130,7 @@ class ParallelGHistBuilder {
/*! \brief Buffer for additional histograms for Parallel processing */
GHistRowT hist_device_buffer_;

::sycl::queue qu_;
::sycl::queue* qu_;
};

/*!
Expand All @@ -143,7 +143,7 @@ class GHistBuilder {
using GHistRowT = GHistRow<GradientSumT, memory_type>;

GHistBuilder() = default;
GHistBuilder(::sycl::queue qu, uint32_t nbins) : qu_{qu}, nbins_{nbins} {}
GHistBuilder(::sycl::queue* qu, uint32_t nbins) : qu_{qu}, nbins_{nbins} {}

// Construct a histogram via histogram aggregation
::sycl::event BuildHist(const USMVector<GradientPair, MemoryType::on_device>& gpair_device,
Expand All @@ -168,7 +168,7 @@ class GHistBuilder {
/*! \brief Number of all bins over all features */
uint32_t nbins_ { 0 };

::sycl::queue qu_;
::sycl::queue* qu_;
};
} // namespace common
} // namespace sycl
Expand Down
Loading

0 comments on commit 1a7b341

Please sign in to comment.