[SYCL] Add split evaluation (#10119)

---------

Co-authored-by: Dmitry Razdoburdin <>

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 <cmath>
+#include <cstring>
+#include <limits>
+#include <string>
+#include <vector>
+
+
+#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 <CL/sycl.hpp>
+
+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 <typename GradType>
+using GradStats = xgboost::detail::GradientPairInternal<GradType>;
+
+}  // namespace tree
+}  // namespace sycl
+}  // namespace xgboost
+#endif  // PLUGIN_SYCL_TREE_PARAM_H_

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 <dmlc/registry.h>
+#include <xgboost/base.h>
+#include <utility>
+#include <vector>
+#include <limits>
+
+#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 <CL/sycl.hpp>
+
+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<typename GradType>
+class TreeEvaluator {
+  // hist and exact use parent id to calculate constraints.
+  static constexpr bst_node_t kRootParentId =
+      (-1 & static_cast<bst_node_t>((1U << 31) - 1));
+
+  USMVector<GradType> lower_bounds_;
+  USMVector<GradType> upper_bounds_;
+  USMVector<int> 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<GradType>::lowest());
+      upper_bounds_.Resize(&qu_, p.MaxNodes(), std::numeric_limits<GradType>::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<GradType>& left,
+                        const GradStats<GradType>& right) const {
+      const GradType negative_infinity = -std::numeric_limits<GradType>::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<GradType>& 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<GradType>& 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<GradType>& 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_

tests/cpp/plugin/test_sycl_split_evaluator.cc

@@ -0,0 +1,134 @@
+/**
+ * Copyright 2020-2024 by XGBoost contributors
+ */
+#include <gtest/gtest.h>
+#include <vector>
+
+#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<typename GradientSumT>
+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<GradientSumT> 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<typename GradientSumT>
+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<GradientSumT> 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<GradientSumT> left(0.1, 0.2);
+    GradStats<GradientSumT> 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<GradientSumT>::infinity());
+    }
+  }
+}
+
+TEST(SyclSplitEvaluator, BasicTest) {
+  BasicTestSplitEvaluator<float>("( 0,  0)", false);
+  BasicTestSplitEvaluator<float>("( 1,  0)", true);
+  BasicTestSplitEvaluator<float>("( 0,  1)", true);
+  BasicTestSplitEvaluator<float>("(-1,  0)", true);
+  BasicTestSplitEvaluator<float>("( 0, -1)", true);
+  BasicTestSplitEvaluator<float>("( 1,  1)", true);
+  BasicTestSplitEvaluator<float>("(-1, -1)", true);
+  BasicTestSplitEvaluator<float>("( 1, -1)", true);
+  BasicTestSplitEvaluator<float>("(-1,  1)", true);
+}
+
+TEST(SyclSplitEvaluator, TestMath) {
+  // Without constraints
+  TestSplitEvaluator<float>("( 0,  0)");
+  // With constraints
+  TestSplitEvaluator<float>("( 1,  0)");
+}
+
+}  // namespace xgboost::sycl::tree