diff --git a/.bazelrc b/.bazelrc index 21ca904cfe2..3226dabcd70 100644 --- a/.bazelrc +++ b/.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. diff --git a/INSTALL.md b/INSTALL.md index cd59c537a80..d3ea6a4aecd 100644 --- a/INSTALL.md +++ b/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. diff --git a/MODULE.bazel b/MODULE.bazel index c55fe156563..7bcc3ea10f8 100644 --- a/MODULE.bazel +++ b/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", diff --git a/cpp/oneapi/dal/BUILD b/cpp/oneapi/dal/BUILD index d161a82843d..af0026137f7 100644 --- a/cpp/oneapi/dal/BUILD +++ b/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", diff --git a/cpp/oneapi/dal/algo/kmeans/test/batch.cpp b/cpp/oneapi/dal/algo/kmeans/test/batch.cpp index 890cf382a2d..89c4d531b59 100644 --- a/cpp/oneapi/dal/algo/kmeans/test/batch.cpp +++ b/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>; diff --git a/cpp/oneapi/dal/algo/kmeans/test/fixture.hpp b/cpp/oneapi/dal/algo/kmeans/test/fixture.hpp index f646af7c1bb..27406eeaa8a 100644 --- a/cpp/oneapi/dal/algo/kmeans/test/fixture.hpp +++ b/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)); diff --git a/cpp/oneapi/dal/algo/logistic_regression/test/fixture.hpp b/cpp/oneapi/dal/algo/logistic_regression/test/fixture.hpp index 7fd7ed2edda..54a8d5ef52e 100644 --- a/cpp/oneapi/dal/algo/logistic_regression/test/fixture.hpp +++ b/cpp/oneapi/dal/algo/logistic_regression/test/fixture.hpp @@ -262,9 +262,17 @@ class log_reg_test : public te::crtp_algo_fixture { 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 diff --git a/cpp/oneapi/dal/algo/pca/backend/gpu/misc.hpp b/cpp/oneapi/dal/algo/pca/backend/gpu/misc.hpp index e503698ac7d..b547868abcf 100644 --- a/cpp/oneapi/dal/algo/pca/backend/gpu/misc.hpp +++ b/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(queue, - column_count, - corr, - lda, - eigenvalues, - { deps }); + syevd_event = pr::syevd(queue, + column_count, + corr, + lda, + eigenvalues, + { deps }); } syevd_event.wait_and_throw(); return std::make_tuple(eigenvalues, syevd_event); diff --git a/cpp/oneapi/dal/algo/pca/backend/gpu/train_kernel_svd_impl_dpc.cpp b/cpp/oneapi/dal/algo/pca/backend/gpu/train_kernel_svd_impl_dpc.cpp index 4cb2e739998..e6724d1f58d 100644 --- a/cpp/oneapi/dal/algo/pca/backend/gpu/train_kernel_svd_impl_dpc.cpp +++ b/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; diff --git a/cpp/oneapi/dal/backend/math_backend.hpp b/cpp/oneapi/dal/backend/math_backend.hpp new file mode 100644 index 00000000000..fc0fb88287b --- /dev/null +++ b/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 (``, namespace `oneapi::mkl`) — the default. Its +/// device backend targets Intel GPUs only. +/// * oneMath (``, 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 +#else +#include +#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 diff --git a/cpp/oneapi/dal/backend/primitives/blas/gemm_dpc.cpp b/cpp/oneapi/dal/backend/primitives/blas/gemm_dpc.cpp index 156603d2905..a3fb61faa5a 100644 --- a/cpp/oneapi/dal/backend/primitives/blas/gemm_dpc.cpp +++ b/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 +#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); } } diff --git a/cpp/oneapi/dal/backend/primitives/blas/gemv_dpc.cpp b/cpp/oneapi/dal/backend/primitives/blas/gemv_dpc.cpp index d13e51e1e00..8bbb9d0445f 100644 --- a/cpp/oneapi/dal/backend/primitives/blas/gemv_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/blas/gemv_dpc.cpp @@ -18,7 +18,7 @@ #include "oneapi/dal/backend/primitives/blas/gemv.hpp" #include "oneapi/dal/backend/primitives/blas/misc.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { @@ -45,35 +45,35 @@ sycl::event gemv(sycl::queue& queue, if constexpr (ao == ndorder::c) { ONEDAL_ASSERT(lda >= n); - return mkl::blas::gemv(queue, - mkl::transpose::trans, - n, - m, - alpha, - a.get_data(), - lda, - x.get_data(), - std::int64_t(1), - beta, - y.get_mutable_data(), - std::int64_t(1), - deps); + return mkl::blas::column_major::gemv(queue, + mkl::transpose::trans, + n, + m, + alpha, + a.get_data(), + lda, + x.get_data(), + std::int64_t(1), + beta, + y.get_mutable_data(), + std::int64_t(1), + deps); } else { ONEDAL_ASSERT(lda >= m); - return mkl::blas::gemv(queue, - mkl::transpose::nontrans, - m, - n, - alpha, - a.get_data(), - lda, - x.get_data(), - std::int64_t(1), - beta, - y.get_mutable_data(), - std::int64_t(1), - deps); + return mkl::blas::column_major::gemv(queue, + mkl::transpose::nontrans, + m, + n, + alpha, + a.get_data(), + lda, + x.get_data(), + std::int64_t(1), + beta, + y.get_mutable_data(), + std::int64_t(1), + deps); } } diff --git a/cpp/oneapi/dal/backend/primitives/blas/misc.hpp b/cpp/oneapi/dal/backend/primitives/blas/misc.hpp index 64f9b70fece..a6f2a5844fb 100644 --- a/cpp/oneapi/dal/backend/primitives/blas/misc.hpp +++ b/cpp/oneapi/dal/backend/primitives/blas/misc.hpp @@ -18,11 +18,11 @@ #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; #ifdef ONEDAL_DATA_PARALLEL /// Convert oneDAL `ndorder` to oneMKL `layout` diff --git a/cpp/oneapi/dal/backend/primitives/blas/syrk.hpp b/cpp/oneapi/dal/backend/primitives/blas/syrk.hpp index c254eddaadd..9e3f876ea3b 100644 --- a/cpp/oneapi/dal/backend/primitives/blas/syrk.hpp +++ b/cpp/oneapi/dal/backend/primitives/blas/syrk.hpp @@ -23,7 +23,7 @@ namespace oneapi::dal::backend::primitives { #ifdef ONEDAL_DATA_PARALLEL -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; template sycl::event syrk(sycl::queue& queue, diff --git a/cpp/oneapi/dal/backend/primitives/blas/syrk_dpc.cpp b/cpp/oneapi/dal/backend/primitives/blas/syrk_dpc.cpp index dc883a3e77f..3e25324142d 100644 --- a/cpp/oneapi/dal/backend/primitives/blas/syrk_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/blas/syrk_dpc.cpp @@ -18,7 +18,7 @@ #include "oneapi/dal/backend/primitives/blas/syrk.hpp" #include "oneapi/dal/backend/primitives/blas/misc.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { @@ -40,7 +40,18 @@ static sycl::event syrk_wrapper(sycl::queue& queue, ONEDAL_ASSERT(is_trans || lda >= n); ONEDAL_ASSERT(!is_trans || lda >= k); - return mkl::blas::syrk(queue, uplo, trans, n, k, alpha, a, lda, beta, c, ldc, deps); + return mkl::blas::column_major::syrk(queue, + uplo, + trans, + n, + k, + alpha, + a, + lda, + beta, + c, + ldc, + deps); } template diff --git a/cpp/oneapi/dal/backend/primitives/lapack/gesvd.hpp b/cpp/oneapi/dal/backend/primitives/lapack/gesvd.hpp index 311f2a5a3c2..f28714ab609 100644 --- a/cpp/oneapi/dal/backend/primitives/lapack/gesvd.hpp +++ b/cpp/oneapi/dal/backend/primitives/lapack/gesvd.hpp @@ -24,7 +24,7 @@ namespace oneapi::dal::backend::primitives { #ifdef ONEDAL_DATA_PARALLEL -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; template sycl::event gesvd(sycl::queue& queue, diff --git a/cpp/oneapi/dal/backend/primitives/lapack/gesvd_dpc.cpp b/cpp/oneapi/dal/backend/primitives/lapack/gesvd_dpc.cpp index 9602d5afeb6..4f0a8cc4b83 100644 --- a/cpp/oneapi/dal/backend/primitives/lapack/gesvd_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/lapack/gesvd_dpc.cpp @@ -18,7 +18,7 @@ #include "oneapi/dal/backend/primitives/lapack/gesvd.hpp" #include "oneapi/dal/backend/primitives/blas/misc.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { diff --git a/cpp/oneapi/dal/backend/primitives/lapack/misc.hpp b/cpp/oneapi/dal/backend/primitives/lapack/misc.hpp index 7893d7523b9..cefffdf4cf3 100644 --- a/cpp/oneapi/dal/backend/primitives/lapack/misc.hpp +++ b/cpp/oneapi/dal/backend/primitives/lapack/misc.hpp @@ -18,11 +18,11 @@ #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; inline constexpr mkl::job ident_job(mkl::job order) { constexpr auto novec = mkl::job::novec; diff --git a/cpp/oneapi/dal/backend/primitives/lapack/syevd.hpp b/cpp/oneapi/dal/backend/primitives/lapack/syevd.hpp index dbed4c9f84c..183a77ea5c0 100644 --- a/cpp/oneapi/dal/backend/primitives/lapack/syevd.hpp +++ b/cpp/oneapi/dal/backend/primitives/lapack/syevd.hpp @@ -24,7 +24,7 @@ namespace oneapi::dal::backend::primitives { #ifdef ONEDAL_DATA_PARALLEL -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; template sycl::event syevd(sycl::queue& queue, diff --git a/cpp/oneapi/dal/backend/primitives/lapack/syevd_dpc.cpp b/cpp/oneapi/dal/backend/primitives/lapack/syevd_dpc.cpp index cfd0be634de..6f3ddd772a1 100644 --- a/cpp/oneapi/dal/backend/primitives/lapack/syevd_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/lapack/syevd_dpc.cpp @@ -18,7 +18,7 @@ #include "oneapi/dal/backend/primitives/lapack/syevd.hpp" #include "oneapi/dal/backend/primitives/blas/misc.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { diff --git a/cpp/oneapi/dal/backend/primitives/objective_function/test/logloss_dpc.cpp b/cpp/oneapi/dal/backend/primitives/objective_function/test/logloss_dpc.cpp index 1bd51dfc14b..328c01f5462 100644 --- a/cpp/oneapi/dal/backend/primitives/objective_function/test/logloss_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/objective_function/test/logloss_dpc.cpp @@ -48,6 +48,7 @@ TEMPLATE_LIST_TEST_M(logloss_test, "test random input with L1", "[logloss]", log TEMPLATE_LIST_TEST_M(logloss_test, "sparse data test without L2", "[logloss]", logloss_types) { SKIP_IF(this->not_float64_friendly()); SKIP_IF(this->get_policy().is_cpu()); + SKIP_IF(!te::device_sparse_blas_supported()); this->generate_sparse_input(); this->run_sparse_test(0.0f, this->fit_intercept_); } @@ -55,6 +56,7 @@ TEMPLATE_LIST_TEST_M(logloss_test, "sparse data test without L2", "[logloss]", l TEMPLATE_LIST_TEST_M(logloss_test, "sparse data test", "[logloss]", logloss_types) { SKIP_IF(this->not_float64_friendly()); SKIP_IF(this->get_policy().is_cpu()); + SKIP_IF(!te::device_sparse_blas_supported()); this->generate_sparse_input(); this->run_sparse_test(1.3f, this->fit_intercept_); } diff --git a/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp b/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp index a49faaa3803..45d113f7fcd 100644 --- a/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp +++ b/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp @@ -27,13 +27,25 @@ #include #include -#include +#include "oneapi/dal/backend/math_backend.hpp" -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; namespace oneapi::dal::backend::primitives { #ifdef ONEDAL_DATA_PARALLEL +/// Not every math backend provides every engine below: oneMath declares +/// `philox4x32x10` and `mrg32k3a` only, so under `ONEDAL_MATH_BACKEND_ONEMATH` +/// the `mt2203`, `mt19937` and `mcg59` wrappers hold no device generator. +/// +/// They stay constructible rather than being compiled out, because +/// `device_engine` always builds a host engine and a device engine as a pair, +/// and the host-side entry points -- `shuffle`, `uniform_without_replacement` +/// and `partial_fisher_yates_shuffle`, all of which draw from the DAAL engine -- +/// never touch the device one. Keeping the object lets those keep working with +/// any engine type; only a device-side `generate` on an absent engine throws, +/// from `generate_rng` in device_engine_dpc.cpp. + /// Abstract base class for all random number generators (RNGs). /// It defines a common interface for working with different types of RNGs, including methods /// for retrieving the engine method and skipping ahead in the random number sequence. @@ -60,8 +72,12 @@ class gen_mt2203 : public gen_base { /// Constructor that initializes the mt2203 generator for use on the GPU. /// @param[in] queue The SYCL queue to manage device operations. /// @param[in] seed The initial seed for the generator. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + gen_mt2203(sycl::queue, std::int64_t, std::int64_t = 0) {} +#else gen_mt2203(sycl::queue queue, std::int64_t seed, std::int64_t engine_idx = 0) : _gen(queue, seed, engine_idx) {} +#endif /// Returns the engine method for mt2203. /// @return The `mt2203` engine method as an enum value of `engine_type_internal`. @@ -76,14 +92,16 @@ class gen_mt2203 : public gen_base { //skip; } +#ifndef ONEDAL_MATH_BACKEND_ONEMATH /// Retrieves a pointer to the underlying mt2203 generator. /// @return A pointer to the `mt2203` RNG. - oneapi::mkl::rng::mt2203* get() { + mkl::rng::mt2203* get() { return &_gen; } protected: - oneapi::mkl::rng::mt2203 _gen; + mkl::rng::mt2203 _gen; +#endif }; /// Implementation of the philox4x32x10 random number generator for GPU. @@ -112,12 +130,12 @@ class gen_philox : public gen_base { /// Retrieves a pointer to the underlying philox4x32x10 generator. /// @return A pointer to the `philox4x32x10` RNG. - oneapi::mkl::rng::philox4x32x10* get() { + mkl::rng::philox4x32x10* get() { return &_gen; } protected: - oneapi::mkl::rng::philox4x32x10 _gen; + mkl::rng::philox4x32x10 _gen; }; /// Implementation of the mrg32k3a random number generator for GPU. @@ -146,12 +164,12 @@ class gen_mrg32k : public gen_base { /// Retrieves a pointer to the underlying mrg32k3a generator. /// @return A pointer to the `mrg32k3a` RNG. - oneapi::mkl::rng::mrg32k3a* get() { + mkl::rng::mrg32k3a* get() { return &_gen; } protected: - oneapi::mkl::rng::mrg32k3a _gen; + mkl::rng::mrg32k3a _gen; }; /// Implementation of the mt19937 random number generator for GPU. @@ -164,7 +182,11 @@ class gen_mt19937 : public gen_base { /// Constructor that initializes the mt19937 generator for use on the GPU. /// @param[in] queue The SYCL queue to manage device operations. /// @param[in] seed The initial seed for the generator. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + gen_mt19937(sycl::queue, std::int64_t) {} +#else gen_mt19937(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} +#endif /// Returns the engine method for mt19937. /// @return The `mt19937` engine method as an enum value of `engine_type_internal`. @@ -174,18 +196,22 @@ class gen_mt19937 : public gen_base { /// Skips ahead in the random number sequence for mt19937 on the GPU. /// @param[in] nSkip The number of steps to skip in the sequence. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + void skip_ahead_gpu(std::int64_t) override {} +#else void skip_ahead_gpu(std::int64_t nSkip) override { skip_ahead(_gen, nSkip); } /// Retrieves a pointer to the underlying mt19937 generator. /// @return A pointer to the `mt19937` RNG. - oneapi::mkl::rng::mt19937* get() { + mkl::rng::mt19937* get() { return &_gen; } protected: - oneapi::mkl::rng::mt19937 _gen; + mkl::rng::mt19937 _gen; +#endif }; /// Implementation of the mcg59 random number generator for GPU. @@ -198,7 +224,11 @@ class gen_mcg59 : public gen_base { /// Constructor that initializes the mcg59 generator for use on the GPU. /// @param[in] queue The SYCL queue to manage device operations. /// @param[in] seed The initial seed for the generator. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + gen_mcg59(sycl::queue, std::int64_t) {} +#else gen_mcg59(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} +#endif /// Returns the engine method for mcg59. /// @return The `mcg59` engine method as an enum value of `engine_type_internal`. @@ -208,18 +238,22 @@ class gen_mcg59 : public gen_base { /// Skips ahead in the random number sequence for mcg59 on the GPU. /// @param[in] nSkip The number of steps to skip in the sequence. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + void skip_ahead_gpu(std::int64_t) override {} +#else void skip_ahead_gpu(std::int64_t nSkip) override { skip_ahead(_gen, nSkip); } /// Retrieves a pointer to the underlying mcg59 generator. /// @return A pointer to the `mcg59` RNG. - oneapi::mkl::rng::mcg59* get() { + mkl::rng::mcg59* get() { return &_gen; } protected: - oneapi::mkl::rng::mcg59 _gen; + mkl::rng::mcg59 _gen; +#endif }; /// A class that provides a unified interface for random number generation on both CPU and GPU devices. diff --git a/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp b/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp index 88b70de7d06..3b66ad84a04 100644 --- a/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp @@ -17,7 +17,7 @@ #include "oneapi/dal/backend/primitives/rng/device_engine.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { @@ -34,28 +34,40 @@ sycl::event generate_rng(Distribution& distr, if (engine_type == engine_type_internal::philox4x32x10) { auto& device_engine = *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); - return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); - } - else if (engine_type == engine_type_internal::mt19937) { - auto& device_engine = - *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); - return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + return mkl::rng::generate(distr, device_engine, count, dst, deps); } else if (engine_type == engine_type_internal::mrg32k3a) { auto& device_engine = *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); - return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + return mkl::rng::generate(distr, device_engine, count, dst, deps); + } +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + // The wrappers for these three exist -- so a `device_engine` carrying one is + // still usable through the host-side rng paths -- but oneMath declares no + // such device engine to hand to `generate`. See device_engine.hpp. + else if (engine_type == engine_type_internal::mt19937 || + engine_type == engine_type_internal::mcg59 || + engine_type == engine_type_internal::mt2203) { + throw unimplemented( + dal::detail::error_messages::rng_engine_not_available_in_math_backend()); + } +#else + else if (engine_type == engine_type_internal::mt19937) { + auto& device_engine = + *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return mkl::rng::generate(distr, device_engine, count, dst, deps); } else if (engine_type == engine_type_internal::mcg59) { auto& device_engine = *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); - return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + return mkl::rng::generate(distr, device_engine, count, dst, deps); } else if (engine_type == engine_type_internal::mt2203) { auto& device_engine = *(static_cast(engine_.get_device_engine_base_ptr().get()))->get(); - return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + return mkl::rng::generate(distr, device_engine, count, dst, deps); } +#endif else { throw std::runtime_error("Unsupported engine type in generate_rng"); } @@ -81,7 +93,7 @@ sycl::event uniform(sycl::queue& queue, if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == sycl::usm::alloc::host) { throw domain_error(dal::detail::error_messages::unsupported_data_type()); } - oneapi::mkl::rng::uniform distr(a, b); + mkl::rng::uniform distr(a, b); engine_.skip_ahead_cpu(count); auto event = generate_rng(distr, engine_, count, dst, deps); return event; diff --git a/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp b/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp index f1e08226b9a..bb4c5b8a7de 100644 --- a/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp @@ -110,7 +110,16 @@ class rng_test : public te::float_algo_fixture } }; +// Every test below draws on the device, and oneMath declares no `mt2203`, +// `mt19937` or `mcg59` device engine, so under that backend those three throw +// `unimplemented` out of `generate_rng`. Narrow the lists rather than skipping +// the whole suite: `philox4x32x10` and `mrg32k3a` are still worth covering, and +// they are the engines the algorithms actually default to. +#ifdef ONEDAL_MATH_BACKEND_ONEMATH +using rng_types = COMBINE_TYPES((float, double), (mrg32k3a, philox4x32x10)); +#else using rng_types = COMBINE_TYPES((float, double), (mt2203, mt19937, mcg59, mrg32k3a, philox4x32x10)); +#endif TEMPLATE_LIST_TEST_M(rng_test, "rng cpu vs gpu", "[rng]", rng_types) { SKIP_IF(this->get_policy().is_cpu()); @@ -135,8 +144,12 @@ TEMPLATE_LIST_TEST_M(rng_test, "rng cpu vs gpu", "[rng]", rng_types) { this->check_results(arr_gpu, arr_host); } +#ifdef ONEDAL_MATH_BACKEND_ONEMATH +using rng_types_skip_ahead_support = COMBINE_TYPES((float, double), (mrg32k3a, philox4x32x10)); +#else using rng_types_skip_ahead_support = COMBINE_TYPES((float, double), (mt19937, mcg59, mrg32k3a, philox4x32x10)); +#endif TEMPLATE_LIST_TEST_M(rng_test, "mixed rng cpu skip", "[rng]", rng_types_skip_ahead_support) { SKIP_IF(this->get_policy().is_cpu()); diff --git a/cpp/oneapi/dal/backend/primitives/sparse_blas/BUILD b/cpp/oneapi/dal/backend/primitives/sparse_blas/BUILD index bcd89758acc..cf9212a046c 100644 --- a/cpp/oneapi/dal/backend/primitives/sparse_blas/BUILD +++ b/cpp/oneapi/dal/backend/primitives/sparse_blas/BUILD @@ -14,11 +14,20 @@ dal_module( ], ) +# Deselected under `--dpc_math_backend=onemath`, where the whole domain throws +# `unimplemented`: oneMath ships the newer sparse specification (`init_csr_matrix` +# plus descriptor-driven `spmv`/`spmm`) rather than the handle API these +# primitives are written against. `@platforms//:incompatible` makes Bazel skip +# the targets instead of running them to a certain failure. dal_test_suite( name = "tests", framework = "catch2", compile_as = [ "dpc++" ], private = True, + target_compatible_with = select({ + "@config//:dpc_math_backend_onemath": ["@platforms//:incompatible"], + "//conditions:default": [], + }), hdrs = glob([ "test/*.hpp", ]), diff --git a/cpp/oneapi/dal/backend/primitives/sparse_blas/gemm_dpc.cpp b/cpp/oneapi/dal/backend/primitives/sparse_blas/gemm_dpc.cpp index 48965de5c5a..5f9dc58120e 100644 --- a/cpp/oneapi/dal/backend/primitives/sparse_blas/gemm_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/sparse_blas/gemm_dpc.cpp @@ -31,6 +31,11 @@ sycl::event gemm(sycl::queue& queue, ONEDAL_ASSERT(b.get_dimension(1) == c.get_dimension(1)); ONEDAL_ASSERT(c.has_mutable_data()); +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + // See the note in gemv_dpc.cpp: oneMath's `spmm` is a different interface + // generation, not a different spelling of this call. + throw unimplemented(dal::detail::error_messages::sparse_blas_not_available_in_math_backend()); +#else if (co == ndorder::c) { return mkl::sparse::gemm(queue, order_as_layout(co), @@ -63,6 +68,7 @@ sycl::event gemm(sycl::queue& queue, } ONEDAL_ASSERT(false); return sycl::event(); +#endif // ONEDAL_MATH_BACKEND_ONEMATH } #define INSTANTIATE(F, bo, co) \ diff --git a/cpp/oneapi/dal/backend/primitives/sparse_blas/gemv_dpc.cpp b/cpp/oneapi/dal/backend/primitives/sparse_blas/gemv_dpc.cpp index 74fb48d5035..b563a07c89a 100644 --- a/cpp/oneapi/dal/backend/primitives/sparse_blas/gemv_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/sparse_blas/gemv_dpc.cpp @@ -31,6 +31,13 @@ sycl::event gemv(sycl::queue& queue, ONEDAL_ASSERT(x.has_data()); ONEDAL_ASSERT(y.has_mutable_data()); +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + // oneMath implements the newer sparse specification -- `spmv` over a + // descriptor with separate buffer-size, optimize and execute stages -- which + // is a different shape from the handle API used here, not a renaming. The + // port is deliberately left out of the change that added the backend. + throw unimplemented(dal::detail::error_messages::sparse_blas_not_available_in_math_backend()); +#else return mkl::sparse::gemv(queue, transpose_to_mkl(transpose_a), alpha, @@ -39,6 +46,7 @@ sycl::event gemv(sycl::queue& queue, beta, y.get_mutable_data(), dependencies); +#endif // ONEDAL_MATH_BACKEND_ONEMATH } #define INSTANTIATE(F) \ diff --git a/cpp/oneapi/dal/backend/primitives/sparse_blas/misc.hpp b/cpp/oneapi/dal/backend/primitives/sparse_blas/misc.hpp index bcc373a12f4..34943b9ebea 100644 --- a/cpp/oneapi/dal/backend/primitives/sparse_blas/misc.hpp +++ b/cpp/oneapi/dal/backend/primitives/sparse_blas/misc.hpp @@ -18,12 +18,11 @@ #include "oneapi/dal/table/common.hpp" -#include -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::backend::primitives { -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; /// Convert oneDAL `sparse_indexing` to oneMKL `index_base` inline constexpr mkl::index_base sparse_indexing_to_mkl(const sparse_indexing indexing) { diff --git a/cpp/oneapi/dal/backend/primitives/sparse_blas/set_csr_data_dpc.cpp b/cpp/oneapi/dal/backend/primitives/sparse_blas/set_csr_data_dpc.cpp index 998efac1c2c..c354d9d9ba9 100644 --- a/cpp/oneapi/dal/backend/primitives/sparse_blas/set_csr_data_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/sparse_blas/set_csr_data_dpc.cpp @@ -35,27 +35,32 @@ sycl::event set_csr_data(sycl::queue &queue, ONEDAL_ASSERT(data.get_count()); ONEDAL_ASSERT(column_indices.get_count()); ONEDAL_ASSERT(row_offsets.get_count() == row_count + 1); -#if INTEL_MKL_VERSION >= 20250300 - return oneapi::mkl::sparse::set_csr_data(queue, - dal::detail::get_impl(handle).get(), - row_count, - column_count, - nnz, - sparse_indexing_to_mkl(indexing), - const_cast(row_offsets.get_data()), - const_cast(column_indices.get_data()), - const_cast(data.get_data()), - deps); +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + // oneMath creates the handle from the CSR arrays in one step + // (`init_csr_matrix`) instead of filling an already-created one, so there is + // no equivalent of this call. See the note in gemv_dpc.cpp. + throw unimplemented(dal::detail::error_messages::sparse_blas_not_available_in_math_backend()); +#elif INTEL_MKL_VERSION >= 20250300 + return mkl::sparse::set_csr_data(queue, + dal::detail::get_impl(handle).get(), + row_count, + column_count, + nnz, + sparse_indexing_to_mkl(indexing), + const_cast(row_offsets.get_data()), + const_cast(column_indices.get_data()), + const_cast(data.get_data()), + deps); #else - return oneapi::mkl::sparse::set_csr_data(queue, - dal::detail::get_impl(handle).get(), - row_count, - column_count, - sparse_indexing_to_mkl(indexing), - const_cast(row_offsets.get_data()), - const_cast(column_indices.get_data()), - const_cast(data.get_data()), - deps); + return mkl::sparse::set_csr_data(queue, + dal::detail::get_impl(handle).get(), + row_count, + column_count, + sparse_indexing_to_mkl(indexing), + const_cast(row_offsets.get_data()), + const_cast(column_indices.get_data()), + const_cast(data.get_data()), + deps); #endif } @@ -73,27 +78,29 @@ sycl::event set_csr_data(sycl::queue &queue, ONEDAL_ASSERT(data); ONEDAL_ASSERT(column_indices); ONEDAL_ASSERT(row_offsets); -#if INTEL_MKL_VERSION >= 20250300 - return oneapi::mkl::sparse::set_csr_data(queue, - dal::detail::get_impl(handle).get(), - row_count, - column_count, - nnz, - sparse_indexing_to_mkl(indexing), - const_cast(row_offsets), - const_cast(column_indices), - const_cast(data), - deps); +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + throw unimplemented(dal::detail::error_messages::sparse_blas_not_available_in_math_backend()); +#elif INTEL_MKL_VERSION >= 20250300 + return mkl::sparse::set_csr_data(queue, + dal::detail::get_impl(handle).get(), + row_count, + column_count, + nnz, + sparse_indexing_to_mkl(indexing), + const_cast(row_offsets), + const_cast(column_indices), + const_cast(data), + deps); #else - return oneapi::mkl::sparse::set_csr_data(queue, - dal::detail::get_impl(handle).get(), - row_count, - column_count, - sparse_indexing_to_mkl(indexing), - const_cast(row_offsets), - const_cast(column_indices), - const_cast(data), - deps); + return mkl::sparse::set_csr_data(queue, + dal::detail::get_impl(handle).get(), + row_count, + column_count, + sparse_indexing_to_mkl(indexing), + const_cast(row_offsets), + const_cast(column_indices), + const_cast(data), + deps); #endif } diff --git a/cpp/oneapi/dal/detail/error_messages.cpp b/cpp/oneapi/dal/detail/error_messages.cpp index aa586a42bfb..822a3240477 100644 --- a/cpp/oneapi/dal/detail/error_messages.cpp +++ b/cpp/oneapi/dal/detail/error_messages.cpp @@ -55,6 +55,14 @@ MSG(invalid_column_indices_block_size, "Invalid column indices block size") MSG(method_not_implemented, "Method is not implemented") MSG(negative_integral_value_conversion_to_unsigned, "Negative integral value conversion to unsigned") +MSG(rng_engine_not_available_in_math_backend, + "The requested RNG engine is not provided by the math backend this build of " + "oneDAL was linked against. Consider philox4x32x10 or mrg32k3a, or rebuild " + "with the default oneMKL device backend.") +MSG(sparse_blas_not_available_in_math_backend, + "Sparse BLAS is not provided by the math backend this build of oneDAL was " + "linked against, so sparse input cannot be processed on the device. Rebuild " + "with the default oneMKL device backend.") MSG(unsupported_feature_type, "Feature type is not supported") MSG(unknown_memcpy_error, "Unknown error during memory copying") MSG(unknown_usm_pointer_type, "USM pointer type is unknown in the current context") diff --git a/cpp/oneapi/dal/detail/error_messages.hpp b/cpp/oneapi/dal/detail/error_messages.hpp index 36d0967e1e7..428a3e8c666 100644 --- a/cpp/oneapi/dal/detail/error_messages.hpp +++ b/cpp/oneapi/dal/detail/error_messages.hpp @@ -72,6 +72,8 @@ class ONEDAL_EXPORT error_messages { MSG(invalid_column_indices_block_size); MSG(method_not_implemented); MSG(negative_integral_value_conversion_to_unsigned); + MSG(rng_engine_not_available_in_math_backend); + MSG(sparse_blas_not_available_in_math_backend); MSG(only_homogen_table_is_supported); MSG(overflow_found_in_multiplication_of_two_values); MSG(overflow_found_in_sum_of_two_values); diff --git a/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.cpp b/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.cpp index 1c33c352c21..3a4408472a3 100644 --- a/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.cpp +++ b/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.cpp @@ -22,6 +22,20 @@ namespace v1 { #ifdef ONEDAL_DATA_PARALLEL +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + +// oneMath has no empty-handle-then-fill flow: its `init_csr_matrix` takes the +// CSR arrays and produces a handle in one step, so there is nothing to create +// here and nothing to release. The handle stays null, and every sparse entry +// point that would use it throws -- see the `sparse_blas` primitives. +sparse_matrix_handle_impl::sparse_matrix_handle_impl(sycl::queue& queue) + : handle_(nullptr), + queue_(queue) {} + +sparse_matrix_handle_impl::~sparse_matrix_handle_impl() = default; + +#else + sparse_matrix_handle_impl::sparse_matrix_handle_impl(sycl::queue& queue) : queue_(queue) { mkl::sparse::init_matrix_handle(&handle_); } @@ -30,6 +44,8 @@ sparse_matrix_handle_impl::~sparse_matrix_handle_impl() { mkl::sparse::release_matrix_handle(queue_, &handle_, {}).wait(); } +#endif // ONEDAL_MATH_BACKEND_ONEMATH + #endif // ONEDAL_DATA_PARALLEL } // namespace v1 diff --git a/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.hpp b/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.hpp index d644d89dae5..e14d2d17254 100644 --- a/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.hpp +++ b/cpp/oneapi/dal/detail/sparse_matrix_handle_impl.hpp @@ -18,13 +18,13 @@ #ifdef ONEDAL_DATA_PARALLEL -#include +#include "oneapi/dal/backend/math_backend.hpp" namespace oneapi::dal::detail { namespace v1 { -namespace mkl = oneapi::mkl; +namespace mkl = oneapi::dal::backend::math; /// Class that hides the implementation details of the `backend::primitives::sparse_matrix_handle` class class sparse_matrix_handle_impl { @@ -46,7 +46,13 @@ class sparse_matrix_handle_impl { private: mkl::sparse::matrix_handle_t handle_; +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + // Kept so the class has one shape for both backends, but the oneMath path + // neither creates nor releases a handle, so nothing there reads it. + [[maybe_unused]] sycl::queue& queue_; +#else sycl::queue& queue_; +#endif }; } // namespace v1 diff --git a/cpp/oneapi/dal/test/engine/fixtures.hpp b/cpp/oneapi/dal/test/engine/fixtures.hpp index 4516f7ae349..b761ffdb9c6 100644 --- a/cpp/oneapi/dal/test/engine/fixtures.hpp +++ b/cpp/oneapi/dal/test/engine/fixtures.hpp @@ -23,6 +23,25 @@ namespace oneapi::dal::test::engine { +/// Whether the device math backend provides the sparse BLAS domain. +/// +/// oneMath ships the newer sparse specification (`init_csr_matrix` plus +/// descriptor-driven `spmv`/`spmm`) rather than the handle API oneDAL's sparse +/// primitives are written against, so under that backend every device-side +/// sparse path throws `unimplemented`. Tests that hand a `csr_table` to the +/// device therefore have to be skipped rather than run to a certain failure. +/// The define is set only for the device translation units, so the host +/// versions of the same tests keep running. +/// +/// @return True unless the library was built against oneMath. +inline constexpr bool device_sparse_blas_supported() { +#ifdef ONEDAL_MATH_BACKEND_ONEMATH + return false; +#else + return true; +#endif +} + class policy_fixture { public: auto& get_policy() { diff --git a/dev/AGENTS.md b/dev/AGENTS.md index c29833a485b..b92964dc6d1 100644 --- a/dev/AGENTS.md +++ b/dev/AGENTS.md @@ -26,7 +26,8 @@ make -f makefile daal oneapi_c PLAT=lnx32e -j$(nproc) - `PLAT`: `lnx32e`, `win32e`, `mac32e`, `lnxarm`, `winarm`, `lnxriscv64`. - `COMPILER`: `icx` (x86-64 default), `gnu`, `clang`, `vc`; the allowed set per platform is in `make/function_definitions/`. - `REQCPU`: subset of `sse2 avx2 avx512` on x86-64, `sve` on ARM, `rv64` on RISC-V. -- `BACKEND_CONFIG`: `mkl` (default on x86-64) or `ref` (OpenBLAS; default on ARM and RISC-V). +- `BACKEND_CONFIG`: `mkl` (default on x86-64) or `ref` (OpenBLAS; default on ARM and RISC-V). Names the host math library only. +- `DPC_MATH_BACKEND`: `mkl` (default) or `onemath`, the math library behind the DPC++ device code. `onemath` additionally reaches NVIDIA GPUs and needs `ONEMATHROOT`; it mirrors Bazel's `--dpc_math_backend` and is independent of `BACKEND_CONFIG`. ## Bazel diff --git a/dev/bazel/README.md b/dev/bazel/README.md index 9c0a86957b5..09278cff722 100644 --- a/dev/bazel/README.md +++ b/dev/bazel/README.md @@ -823,3 +823,247 @@ validated on ARM), also set `--rng_backend=openrng` and point `OPENRNGROOT` at a OpenRNG install built with `.ci/env/openrng.sh`. `--rng_backend=openrng` only takes effect together with `--backend_config=ref`; it is silently ignored on the MKL backend, whose build never compiles the ref RNG shim that OpenRNG replaces. + +## NVIDIA GPUs through oneMath (experimental) + +The DPC++ device sources are written against the oneMKL DPC++ interface but use +no Intel-specific SYCL extension: sub-group sizes are queried at run time and +there is no ESIMD, no `sycl::ext::intel`, and no hard-coded device selector. +The only Intel-specific piece is the math library itself, so pointing the build +at [oneMath](https://github.com/uxlfoundation/oneMath) — the open-source +implementation of the same interface, which additionally dispatches to cuBLAS, +cuSOLVER, cuSPARSE and cuRAND — lets the same kernels run on an NVIDIA GPU, as +far as oneMath implements the interface the kernels use. That is BLAS, LAPACK +and two of the five device RNG engines, but not sparse BLAS; see the table below +for what that leaves out. + +Which library is used is decided by `--dpc_math_backend`: + +| Value | Library | Device support | +|---------------------|---------|--------------------------------------| +| `mkl` _(default)_ | oneMKL | Intel GPUs | +| `onemath` | oneMath | Intel GPUs plus NVIDIA GPUs, without sparse BLAS | + +Note that `onemath` is not an NVIDIA-only setting: oneMath's own `mklgpu` +backend targets Intel GPUs, and a oneMath build with no CUDA backend and no +`ONEDAL_SYCL_TARGETS` is a perfectly ordinary Intel GPU build. That is how this +configuration is tested here, since the two gaps below are interface gaps rather +than device ones and show up on any target. + +This is orthogonal to `--backend_config`, which selects the *host* math library +(`mkl` or `ref`); a DPC++ build keeps using whatever that says on the CPU side. +Only `cpp/oneapi/dal/backend/math_backend.hpp` knows the difference — it maps +the `mkl::` namespace used by ~150 call sites onto `oneapi::mkl` or +`oneapi::math`, driven by the `ONEDAL_MATH_BACKEND_ONEMATH` define that the +selected dependency carries. + +Two things have to be provided: + +1. **oneMath**, built with the backends you want, pointed at by `ONEMATHROOT`. + There is no redistributable to download because the backend set is fixed at + oneMath configure time: + + ```sh + cmake -B build -S . -GNinja \ + -DCMAKE_CXX_COMPILER=icpx \ + -DENABLE_MKLGPU_BACKEND=OFF \ + -DENABLE_CUBLAS_BACKEND=ON \ + -DENABLE_CUSOLVER_BACKEND=ON \ + -DENABLE_CURAND_BACKEND=ON \ + -DENABLE_CUSPARSE_BACKEND=ON \ + -DCMAKE_INSTALL_PREFIX=$HOME/onemath + cmake --build build --target install + export ONEMATHROOT=$HOME/onemath + ``` + + oneDAL links only the run-time dispatching `libonemath.so`; the per-domain + backend libraries next to it are loaded by the dispatcher. + +2. **An NVPTX target for the compiler.** `ONEDAL_SYCL_TARGETS` is appended to + both the DPC++ compile and link actions as `-fsycl-targets=`; device code is + produced at both, so it has to reach both. Under Bazel it is an environment + variable rather than a build flag because the DPC++ flag sets are baked in + when the toolchain repository is configured, before build flags are visible; + changing it refetches the toolchain. Make reads the same variable, and takes + it as a make variable too. + + ```sh + export ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda + export LD_LIBRARY_PATH=$ONEMATHROOT/lib:$LD_LIBRARY_PATH + bazel test --config=nvidia-gpu //cpp/oneapi/dal/backend/primitives/blas:tests + ``` + +`--config=nvidia-gpu` is shorthand for `--dpc_math_backend=onemath`, plus +`--test_env=LD_LIBRARY_PATH` on the test lane. That last part is needed because +oneMath's dispatcher `dlopen`s its per-domain backends through a `$ORIGIN` +RUNPATH: under Bazel `$ORIGIN` is the `_solib` directory the dispatcher is staged +into, not the oneMath install, so the backends are not next to it and without a +search path every test aborts with `oneapi::math::backend_not_found`. Hence +`$ONEMATHROOT/lib` on `LD_LIBRARY_PATH` above. The Make build needs none of this, +because it records the install directory as an rpath. + +The Make build has the same switch, spelled `DPC_MATH_BACKEND` and reading the +same two environment variables: + +```sh +export ONEMATHROOT=/path/to/onemath +export ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda +make -f makefile oneapi_dpc PLAT=lnx32e DPC_MATH_BACKEND=onemath +``` + +`DPC_MATH_BACKEND` is independent of `BACKEND_CONFIG` there too, and the +non-default value gets its own `__work_onemath` directory: the incremental build +only re-examines a command line when a makefile is newer than the object, so a +shared working directory would mix objects compiled against two different math +libraries. `$ONEMATHROOT/lib` is recorded as an rpath on `libonedal_dpc.so`, so +consumers link and run without `LD_LIBRARY_PATH`; an absolute rpath is warranted +here because oneMath has no redistributable and no standard prefix, and without +it `ld` reports every oneMath symbol undefined when a consumer links against the +result. Only `PLAT=lnx32e` is wired up; anything else is rejected rather than +linked against a guessed library name. + +What works today, measured through both build systems against oneMath built from +source at commit `3273ca2` with the `mklcpu` and `mklgpu` backends, running on an +Intel Data Center GPU Max 1100: + +| domain | status | +|---|---| +| BLAS | works | +| LAPACK | works | +| RNG | `philox4x32x10` and `mrg32k3a` work; `mt2203`, `mt19937` and `mcg59` throw `unimplemented` | +| sparse BLAS | throws `unimplemented` | + +The library builds and links as a whole either way — the two gaps below are +compiled out rather than left to fail the build, so a `DPC_MATH_BACKEND=onemath` +build produces a `libonedal_dpc.so` that examples and tests can be linked and run +against. What is not there reports itself as `unimplemented` at the point of use +instead of being silently substituted, which matters because the missing RNG +engines have no numerically equivalent stand-in. + +Concretely: all 45 DPC++ examples build against such a library, and 44 of them +run and produce the same output as the oneMKL build. The one that does not is +`kmeans_lloyd_csr_batch`, which throws `unimplemented` out of the sparse BLAS +layer. + +The Bazel test suites run too, not just build. Across `primitives/{blas,lapack, +rng,objective_function}` and `algo/{pca,kmeans,decision_forest,covariance, +linear_regression,basic_statistics,logistic_regression}`, 67 targets pass under +`--config=nvidia-gpu` on the same GPU, the three `sparse_blas` targets are +skipped as incompatible, and `decision_forest:test_spmd_dpc` fails exactly as it +does on the default backend there. + +Neither gap is as wide as the domain table suggests, which is worth spelling out +so the table is not read as a list of unusable algorithms: + +- The three absent engines are reachable only on request. `decision_forest` is + the only oneAPI descriptor that exposes `engine_type`, its default is + `philox4x32x10`, and that default is what the one device-side draw + (`pr::uniform` on the bootstrap indices) ends up using — so decision forest + runs unless the caller explicitly asks for `mt2203`, `mt19937` or `mcg59`. +- The host-side RNG paths — `shuffle`, `uniform_without_replacement` and the + Fisher-Yates helper, which is how `kmeans_init` draws its random centroids — + take their numbers from the DAAL engines and never touch the device engine, so + they work with any engine type. +- Sparse BLAS is only entered for a `csr_table` on the device, so it bites + exactly the CSR-input algorithms. + +The two gaps are not a matter of naming and cannot be closed by this +configuration alone: + +- **RNG.** oneMath's `oneapi::math::rng` provides `philox4x32x10` and + `mrg32k3a` only. `primitives/rng/device_engine.hpp` also instantiates + `mt2203`, `mt19937` and `mcg59`, which the interface does not declare at + all. Closing this means either contributing the engines upstream to oneMath + or giving those three a fallback — and `mt2203` in particular has no GPU + `skip_ahead`, so it cannot simply be mapped onto a counter-based engine. +- **Sparse BLAS.** The two libraries expose disjoint generations of the + interface. oneDAL uses oneMKL's handle API (`init_matrix_handle`, + `set_csr_data`, `gemv`, `gemm`, `optimize_gemv`); oneMath ships the newer + specification (`init_csr_matrix`, `release_sparse_matrix`, and `spmv` / + `spmm` driven by descriptors with separate buffer-size, optimize and execute + stages). oneMKL 2026 does not offer the newer form, so there is no spelling + that satisfies both and the layer needs a per-backend implementation. + +Until either is closed, `oneapi::math` is reached only for BLAS and LAPACK, and +the `sparse_blas` primitives and three of the five device RNG engines are +`unimplemented` under `onemath`. Both gaps live behind +`ONEDAL_MATH_BACKEND_ONEMATH`, so the default oneMKL build compiles exactly the +code it did before. + +### What the CUDA backends cover + +The two gaps above are interface gaps and apply to any oneMath build. On top of +them, oneMath's CUDA backends do not implement everything oneMath declares, so a +`--dpc_math_backend=onemath` build with CUDA backends has a second, narrower set +of restrictions. The table is derived from reading the oneMath sources at commit +`3273ca2` — specifically the `unimplemented` throw sites in +`src/{blas,lapack,rng,sparse_blas}/backends/cu*` — for exactly the entry points +oneDAL's device code calls. **It is not measured**: there is no NVIDIA hardware +here, so treat it as what the backend source says it will do. + +| oneDAL uses | CUDA backend | covered? | +|---|---|---| +| `blas::column_major::{gemm,gemv,syrk}`, float and double | cuBLAS | yes — only `bfloat16` `gemm` and the `row_major` forms are `unimplemented` | +| `lapack::{potrf,potrs,syevd}` and their scratchpad sizes | cuSOLVER | yes, with no restriction | +| `lapack::gesvd` | cuSOLVER | **only for m ≥ n** | +| `rng::uniform` over `philox4x32x10` / `mrg32k3a`, USM API, `float` / `double` / `std::int32_t` | cuRAND | yes, including scalar `skip_ahead` — but see the semantics note below | +| `sparse::*` | cuSPARSE | n/a — blocked by the interface-generation gap above, not by the backend | + +Two consequences worth calling out, because they are not visible from the domain +table: + +- **PCA `method::svd` cannot run on NVIDIA.** `train_kernel_svd_impl_dpc.cpp` + calls `pr::gesvd(queue, column_count, row_count, ...)`, i.e. + m = columns and n = rows, so m < n for any dataset with more rows than + columns — the normal case. cuSOLVER's `gesvd` throws + `unimplemented("lapack", "gesvd", "cusolver gesvd does not support m < n")` + there. PCA `method::cov` goes through `syevd` and is unaffected, so PCA as a + whole is usable; only the SVD method is not. +- **Random forest works, and should be left on its default engine.** + `decision_forest` is the only descriptor exposing `engine_type`, its default + is `philox4x32x10`, and cuRAND implements everything the one device-side draw + needs. `mrg32k3a` is the other engine cuRAND covers, so either of those two is + a valid explicit choice; `mt2203`, `mt19937` and `mcg59` are not declared by + oneMath at all. + + One caveat on repeated draws: oneDAL advances an engine with + `skip_ahead_gpu(count)` after each device draw, expecting oneMKL's *relative* + advance. cuRAND's backend implements `skip_ahead(n)` as + `curandSetGeneratorOffset(engine_, n)`, an *absolute* offset, so successive + draws from one engine would restart from the same place rather than continue. + Single-draw use — which is what decision forest does per tree block — is + unaffected, but this is a real divergence and is part of the upstream report. + +### Deselecting what the backend cannot run + +Rather than letting these paths build and then throw, the tests and examples that +reach them are deselected, the same way the reference host backend deselects what +it does not implement: + +- `te::device_sparse_blas_supported()` (in `test/engine/fixtures.hpp`) is + `false` under `ONEDAL_MATH_BACKEND_ONEMATH`. The CSR-only test cases in + `kmeans` and `objective_function` check it with `SKIP_IF`. The define is set + only for device translation units, so the host versions of the same tests keep + running. +- Type lists that mix a dense and a sparse method — `kmeans_types`, + `log_reg_types` — drop the sparse method under the same define, and the device + RNG test's engine lists drop the three absent engines. A `COMBINE_TYPES` list + cannot be empty, which is why the wholly-sparse cases use `SKIP_IF` instead. +- `//cpp/oneapi/dal/backend/primitives/sparse_blas:tests` is + `target_compatible_with` `@platforms//:incompatible` under + `--dpc_math_backend=onemath`, so Bazel skips it instead of building a test + whose every assertion would throw. +- The DPC++ examples take two CMake options, mirroring `REF_BACKEND` / + `OPENRNG_BACKEND` on the host side: `-DONEMATH_BACKEND=ON` excludes + `kmeans_lloyd_csr_batch`, and adding `-DCUDA_BACKEND=ON` additionally excludes + `pca_svd_dense_batch` for the `gesvd` restriction above. There is no Bazel + equivalent of the second one: which backends a `libonemath.so` dispatches to is + fixed when oneMath is configured and is not visible to the build graph. + +Other known limitations: + +- Not covered by CI, and no NVIDIA hardware is validated. The oneMath side was + exercised with the `mklgpu` backend on an Intel GPU, which checks the + interface but none of the CUDA code paths. +- Devices without `aspect::fp64` need `--test_disable_fp64=yes`; that covers the + tests, not the algorithms themselves. diff --git a/dev/bazel/config/config.tpl.BUILD b/dev/bazel/config/config.tpl.BUILD index a7b2a02a358..d7b29e63234 100644 --- a/dev/bazel/config/config.tpl.BUILD +++ b/dev/bazel/config/config.tpl.BUILD @@ -120,6 +120,27 @@ config_flag( ], ) +# SYCL math library behind the DPC++ device primitives (BLAS, LAPACK, sparse +# BLAS, RNG). Orthogonal to `backend_config`, which picks the *host* math +# library: a DPC++ build keeps using MKL or the reference backend on the CPU +# side either way. `config_flag` auto-generates the matching +# `:dpc_math_backend_mkl` / `:dpc_math_backend_onemath` config_setting targets. +# +# * "mkl" — oneMKL, Intel GPUs only. The default, and what Make builds. +# * "onemath" — oneMath (https://github.com/uxlfoundation/oneMath), the +# open-source implementation of the same DPC++ interface. Adds cuBLAS, +# cuSOLVER, cuSPARSE and cuRAND backends, so the same sources run on NVIDIA +# GPUs. Requires ONEMATHROOT and a compiler asked for an NVPTX target via +# ONEDAL_SYCL_TARGETS; see `--config=nvidia-gpu` in .bazelrc. +config_flag( + name = "dpc_math_backend", + build_setting_default = "mkl", + allowed_build_setting_values = [ + "mkl", + "onemath", + ], +) + # Use this, not the bare `:rng_backend_openrng`, to gate the OpenRNG define and # dependency: RNG_OPENRNG lives in dev/make/deps.ref.mk, so Make can only reach # it on the ref backend, and .ci/scripts/build.sh guards `--use-openrng yes` diff --git a/dev/bazel/dal.bzl b/dev/bazel/dal.bzl index 7f6735a3294..bfae1d34e4b 100644 --- a/dev/bazel/dal.bzl +++ b/dev/bazel/dal.bzl @@ -170,7 +170,7 @@ def dal_test(name, hdrs=[], srcs=[], dal_deps=[], dal_test_deps=[], dpc_hdrs=[], dpc_srcs=[], dpc_deps=[], compile_as=[ "c++", "dpc++" ], framework="catch2", data=[], tags=[], private=False, mpi=False, ccl=False, mpi_ranks=0, args=[], - use_onedal_release_libs=True, **kwargs): + use_onedal_release_libs=True, target_compatible_with=None, **kwargs): # TODO: Check `compile_as` parameter # TODO: Refactor this rule once decision on the tests structure is made if not framework in ["catch2", "none"]: @@ -226,6 +226,7 @@ def dal_test(name, hdrs=[], srcs=[], dal_deps=[], dal_test_deps=[], data = _expand_select(data + _test_runtime_data()), tags = common_tags + tags + ["host", iface_access_tag], args = test_args, + target_compatible_with = target_compatible_with, ) tests_for_test_suite.append(name + "_host") if "dpc++" in compile_as: @@ -241,6 +242,7 @@ def dal_test(name, hdrs=[], srcs=[], dal_deps=[], dal_test_deps=[], data = _expand_select(data + _test_runtime_data()), tags = common_tags + tags + ["dpc", iface_access_tag], args = test_args, + target_compatible_with = target_compatible_with, ) tests_for_test_suite.append(name + "_dpc") native.test_suite( diff --git a/dev/bazel/deps/onemath.bzl b/dev/bazel/deps/onemath.bzl new file mode 100644 index 00000000000..cf0c40d6973 --- /dev/null +++ b/dev/bazel/deps/onemath.bzl @@ -0,0 +1,51 @@ +#=============================================================================== +# 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. +#=============================================================================== + +load("@onedal//dev/bazel:repos.bzl", "repos") + +# oneMath (https://github.com/uxlfoundation/oneMath) implements the same DPC++ +# interface as oneMKL and additionally dispatches to cuBLAS, cuSOLVER, cuSPARSE +# and cuRAND, which is what lets the oneDAL device sources run on NVIDIA GPUs. +# +# Unlike TBB or oneMKL there is no redistributable package to download: the +# backends a build needs are chosen at oneMath configure time +# (`-DENABLE_CUBLAS_BACKEND=ON` and friends), so the library has to be built +# locally and pointed at through ONEMATHROOT. Consequently this repository has +# no `urls` and only resolves when something actually depends on it, i.e. under +# `--dpc_math_backend=onemath`. +# +# oneDAL links only the run-time dispatching library, `libonemath.so`. The +# per-backend shared objects are `dlopen`ed by the dispatcher through a +# `$ORIGIN` RUNPATH, which under Bazel is the `_solib` directory and not the +# oneMath install, so they are found through `LD_LIBRARY_PATH` (see +# `--config=nvidia-gpu` in dev/bazel/README.md) and are not staged here. +# +# Every pattern is globbed the way `mkl.bzl` globs its shared objects, because +# oneMath installs `libonemath.so` as a symlink to the SONAME'd +# `libonemath.so.`. Staging only the symlink links fine -- `ld` follows it +# -- but records a `DT_NEEDED` on the SONAME that is then absent from the test +# runfiles, and every binary dies at startup with `libonemath.so.0: cannot open +# shared object file`. Bazel treats the versioned file as a runtime-only input, +# so naming both is exactly what is wanted. +onemath_repo = repos.prebuilt_libs_repo_rule( + includes = [ + "include", + ], + libs = [ + "lib/libonemath.so*", + ], + build_template = "@onedal//dev/bazel/deps:onemath.tpl.BUILD", +) diff --git a/dev/bazel/deps/onemath.tpl.BUILD b/dev/bazel/deps/onemath.tpl.BUILD new file mode 100644 index 00000000000..1dd26419364 --- /dev/null +++ b/dev/bazel/deps/onemath.tpl.BUILD @@ -0,0 +1,45 @@ +package(default_visibility = ["//visibility:public"]) +load("@rules_cc//cc:defs.bzl", "cc_library") + +cc_library( + name = "headers", + # oneMath ships `.hpp` headers plus the `.hxx` files they include; it ships + # no `.h` at all, so that pattern has to tolerate matching nothing rather + # than fail the glob. + hdrs = glob( + [ + "include/**/*.h", + "include/**/*.hpp", + "include/**/*.hxx", + ], + allow_empty = True, + ), + includes = [ + "include", + ], +) + +cc_library( + name = "onemath_dpc", + # Only the run-time dispatching library is linked; the dispatcher loads the + # per-domain backend libraries itself (see onemath.bzl). + srcs = glob([ + "lib/libonemath.so*", + ]), + linkopts = [ + # Same cap as `@mkl//:mkl_dpc`: a fixed 16 gives the best trade-off + # between device-code link speedup and memory used, and is clamped to + # `nproc` on smaller machines. + "-fsycl-max-parallel-link-jobs=16", + ], + deps = [ + ":headers", + "@opencl//:opencl_binary", + ], + # Switches `dal/backend/math_backend.hpp` from to + # . Carried by the dependency rather than by the algorithm + # targets so the define cannot drift from the library being linked. + defines = [ + "ONEDAL_MATH_BACKEND_ONEMATH", + ], +) diff --git a/dev/bazel/repos.bzl b/dev/bazel/repos.bzl index 8d72461b175..a9bbb9978cd 100644 --- a/dev/bazel/repos.bzl +++ b/dev/bazel/repos.bzl @@ -154,9 +154,29 @@ def _create_optional_symlinks(repo_ctx, root, entries, substitutions=None, mappi present = [] for entry in entries: entry_fmt = utils.substitute(entry, substitutions) - src_entry_path = utils.substitute(paths.join(root, entry_fmt), mapping) - if repo_ctx.path(src_entry_path).exists: - present.append(entry) + if "*" in entry_fmt: + # `exists` on a pattern is always false, so a globbed entry has to be + # resolved to the concrete names before the presence test -- without + # this, a pattern is silently dropped even when the package ships + # files matching it. Resolved here rather than handed to + # `_create_symlinks`, which fails an unmatched pattern by design. + pattern = entry_fmt.split("/")[-1] + dir_part = entry_fmt[:entry_fmt.rfind("/")] if "/" in entry_fmt else "" + root_with_dir = utils.substitute( + paths.join(root, dir_part) if dir_part else root, + mapping + ) + dir_path = repo_ctx.path(root_with_dir) + if not dir_path.exists: + continue + for fs_entry in dir_path.readdir(): + if _matches_glob(fs_entry.basename, pattern): + present.append(paths.join(dir_part, fs_entry.basename) + if dir_part else fs_entry.basename) + else: + src_entry_path = utils.substitute(paths.join(root, entry_fmt), mapping) + if repo_ctx.path(src_entry_path).exists: + present.append(entry) _create_symlinks(repo_ctx, root, present, substitutions, mapping) def _matches_glob(name, pattern): diff --git a/dev/bazel/toolchains/cc_toolchain.bzl b/dev/bazel/toolchains/cc_toolchain.bzl index 8815073f946..42686bedb5a 100644 --- a/dev/bazel/toolchains/cc_toolchain.bzl +++ b/dev/bazel/toolchains/cc_toolchain.bzl @@ -76,5 +76,9 @@ onedal_cc_toolchain = repository_rule( # ONEDAL_WIN_COMPILER=icx selects icx/icpx; anything else falls # back to the rules_cc MSVC cl auto-config. "ONEDAL_WIN_COMPILER", + # Comma-separated SYCL target triples for the DPC++ compile and link + # actions, e.g. ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda. Empty leaves + # the compiler at its default (Intel SPIR-V). + "ONEDAL_SYCL_TARGETS", ], ) diff --git a/dev/bazel/toolchains/cc_toolchain_lnx.bzl b/dev/bazel/toolchains/cc_toolchain_lnx.bzl index 55a171dc2cc..ef1b807a3ab 100644 --- a/dev/bazel/toolchains/cc_toolchain_lnx.bzl +++ b/dev/bazel/toolchains/cc_toolchain_lnx.bzl @@ -179,6 +179,22 @@ def _add_sycl_linkage(repo_ctx, cc): else: return [] +def _get_sycl_target_flags(repo_ctx): + """Retarget SYCL device code, for example to NVIDIA NVPTX. + + oneDAL's device sources are target-agnostic -- they use no Intel SYCL + extensions and query sub-group sizes at run time -- so changing GPU vendor + is purely a toolchain concern: `-fsycl-targets` has to reach both the DPC++ + compile and the DPC++ link action, and nothing else in the build changes. + + Read from ONEDAL_SYCL_TARGETS rather than a `--@config//` build flag + because the DPC++ flag sets are baked into the toolchain when this + repository is configured, before any build flag is visible. Unset leaves + the compiler at its default target (Intel SPIR-V). + """ + targets = repo_ctx.os.environ.get("ONEDAL_SYCL_TARGETS", "").strip() + return ["-fsycl-targets={}".format(targets)] if targets else [] + def configure_cc_toolchain_lnx(repo_ctx, reqs): if reqs.os_id != "lnx": auto_configure_fail("Cannot configure Linux toolchain for '{}'".format(reqs.os_id)) @@ -205,6 +221,9 @@ def configure_cc_toolchain_lnx(repo_ctx, reqs): # DPC++ kernel code split option dpcc_code_split = "per_kernel" + # Extra SYCL target triples, e.g. `nvptx64-nvidia-cuda` for NVIDIA GPUs. + sycl_target_flags = _get_sycl_target_flags(repo_ctx) + repo_ctx.template( "BUILD", Label("@onedal//dev/bazel/toolchains:cc_toolchain_lnx.tpl.BUILD"), @@ -277,6 +296,7 @@ def configure_cc_toolchain_lnx(repo_ctx, reqs): tools.dpcc, "-fsycl-device-code-split={}".format(dpcc_code_split), ) + + sycl_target_flags + add_compiler_option_if_supported( repo_ctx, tools.cc, @@ -339,6 +359,7 @@ def configure_cc_toolchain_lnx(repo_ctx, reqs): tools.dpcc, "-fsycl-device-code-split={}".format(dpcc_code_split), ) + + sycl_target_flags + add_linker_option_if_supported( repo_ctx, tools.dpcc, diff --git a/dev/make/compiler_definitions/dpcpp.mk b/dev/make/compiler_definitions/dpcpp.mk index 70f3ac204f7..7934ce3f8c1 100644 --- a/dev/make/compiler_definitions/dpcpp.mk +++ b/dev/make/compiler_definitions/dpcpp.mk @@ -73,18 +73,31 @@ endif -asanstatic.dpcpp = -static-libasan -asanshared.dpcpp = -shared-libasan +# Offload targets for the device code. Unset leaves the compiler at its default, +# which is Intel SPIR-V; ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda retargets it at +# NVIDIA GPUs. oneDAL's device sources are target-agnostic -- no Intel SYCL +# extension, sub-group sizes queried at run time -- so changing GPU vendor is a +# toolchain concern and nothing else in the build reacts to this. The value is +# passed through verbatim, so a comma-separated list of triples works, and it has +# to reach both the compile and the link because device code is produced at both. +-fsycl-targets.dpcpp = $(if $(ONEDAL_SYCL_TARGETS),-fsycl-targets=$(ONEDAL_SYCL_TARGETS)) + COMPILER.lnx.dpcpp = icpx -fsycl -m64 -stdlib=libstdc++ -fgnu-runtime -fwrapv \ - -Werror -Wreturn-type -fsycl-device-code-split=per_kernel + -Werror -Wreturn-type -fsycl-device-code-split=per_kernel \ + $(-fsycl-targets.dpcpp) COMPILER.win.dpcpp = icx -fsycl $(if $(MSVC_RT_is_release),-MD, -MDd /debug:none) -nologo -WX \ - -Wno-deprecated-declarations -Wno-ignored-attributes -fsycl-device-code-split=per_kernel + -Wno-deprecated-declarations -Wno-ignored-attributes -fsycl-device-code-split=per_kernel \ + $(-fsycl-targets.dpcpp) linker.ld.flag := $(if $(LINKER),-fuse-ld=$(LINKER),) link.dynamic.lnx.dpcpp = icpx $(linker.ld.flag) -fsycl -m64 -lgomp \ - -fsycl-device-code-split=per_kernel -fsycl-max-parallel-link-jobs=$(SYCL_LINK_PRL) + -fsycl-device-code-split=per_kernel -fsycl-max-parallel-link-jobs=$(SYCL_LINK_PRL) \ + $(-fsycl-targets.dpcpp) link.dynamic.lnx.dpcpp += $(if $(filter yes,$(GCOV_ENABLED)),-Xscoverage,) link.dynamic.win.dpcpp = icx $(linker.ld.flag) -fsycl -m64 \ - -fsycl-device-code-split=per_kernel -fsycl-max-parallel-link-jobs=$(SYCL_LINK_PRL) + -fsycl-device-code-split=per_kernel -fsycl-max-parallel-link-jobs=$(SYCL_LINK_PRL) \ + $(-fsycl-targets.dpcpp) pedantic.opts.lnx.dpcpp = -pedantic \ -Wall \ diff --git a/dev/make/deps.dpc.onemath.mk b/dev/make/deps.dpc.onemath.mk new file mode 100644 index 00000000000..bd04b801e10 --- /dev/null +++ b/dev/make/deps.dpc.onemath.mk @@ -0,0 +1,88 @@ +#=============================================================================== +# 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. +#=============================================================================== + +#++ +# Device-side math backend (oneMath) definitions for makefile +# +# oneMath (https://github.com/uxlfoundation/oneMath) implements the same DPC++ +# interface as oneMKL and additionally dispatches to cuBLAS, cuSOLVER, cuSPARSE +# and cuRAND, which is what lets the unmodified oneDAL device sources run on +# NVIDIA GPUs. Selected with DPC_MATH_BACKEND=onemath. +# +# Only the device-side (`_dpc`) pieces are replaced here. The host math library +# stays whatever BACKEND_CONFIG selected, so a build can take oneMKL on the CPU +# and oneMath on the GPU. +# +# Not every domain oneDAL uses is available in oneMath: BLAS and LAPACK are, +# sparse BLAS is not, and RNG only in part. The gaps are compiled out behind +# ONEDAL_MATH_BACKEND_ONEMATH and throw `unimplemented` at the point of use, so +# the library still builds and links as a whole. See the "NVIDIA GPUs through +# oneMath" section of dev/bazel/README.md for what that means in practice. +#-- + +ifeq (,$(ONEMATHROOT)) + $(error DPC_MATH_BACKEND=onemath requires ONEMATHROOT to point at a oneMath install) +endif + +# There is no redistributable to download and no default install location: which +# CUDA backends a libonemath.so contains is fixed when oneMath itself is +# configured, so the library has to be built locally and named explicitly. +ONEMATHDIR := $(subst \,/,$(ONEMATHROOT)) +ONEMATHDIR.include := $(ONEMATHDIR)/include +ONEMATHDIR.lib := $(ONEMATHDIR)/lib + +ifeq (,$(wildcard $(ONEMATHDIR.include)/oneapi/math.hpp)) + $(error $(ONEMATHDIR.include)/oneapi/math.hpp not found; ONEMATHROOT=$(ONEMATHROOT) does not look like a oneMath install) +endif + +# Windows and the non-x86 platforms are deliberately not wired up: the oneMath +# import-library naming and the NVPTX toolchain setup there are unverified, and a +# silently wrong link line is worse than a refusal. +ifneq ($(PLAT),lnx32e) + $(error DPC_MATH_BACKEND=onemath is only supported for PLAT=lnx32e, got '$(PLAT)') +endif + +# Appended rather than assigned: the oneAPI sources also include DAAL headers, +# which need the host backend's include directory whatever it is. Adding oneMath +# alongside is safe because and do not +# collide, and math_backend.hpp includes exactly one of them. +daaldep.math_backend_oneapi.incdir += $(ONEMATHDIR.include) + +# Switches cpp/oneapi/dal/backend/math_backend.hpp from to +# . Defined next to the link line it belongs with, so the two +# cannot drift apart. +daaldep.math_backend.dpc_defines := -DONEDAL_MATH_BACKEND_ONEMATH + +# Only the run-time dispatching library is linked. The per-domain backend shared +# objects that sit next to it (libonemath_blas_cublas.so and friends) are +# dlopened by the dispatcher, so they have to be findable at run time, but they +# are not link inputs and which of them exist depends on how oneMath was +# configured. +# +# The rpath is what makes the result usable rather than merely linkable. oneMath +# has no redistributable and no standard prefix, so unlike oneMKL it cannot be +# reached through the release tree's $ORIGIN-relative rpath. Without it, +# libonedal_dpc.so records a DT_NEEDED on libonemath.so.0 that nothing can +# resolve: linking any consumer against it fails with `libonemath.so.0 ... not +# found` followed by every oneMath symbol reported undefined, because ld does not +# consult LD_LIBRARY_PATH when resolving a dependency of a shared library. +# Recording the absolute path is acceptable here in a way it would not be for a +# redistributable build -- this configuration is experimental and built against +# one specific local install by definition. +onemath_libs.lnx32e := -L$(ONEMATHDIR.lib) -Wl,-rpath,$(ONEMATHDIR.lib) \ + -lonemath -lsycl -lm -ldl + +daaldep.math_backend.dpc_link_deps := $(onemath_libs.$(PLAT)) diff --git a/examples/oneapi/dpc/CMakeLists.txt b/examples/oneapi/dpc/CMakeLists.txt index 3e7f0c64ef6..4185948e350 100644 --- a/examples/oneapi/dpc/CMakeLists.txt +++ b/examples/oneapi/dpc/CMakeLists.txt @@ -32,12 +32,38 @@ set_common_compiler_options() include_directories(source) +option(ONEMATH_BACKEND "Exclude not supported files for the oneMath device backend" OFF) +option(CUDA_BACKEND "Exclude not supported files for oneMath dispatching to CUDA" OFF) + # Initialize the EXCLUDE_LIST variable set(EXCLUDE_LIST "source/utils/*" "source/example_util/*" ) +# A library built with DPC_MATH_BACKEND=onemath / --dpc_math_backend=onemath does +# not cover every domain oneDAL uses; see the "NVIDIA GPUs through oneMath" +# section of dev/bazel/README.md. What is missing throws `unimplemented` at the +# point of use, so the examples below would build and then fail at run time -- +# exclude them instead, the same way the reference host backend does. +if(ONEMATH_BACKEND) + # The only example that hands a csr_table to the device, which is the one + # entry into the sparse BLAS primitives. + set(EXCLUDE_LIST + ${EXCLUDE_LIST} + "source/kmeans/kmeans_lloyd_csr_batch.cpp" + ) + if (CUDA_BACKEND) + # cuSOLVER's gesvd requires m >= n, and the PCA svd kernel calls it with + # m = column_count, n = row_count, so it throws for any dataset with + # more rows than columns. PCA cov goes through syevd and is unaffected. + set(EXCLUDE_LIST + ${EXCLUDE_LIST} + "source/pca/pca_svd_dense_batch.cpp" + ) + endif() +endif() + # Define variable to specify the examples or directories to include or exclude option(EXAMPLES_LIST "") diff --git a/makefile b/makefile index bcee9a445e1..78c8272d50f 100644 --- a/makefile +++ b/makefile @@ -205,11 +205,25 @@ USECPUS.out.defs.filter := $(if $(USECPUS.out.defs),sed $(sed.-b) $(sed.-i) -E - DAALTHRS ?= tbb DAALAY ?= a y +# Math library behind the DPC++ device primitives, chosen independently of +# BACKEND_CONFIG: that one names the host math library, which a DPC++ build keeps +# using on the CPU side either way. Defined here rather than next to the rest of +# the math backend handling because WORKDIR below is keyed on it. +DPC_MATH_BACKEND ?= mkl +$(if $(filter $(DPC_MATH_BACKEND),mkl onemath),,$(error unknown dpc math backend $(DPC_MATH_BACKEND))) + +# A non-default device math backend gets its own working directory. The +# incremental build only re-examines a command line when a makefile is newer than +# the object, so sharing one directory would otherwise mix objects compiled +# against two different math libraries and fail at link. The default path keeps +# the directory name it has always had. +WORKDIR.dpc_math_suffix := $(if $(filter-out mkl,$(DPC_MATH_BACKEND)),_$(DPC_MATH_BACKEND)) + DIR:=. CPPDIR:=$(DIR)/cpp CPPDIR.daal:=$(CPPDIR)/daal CPPDIR.onedal:=$(CPPDIR)/oneapi/dal -WORKDIR ?= $(DIR)/__work$(CMPLRDIRSUFF.$(COMPILER))/$(if $(MSVC_RT_is_release),md,mdd)/$(PLAT) +WORKDIR ?= $(DIR)/__work$(CMPLRDIRSUFF.$(COMPILER))$(WORKDIR.dpc_math_suffix)/$(if $(MSVC_RT_is_release),md,mdd)/$(PLAT) RELEASEDIR ?= $(DIR)/__release_$(_OS)$(CMPLRDIRSUFF.$(COMPILER)) RELEASEDIR.daal := $(RELEASEDIR)/daal/latest RELEASEDIR.data := $(RELEASEDIR.daal)/data @@ -306,6 +320,14 @@ endif include dev/make/deps.$(BACKEND_CONFIG).mk +# Device-side math library (DPC_MATH_BACKEND, validated near WORKDIR above). +# `mkl` needs no file of its own, because the device link line for oneMKL is +# already part of deps.mkl.mk; only a device backend that differs from the host +# one has something to add. +ifneq ($(DPC_MATH_BACKEND), mkl) + include dev/make/deps.dpc.$(DPC_MATH_BACKEND).mk +endif + #=============================== VTune SDK folders ====================================== ifeq ($(REQPROFILE), yes) @@ -722,6 +744,7 @@ $(ONEAPI.objs_y.dpc): COPT += $(-fPIC) $(-cxx17) $(-optlevel.dpcpp) $(-Zl_DPCPP) -DDAAL_NOTHROW_EXCEPTIONS \ -DDAAL_HIDE_DEPRECATED \ -DONEDAL_DATA_PARALLEL \ + $(daaldep.math_backend.dpc_defines) \ -D_ENABLE_ATOMIC_ALIGNMENT_FIX \ $(if $(CHECK_DLL_SIG),-DDAAL_CHECK_DLL_SIG) \ -D__ONEDAL_ENABLE_EXPORT__ \ @@ -1147,6 +1170,17 @@ Flags: assertions and adds debug symbols. Any value will enable debug mode, except value "symbols" which will only add debug symbols. special value: symbols + DPC_MATH_BACKEND - math library behind the DPC++ device primitives, chosen + independently of BACKEND_CONFIG (which names the host one) + mkl - oneMKL, Intel GPUs [default] + onemath - oneMath, which also dispatches to cuBLAS/cuSOLVER/cuSPARSE/ + cuRAND and so runs the device code on NVIDIA GPUs. Experimental: + only the BLAS and LAPACK domains build against it. Requires + ONEMATHROOT, and PLAT=lnx32e. Combine with + ONEDAL_SYCL_TARGETS=nvptx64-nvidia-cuda to build for NVPTX. + See the "NVIDIA GPUs through oneMath" section of dev/bazel/README.md. + ONEDAL_SYCL_TARGETS - value for the compiler's -fsycl-targets, applied to both + the DPC++ compile and link. Empty leaves the compiler at its default target REQPROFILE - flag that enables kernel profiling using STDALLOC - flag that makes the library use stdlib allocators instead of MKL's (default: no)