diff --git a/plugin/sycl/tree/param.h b/plugin/sycl/tree/param.h new file mode 100644 index 000000000000..1b47d83a4035 --- /dev/null +++ b/plugin/sycl/tree/param.h @@ -0,0 +1,55 @@ +/*! + * Copyright 2014-2024 by Contributors + */ +#ifndef PLUGIN_SYCL_TREE_PARAM_H_ +#define PLUGIN_SYCL_TREE_PARAM_H_ + + +#include +#include +#include +#include +#include + + +#include "xgboost/parameter.h" +#include "xgboost/data.h" +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#include "../src/tree/param.h" +#pragma GCC diagnostic pop + +#include + +namespace xgboost { +namespace sycl { +namespace tree { + + +/*! \brief Wrapper for necessary training parameters for regression tree to access on device */ +/* The original structure xgboost::tree::TrainParam can't be used, + * since std::vector are not copyable on sycl-devices. + */ +struct TrainParam { + float min_child_weight; + float reg_lambda; + float reg_alpha; + float max_delta_step; + + TrainParam() {} + + explicit TrainParam(const xgboost::tree::TrainParam& param) { + reg_lambda = param.reg_lambda; + reg_alpha = param.reg_alpha; + min_child_weight = param.min_child_weight; + max_delta_step = param.max_delta_step; + } +}; + +template +using GradStats = xgboost::detail::GradientPairInternal; + +} // namespace tree +} // namespace sycl +} // namespace xgboost +#endif // PLUGIN_SYCL_TREE_PARAM_H_ diff --git a/plugin/sycl/tree/split_evaluator.h b/plugin/sycl/tree/split_evaluator.h new file mode 100644 index 000000000000..2f1e8c7c4e66 --- /dev/null +++ b/plugin/sycl/tree/split_evaluator.h @@ -0,0 +1,208 @@ +/*! + * Copyright 2018-2024 by Contributors + */ + +#ifndef PLUGIN_SYCL_TREE_SPLIT_EVALUATOR_H_ +#define PLUGIN_SYCL_TREE_SPLIT_EVALUATOR_H_ + +#include +#include +#include +#include +#include + +#include "param.h" +#include "../data.h" + +#include "xgboost/tree_model.h" +#include "xgboost/host_device_vector.h" +#include "xgboost/context.h" +#include "../../src/common/transform.h" +#include "../../src/common/math.h" +#include "../../src/tree/param.h" + +#include + +namespace xgboost { +namespace sycl { +namespace tree { + +/*! \brief SYCL implementation of TreeEvaluator, with USM memory for temporary buffer to access on device. + * It also contains own implementation of SplitEvaluator for device compilation, because some of the + functions from the original SplitEvaluator are currently not supported + */ + +template +class TreeEvaluator { + // hist and exact use parent id to calculate constraints. + static constexpr bst_node_t kRootParentId = + (-1 & static_cast((1U << 31) - 1)); + + USMVector lower_bounds_; + USMVector upper_bounds_; + USMVector monotone_; + TrainParam param_; + ::sycl::queue qu_; + bool has_constraint_; + + public: + void Reset(::sycl::queue qu, xgboost::tree::TrainParam const& p, bst_feature_t n_features) { + qu_ = qu; + + has_constraint_ = false; + for (const auto& constraint : p.monotone_constraints) { + if (constraint != 0) { + has_constraint_ = true; + break; + } + } + + if (has_constraint_) { + monotone_.Resize(&qu_, n_features, 0); + qu_.memcpy(monotone_.Data(), p.monotone_constraints.data(), + sizeof(int) * p.monotone_constraints.size()); + qu_.wait(); + + lower_bounds_.Resize(&qu_, p.MaxNodes(), std::numeric_limits::lowest()); + upper_bounds_.Resize(&qu_, p.MaxNodes(), std::numeric_limits::max()); + } + param_ = TrainParam(p); + } + + bool HasConstraint() const { + return has_constraint_; + } + + TreeEvaluator(::sycl::queue qu, xgboost::tree::TrainParam const& p, bst_feature_t n_features) { + Reset(qu, p, n_features); + } + + struct SplitEvaluator { + const int* constraints; + const GradType* lower; + const GradType* upper; + bool has_constraint; + TrainParam param; + + GradType CalcSplitGain(bst_node_t nidx, + bst_feature_t fidx, + const GradStats& left, + const GradStats& right) const { + const GradType negative_infinity = -std::numeric_limits::infinity(); + GradType wleft = this->CalcWeight(nidx, left); + GradType wright = this->CalcWeight(nidx, right); + + GradType gain = this->CalcGainGivenWeight(nidx, left, wleft) + + this->CalcGainGivenWeight(nidx, right, wright); + if (!has_constraint) { + return gain; + } + + int constraint = constraints[fidx]; + if (constraint == 0) { + return gain; + } else if (constraint > 0) { + return wleft <= wright ? gain : negative_infinity; + } else { + return wleft >= wright ? gain : negative_infinity; + } + } + + inline static GradType ThresholdL1(GradType w, float alpha) { + if (w > + alpha) { + return w - alpha; + } + if (w < - alpha) { + return w + alpha; + } + return 0.0; + } + + inline GradType CalcWeight(GradType sum_grad, GradType sum_hess) const { + if (sum_hess < param.min_child_weight || sum_hess <= 0.0) { + return 0.0; + } + GradType dw = -this->ThresholdL1(sum_grad, param.reg_alpha) / (sum_hess + param.reg_lambda); + if (param.max_delta_step != 0.0f && std::abs(dw) > param.max_delta_step) { + dw = ::sycl::copysign((GradType)param.max_delta_step, dw); + } + return dw; + } + + inline GradType CalcWeight(bst_node_t nodeid, const GradStats& stats) const { + GradType w = this->CalcWeight(stats.GetGrad(), stats.GetHess()); + if (!has_constraint) { + return w; + } + + if (nodeid == kRootParentId) { + return w; + } else if (w < lower[nodeid]) { + return lower[nodeid]; + } else if (w > upper[nodeid]) { + return upper[nodeid]; + } else { + return w; + } + } + + inline GradType CalcGainGivenWeight(GradType sum_grad, GradType sum_hess, GradType w) const { + return -(2.0f * sum_grad * w + (sum_hess + param.reg_lambda) * xgboost::common::Sqr(w)); + } + + inline GradType CalcGainGivenWeight(bst_node_t nid, const GradStats& stats, + GradType w) const { + if (stats.GetHess() <= 0) { + return .0f; + } + // Avoiding tree::CalcGainGivenWeight can significantly reduce avg floating point error. + if (param.max_delta_step == 0.0f && has_constraint == false) { + return xgboost::common::Sqr(this->ThresholdL1(stats.GetGrad(), param.reg_alpha)) / + (stats.GetHess() + param.reg_lambda); + } + return this->CalcGainGivenWeight(stats.GetGrad(), stats.GetHess(), w); + } + + GradType CalcGain(bst_node_t nid, const GradStats& stats) const { + return this->CalcGainGivenWeight(nid, stats, this->CalcWeight(nid, stats)); + } + }; + + public: + /* Get a view to the evaluator that can be passed down to device. */ + auto GetEvaluator() const { + return SplitEvaluator{monotone_.DataConst(), + lower_bounds_.DataConst(), + upper_bounds_.DataConst(), + has_constraint_, + param_}; + } + + void AddSplit(bst_node_t nodeid, bst_node_t leftid, bst_node_t rightid, + bst_feature_t f, GradType left_weight, GradType right_weight) { + if (!has_constraint_) { + return; + } + + lower_bounds_[leftid] = lower_bounds_[nodeid]; + upper_bounds_[leftid] = upper_bounds_[nodeid]; + + lower_bounds_[rightid] = lower_bounds_[nodeid]; + upper_bounds_[rightid] = upper_bounds_[nodeid]; + int32_t c = monotone_[f]; + GradType mid = (left_weight + right_weight) / 2; + + if (c < 0) { + lower_bounds_[leftid] = mid; + upper_bounds_[rightid] = mid; + } else if (c > 0) { + upper_bounds_[leftid] = mid; + lower_bounds_[rightid] = mid; + } + } +}; +} // namespace tree +} // namespace sycl +} // namespace xgboost + +#endif // PLUGIN_SYCL_TREE_SPLIT_EVALUATOR_H_ diff --git a/tests/cpp/plugin/test_sycl_split_evaluator.cc b/tests/cpp/plugin/test_sycl_split_evaluator.cc new file mode 100644 index 000000000000..507490fd17e1 --- /dev/null +++ b/tests/cpp/plugin/test_sycl_split_evaluator.cc @@ -0,0 +1,134 @@ +/** + * Copyright 2020-2024 by XGBoost contributors + */ +#include +#include + +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wtautological-constant-compare" +#pragma GCC diagnostic ignored "-W#pragma-messages" +#include "../../../plugin/sycl/tree/split_evaluator.h" +#pragma GCC diagnostic pop + +#include "../../../plugin/sycl/device_manager.h" +#include "../helpers.h" + +namespace xgboost::sycl::tree { + +template +void BasicTestSplitEvaluator(const std::string& monotone_constraints, bool has_constrains) { + const size_t n_columns = 2; + + xgboost::tree::TrainParam param; + param.UpdateAllowUnknown(Args{{"min_child_weight", "0"}, + {"reg_lambda", "0"}, + {"monotone_constraints", monotone_constraints}}); + + DeviceManager device_manager; + auto qu = device_manager.GetQueue(DeviceOrd::SyclDefault()); + + TreeEvaluator tree_evaluator(qu, param, n_columns); + { + // Check correctness of has_constrains flag + ASSERT_EQ(tree_evaluator.HasConstraint(), has_constrains); + } + + auto split_evaluator = tree_evaluator.GetEvaluator(); + { + // Check if params were inititialised correctly + ASSERT_EQ(split_evaluator.param.min_child_weight, param.min_child_weight); + ASSERT_EQ(split_evaluator.param.reg_lambda, param.reg_lambda); + ASSERT_EQ(split_evaluator.param.reg_alpha, param.reg_alpha); + ASSERT_EQ(split_evaluator.param.max_delta_step, param.max_delta_step); + } +} + +template +void TestSplitEvaluator(const std::string& monotone_constraints) { + const size_t n_columns = 2; + + xgboost::tree::TrainParam param; + param.UpdateAllowUnknown(Args{{"min_child_weight", "0"}, + {"reg_lambda", "0"}, + {"monotone_constraints", monotone_constraints}}); + + DeviceManager device_manager; + auto qu = device_manager.GetQueue(DeviceOrd::SyclDefault()); + + TreeEvaluator tree_evaluator(qu, param, n_columns); + auto split_evaluator = tree_evaluator.GetEvaluator(); + { + // Test ThresholdL1 + const GradientSumT alpha = 0.5; + { + const GradientSumT val = 0.0; + const auto trh = split_evaluator.ThresholdL1(val, alpha); + ASSERT_EQ(trh, 0.0); + } + + { + const GradientSumT val = 1.0; + const auto trh = split_evaluator.ThresholdL1(val, alpha); + ASSERT_EQ(trh, val - alpha); + } + + { + const GradientSumT val = -1.0; + const auto trh = split_evaluator.ThresholdL1(val, alpha); + ASSERT_EQ(trh, val + alpha); + } + } + + { + constexpr float eps = 1e-8; + tree_evaluator.AddSplit(0, 1, 2, 0, 0.3, 0.7); + + GradStats left(0.1, 0.2); + GradStats right(0.3, 0.4); + bst_node_t nidx = 0; + bst_feature_t fidx = 0; + + GradientSumT wleft = split_evaluator.CalcWeight(nidx, left); + // wleft = -grad/hess = -0.1/0.2 + EXPECT_NEAR(wleft, -0.5, eps); + GradientSumT wright = split_evaluator.CalcWeight(nidx, right); + // wright = -grad/hess = -0.3/0.4 + EXPECT_NEAR(wright, -0.75, eps); + + GradientSumT gweight_left = split_evaluator.CalcGainGivenWeight(nidx, left, wleft); + // gweight_left = left.grad**2 / left.hess = 0.1*0.1/0.2 = 0.05 + EXPECT_NEAR(gweight_left, 0.05, eps); + // gweight_left = right.grad**2 / right.hess = 0.3*0.3/0.4 = 0.225 + GradientSumT gweight_right = split_evaluator.CalcGainGivenWeight(nidx, right, wright); + EXPECT_NEAR(gweight_right, 0.225, eps); + + GradientSumT split_gain = split_evaluator.CalcSplitGain(nidx, fidx, left, right); + if (!tree_evaluator.HasConstraint()) { + EXPECT_NEAR(split_gain, gweight_left + gweight_right, eps); + } else { + // Parameters are chosen to have -inf here + ASSERT_EQ(split_gain, -std::numeric_limits::infinity()); + } + } +} + +TEST(SyclSplitEvaluator, BasicTest) { + BasicTestSplitEvaluator("( 0, 0)", false); + BasicTestSplitEvaluator("( 1, 0)", true); + BasicTestSplitEvaluator("( 0, 1)", true); + BasicTestSplitEvaluator("(-1, 0)", true); + BasicTestSplitEvaluator("( 0, -1)", true); + BasicTestSplitEvaluator("( 1, 1)", true); + BasicTestSplitEvaluator("(-1, -1)", true); + BasicTestSplitEvaluator("( 1, -1)", true); + BasicTestSplitEvaluator("(-1, 1)", true); +} + +TEST(SyclSplitEvaluator, TestMath) { + // Without constraints + TestSplitEvaluator("( 0, 0)"); + // With constraints + TestSplitEvaluator("( 1, 0)"); +} + +} // namespace xgboost::sycl::tree