feat(bazel): run the DPC++ device code on NVIDIA GPUs through oneMath

.bazelrc

@@ -20,6 +20,7 @@ build --flag_alias=cpu=@config//:cpu
 build --flag_alias=enable_assert=@config//:enable_assert
 build --flag_alias=stdalloc=@config//:stdalloc
 build --flag_alias=rng_backend=@config//:rng_backend
+build --flag_alias=dpc_math_backend=@config//:dpc_math_backend
 build --flag_alias=build_parameters_lib=@config//:build_parameters_lib
 build --flag_alias=msvc_runtime=@config//:msvc_runtime
 
@@ -149,6 +150,31 @@ test:dpc-private \
 build:release-dpc \
     --release_dpc=true
 
+# Configuration: 'nvidia-gpu'
+# Experimental: run the DPC++ device code on NVIDIA GPUs.
+#
+# Swaps oneMKL for oneMath, which additionally dispatches to cuBLAS, cuSOLVER,
+# cuSPARSE and cuRAND. The device sources themselves are unchanged. Needs
+#   * oneMath built with the CUDA backends enabled, pointed at by ONEMATHROOT;
+#   * a DPC++ compiler with CUDA support, asked for the NVPTX target through
+#     ONEDAL_SYCL_TARGETS (an env var, because the DPC++ flag sets are fixed
+#     when the toolchain repository is configured):
+#
+#   export ONEMATHROOT=/path/to/onemath
+#   export ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda
+#   bazel test --config=nvidia-gpu //cpp/oneapi/dal/backend/primitives/blas:tests
+#
+# The test lane also forwards LD_LIBRARY_PATH. oneMath's dispatcher `dlopen`s its
+# per-domain backends through a `$ORIGIN` RUNPATH, and under Bazel `$ORIGIN` is
+# the `_solib` directory the dispatcher is staged into rather than the oneMath
+# install, so without a search path every test aborts with
+# `oneapi::math::backend_not_found`. Put $ONEMATHROOT/lib on LD_LIBRARY_PATH.
+build:nvidia-gpu \
+    --dpc_math_backend=onemath
+test:nvidia-gpu \
+    --dpc_math_backend=onemath \
+    --test_env=LD_LIBRARY_PATH
+
 # Configuration: 'dev'
 # Fast local development build — compiles only for the host machine's
 # highest ISA (auto-detected). Use this to speed up iterative builds.

INSTALL.md

@@ -281,6 +281,24 @@ It is possible to integrate various sanitizers by specifying the REQSAN flag, av
 
             make -f makefile daal oneapi_c PLAT=lnx32e REQPROFILE=yes
 
+- To build the DPC++ device code against [oneMath](https://github.com/uxlfoundation/oneMath) instead of oneMKL, so that it can run on NVIDIA GPUs, add `DPC_MATH_BACKEND=onemath`:
+
+    _Note: experimental, and Linux x86-64 only. The library builds and links as a whole, but oneMath does not cover every domain oneDAL uses: sparse BLAS, and three of the five device RNG engines, throw `unimplemented` rather than running. See the "NVIDIA GPUs through oneMath" section of [the Bazel docs](https://github.com/uxlfoundation/oneDAL/tree/main/dev/bazel) for exactly what that affects and for the oneMath cmake recipe._
+
+    - Point `ONEMATHROOT` at a oneMath install built with the backends you need. Its `lib` directory is recorded as an rpath on `libonedal_dpc.so`, so no `LD_LIBRARY_PATH` is needed to link or run against the result:
+
+            export ONEMATHROOT=/path/to/onemath
+
+    - Ask the compiler for an NVPTX device target. `ONEDAL_SYCL_TARGETS` is passed to both the DPC++ compile and link as `-fsycl-targets=`; leaving it unset keeps the compiler's default (Intel SPIR-V), which is useful for checking the build without NVIDIA hardware:
+
+            export ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda
+
+    - Run `make` to build oneDAL:
+
+            make -f makefile oneapi_dpc PLAT=lnx32e DPC_MATH_BACKEND=onemath
+
+    `DPC_MATH_BACKEND` is independent of `BACKEND_CONFIG`, which selects the *host* math library: a build can use oneMKL on the CPU and oneMath on the GPU. Intermediate objects go to `__work_onemath` rather than `__work`, so switching back and forth does not need a `clean`.
+
 ---
 **NOTE:** Built libraries are located in the `__release_{os_name}[_{compiler_name}]/daal` directory.
 

MODULE.bazel

@@ -72,6 +72,12 @@ openblas_repo(name = "openblas", root_env_var = "OPENBLASROOT",)
 openrng_repo = use_repo_rule("@onedal//dev/bazel/deps:openrng.bzl", "openrng_repo")
 openrng_repo(name = "openrng", root_env_var = "OPENRNGROOT",)
 
+# Alternative SYCL math library, selected with `--dpc_math_backend=onemath`.
+# Has to be built locally with the wanted backends enabled, so there is no URL
+# to download from; the repository is only resolved when that flag is set.
+onemath_repo = use_repo_rule("@onedal//dev/bazel/deps:onemath.bzl", "onemath_repo")
+onemath_repo(name = "onemath", root_env_var = "ONEMATHROOT",)
+
 tbb_repo = use_repo_rule("@onedal//dev/bazel/deps:tbb.bzl", "tbb_repo")
 tbb_repo(
     name = "tbb",

cpp/oneapi/dal/BUILD

@@ -48,11 +48,27 @@ dal_module(
         "@onedal//cpp/daal:data_management",
     ],
     dpc_deps = [
-        "@mkl//:mkl_dpc",
+        ":math_backend_dpc",
         "@dpl//:headers",
     ],
 )
 
+# SYCL math library the device primitives are linked against, chosen by
+# `--dpc_math_backend`. Both targets expose the same oneMKL DPC++ interface and
+# carry the `defines` that `dal/backend/math_backend.hpp` switches on, so the
+# device sources are identical for either library.
+#
+# An `alias` rather than a `select()` in `dpc_deps` above: `dal_module` walks
+# its dependency lists at macro-evaluation time, where a `select` is still
+# opaque.
+alias(
+    name = "math_backend_dpc",
+    actual = select({
+        "@config//:dpc_math_backend_onemath": "@onemath//:onemath_dpc",
+        "//conditions:default": "@mkl//:mkl_dpc",
+    }),
+)
+
 dal_collect_modules(
     name = "core",
     root = "@onedal//cpp/oneapi/dal",

cpp/oneapi/dal/algo/kmeans/test/batch.cpp

@@ -318,6 +318,7 @@ TEMPLATE_LIST_TEST_M(kmeans_batch_test,
                      "[kmeans][batch][external-dataset]",
                      kmeans_types_csr) {
     SKIP_IF(!this->is_sparse_method());
+    SKIP_IF(!te::device_sparse_blas_supported());
     SKIP_IF(this->not_float64_friendly());
 
     using Float = std::tuple_element_t<0, TestType>;
@@ -393,6 +394,7 @@ TEMPLATE_LIST_TEST_M(kmeans_batch_test,
                      "[kmeans][batch]",
                      kmeans_types_csr) {
     SKIP_IF(!this->is_sparse_method());
+    SKIP_IF(!te::device_sparse_blas_supported());
     SKIP_IF(this->not_float64_friendly());
     SKIP_IF(this->get_policy().is_gpu());
     using Float = std::tuple_element_t<0, TestType>;
@@ -444,6 +446,7 @@ TEMPLATE_LIST_TEST_M(kmeans_batch_test,
                      "[kmeans][batch]",
                      kmeans_types_csr) {
     SKIP_IF(!this->is_sparse_method());
+    SKIP_IF(!te::device_sparse_blas_supported());
     SKIP_IF(this->not_float64_friendly());
     using Float = std::tuple_element_t<0, TestType>;
 
@@ -526,6 +529,7 @@ TEMPLATE_LIST_TEST_M(kmeans_batch_test,
                      "[kmeans][batch]",
                      kmeans_types_csr) {
     SKIP_IF(!this->is_sparse_method());
+    SKIP_IF(!te::device_sparse_blas_supported());
     SKIP_IF(this->not_float64_friendly());
     using Float = std::tuple_element_t<0, TestType>;
 
@@ -665,6 +669,7 @@ TEMPLATE_LIST_TEST_M(kmeans_batch_test,
                      kmeans_types_csr) {
     SKIP_IF(this->get_policy().is_cpu());
     SKIP_IF(!this->is_sparse_method());
+    SKIP_IF(!te::device_sparse_blas_supported());
     SKIP_IF(this->not_float64_friendly());
     using Float = std::tuple_element_t<0, TestType>;
 

cpp/oneapi/dal/algo/kmeans/test/fixture.hpp

@@ -35,8 +35,17 @@ namespace oneapi::dal::kmeans::test {
 namespace te = dal::test::engine;
 namespace la = dal::test::engine::linalg;
 
+// `lloyd_csr` on the device goes through the sparse BLAS primitives, which the
+// oneMath backend does not provide -- see `te::device_sparse_blas_supported()`.
+// The method is dropped from the mixed list rather than skipped case by case;
+// the CSR-only cases below carry a `SKIP_IF` instead, because a
+// `COMBINE_TYPES` list cannot be empty.
+#ifdef ONEDAL_MATH_BACKEND_ONEMATH
+using kmeans_types = COMBINE_TYPES((float, double), (kmeans::method::lloyd_dense));
+#else
 using kmeans_types = COMBINE_TYPES((float, double),
                                    (kmeans::method::lloyd_dense, kmeans::method::lloyd_csr));
+#endif
 
 using kmeans_types_csr = COMBINE_TYPES((float, double), (kmeans::method::lloyd_csr));
 

cpp/oneapi/dal/algo/logistic_regression/test/fixture.hpp

@@ -262,9 +262,17 @@ class log_reg_test : public te::crtp_algo_fixture<TestType, Derived> {
     table X_test_;
 };
 
+// `method::sparse` reaches the sparse BLAS primitives, which the oneMath
+// backend does not provide -- see `te::device_sparse_blas_supported()`.
+#ifdef ONEDAL_MATH_BACKEND_ONEMATH
+using log_reg_types = COMBINE_TYPES((float, double),
+                                    (logistic_regression::method::dense_batch),
+                                    (logistic_regression::task::classification));
+#else
 using log_reg_types = COMBINE_TYPES((float, double),
                                     (logistic_regression::method::dense_batch,
                                      logistic_regression::method::sparse),
                                     (logistic_regression::task::classification));
+#endif
 
 } // namespace oneapi::dal::logistic_regression::test

cpp/oneapi/dal/algo/pca/backend/gpu/misc.hpp

@@ -53,12 +53,12 @@ auto syevd_computation(sycl::queue& queue,
 
     sycl::event syevd_event;
     {
-        syevd_event = pr::syevd<mkl::job::vec, mkl::uplo::upper>(queue,
-                                                                 column_count,
-                                                                 corr,
-                                                                 lda,
-                                                                 eigenvalues,
-                                                                 { deps });
+        syevd_event = pr::syevd<pr::mkl::job::vec, pr::mkl::uplo::upper>(queue,
+                                                                         column_count,
+                                                                         corr,
+                                                                         lda,
+                                                                         eigenvalues,
+                                                                         { deps });
     }
     syevd_event.wait_and_throw();
     return std::make_tuple(eigenvalues, syevd_event);

cpp/oneapi/dal/algo/pca/backend/gpu/train_kernel_svd_impl_dpc.cpp

@@ -32,7 +32,7 @@ namespace oneapi::dal::pca::backend {
 
 namespace bk = dal::backend;
 namespace pr = dal::backend::primitives;
-namespace mkl = oneapi::mkl;
+namespace mkl = oneapi::dal::backend::math;
 using alloc = sycl::usm::alloc;
 
 using bk::context_gpu;

cpp/oneapi/dal/backend/math_backend.hpp

@@ -0,0 +1,60 @@
+/*******************************************************************************
+* Copyright contributors to the oneDAL project
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+*     http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing, software
+* distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions and
+* limitations under the License.
+*******************************************************************************/
+
+#pragma once
+
+/// Single point where the SYCL math library behind oneDAL's device primitives
+/// is chosen.
+///
+/// Two libraries implement the same oneMKL DPC++ interface that the BLAS,
+/// LAPACK, sparse BLAS and RNG primitives are written against:
+///
+///  * oneMKL (`<oneapi/mkl.hpp>`, namespace `oneapi::mkl`) — the default. Its
+///    device backend targets Intel GPUs only.
+///  * oneMath (`<oneapi/math.hpp>`, namespace `oneapi::math`) — the
+///    open-source implementation of the same specification. Besides the Intel
+///    backends it dispatches to cuBLAS, cuSOLVER, cuSPARSE and cuRAND, so the
+///    unmodified oneDAL device sources also run on NVIDIA GPUs once the
+///    compiler is asked for an `nvptx64-nvidia-cuda` target.
+///
+/// Both are reached through the `oneapi::dal::backend::math` alias, so the
+/// ~150 `mkl::` call sites stay untouched and the decision is made once, at
+/// build time, by defining `ONEDAL_MATH_BACKEND_ONEMATH`. With Bazel that
+/// define arrives from the `@onemath//:onemath_dpc` dependency selected by
+/// `--dpc_math_backend=onemath`.
+///
+/// This header intentionally contains no oneDAL types, so it can be included
+/// from `dal/detail` as well as from `dal/backend/primitives`.
+
+#ifdef ONEDAL_DATA_PARALLEL
+
+#ifdef ONEDAL_MATH_BACKEND_ONEMATH
+#include <oneapi/math.hpp>
+#else
+#include <oneapi/mkl.hpp>
+#endif
+
+namespace oneapi::dal::backend {
+
+#ifdef ONEDAL_MATH_BACKEND_ONEMATH
+namespace math = ::oneapi::math;
+#else
+namespace math = ::oneapi::mkl;
+#endif
+
+} // namespace oneapi::dal::backend
+
+#endif // ONEDAL_DATA_PARALLEL

cpp/oneapi/dal/backend/primitives/blas/gemm_dpc.cpp

@@ -18,7 +18,7 @@
 #include "oneapi/dal/backend/primitives/blas/gemm.hpp"
 #include "oneapi/dal/backend/primitives/blas/misc.hpp"
 
-#include <oneapi/mkl.hpp>
+#include "oneapi/dal/backend/math_backend.hpp"
 
 namespace oneapi::dal::backend::primitives {
 
@@ -43,38 +43,38 @@ sycl::event gemm(sycl::queue& queue,
 
     constexpr bool is_c_trans = (co == ndorder::c);
     if constexpr (is_c_trans) {
-        return mkl::blas::gemm(queue,
-                               f_order_as_transposed(bo),
-                               f_order_as_transposed(ao),
-                               c.get_dimension(1),
-                               c.get_dimension(0),
-                               a.get_dimension(1),
-                               alpha,
-                               b.get_data(),
-                               b.get_leading_stride(),
-                               a.get_data(),
-                               a.get_leading_stride(),
-                               beta,
-                               c.get_mutable_data(),
-                               c.get_leading_stride(),
-                               deps);
+        return mkl::blas::column_major::gemm(queue,
+                                             f_order_as_transposed(bo),
+                                             f_order_as_transposed(ao),
+                                             c.get_dimension(1),
+                                             c.get_dimension(0),
+                                             a.get_dimension(1),
+                                             alpha,
+                                             b.get_data(),
+                                             b.get_leading_stride(),
+                                             a.get_data(),
+                                             a.get_leading_stride(),
+                                             beta,
+                                             c.get_mutable_data(),
+                                             c.get_leading_stride(),
+                                             deps);
     }
     else {
-        return mkl::blas::gemm(queue,
-                               c_order_as_transposed(ao),
-                               c_order_as_transposed(bo),
-                               c.get_dimension(0),
-                               c.get_dimension(1),
-                               a.get_dimension(1),
-                               alpha,
-                               a.get_data(),
-                               a.get_leading_stride(),
-                               b.get_data(),
-                               b.get_leading_stride(),
-                               beta,
-                               c.get_mutable_data(),
-                               c.get_leading_stride(),
-                               deps);
+        return mkl::blas::column_major::gemm(queue,
+                                             c_order_as_transposed(ao),
+                                             c_order_as_transposed(bo),
+                                             c.get_dimension(0),
+                                             c.get_dimension(1),
+                                             a.get_dimension(1),
+                                             alpha,
+                                             a.get_data(),
+                                             a.get_leading_stride(),
+                                             b.get_data(),
+                                             b.get_leading_stride(),
+                                             beta,
+                                             c.get_mutable_data(),
+                                             c.get_leading_stride(),
+                                             deps);
     }
 }