feature(rf optimizations): enabling oneDPL and sort primitive refactoring

.ci/env/apt.sh

@@ -38,9 +38,14 @@ function install_tbb {
     sudo apt-get install -y intel-oneapi-tbb-devel-2022.0
 }
 
+function install_dpl {
+    sudo apt-get install -y intel-oneapi-libdpstd-devel
+}
+
 function install_mkl {
     sudo apt-get install -y intel-oneapi-mkl-devel-2025.0
     install_tbb
+    install_dpl
 }
 
 function install_clang-format {
@@ -129,6 +134,9 @@ elif [ "${component}" == "tbb" ]; then
 elif [ "${component}" == "mkl" ]; then
     add_repo
     install_mkl
+elif [ "${component}" == "dpl" ]; then
+    add_repo
+    install_dpl
 elif [ "${component}" == "gnu-cross-compilers" ]; then
     update
     install_gnu-cross-compilers "$2"

.ci/pipeline/ci.yml

@@ -29,7 +29,7 @@ variables:
   VM_IMAGE : 'ubuntu-24.04'
   SYSROOT_OS: 'noble'
   WINDOWS_BASEKIT_URL: 'https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b380d914-366b-4b77-a74a-05e3c38b3514/intel-oneapi-base-toolkit-2025.0.0.882_offline.exe'
-  WINDOWS_DPCPP_COMPONENTS: 'intel.oneapi.win.mkl.devel:intel.oneapi.win.tbb.devel'
+  WINDOWS_DPCPP_COMPONENTS: 'intel.oneapi.win.mkl.devel:intel.oneapi.win.tbb.devel:intel.oneapi.win.dpl'
 
 resources:
   repositories:

INSTALL.md

@@ -23,6 +23,7 @@ Required Software:
 * BLAS and LAPACK libraries - both provided by oneMKL
 * Python version 3.9 or higher
 * oneTBB library (repository contains script to download it)
+* oneDPL library
 * Microsoft Visual Studio\* (Windows\* only)
 * [MSYS2](http://msys2.github.io) (Windows\* only)
 * `make` and `dos2unix` tools; install these packages using MSYS2 on Windows\* as follows:
@@ -105,9 +106,24 @@ is available as an alternative to the manual setup.
 
             ./dev/download_tbb.sh
 
-6. Download and install Python (version 3.9 or higher).
+6. Set up oneDPL
+  _Note: if you used the general oneAPI setvars script from a Base Toolkit installation, this step will not be necessary as oneDPL will already have been set up._
 
-7. Build oneDAL via command-line interface. Choose the appropriate commands based on the interface, platform, and the compiler you use. Interface and platform are required arguments of makefile while others are optional. Below you can find the set of examples for building oneDAL. You may use a combination of them to get the desired build configuration:
+    Download and install [Intel(R) oneDPL](https://www.intel.com/content/www/us/en/developer/tools/oneapi/dpc-library.html).
+    Set the environment variables for for Intel(R) oneDPL. For example:
+
+    - oneDPL (Windows\*):
+
+            call "C:\Program Files (x86)\Intel\oneAPI\dpl\latest\env\vars.bat" intel64
+
+    - oneDPL (Linux\*):
+
+            source /opt/intel/oneapi/dpl/latest/env/vars.sh intel64
+
+
+7. Download and install Python (version 3.9 or higher).
+
+8. Build oneDAL via command-line interface. Choose the appropriate commands based on the interface, platform, and the compiler you use. Interface and platform are required arguments of makefile while others are optional. Below you can find the set of examples for building oneDAL. You may use a combination of them to get the desired build configuration:
 
     - DAAL interfaces on **Linux\*** using **Intel(R) C++ Compiler**:
 

MODULE.bazel

@@ -90,6 +90,21 @@ ccl_repo(
     ]
 )
 
+dpl_repo = use_repo_rule("@onedal//dev/bazel/deps:dpl.bzl", "dpl_repo")
+dpl_repo(
+    name = "dpl",
+    root_env_var = "DPL_ROOT",
+    urls = [
+        "https://files.pythonhosted.org/packages/95/f6/18f78cb933e01ecd9e99d37a10da4971a795fcfdd1d24640799b4050fdbb/onedpl_devel-2022.7.1-py2.py3-none-manylinux_2_28_x86_64.whl",
+    ],
+    sha256s = [
+        "3b270999d2464c5151aa0e7995dda9e896d072c75069ccee1efae9dc56bdc417",
+    ],
+    strip_prefixes = [
+        "onedpl_devel-2022.7.1.data/data",
+    ],
+)
+
 mkl_repo = use_repo_rule("@onedal//dev/bazel/deps:mkl.bzl", "mkl_repo")
 mkl_repo(
     name = "mkl",

cpp/oneapi/dal.hpp

@@ -24,6 +24,7 @@
 #include "oneapi/dal/exceptions.hpp"
 #include "oneapi/dal/infer.hpp"
 #include "oneapi/dal/read.hpp"
+#include "oneapi/dal/rng.hpp"
 #include "oneapi/dal/train.hpp"
 #include "oneapi/dal/partial_compute.hpp"
 #include "oneapi/dal/finalize_compute.hpp"

cpp/oneapi/dal/BUILD

@@ -31,6 +31,7 @@ dal_module(
     ],
     dpc_deps = [
         "@mkl//:mkl_dpc",
+        "@dpl//:headers",
     ],
 )
 

cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_feature_type_dpc.cpp

@@ -20,6 +20,7 @@
 #include "oneapi/dal/table/row_accessor.hpp"
 #include "oneapi/dal/backend/memory.hpp"
 #include "oneapi/dal/detail/profiler.hpp"
+#include <string>
 
 #ifdef ONEDAL_DATA_PARALLEL
 
@@ -34,8 +35,15 @@ inline sycl::event sort_inplace(sycl::queue& queue_,
                                 pr::ndarray<Float, 1>& src,
                                 const bk::event_vector& deps = {}) {
     ONEDAL_ASSERT(src.get_count() > 0);
+    auto device = queue_.get_device();
+    std::string device_name = device.get_info<sycl::info::device::name>();
     auto src_ind = pr::ndarray<Index, 1>::empty(queue_, { src.get_count() });
-    return pr::radix_sort_indices_inplace<Float, Index>{ queue_ }(src, src_ind, deps);
+    if (device_name.find("Data Center GPU Max") != std::string::npos) {
+        return pr::radix_sort_indices_inplace_dpl<Float, Index>(queue_, src, src_ind, deps);
+    }
+    else {
+        return pr::radix_sort_indices_inplace<Float, Index>{ queue_ }(src, src_ind, deps);
+    }
 }
 
 template <typename Float, typename Bin, typename Index>
@@ -429,15 +437,36 @@ sycl::event indexed_features<Float, Bin, Index>::operator()(const table& tbl,
             pr::ndarray<Bin, 1>::empty(queue_, { row_count_ }, sycl::usm::alloc::device);
     }
 
-    pr::radix_sort_indices_inplace<Float, Index> sort{ queue_ };
-
     sycl::event last_event;
-
-    for (Index i = 0; i < column_count_; i++) {
-        last_event = extract_column(data_nd_, values_nd, indices_nd, i, { last_event });
-        last_event = sort(values_nd, indices_nd, { last_event });
-        last_event =
-            compute_bins(values_nd, indices_nd, column_bin_vec_[i], entries_[i], i, { last_event });
+    auto device = queue_.get_device();
+    std::string device_name = device.get_info<sycl::info::device::name>();
+    if (device_name.find("Data Center GPU Max") != std::string::npos) {
+        for (Index i = 0; i < column_count_; i++) {
+            last_event = extract_column(data_nd_, values_nd, indices_nd, i, { last_event });
+            last_event = pr::radix_sort_indices_inplace_dpl<Float, Index>(queue_,
+                                                                          values_nd,
+                                                                          indices_nd,
+                                                                          { last_event });
+            last_event = compute_bins(values_nd,
+                                      indices_nd,
+                                      column_bin_vec_[i],
+                                      entries_[i],
+                                      i,
+                                      { last_event });
+        }
+    }
+    else {
+        pr::radix_sort_indices_inplace<Float, Index> sort{ queue_ };
+        for (Index i = 0; i < column_count_; i++) {
+            last_event = extract_column(data_nd_, values_nd, indices_nd, i, { last_event });
+            last_event = sort(values_nd, indices_nd, { last_event });
+            last_event = compute_bins(values_nd,
+                                      indices_nd,
+                                      column_bin_vec_[i],
+                                      entries_[i],
+                                      i,
+                                      { last_event });
+        }
     }
 
     last_event.wait_and_throw();

cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl.hpp

@@ -21,7 +21,7 @@
 #include "oneapi/dal/backend/primitives/utils.hpp"
 #include "oneapi/dal/algo/decision_forest/train_types.hpp"
 
-#include "oneapi/dal/backend/primitives/rng/host_engine_collection.hpp"
+#include "oneapi/dal/backend/primitives/rng/device_engine.hpp"
 
 #include "oneapi/dal/algo/decision_forest/backend/gpu/train_misc_structs.hpp"
 #include "oneapi/dal/algo/decision_forest/backend/gpu/train_impurity_data.hpp"
@@ -50,8 +50,7 @@ class train_kernel_hist_impl {
     using model_manager_t = train_model_manager<Float, Index, Task>;
     using train_context_t = train_context<Float, Index, Task>;
     using imp_data_t = impurity_data<Float, Index, Task>;
-    using rng_engine_t = pr::host_engine;
-    using rng_engine_list_t = std::vector<rng_engine_t>;
+    using rng_engine_t = pr::device_engine;
     using msg = dal::detail::error_messages;
     using comm_t = bk::communicator<spmd::device_memory_access::usm>;
     using node_t = node<Index>;
@@ -79,7 +78,7 @@ class train_kernel_hist_impl {
                                           Index class_count) const;
 
     sycl::event gen_initial_tree_order(train_context_t& ctx,
-                                       rng_engine_list_t& rng_engine_list,
+                                       rng_engine_t& rng_engine,
                                        pr::ndarray<Index, 1>& node_list,
                                        pr::ndarray<Index, 1>& tree_order_level,
                                        Index engine_offset,
@@ -103,6 +102,7 @@ class train_kernel_hist_impl {
                      const table& data,
                      const table& labels,
                      const table& weights);
+
     /// Allocates all buffers that are used for training.
     /// @param[in] ctx  a training context structure for a GPU backend
     void allocate_buffers(const train_context_t& ctx);
@@ -115,12 +115,12 @@ class train_kernel_hist_impl {
     /// @param[in] ctx              a training context structure for a GPU backend
     /// @param[in] node_count       number of nodes on the current level
     /// @param[in] node_vs_tree_map an initial tree order
-    /// @param[in] rng_engine_list  a list of random generator engines
+    /// @param[in] rng_engine  a random generator engine
     std::tuple<pr::ndarray<Index, 1>, sycl::event> gen_feature_list(
         const train_context_t& ctx,
         Index node_count,
         const pr::ndarray<Index, 1>& node_vs_tree_map,
-        rng_engine_list_t& rng_engine_list);
+        rng_engine_t& rng_engine);
 
     /// Generates random thresholds for each node and for each selected feature for node.
     /// Thresholds are used for a random splitter kernel to split each node.
@@ -129,12 +129,12 @@ class train_kernel_hist_impl {
     /// @param[in] ctx              a training context structure for a GPU backend
     /// @param[in] node_count       number of nodes on the current level
     /// @param[in] node_vs_tree_map an initial tree order
-    /// @param[in] rng_engine_list  a list of random generator engines
+    /// @param[in] rng_engine  a random generator engine
     std::tuple<pr::ndarray<Float, 1>, sycl::event> gen_random_thresholds(
         const train_context_t& ctx,
         Index node_count,
         const pr::ndarray<Index, 1>& node_vs_tree_map,
-        rng_engine_list_t& rng_engine_list);
+        rng_engine_t& rng_engine);
 
     /// Computes initial impurity for each node.
     ///
@@ -561,7 +561,7 @@ class train_kernel_hist_impl {
     /// @param[in] oob_per_obs_list     an array of OOB values per observation
     /// @param[in] var_imp              variable importance values
     /// @param[in] var_imp_variance     variable importance variance values
-    /// @param[in] rng_engine_arr       a list of random generator engines
+    /// @param[in] rng_engine       a random generator engine
     /// @param[in] tree_idx             a tree index
     /// @param[in] tree_in_block        number of trees in the computational block
     /// @param[in] built_tree_count     number of built trees
@@ -575,7 +575,7 @@ class train_kernel_hist_impl {
                                 pr::ndarray<hist_type_t, 1>& oob_per_obs_list,
                                 pr::ndarray<Float, 1>& var_imp,
                                 pr::ndarray<Float, 1>& var_imp_variance,
-                                const rng_engine_list_t& rng_engine_arr,
+                                rng_engine_t& rng_engine,
                                 Index tree_idx,
                                 Index tree_in_block,
                                 Index built_tree_count,