initial

Author: Dmitry Razdoburdin

plugin/sycl/common/optional_weight.cc

@@ -8,14 +8,13 @@
 #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());
 
+template <typename T>
+T ElementWiseSum(::sycl::queue* qu, OptionalWeights const& weights) {
   const auto* data = weights.Data();
-  double result = 0;
+  T result = 0;
   {
-    ::sycl::buffer<double> buff(&result, 1);
+    ::sycl::buffer<T> 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,
@@ -28,4 +27,16 @@ double SumOptionalWeights(Context const* ctx, OptionalWeights const& weights) {
 
   return result;
 }
+
+double SumOptionalWeights(Context const* ctx, OptionalWeights const& weights) {
+  sycl::DeviceManager device_manager;
+  auto* qu = device_manager.GetQueue(ctx->Device());
+
+  bool has_fp64_support = qu->get_device().has(::sycl::aspect::fp64);
+  if (has_fp64_support) {
+    return ElementWiseSum<double>(qu, weights);
+  } else {
+    return ElementWiseSum<float>(qu, weights);
+  }
+}
 }  // namespace xgboost::common::sycl_impl

plugin/sycl/common/stats.cc

@@ -0,0 +1,26 @@
+/*!
+ * Copyright by Contributors 2017-2025
+ */
+#include <sycl/sycl.hpp>
+
+#include "../../../src/common/stats.h"
+
+#include "../device_manager.h"
+
+namespace xgboost::common::sycl_impl {
+    void Mean(Context const* ctx, linalg::VectorView<float const> v, linalg::VectorView<float> out) {
+        sycl::DeviceManager device_manager;
+        auto* qu = device_manager.GetQueue(ctx->Device());
+
+        qu->submit([&](::sycl::handler& cgh) {
+            auto reduction =
+                ::sycl::reduction(&(out(0)), 0.0f, ::sycl::plus<float>(),
+                                  ::sycl::property::reduction::initialize_to_identity());
+            cgh.parallel_for<>(::sycl::range<1>(v.Size()), reduction,
+                               [=](::sycl::id<1> pid, auto& sum) {
+                size_t i = pid[0];
+                sum += v(i);
+            });
+        }).wait_and_throw();
+    }
+}

plugin/sycl/context_helper.cc

@@ -14,7 +14,10 @@ namespace sycl {
 
 DeviceOrd DeviceFP64(const DeviceOrd& device) {
   DeviceManager device_manager;
-  bool support_fp64 = device_manager.GetQueue(device)->get_device().has(::sycl::aspect::fp64);
+  bool support_fp64 = true;
+  if (device.IsSycl()) {
+    support_fp64 = device_manager.GetQueue(device)->get_device().has(::sycl::aspect::fp64);
+  }
   if (support_fp64) {
     return device;
   } else {

plugin/sycl/device_manager.cc

@@ -21,18 +21,25 @@ ::sycl::queue* DeviceManager::GetQueue(const DeviceOrd& device_spec) const {
                                     (collective::IsDistributed());
     DeviceRegister& device_register = GetDevicesRegister();
     if (not_use_default_selector) {
-        const int device_idx =
-            collective::IsDistributed() ? collective::GetRank() : device_spec.ordinal;
         if (device_spec.IsSyclDefault()) {
             auto& devices = device_register.devices;
+            const int device_idx = collective::IsDistributed()
+                                   ? collective::GetRank() % devices.size()
+                                   : device_spec.ordinal;
             CHECK_LT(device_idx, devices.size());
             queue_idx = device_idx;
         } else if (device_spec.IsSyclCPU()) {
             auto& cpu_devices_idxes = device_register.cpu_devices_idxes;
+            const int device_idx = collective::IsDistributed()
+                                   ? collective::GetRank() % cpu_devices_idxes.size()
+                                   : device_spec.ordinal;
             CHECK_LT(device_idx, cpu_devices_idxes.size());
             queue_idx = cpu_devices_idxes[device_idx];
         } else if (device_spec.IsSyclGPU()) {
             auto& gpu_devices_idxes = device_register.gpu_devices_idxes;
+            const int device_idx = collective::IsDistributed()
+                                   ? collective::GetRank() % gpu_devices_idxes.size()
+                                   : device_spec.ordinal;
             CHECK_LT(device_idx, gpu_devices_idxes.size());
             queue_idx = gpu_devices_idxes[device_idx];
         } else {

src/common/linalg_op.h

@@ -159,9 +159,13 @@ void ElementWiseKernel(Context const* ctx, TensorView<T, D> t, Fn&& fn) {
 #elif defined(SYCL_LANGUAGE_VERSION)
 template <typename T, std::int32_t D, typename Fn, auto _tag = detail::SysTag()>
 void ElementWiseKernel(Context const* ctx, TensorView<T, D> t, Fn&& fn) {
-  ctx->DispatchDevice([&] { cpu_impl::ElementWiseKernel(t, ctx->Threads(), std::forward<Fn>(fn)); },
-                      [&] { LOG(FATAL) << "Invalid TU"; },
-                      [&] { ::xgboost::sycl::linalg::ElementWiseKernel(t, std::forward<Fn>(fn)); });
+  if (t.Device().IsCPU()) {
+    cpu_impl::ElementWiseKernel(t, ctx->Threads(), std::forward<Fn>(fn));
+  } else {
+    ctx->DispatchDevice([&] { cpu_impl::ElementWiseKernel(t, ctx->Threads(), std::forward<Fn>(fn)); },
+                        [&] { LOG(FATAL) << "Invalid TU"; },
+                        [&] { ::xgboost::sycl::linalg::ElementWiseKernel(t, std::forward<Fn>(fn)); });
+  }
 }
 #else
 template <typename T, std::int32_t D, typename Fn, auto _tag = detail::SysTag()>

src/common/stats.cc

@@ -51,6 +51,8 @@ void Mean(Context const* ctx, linalg::VectorView<float const> v, linalg::Vector<
 
   if (ctx->IsCUDA()) {
     cuda_impl::Mean(ctx, v, out->View(ctx->Device()));
+  } else if (ctx->IsSycl()) {
+    sycl_impl::Mean(ctx, v, out->View(ctx->Device()));
   } else {
     auto h_v = v;
     float n = v.Size();

src/common/stats.h

@@ -14,7 +14,7 @@
 #include "xgboost/linalg.h"   // TensorView,VectorView
 #include "xgboost/logging.h"  // CHECK_GE
 
-#if !defined(XGBOOST_USE_CUDA)
+#if !defined(XGBOOST_USE_CUDA) && !defined(XGBOOST_USE_SYCL)
 #include "common.h"  // AssertGPUSupport
 #endif
 
@@ -140,6 +140,17 @@ inline void WeightedSampleMean(Context const*, bool, linalg::MatrixView<float co
 #endif  // !defined(XGBOOST_USE_CUDA)
 }  // namespace cuda_impl
 
+namespace sycl_impl {
+void Mean(Context const* ctx, linalg::VectorView<float const> v, linalg::VectorView<float> out);
+
+#if !defined(XGBOOST_USE_SYCL)
+inline void Mean(Context const*, linalg::VectorView<float const>, linalg::VectorView<float>) {
+  common::AssertGPUSupport();
+}
+
+#endif  // !defined(XGBOOST_USE_SYCL)
+}  // namespace sycl_impl
+
 /**
  * @brief Calculate medians for each column of the input matrix.
  */

src/objective/multiclass_obj.cu

@@ -106,10 +106,13 @@ class SoftmaxMultiClassObj : public ObjFunction {
           << "Number of weights should be equal to number of data points.";
     }
     info.weights_.SetDevice(device);
-    auto weights = common::MakeOptionalWeights(this->ctx_->Device(), info.weights_);
+    auto weights = common::MakeOptionalWeights(device, info.weights_);
 
     preds.SetDevice(device);
-    auto predt = linalg::MakeTensorView(this->ctx_, &preds, n_samples, n_classes);
+    Context cpu_context = Context();
+    auto predt = linalg::MakeTensorView(
+      device == ctx_->Device() ? this->ctx_ : &cpu_context,
+      &preds, n_samples, n_classes);
     CHECK_EQ(labels.Shape(1), 1);
     auto y1d = labels.Slice(linalg::All(), 0);
     CHECK_EQ(y1d.Shape(0), info.num_row_);