SYCL. Add sycl-specific implementation for absoluteerror

plugin/sycl/objective/regression_obj.cc

@@ -26,6 +26,9 @@
 #include "../../../src/objective/regression_loss.h"
 #pragma GCC diagnostic pop
 #include "../../../src/objective/regression_param.h"
+#include "../../../src/objective/init_estimation.h"
+#include "../../../src/objective/adaptive.h"
+#include "../../../src/common/optional_weight.h"  // OptionalWeights
 
 #include "../common/linalg_op.h"
 
@@ -192,6 +195,86 @@ XGBOOST_REGISTER_OBJECTIVE(LogisticRaw,
           "before logistic transformation with SYCL backend.")
 .set_body([]() { return new RegLossObj<xgboost::obj::LogisticRaw>(); });
 
+class MeanAbsoluteError : public ObjFunction {
+ public:
+  void Configure(Args const&) override {}
+
+  ObjInfo Task() const override {
+    return {ObjInfo::kRegression, true, true};
+  }
+
+  bst_target_t Targets(MetaInfo const& info) const override {
+    return std::max(static_cast<std::size_t>(1), info.labels.Shape(1));
+  }
+
+  void GetGradient(HostDeviceVector<float> const& preds, const MetaInfo& info,
+                   std::int32_t, xgboost::linalg::Matrix<GradientPair>* out_gpair) override {
+    if (qu_ == nullptr) {
+      qu_ = device_manager.GetQueue(ctx_->Device());
+    }
+
+    size_t const ndata = preds.Size();
+    auto const n_targets = this->Targets(info);
+
+    xgboost::obj::CheckInitInputs(info);
+    CHECK_EQ(info.labels.Size(), preds.Size()) << "Invalid shape of labels.";
+    const bst_float* label_ptr = info.labels.Data()->ConstDevicePointer();
+
+    out_gpair->SetDevice(ctx_->Device());
+    out_gpair->Reshape(info.num_row_, this->Targets(info));
+    GradientPair* out_gpair_ptr  = out_gpair->Data()->DevicePointer();
+
+    preds.SetDevice(ctx_->Device());
+    const bst_float* preds_ptr = preds.ConstDevicePointer();
+    auto predt = xgboost::linalg::MakeTensorView(ctx_, &preds, info.num_row_, this->Targets(info));
+    info.weights_.SetDevice(ctx_->Device());
+    common::OptionalWeights weight{ctx_->IsCPU() ? info.weights_.ConstHostSpan()
+                                                 : info.weights_.ConstDeviceSpan()};
+
+    qu_->submit([&](::sycl::handler& cgh) {
+      cgh.parallel_for<>(::sycl::range<1>(ndata), [=](::sycl::id<1> pid) {
+        int idx = pid[0];
+        auto sign = [](auto x) {
+          return (x > static_cast<decltype(x)>(0)) - (x < static_cast<decltype(x)>(0));
+        };
+        const bst_float pred = preds_ptr[idx];
+        const bst_float label = label_ptr[idx];
+
+        bst_float hess = weight[idx/n_targets];
+        bst_float grad = sign(pred - label) * hess;
+        out_gpair_ptr[idx] = GradientPair{grad, hess};
+      });
+    });
+    qu_->wait_and_throw();
+  }
+
+  void UpdateTreeLeaf(HostDeviceVector<bst_node_t> const& position, MetaInfo const& info,
+                      float learning_rate, HostDeviceVector<float> const& prediction,
+                      std::int32_t group_idx, RegTree* p_tree) const override {
+    ::xgboost::obj::UpdateTreeLeaf(ctx_, position, group_idx, info, learning_rate, prediction, 0.5,
+                                   p_tree);
+  }
+
+  const char* DefaultEvalMetric() const override { return "mae"; }
+
+  void SaveConfig(Json* p_out) const override {
+    auto& out = *p_out;
+    out["name"] = String("reg:absoluteerror");
+  }
+
+  void LoadConfig(Json const& in) override {
+    CHECK_EQ(StringView{get<String const>(in["name"])}, StringView{"reg:absoluteerror"});
+  }
+
+ protected:
+  sycl::DeviceManager device_manager;
+  mutable ::sycl::queue* qu_ = nullptr;
+};
+
+XGBOOST_REGISTER_OBJECTIVE(MeanAbsoluteError, "reg:absoluteerror_sycl")
+    .describe("Mean absoluate error.")
+    .set_body([]() { return new MeanAbsoluteError(); });
+
 }  // namespace obj
 }  // namespace sycl
 }  // namespace xgboost

src/common/algorithm.h

@@ -75,7 +75,7 @@ void Sort(Context const *ctx, Iter begin, Iter end, Comp comp) {
 template <typename Idx, typename Iter, typename V = typename std::iterator_traits<Iter>::value_type,
           typename Comp = std::less<V>>
 std::vector<Idx> ArgSort(Context const *ctx, Iter begin, Iter end, Comp comp = std::less<V>{}) {
-  CHECK(ctx->IsCPU());
+  CHECK(!ctx->IsCUDA());
   auto n = std::distance(begin, end);
   std::vector<Idx> result(n);
   Iota(ctx, result.begin(), result.end(), 0);

tests/cpp/objective/test_regression_obj.cc

@@ -122,4 +122,103 @@ void TestsLogisticRawGPair(const Context* ctx) {
                    {0.25f, 0.24f, 0.20f, 0.19f, 0.25f,  0.24f,  0.20f,  0.19f});
 }
 
+void TestAbsoluteError(const Context* ctx) {
+  std::unique_ptr<ObjFunction> obj{ObjFunction::Create("reg:absoluteerror", ctx)};
+  obj->Configure({});
+  CheckConfigReload(obj, "reg:absoluteerror");
+
+  MetaInfo info;
+  std::vector<float> labels{0.f, 3.f, 2.f, 5.f, 4.f, 7.f};
+  info.labels.Reshape(6, 1);
+  info.labels.Data()->HostVector() = labels;
+  info.num_row_ = labels.size();
+  HostDeviceVector<float> predt{1.f, 2.f, 3.f, 4.f, 5.f, 6.f};
+  info.weights_.HostVector() = {1.f, 1.f, 1.f, 1.f, 1.f, 1.f};
+
+  CheckObjFunction(obj, predt.HostVector(), labels, info.weights_.HostVector(),
+                   {1.f, -1.f, 1.f, -1.f, 1.f, -1.f}, info.weights_.HostVector());
+
+  RegTree tree;
+  tree.ExpandNode(0, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
+
+  HostDeviceVector<bst_node_t> position(labels.size(), 0);
+  auto& h_position = position.HostVector();
+  for (size_t i = 0; i < labels.size(); ++i) {
+    if (i < labels.size() / 2) {
+      h_position[i] = 1;  // left
+    } else {
+      h_position[i] = 2;  // right
+    }
+  }
+
+  auto& h_predt = predt.HostVector();
+  for (size_t i = 0; i < h_predt.size(); ++i) {
+    h_predt[i] = labels[i] + i;
+  }
+
+  tree::TrainParam param;
+  param.Init(Args{});
+  auto lr = param.learning_rate;
+
+  obj->UpdateTreeLeaf(position, info, param.learning_rate, predt, 0, &tree);
+  ASSERT_EQ(tree[1].LeafValue(), -1.0f * lr);
+  ASSERT_EQ(tree[2].LeafValue(), -4.0f * lr);
+}
+
+void TestAbsoluteErrorLeaf(const Context* ctx) {
+  bst_target_t constexpr kTargets = 3, kRows = 16;
+  std::unique_ptr<ObjFunction> obj{ObjFunction::Create("reg:absoluteerror", ctx)};
+  obj->Configure({});
+
+  MetaInfo info;
+  info.num_row_ = kRows;
+  info.labels.Reshape(16, kTargets);
+  HostDeviceVector<float> predt(info.labels.Size());
+
+  for (bst_target_t t{0}; t < kTargets; ++t) {
+    auto h_labels = info.labels.HostView().Slice(linalg::All(), t);
+    std::iota(linalg::begin(h_labels), linalg::end(h_labels), 0);
+
+    auto h_predt =
+        linalg::MakeTensorView(ctx, predt.HostSpan(), kRows, kTargets).Slice(linalg::All(), t);
+    for (size_t i = 0; i < h_predt.Size(); ++i) {
+      h_predt(i) = h_labels(i) + i;
+    }
+
+    HostDeviceVector<bst_node_t> position(h_labels.Size(), 0);
+    auto& h_position = position.HostVector();
+    for (int32_t i = 0; i < 3; ++i) {
+      h_position[i] = ~i;  // negation for sampled nodes.
+    }
+    for (size_t i = 3; i < 8; ++i) {
+      h_position[i] = 3;
+    }
+    // empty leaf for node 4
+    for (size_t i = 8; i < 13; ++i) {
+      h_position[i] = 5;
+    }
+    for (size_t i = 13; i < h_labels.Size(); ++i) {
+      h_position[i] = 6;
+    }
+
+    RegTree tree;
+    tree.ExpandNode(0, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
+    tree.ExpandNode(1, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
+    tree.ExpandNode(2, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
+    ASSERT_EQ(tree.GetNumLeaves(), 4);
+
+    auto empty_leaf = tree[4].LeafValue();
+
+    tree::TrainParam param;
+    param.Init(Args{});
+    auto lr = param.learning_rate;
+
+    obj->UpdateTreeLeaf(position, info, lr, predt, t, &tree);
+    ASSERT_EQ(tree[3].LeafValue(), -5.0f * lr);
+    ASSERT_EQ(tree[4].LeafValue(), empty_leaf * lr);
+    ASSERT_EQ(tree[5].LeafValue(), -10.0f * lr);
+    ASSERT_EQ(tree[6].LeafValue(), -14.0f * lr);
+  }
+}
+
 }  // namespace xgboost

tests/cpp/objective/test_regression_obj.h

@@ -18,6 +18,10 @@ void TestLogisticRegressionBasic(const Context* ctx);
 
 void TestsLogisticRawGPair(const Context* ctx);
 
+void TestAbsoluteError(const Context* ctx);
+
+void TestAbsoluteErrorLeaf(const Context* ctx);
+
 }  // namespace xgboost
 
 #endif  // XGBOOST_TEST_REGRESSION_OBJ_H_

tests/cpp/objective/test_regression_obj_cpu.cc

@@ -290,103 +290,12 @@ TEST(Objective, CoxRegressionGPair) {
 
 TEST(Objective, DeclareUnifiedTest(AbsoluteError)) {
   Context ctx = MakeCUDACtx(GPUIDX);
-  std::unique_ptr<ObjFunction> obj{ObjFunction::Create("reg:absoluteerror", &ctx)};
-  obj->Configure({});
-  CheckConfigReload(obj, "reg:absoluteerror");
-
-  MetaInfo info;
-  std::vector<float> labels{0.f, 3.f, 2.f, 5.f, 4.f, 7.f};
-  info.labels.Reshape(6, 1);
-  info.labels.Data()->HostVector() = labels;
-  info.num_row_ = labels.size();
-  HostDeviceVector<float> predt{1.f, 2.f, 3.f, 4.f, 5.f, 6.f};
-  info.weights_.HostVector() = {1.f, 1.f, 1.f, 1.f, 1.f, 1.f};
-
-  CheckObjFunction(obj, predt.HostVector(), labels, info.weights_.HostVector(),
-                   {1.f, -1.f, 1.f, -1.f, 1.f, -1.f}, info.weights_.HostVector());
-
-  RegTree tree;
-  tree.ExpandNode(0, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
-
-  HostDeviceVector<bst_node_t> position(labels.size(), 0);
-  auto& h_position = position.HostVector();
-  for (size_t i = 0; i < labels.size(); ++i) {
-    if (i < labels.size() / 2) {
-      h_position[i] = 1;  // left
-    } else {
-      h_position[i] = 2;  // right
-    }
-  }
-
-  auto& h_predt = predt.HostVector();
-  for (size_t i = 0; i < h_predt.size(); ++i) {
-    h_predt[i] = labels[i] + i;
-  }
-
-  tree::TrainParam param;
-  param.Init(Args{});
-  auto lr = param.learning_rate;
-
-  obj->UpdateTreeLeaf(position, info, param.learning_rate, predt, 0, &tree);
-  ASSERT_EQ(tree[1].LeafValue(), -1.0f * lr);
-  ASSERT_EQ(tree[2].LeafValue(), -4.0f * lr);
+  TestAbsoluteError(&ctx);
 }
 
 TEST(Objective, DeclareUnifiedTest(AbsoluteErrorLeaf)) {
   Context ctx = MakeCUDACtx(GPUIDX);
-  bst_target_t constexpr kTargets = 3, kRows = 16;
-  std::unique_ptr<ObjFunction> obj{ObjFunction::Create("reg:absoluteerror", &ctx)};
-  obj->Configure({});
-
-  MetaInfo info;
-  info.num_row_ = kRows;
-  info.labels.Reshape(16, kTargets);
-  HostDeviceVector<float> predt(info.labels.Size());
-
-  for (bst_target_t t{0}; t < kTargets; ++t) {
-    auto h_labels = info.labels.HostView().Slice(linalg::All(), t);
-    std::iota(linalg::begin(h_labels), linalg::end(h_labels), 0);
-
-    auto h_predt =
-        linalg::MakeTensorView(&ctx, predt.HostSpan(), kRows, kTargets).Slice(linalg::All(), t);
-    for (size_t i = 0; i < h_predt.Size(); ++i) {
-      h_predt(i) = h_labels(i) + i;
-    }
-
-    HostDeviceVector<bst_node_t> position(h_labels.Size(), 0);
-    auto& h_position = position.HostVector();
-    for (int32_t i = 0; i < 3; ++i) {
-      h_position[i] = ~i;  // negation for sampled nodes.
-    }
-    for (size_t i = 3; i < 8; ++i) {
-      h_position[i] = 3;
-    }
-    // empty leaf for node 4
-    for (size_t i = 8; i < 13; ++i) {
-      h_position[i] = 5;
-    }
-    for (size_t i = 13; i < h_labels.Size(); ++i) {
-      h_position[i] = 6;
-    }
-
-    RegTree tree;
-    tree.ExpandNode(0, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
-    tree.ExpandNode(1, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
-    tree.ExpandNode(2, /*split_index=*/1, 2, true, 0.0f, 2.f, 3.f, 4.f, 2.f, 1.f, 1.f);
-    ASSERT_EQ(tree.GetNumLeaves(), 4);
-
-    auto empty_leaf = tree[4].LeafValue();
-
-    tree::TrainParam param;
-    param.Init(Args{});
-    auto lr = param.learning_rate;
-
-    obj->UpdateTreeLeaf(position, info, lr, predt, t, &tree);
-    ASSERT_EQ(tree[3].LeafValue(), -5.0f * lr);
-    ASSERT_EQ(tree[4].LeafValue(), empty_leaf * lr);
-    ASSERT_EQ(tree[5].LeafValue(), -10.0f * lr);
-    ASSERT_EQ(tree[6].LeafValue(), -14.0f * lr);
-  }
+  TestAbsoluteErrorLeaf(&ctx);
 }
 
 TEST(Adaptive, DeclareUnifiedTest(MissingLeaf)) {

tests/cpp/plugin/test_sycl_regression_obj.cc

@@ -45,11 +45,23 @@ TEST(SyclObjective, LogisticRawGPair) {
   TestsLogisticRawGPair(&ctx);
 }
 
+TEST(SyclObjective, AbsoluteError) {
+  Context ctx;
+  ctx.UpdateAllowUnknown(Args{{"device", "sycl"}});
+  TestAbsoluteError(&ctx);
+}
+
+TEST(SyclObjective, AbsoluteErrorLeaf) {
+  Context ctx;
+  ctx.UpdateAllowUnknown(Args{{"device", "sycl"}});
+  TestAbsoluteErrorLeaf(&ctx);
+}
+
 TEST(SyclObjective, CPUvsSycl) {
   Context ctx_sycl;
   ctx_sycl.UpdateAllowUnknown(Args{{"device", "sycl"}});
   ObjFunction * obj_sycl =
-      ObjFunction::Create("reg:squarederror_sycl", &ctx_sycl);
+      ObjFunction::Create("reg:squarederror", &ctx_sycl);
 
   Context ctx_cpu;
   ctx_cpu.UpdateAllowUnknown(Args{{"device", "cpu"}});