Add support inference on SYCL devices (#9800)

---------

Co-authored-by: Dmitry Razdoburdin <>
Co-authored-by: Nikolay Petrov <nikolay.a.petrov@intel.com>
Co-authored-by: Alexandra <alexandra.epanchinzeva@intel.com>

plugin/sycl/README.md

@@ -0,0 +1,40 @@
+<!--
+******************************************************************************
+* Copyright by Contributors 2017-2023
+*******************************************************************************/-->
+
+# SYCL-based Algorithm for Tree Construction
+This plugin adds support of SYCL programming model for prediction algorithms to XGBoost.
+
+## Usage
+Specify the 'device' parameter as described in the table below to offload model training and inference on SYCL device.
+
+### Algorithms
+| device | Description |
+| --- | --- |
+sycl | use default sycl device  |
+sycl:gpu | use default sycl gpu  |
+sycl:cpu | use default sycl cpu  |
+sycl:gpu:N | use sycl gpu number N |
+sycl:cpu:N | use sycl cpu number N |
+
+Python example:
+```python
+param['device'] = 'sycl:gpu:0'
+```
+Note: 'sycl:cpu' devices have full functional support but can't provide good enough performance. We recommend use 'sycl:cpu' devices only for test purposes.
+Note: if device is specified to be 'sycl', device type will be automatically chosen. In case the system has both sycl GPU and sycl CPU, GPU will on use.
+
+## Dependencies
+To build and use the plugin, install [Intel® oneAPI DPC++/C++ Compiler](https://www.intel.com/content/www/us/en/developer/tools/oneapi/dpc-compiler.html).
+See also [Intel® oneAPI Programming Guide](https://www.intel.com/content/www/us/en/docs/oneapi/programming-guide/2024-0/overview.html).
+
+## Build
+From the ``xgboost`` directory, run:
+
+```bash
+$ mkdir build
+$ cd build
+$ cmake .. -DPLUGIN_SYCL=ON
+$ make -j
+```

plugin/sycl/data.h

@@ -0,0 +1,256 @@
+/*!
+ * Copyright by Contributors 2017-2023
+ */
+#ifndef PLUGIN_SYCL_DATA_H_
+#define PLUGIN_SYCL_DATA_H_
+
+#include <cstddef>
+#include <limits>
+#include <mutex>
+#include <vector>
+#include <memory>
+#include <algorithm>
+
+#include "xgboost/base.h"
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wtautological-constant-compare"
+#pragma GCC diagnostic ignored "-W#pragma-messages"
+#include "xgboost/data.h"
+#pragma GCC diagnostic pop
+#include "xgboost/logging.h"
+#include "xgboost/host_device_vector.h"
+
+#include "../../src/common/threading_utils.h"
+
+#include "CL/sycl.hpp"
+
+namespace xgboost {
+namespace sycl {
+enum class MemoryType { shared, on_device};
+
+
+template <typename T>
+class USMDeleter {
+ public:
+  explicit USMDeleter(::sycl::queue qu) : qu_(qu) {}
+
+  void operator()(T* data) const {
+    ::sycl::free(data, qu_);
+  }
+
+ private:
+  ::sycl::queue qu_;
+};
+
+template <typename T, MemoryType memory_type = MemoryType::shared>
+class USMVector {
+  static_assert(std::is_standard_layout<T>::value, "USMVector admits only POD types");
+
+  std::shared_ptr<T> allocate_memory_(::sycl::queue* qu, size_t size) {
+    if constexpr (memory_type == MemoryType::shared) {
+      return std::shared_ptr<T>(::sycl::malloc_shared<T>(size_, *qu), USMDeleter<T>(*qu));
+    } else {
+      return std::shared_ptr<T>(::sycl::malloc_device<T>(size_, *qu), USMDeleter<T>(*qu));
+    }
+  }
+
+  void copy_vector_to_memory_(::sycl::queue* qu, const std::vector<T> &vec) {
+    if constexpr (memory_type == MemoryType::shared) {
+      std::copy(vec.begin(), vec.end(), data_.get());
+    } else {
+      qu->memcpy(data_.get(), vec.data(), size_ * sizeof(T));
+    }
+  }
+
+
+ public:
+  USMVector() : size_(0), capacity_(0), data_(nullptr) {}
+
+  USMVector(::sycl::queue& qu, size_t size) : size_(size), capacity_(size) {
+    data_ = allocate_memory_(qu, size_);
+  }
+
+  USMVector(::sycl::queue& qu, size_t size, T v) : size_(size), capacity_(size) {
+    data_ = allocate_memory_(qu, size_);
+    qu.fill(data_.get(), v, size_).wait();
+  }
+
+  USMVector(::sycl::queue* qu, const std::vector<T> &vec) {
+    size_ = vec.size();
+    capacity_ = size_;
+    data_ = allocate_memory_(qu, size_);
+    copy_vector_to_memory_(qu, vec);
+  }
+
+  ~USMVector() {
+  }
+
+  USMVector<T>& operator=(const USMVector<T>& other) {
+    size_ = other.size_;
+    capacity_ = other.capacity_;
+    data_ = other.data_;
+    return *this;
+  }
+
+  T* Data() { return data_.get(); }
+  const T* DataConst() const { return data_.get(); }
+
+  size_t Size() const { return size_; }
+
+  size_t Capacity() const { return capacity_; }
+
+  T& operator[] (size_t i) { return data_.get()[i]; }
+  const T& operator[] (size_t i) const { return data_.get()[i]; }
+
+  T* Begin () const { return data_.get(); }
+  T* End () const { return data_.get() + size_; }
+
+  bool Empty() const { return (size_ == 0); }
+
+  void Clear() {
+    data_.reset();
+    size_ = 0;
+    capacity_ = 0;
+  }
+
+  void Resize(::sycl::queue* qu, size_t size_new) {
+    if (size_new <= capacity_) {
+      size_ = size_new;
+    } else {
+      size_t size_old = size_;
+      auto data_old = data_;
+      size_ = size_new;
+      capacity_ = size_new;
+      data_ = allocate_memory_(qu, size_);;
+      if (size_old > 0) {
+        qu->memcpy(data_.get(), data_old.get(), sizeof(T) * size_old).wait();
+      }
+    }
+  }
+
+  void Resize(::sycl::queue* qu, size_t size_new, T v) {
+    if (size_new <= size_) {
+      size_ = size_new;
+    } else if (size_new <= capacity_) {
+      qu->fill(data_.get() + size_, v, size_new - size_).wait();
+      size_ = size_new;
+    } else {
+      size_t size_old = size_;
+      auto data_old = data_;
+      size_ = size_new;
+      capacity_ = size_new;
+      data_ = allocate_memory_(qu, size_);
+      if (size_old > 0) {
+        qu->memcpy(data_.get(), data_old.get(), sizeof(T) * size_old).wait();
+      }
+      qu->fill(data_.get() + size_old, v, size_new - size_old).wait();
+    }
+  }
+
+  ::sycl::event ResizeAsync(::sycl::queue* qu, size_t size_new, T v) {
+    if (size_new <= size_) {
+      size_ = size_new;
+      return ::sycl::event();
+    } else if (size_new <= capacity_) {
+      auto event = qu->fill(data_.get() + size_, v, size_new - size_);
+      size_ = size_new;
+      return event;
+    } else {
+      size_t size_old = size_;
+      auto data_old = data_;
+      size_ = size_new;
+      capacity_ = size_new;
+      data_ = allocate_memory_(qu, size_);
+      ::sycl::event event;
+      if (size_old > 0) {
+        event = qu->memcpy(data_.get(), data_old.get(), sizeof(T) * size_old);
+      }
+      return qu->fill(data_.get() + size_old, v, size_new - size_old, event);
+    }
+  }
+
+  ::sycl::event ResizeAndFill(::sycl::queue* qu, size_t size_new, int v) {
+    if (size_new <= size_) {
+      size_ = size_new;
+      return qu->memset(data_.get(), v, size_new * sizeof(T));
+    } else if (size_new <= capacity_) {
+      size_ = size_new;
+      return qu->memset(data_.get(), v, size_new * sizeof(T));
+    } else {
+      size_t size_old = size_;
+      auto data_old = data_;
+      size_ = size_new;
+      capacity_ = size_new;
+      data_ = allocate_memory_(qu, size_);
+      return qu->memset(data_.get(), v, size_new * sizeof(T));
+    }
+  }
+
+  ::sycl::event Fill(::sycl::queue* qu, T v) {
+    return qu->fill(data_.get(), v, size_);
+  }
+
+  void Init(::sycl::queue* qu, const std::vector<T> &vec) {
+    size_ = vec.size();
+    capacity_ = size_;
+    data_ = allocate_memory_(qu, size_);
+    copy_vector_to_memory_(qu, vec);
+  }
+
+  using value_type = T;  // NOLINT
+
+ private:
+  size_t size_;
+  size_t capacity_;
+  std::shared_ptr<T> data_;
+};
+
+/* Wrapper for DMatrix which stores all batches in a single USM buffer */
+struct DeviceMatrix {
+  DMatrix* p_mat;  // Pointer to the original matrix on the host
+  ::sycl::queue qu_;
+  USMVector<size_t> row_ptr;
+  USMVector<Entry> data;
+  size_t total_offset;
+
+  DeviceMatrix(::sycl::queue qu, DMatrix* dmat) : p_mat(dmat), qu_(qu) {
+    size_t num_row = 0;
+    size_t num_nonzero = 0;
+    for (auto &batch : dmat->GetBatches<SparsePage>()) {
+      const auto& data_vec = batch.data.HostVector();
+      const auto& offset_vec = batch.offset.HostVector();
+      num_nonzero += data_vec.size();
+      num_row += batch.Size();
+    }
+
+    row_ptr.Resize(&qu_, num_row + 1);
+    data.Resize(&qu_, num_nonzero);
+
+    size_t data_offset = 0;
+    for (auto &batch : dmat->GetBatches<SparsePage>()) {
+      const auto& data_vec = batch.data.HostVector();
+      const auto& offset_vec = batch.offset.HostVector();
+      size_t batch_size = batch.Size();
+      if (batch_size > 0) {
+        std::copy(offset_vec.data(), offset_vec.data() + batch_size,
+                  row_ptr.Data() + batch.base_rowid);
+        if (batch.base_rowid > 0) {
+          for (size_t i = 0; i < batch_size; i++)
+            row_ptr[i + batch.base_rowid] += batch.base_rowid;
+        }
+        std::copy(data_vec.data(), data_vec.data() + offset_vec[batch_size],
+                  data.Data() + data_offset);
+        data_offset += offset_vec[batch_size];
+      }
+    }
+    row_ptr[num_row] = data_offset;
+    total_offset = data_offset;
+  }
+
+  ~DeviceMatrix() {
+  }
+};
+}  // namespace sycl
+}  // namespace xgboost
+
+#endif  // PLUGIN_SYCL_DATA_H_

plugin/sycl/device_manager.cc

@@ -0,0 +1,124 @@
+/*!
+ * Copyright 2017-2023 by Contributors
+ * \file device_manager.cc
+ */
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wtautological-constant-compare"
+#pragma GCC diagnostic ignored "-W#pragma-messages"
+#include <rabit/rabit.h>
+#pragma GCC diagnostic pop
+
+#include "../sycl/device_manager.h"
+
+namespace xgboost {
+namespace sycl {
+
+::sycl::device DeviceManager::GetDevice(const DeviceOrd& device_spec) const {
+    if (!device_spec.IsSycl()) {
+        LOG(WARNING) << "Sycl kernel is executed with non-sycl context: "
+                     << device_spec.Name() << ". "
+                     << "Default sycl device_selector will be used.";
+    }
+
+    bool not_use_default_selector = (device_spec.ordinal != kDefaultOrdinal) ||
+                                    (rabit::IsDistributed());
+    if (not_use_default_selector) {
+      DeviceRegister& device_register = GetDevicesRegister();
+      const int device_idx = rabit::IsDistributed() ? rabit::GetRank() : device_spec.ordinal;
+      if (device_spec.IsSyclDefault()) {
+          auto& devices = device_register.devices;
+          CHECK_LT(device_idx, devices.size());
+          return devices[device_idx];
+      } else if (device_spec.IsSyclCPU()) {
+          auto& cpu_devices = device_register.cpu_devices;
+          CHECK_LT(device_idx, cpu_devices.size());
+          return cpu_devices[device_idx];
+      } else {
+          auto& gpu_devices = device_register.gpu_devices;
+          CHECK_LT(device_idx, gpu_devices.size());
+          return gpu_devices[device_idx];
+      }
+    } else {
+        if (device_spec.IsSyclCPU()) {
+            return ::sycl::device(::sycl::cpu_selector_v);
+        } else if (device_spec.IsSyclGPU()) {
+            return ::sycl::device(::sycl::gpu_selector_v);
+        } else {
+            return ::sycl::device(::sycl::default_selector_v);
+        }
+    }
+}
+
+::sycl::queue DeviceManager::GetQueue(const DeviceOrd& device_spec) const {
+    if (!device_spec.IsSycl()) {
+        LOG(WARNING) << "Sycl kernel is executed with non-sycl context: "
+                     << device_spec.Name() << ". "
+                     << "Default sycl device_selector will be used.";
+    }
+
+    QueueRegister_t& queue_register = GetQueueRegister();
+    if (queue_register.count(device_spec.Name()) > 0) {
+        return queue_register.at(device_spec.Name());
+    }
+
+    bool not_use_default_selector = (device_spec.ordinal != kDefaultOrdinal) ||
+                                    (rabit::IsDistributed());
+    std::lock_guard<std::mutex> guard(queue_registering_mutex);
+    if (not_use_default_selector) {
+      DeviceRegister& device_register = GetDevicesRegister();
+      const int device_idx = rabit::IsDistributed() ? rabit::GetRank() : device_spec.ordinal;
+      if (device_spec.IsSyclDefault()) {
+          auto& devices = device_register.devices;
+          CHECK_LT(device_idx, devices.size());
+          queue_register[device_spec.Name()] = ::sycl::queue(devices[device_idx]);
+      } else if (device_spec.IsSyclCPU()) {
+          auto& cpu_devices = device_register.cpu_devices;
+          CHECK_LT(device_idx, cpu_devices.size());
+          queue_register[device_spec.Name()] = ::sycl::queue(cpu_devices[device_idx]);;
+      } else if (device_spec.IsSyclGPU()) {
+          auto& gpu_devices = device_register.gpu_devices;
+          CHECK_LT(device_idx, gpu_devices.size());
+          queue_register[device_spec.Name()] = ::sycl::queue(gpu_devices[device_idx]);
+      }
+    } else {
+        if (device_spec.IsSyclCPU()) {
+            queue_register[device_spec.Name()] = ::sycl::queue(::sycl::cpu_selector_v);
+        } else if (device_spec.IsSyclGPU()) {
+            queue_register[device_spec.Name()] = ::sycl::queue(::sycl::gpu_selector_v);
+        } else {
+            queue_register[device_spec.Name()] = ::sycl::queue(::sycl::default_selector_v);
+        }
+    }
+    return queue_register.at(device_spec.Name());
+}
+
+DeviceManager::DeviceRegister& DeviceManager::GetDevicesRegister() const {
+    static DeviceRegister device_register;
+
+    if (device_register.devices.size() == 0) {
+        std::lock_guard<std::mutex> guard(device_registering_mutex);
+        std::vector<::sycl::device> devices = ::sycl::device::get_devices();
+        for (size_t i = 0; i < devices.size(); i++) {
+            LOG(INFO) << "device_index = " << i << ", name = "
+                      << devices[i].get_info<::sycl::info::device::name>();
+        }
+
+        for (size_t i = 0; i < devices.size(); i++) {
+            device_register.devices.push_back(devices[i]);
+            if (devices[i].is_cpu()) {
+                device_register.cpu_devices.push_back(devices[i]);
+            } else if (devices[i].is_gpu()) {
+                device_register.gpu_devices.push_back(devices[i]);
+            }
+        }
+    }
+    return device_register;
+}
+
+DeviceManager::QueueRegister_t& DeviceManager::GetQueueRegister() const {
+    static QueueRegister_t queue_register;
+    return queue_register;
+}
+
+}  // namespace sycl
+}  // namespace xgboost

plugin/sycl/device_manager.h

@@ -0,0 +1,47 @@
+/*!
+ * Copyright 2017-2023 by Contributors
+ * \file device_manager.h
+ */
+#ifndef PLUGIN_SYCL_DEVICE_MANAGER_H_
+#define PLUGIN_SYCL_DEVICE_MANAGER_H_
+
+#include <vector>
+#include <mutex>
+#include <string>
+#include <unordered_map>
+
+#include <CL/sycl.hpp>
+
+#include "xgboost/context.h"
+
+namespace xgboost {
+namespace sycl {
+
+class DeviceManager {
+ public:
+  ::sycl::queue GetQueue(const DeviceOrd& device_spec) const;
+
+  ::sycl::device GetDevice(const DeviceOrd& device_spec) const;
+
+ private:
+  using QueueRegister_t = std::unordered_map<std::string, ::sycl::queue>;
+  constexpr static int kDefaultOrdinal = -1;
+
+  struct DeviceRegister {
+    std::vector<::sycl::device> devices;
+    std::vector<::sycl::device> cpu_devices;
+    std::vector<::sycl::device> gpu_devices;
+  };
+
+  QueueRegister_t& GetQueueRegister() const;
+
+  DeviceRegister& GetDevicesRegister() const;
+
+  mutable std::mutex queue_registering_mutex;
+  mutable std::mutex device_registering_mutex;
+};
+
+}  // namespace sycl
+}  // namespace xgboost
+
+#endif  // PLUGIN_SYCL_DEVICE_MANAGER_H_

plugin/sycl/predictor/predictor.cc

@@ -0,0 +1,342 @@
+/*!
+ * Copyright by Contributors 2017-2023
+ */
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wtautological-constant-compare"
+#pragma GCC diagnostic ignored "-W#pragma-messages"
+#include <rabit/rabit.h>
+#pragma GCC diagnostic pop
+
+#include <cstddef>
+#include <limits>
+#include <mutex>
+
+#include <CL/sycl.hpp>
+
+#include "../data.h"
+
+#include "dmlc/registry.h"
+
+#include "xgboost/tree_model.h"
+#include "xgboost/predictor.h"
+#include "xgboost/tree_updater.h"
+
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wtautological-constant-compare"
+#include "../../src/data/adapter.h"
+#pragma GCC diagnostic pop
+#include "../../src/common/math.h"
+#include "../../src/gbm/gbtree_model.h"
+
+#include "../device_manager.h"
+
+namespace xgboost {
+namespace sycl {
+namespace predictor {
+
+DMLC_REGISTRY_FILE_TAG(predictor_sycl);
+
+/* Wrapper for descriptor of a tree node */
+struct DeviceNode {
+  DeviceNode()
+      : fidx(-1), left_child_idx(-1), right_child_idx(-1) {}
+
+  union NodeValue {
+    float leaf_weight;
+    float fvalue;
+  };
+
+  int fidx;
+  int left_child_idx;
+  int right_child_idx;
+  NodeValue val;
+
+  explicit DeviceNode(const RegTree::Node& n) {
+    this->left_child_idx = n.LeftChild();
+    this->right_child_idx = n.RightChild();
+    this->fidx = n.SplitIndex();
+    if (n.DefaultLeft()) {
+      fidx |= (1U << 31);
+    }
+
+    if (n.IsLeaf()) {
+      this->val.leaf_weight = n.LeafValue();
+    } else {
+      this->val.fvalue = n.SplitCond();
+    }
+  }
+
+  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 this->left_child_idx;
+    } else {
+      return this->right_child_idx;
+    }
+  }
+
+  float GetFvalue() const { return val.fvalue; }
+
+  float GetWeight() const { return val.leaf_weight; }
+};
+
+/* SYCL implementation of a device model,
+ * storing tree structure in USM buffers to provide access from device kernels
+ */
+class DeviceModel {
+ public:
+  ::sycl::queue qu_;
+  USMVector<DeviceNode> nodes_;
+  USMVector<size_t> tree_segments_;
+  USMVector<int> tree_group_;
+  size_t tree_beg_;
+  size_t tree_end_;
+  int num_group_;
+
+  DeviceModel() {}
+
+  ~DeviceModel() {}
+
+  void Init(::sycl::queue qu, const gbm::GBTreeModel& model, size_t tree_begin, size_t tree_end) {
+    qu_ = qu;
+
+    tree_segments_.Resize(&qu_, (tree_end - tree_begin) + 1);
+    int sum = 0;
+    tree_segments_[0] = sum;
+    for (int tree_idx = tree_begin; tree_idx < tree_end; tree_idx++) {
+      if (model.trees[tree_idx]->HasCategoricalSplit()) {
+        LOG(FATAL) << "Categorical features are not yet supported by sycl";
+      }
+      sum += model.trees[tree_idx]->GetNodes().size();
+      tree_segments_[tree_idx - tree_begin + 1] = sum;
+    }
+
+    nodes_.Resize(&qu_, sum);
+    for (int tree_idx = tree_begin; tree_idx < tree_end; tree_idx++) {
+      auto& src_nodes = model.trees[tree_idx]->GetNodes();
+      for (size_t node_idx = 0; node_idx < src_nodes.size(); node_idx++)
+        nodes_[node_idx + tree_segments_[tree_idx - tree_begin]] =
+          static_cast<DeviceNode>(src_nodes[node_idx]);
+    }
+
+    tree_group_.Resize(&qu_, model.tree_info.size());
+    for (size_t tree_idx = 0; tree_idx < model.tree_info.size(); tree_idx++)
+      tree_group_[tree_idx] = model.tree_info[tree_idx];
+
+    tree_beg_ = tree_begin;
+    tree_end_ = tree_end;
+    num_group_ = model.learner_model_param->num_output_group;
+  }
+};
+
+float GetFvalue(int ridx, int fidx, Entry* data, size_t* row_ptr, bool* is_missing) {
+  // Binary search
+  auto begin_ptr = data + row_ptr[ridx];
+  auto end_ptr = data + row_ptr[ridx + 1];
+  Entry* previous_middle = nullptr;
+  while (end_ptr != begin_ptr) {
+    auto middle = begin_ptr + (end_ptr - begin_ptr) / 2;
+    if (middle == previous_middle) {
+      break;
+    } else {
+      previous_middle = middle;
+    }
+
+    if (middle->index == fidx) {
+      *is_missing = false;
+      return middle->fvalue;
+    } else if (middle->index < fidx) {
+      begin_ptr = middle;
+    } else {
+      end_ptr = middle;
+    }
+  }
+  *is_missing = true;
+  return 0.0;
+}
+
+float GetLeafWeight(int ridx, const DeviceNode* tree, Entry* data, size_t* row_ptr) {
+  DeviceNode n = tree[0];
+  int node_id = 0;
+  bool is_missing;
+  while (!n.IsLeaf()) {
+    float fvalue = GetFvalue(ridx, n.GetFidx(), data, row_ptr, &is_missing);
+    // Missing value
+    if (is_missing) {
+      n = tree[n.MissingIdx()];
+    } else {
+      if (fvalue < n.GetFvalue()) {
+        node_id = n.left_child_idx;
+        n = tree[n.left_child_idx];
+      } else {
+        node_id = n.right_child_idx;
+        n = tree[n.right_child_idx];
+      }
+    }
+  }
+  return n.GetWeight();
+}
+
+void DevicePredictInternal(::sycl::queue qu,
+                           sycl::DeviceMatrix* dmat,
+                           HostDeviceVector<float>* out_preds,
+                           const gbm::GBTreeModel& model,
+                           size_t tree_begin,
+                           size_t tree_end) {
+  if (tree_end - tree_begin == 0) return;
+  if (out_preds->HostVector().size() == 0) return;
+
+  DeviceModel device_model;
+  device_model.Init(qu, model, tree_begin, tree_end);
+
+  auto& out_preds_vec = out_preds->HostVector();
+
+  DeviceNode* nodes = device_model.nodes_.Data();
+  ::sycl::buffer<float, 1> out_preds_buf(out_preds_vec.data(), out_preds_vec.size());
+  size_t* tree_segments = device_model.tree_segments_.Data();
+  int* tree_group = device_model.tree_group_.Data();
+  size_t* row_ptr = dmat->row_ptr.Data();
+  Entry* data = dmat->data.Data();
+  int num_features = dmat->p_mat->Info().num_col_;
+  int num_rows = dmat->row_ptr.Size() - 1;
+  int num_group = model.learner_model_param->num_output_group;
+
+  qu.submit([&](::sycl::handler& cgh) {
+    auto out_predictions = out_preds_buf.template get_access<::sycl::access::mode::read_write>(cgh);
+    cgh.parallel_for<>(::sycl::range<1>(num_rows), [=](::sycl::id<1> pid) {
+      int global_idx = pid[0];
+      if (global_idx >= num_rows) return;
+      if (num_group == 1) {
+        float sum = 0.0;
+        for (int tree_idx = tree_begin; tree_idx < tree_end; tree_idx++) {
+          const DeviceNode* tree = nodes + tree_segments[tree_idx - tree_begin];
+          sum += GetLeafWeight(global_idx, tree, data, row_ptr);
+        }
+        out_predictions[global_idx] += sum;
+      } else {
+        for (int tree_idx = tree_begin; tree_idx < tree_end; tree_idx++) {
+          const DeviceNode* tree = nodes + tree_segments[tree_idx - tree_begin];
+          int out_prediction_idx = global_idx * num_group + tree_group[tree_idx];
+          out_predictions[out_prediction_idx] += GetLeafWeight(global_idx, tree, data, row_ptr);
+        }
+      }
+    });
+  }).wait();
+}
+
+class Predictor : public xgboost::Predictor {
+ protected:
+  void InitOutPredictions(const MetaInfo& info,
+                          HostDeviceVector<bst_float>* out_preds,
+                          const gbm::GBTreeModel& model) const override {
+    CHECK_NE(model.learner_model_param->num_output_group, 0);
+    size_t n = model.learner_model_param->num_output_group * info.num_row_;
+    const auto& base_margin = info.base_margin_.Data()->HostVector();
+    out_preds->Resize(n);
+    std::vector<bst_float>& out_preds_h = out_preds->HostVector();
+    if (base_margin.size() == n) {
+      CHECK_EQ(out_preds->Size(), n);
+      std::copy(base_margin.begin(), base_margin.end(), out_preds_h.begin());
+    } else {
+      auto base_score = model.learner_model_param->BaseScore(ctx_)(0);
+      if (!base_margin.empty()) {
+        std::ostringstream oss;
+        oss << "Ignoring the base margin, since it has incorrect length. "
+            << "The base margin must be an array of length ";
+        if (model.learner_model_param->num_output_group > 1) {
+          oss << "[num_class] * [number of data points], i.e. "
+              << model.learner_model_param->num_output_group << " * " << info.num_row_
+              << " = " << n << ". ";
+        } else {
+          oss << "[number of data points], i.e. " << info.num_row_ << ". ";
+        }
+        oss << "Instead, all data points will use "
+            << "base_score = " << base_score;
+        LOG(WARNING) << oss.str();
+      }
+      std::fill(out_preds_h.begin(), out_preds_h.end(), base_score);
+    }
+  }
+
+ public:
+  explicit Predictor(Context const* context) :
+      xgboost::Predictor::Predictor{context},
+      cpu_predictor(xgboost::Predictor::Create("cpu_predictor", context)) {}
+
+  void PredictBatch(DMatrix *dmat, PredictionCacheEntry *predts,
+                    const gbm::GBTreeModel &model, uint32_t tree_begin,
+                    uint32_t tree_end = 0) const override {
+    ::sycl::queue qu = device_manager.GetQueue(ctx_->Device());
+    // TODO(razdoburdin): remove temporary workaround after cache fix
+    sycl::DeviceMatrix device_matrix(qu, dmat);
+
+    auto* out_preds = &predts->predictions;
+    if (tree_end == 0) {
+      tree_end = model.trees.size();
+    }
+
+    if (tree_begin < tree_end) {
+      DevicePredictInternal(qu, &device_matrix, out_preds, model, tree_begin, tree_end);
+    }
+  }
+
+  bool InplacePredict(std::shared_ptr<DMatrix> p_m,
+                      const gbm::GBTreeModel &model, float missing,
+                      PredictionCacheEntry *out_preds, uint32_t tree_begin,
+                      unsigned tree_end) const override {
+    LOG(WARNING) << "InplacePredict is not yet implemented for SYCL. CPU Predictor is used.";
+    return cpu_predictor->InplacePredict(p_m, model, missing, out_preds, tree_begin, tree_end);
+  }
+
+  void PredictInstance(const SparsePage::Inst& inst,
+                       std::vector<bst_float>* out_preds,
+                       const gbm::GBTreeModel& model, unsigned ntree_limit,
+                       bool is_column_split) const override {
+    LOG(WARNING) << "PredictInstance is not yet implemented for SYCL. CPU Predictor is used.";
+    cpu_predictor->PredictInstance(inst, out_preds, model, ntree_limit, is_column_split);
+  }
+
+  void PredictLeaf(DMatrix* p_fmat, HostDeviceVector<bst_float>* out_preds,
+                   const gbm::GBTreeModel& model, unsigned ntree_limit) const override {
+    LOG(WARNING) << "PredictLeaf is not yet implemented for SYCL. CPU Predictor is used.";
+    cpu_predictor->PredictLeaf(p_fmat, out_preds, model, ntree_limit);
+  }
+
+  void PredictContribution(DMatrix* p_fmat, HostDeviceVector<float>* out_contribs,
+                           const gbm::GBTreeModel& model, uint32_t ntree_limit,
+                           const std::vector<bst_float>* tree_weights,
+                           bool approximate, int condition,
+                           unsigned condition_feature) const override {
+    LOG(WARNING) << "PredictContribution is not yet implemented for SYCL. CPU Predictor is used.";
+    cpu_predictor->PredictContribution(p_fmat, out_contribs, model, ntree_limit, tree_weights,
+                                       approximate, condition, condition_feature);
+  }
+
+  void PredictInteractionContributions(DMatrix* p_fmat, HostDeviceVector<bst_float>* out_contribs,
+                                       const gbm::GBTreeModel& model, unsigned ntree_limit,
+                                       const std::vector<bst_float>* tree_weights,
+                                       bool approximate) const override {
+    LOG(WARNING) << "PredictInteractionContributions is not yet implemented for SYCL. "
+                 << "CPU Predictor is used.";
+    cpu_predictor->PredictInteractionContributions(p_fmat, out_contribs, model, ntree_limit,
+                                                   tree_weights, approximate);
+  }
+
+ private:
+  DeviceManager device_manager;
+
+  std::unique_ptr<xgboost::Predictor> cpu_predictor;
+};
+
+XGBOOST_REGISTER_PREDICTOR(Predictor, "sycl_predictor")
+.describe("Make predictions using SYCL.")
+.set_body([](Context const* ctx) { return new Predictor(ctx); });
+
+}  // namespace predictor
+}  // namespace sycl
+}  // namespace xgboost