SYCL. Fixes for the upcoming v3.1

include/xgboost/linalg.h

@@ -591,13 +591,13 @@ auto MakeTensorView(Context const *ctx, Order order, common::Span<T, ext> data,
 
 template <typename T, typename... S>
 auto MakeTensorView(Context const *ctx, HostDeviceVector<T> *data, S &&...shape) {
-  auto span = ctx->IsCUDA() ? data->DeviceSpan() : data->HostSpan();
+  auto span = ctx->IsCPU() ? data->HostSpan() : data->DeviceSpan();
   return MakeTensorView(ctx->Device(), span, std::forward<S>(shape)...);
 }
 
 template <typename T, typename... S>
 auto MakeTensorView(Context const *ctx, HostDeviceVector<T> const *data, S &&...shape) {
-  auto span = ctx->IsCUDA() ? data->ConstDeviceSpan() : data->ConstHostSpan();
+  auto span = ctx->IsCPU() ? data->ConstHostSpan() : data->ConstDeviceSpan();
   return MakeTensorView(ctx->Device(), span, std::forward<S>(shape)...);
 }
 
@@ -647,13 +647,13 @@ auto MakeVec(T *ptr, size_t s, DeviceOrd device = DeviceOrd::CPU()) {
 
 template <typename T>
 auto MakeVec(HostDeviceVector<T> *data) {
-  return MakeVec(data->Device().IsCUDA() ? data->DevicePointer() : data->HostPointer(),
+  return MakeVec(data->Device().IsCPU() ? data->HostPointer() : data->DevicePointer(),
                  data->Size(), data->Device());
 }
 
 template <typename T>
 auto MakeVec(HostDeviceVector<T> const *data) {
-  return MakeVec(data->Device().IsCUDA() ? data->ConstDevicePointer() : data->ConstHostPointer(),
+  return MakeVec(data->Device().IsCPU() ? data->ConstHostPointer() : data->ConstDevicePointer(),
                  data->Size(), data->Device());
 }
 
@@ -759,7 +759,7 @@ class Tensor {
     for (auto i = D; i < kDim; ++i) {
       shape_[i] = 1;
     }
-    if (device.IsCUDA()) {
+    if (!device.IsCPU()) {
       data_.SetDevice(device);
       data_.ConstDevicePointer();  // Pull to device;
     }
@@ -788,11 +788,11 @@ class Tensor {
       shape_[i] = 1;
     }
     auto size = detail::CalcSize(shape_);
-    if (device.IsCUDA()) {
+    if (!device.IsCPU()) {
       data_.SetDevice(device);
     }
     data_.Resize(size);
-    if (device.IsCUDA()) {
+    if (!device.IsCPU()) {
       data_.DevicePointer();  // Pull to device
     }
   }

plugin/sycl/common/host_device_vector.cc

@@ -16,6 +16,7 @@
 
 #include "../device_manager.h"
 #include "../data.h"
+#include "../predictor/node.h"
 
 namespace xgboost {
 template <typename T>
@@ -405,6 +406,7 @@ template class HostDeviceVector<FeatureType>;
 template class HostDeviceVector<Entry>;
 template class HostDeviceVector<bst_idx_t>;
 template class HostDeviceVector<uint32_t>;  // bst_feature_t
+template class HostDeviceVector<sycl::predictor::Node>;
 
 }  // namespace xgboost
 

plugin/sycl/common/linalg_op.cc

@@ -0,0 +1,52 @@
+/**
+ * Copyright 2021-2025, XGBoost Contributors
+ * \file linalg_op.h
+ */
+
+#include "../data.h"
+#include "../device_manager.h"
+
+#include "../../../src/common/optional_weight.h"  // for OptionalWeights
+#include "xgboost/context.h"  // for Context
+
+#include <sycl/sycl.hpp>
+
+namespace xgboost::sycl::linalg {
+
+void SmallHistogram(Context const* ctx, xgboost::linalg::MatrixView<float const> indices,
+                    xgboost::common::OptionalWeights const& weights,
+                    xgboost::linalg::VectorView<float> bins) {
+  sycl::DeviceManager device_manager;
+  auto* qu = device_manager.GetQueue(ctx->Device());
+
+  qu->submit([&](::sycl::handler& cgh) {
+    cgh.parallel_for<>(::sycl::range<1>(indices.Size()),
+                       [=](::sycl::id<1> pid) {
+      const size_t i = pid[0];
+      auto y = indices(i);
+      auto w = weights[i];
+      AtomicRef<float> bin_val(const_cast<float&>(bins(static_cast<std::size_t>(y))));
+      bin_val += w;
+    });
+  }).wait();
+}
+
+void VecScaMul(Context const* ctx, xgboost::linalg::VectorView<float> x, double mul) {
+  sycl::DeviceManager device_manager;
+  auto* qu = device_manager.GetQueue(ctx->Device());
+
+  qu->submit([&](::sycl::handler& cgh) {
+    cgh.parallel_for<>(::sycl::range<1>(x.Size()),
+                       [=](::sycl::id<1> pid) {
+      const size_t i = pid[0];
+      const_cast<float&>(x(i)) *= mul;
+    });
+  }).wait();
+}
+}  // namespace xgboost::sycl::linalg
+
+namespace xgboost::linalg::sycl_impl {
+void VecScaMul(Context const* ctx, xgboost::linalg::VectorView<float> x, double mul) {
+  xgboost::sycl::linalg::VecScaMul(ctx, x, mul);
+}
+}  // namespace xgboost::linalg::sycl_impl

plugin/sycl/common/optional_weight.cc

@@ -0,0 +1,31 @@
+/*!
+ * Copyright by Contributors 2017-2025
+ */
+#include <sycl/sycl.hpp>
+
+#include "../../../src/common/optional_weight.h"
+
+#include "../device_manager.h"
+
+namespace xgboost::common::sycl_impl {
+double SumOptionalWeights(Context const* ctx, OptionalWeights const& weights) {
+  sycl::DeviceManager device_manager;
+  auto* qu = device_manager.GetQueue(ctx->Device());
+
+  const auto* data = weights.Data();
+  double result = 0;
+  {
+    ::sycl::buffer<double> buff(&result, 1);
+    qu->submit([&](::sycl::handler& cgh) {
+      auto reduction = ::sycl::reduction(buff, cgh, ::sycl::plus<>());
+      cgh.parallel_for<>(::sycl::range<1>(weights.Size()), reduction,
+                        [=](::sycl::id<1> pid, auto& sum) {
+        size_t i = pid[0];
+        sum += data[i];
+      });
+    }).wait_and_throw();
+  }
+
+  return result;
+}
+}  // namespace xgboost::common::sycl_impl

plugin/sycl/device_properties.h

@@ -47,6 +47,9 @@ class DeviceProperties {
   size_t l2_size = 0;
   float l2_size_per_eu = 0;
 
+  DeviceProperties():
+    is_gpu(false) {}
+
   explicit DeviceProperties(const ::sycl::device& device):
     is_gpu(device.is_gpu()),
     usm_host_allocations(device.has(::sycl::aspect::usm_host_allocations)),

plugin/sycl/predictor/node.h

@@ -0,0 +1,69 @@
+/*!
+ * Copyright by Contributors 2017-2025
+ * \file node.h
+ */
+#ifndef PLUGIN_SYCL_PREDICTOR_NODE_H_
+#define PLUGIN_SYCL_PREDICTOR_NODE_H_
+
+#include "../../src/gbm/gbtree_model.h"
+
+namespace xgboost {
+namespace sycl {
+namespace predictor {
+
+union NodeValue {
+  float leaf_weight;
+  float fvalue;
+};
+
+class Node {
+  int fidx;
+  int left_child_idx;
+  int right_child_idx;
+  NodeValue val;
+
+ public:
+  Node() = default;
+
+  explicit Node(const RegTree::Node& n) {
+    left_child_idx = n.LeftChild();
+    right_child_idx = n.RightChild();
+    fidx = n.SplitIndex();
+    if (n.DefaultLeft()) {
+      fidx |= (1U << 31);
+    }
+
+    if (n.IsLeaf()) {
+      val.leaf_weight = n.LeafValue();
+    } else {
+      val.fvalue = n.SplitCond();
+    }
+  }
+
+  int LeftChildIdx() const {return left_child_idx; }
+
+  int RightChildIdx() const {return right_child_idx; }
+
+  bool IsLeaf() const { return left_child_idx == -1; }
+
+  int GetFidx() const { return fidx & ((1U << 31) - 1U); }
+
+  bool MissingLeft() const { return (fidx >> 31) != 0; }
+
+  int MissingIdx() const {
+    if (MissingLeft()) {
+      return left_child_idx;
+    } else {
+      return right_child_idx;
+    }
+  }
+
+  float GetFvalue() const { return val.fvalue; }
+
+  float GetWeight() const { return val.leaf_weight; }
+};
+
+}  // namespace predictor
+}  // namespace sycl
+}  // namespace xgboost
+#endif  // PLUGIN_SYCL_PREDICTOR_NODE_H_