From 9b9548aa98cb23b510a4907786350836430b25a7 Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Thu, 4 Jul 2024 15:59:35 +0200 Subject: [PATCH 01/15] Disable spsv symmetric conjtrans --- docs/domains/sparse_linear_algebra.rst | 3 +++ .../backends/mkl_common/mkl_spmm.cxx | 5 +++- .../backends/mkl_common/mkl_spsv.cxx | 25 ++++++++++++++----- 3 files changed, 26 insertions(+), 7 deletions(-) diff --git a/docs/domains/sparse_linear_algebra.rst b/docs/domains/sparse_linear_algebra.rst index 8798303c9..41b7c9d9e 100644 --- a/docs/domains/sparse_linear_algebra.rst +++ b/docs/domains/sparse_linear_algebra.rst @@ -29,5 +29,8 @@ Currently known limitations: ``oneapi::mkl::transpose::conjtrans`` and has the ``oneapi::mkl::sparse::matrix_property::symmetric`` property will throw an ``oneapi::mkl::unimplemented`` exception. +- Using ``spsv`` on Intel GPU with a sparse matrix that is + ``oneapi::mkl::transpose::conjtrans`` and will throw an + ``oneapi::mkl::unimplemented`` exception. - Scalar parameters ``alpha`` and ``beta`` should be host pointers to prevent synchronizations and copies to the host. diff --git a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx index 7e5ba03b2..aa292921e 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx @@ -67,7 +67,10 @@ void check_valid_spmm(const std::string function_name, sycl::queue &queue, } #if BACKEND == gpu - if (opA == oneapi::mkl::transpose::conjtrans && + detail::data_type data_type = internal_A_handle->get_value_type(); + if ((data_type == detail::data_type::complex_fp32 || + data_type == detail::data_type::complex_fp64) && + opA == oneapi::mkl::transpose::conjtrans && internal_A_handle->has_matrix_property(oneapi::mkl::sparse::matrix_property::symmetric)) { throw mkl::unimplemented( "sparse_blas", function_name, diff --git a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx index f69cca852..ea8410a88 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx @@ -30,7 +30,7 @@ sycl::event release_spsv_descr(sycl::queue &queue, oneapi::mkl::sparse::spsv_des } void check_valid_spsv(const std::string function_name, sycl::queue &queue, - oneapi::mkl::sparse::matrix_view A_view, + oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, oneapi::mkl::sparse::dense_vector_handle_t y_handle, const void *alpha, @@ -47,6 +47,19 @@ void check_valid_spsv(const std::string function_name, sycl::queue &queue, "The backend does not support `no_optimize_alg` unless A_handle has the property `matrix_property::sorted`."); } +#if BACKEND == gpu + detail::data_type data_type = internal_A_handle->get_value_type(); + if ((data_type == detail::data_type::complex_fp32 || + data_type == detail::data_type::complex_fp64) && + opA == oneapi::mkl::transpose::conjtrans) { + throw mkl::unimplemented( + "sparse_blas", function_name, + "The backend does not support spsv using conjtrans."); + } +#else + (void)opA; +#endif // BACKEND + detail::check_all_containers_compatible(function_name, internal_A_handle, x_handle, y_handle); if (A_view.type_view != matrix_descr::triangular) { throw mkl::invalid_argument("sparse_blas", function_name, @@ -58,7 +71,7 @@ void check_valid_spsv(const std::string function_name, sycl::queue &queue, } } -void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose /*opA*/, const void *alpha, +void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, @@ -67,7 +80,7 @@ void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose /*opA*/, const oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, std::size_t &temp_buffer_size) { // TODO: Add support for external workspace once the close-source oneMKL backend supports it. - check_valid_spsv(__func__, queue, A_view, A_handle, x_handle, y_handle, alpha, alg); + check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); temp_buffer_size = 0; } @@ -79,7 +92,7 @@ void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, sycl::buffer /*workspace*/) { - check_valid_spsv(__func__, queue, A_view, A_handle, x_handle, y_handle, alpha, alg); + check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); auto internal_A_handle = detail::get_internal_handle(A_handle); if (!internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -102,7 +115,7 @@ sycl::event spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, void * /*workspace*/, const std::vector &dependencies) { - check_valid_spsv(__func__, queue, A_view, A_handle, x_handle, y_handle, alpha, alg); + check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); auto internal_A_handle = detail::get_internal_handle(A_handle); if (internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -149,7 +162,7 @@ sycl::event spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp oneapi::mkl::sparse::dense_vector_handle_t y_handle, oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t spsv_descr, const std::vector &dependencies) { - check_valid_spsv(__func__, queue, A_view, A_handle, x_handle, y_handle, alpha, alg); + check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); auto value_type = detail::get_internal_handle(A_handle)->get_value_type(); DISPATCH_MKL_OPERATION("spsv", value_type, internal_spsv, queue, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, spsv_descr, dependencies); From 3c7789cf8696976af6a37a0c1144d4372ad6f10d Mon Sep 17 00:00:00 2001 From: HJA Bird Date: Thu, 4 Jul 2024 20:10:07 +0100 Subject: [PATCH 02/15] [Docs] Fix errors in example prose (#526) * Docs incorrectly state that the example builds for Intel, contradicting itself later on. * This error is corrected. --- docs/building_the_project_with_dpcpp.rst | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/docs/building_the_project_with_dpcpp.rst b/docs/building_the_project_with_dpcpp.rst index 6df17f46d..365028237 100644 --- a/docs/building_the_project_with_dpcpp.rst +++ b/docs/building_the_project_with_dpcpp.rst @@ -341,7 +341,7 @@ The following table provides details of CMake options and their default values: CMake invocation examples ########################## -Build oneMKL with support for x86 CPU, Intel GPU, and Nvidia GPUs with tests +Build oneMKL with support for Nvidia GPUs with tests disabled using the Ninja build system: .. code-block:: bash @@ -363,7 +363,7 @@ and Intel GPU (``MKLGPU``) backends are enabled by default, but are disabled here. The backends for Nvidia GPUs must all be explicilty enabled. The tests are disabled, but the examples will still be built. -Building oneMKL with support x86 CPU, Intel GPU, and AMD GPUs with tests +Building oneMKL with support for AMD GPUs with tests disabled: .. code-block:: bash @@ -401,7 +401,7 @@ GPU and Nvidia GPU with testing enabled: Note that this is not a supported configuration, and requires Codeplay's oneAPI for `AMD `_ and `Nvidia `_ GPU -plugins. Like the above example, the MKLCPU and MKLGPU backends are enabled by +plugins. The MKLCPU and MKLGPU backends are enabled by default, with backends for Nvidia GPU and AMD GPU explicitly enabled. ``-DTARGET_DOMAINS=dft`` causes only DFT backends to be built. If this was not set, the backend libraries to enable the use of BLAS, LAPACK and RNG with MKLGPU From ccd72e1403758c685bf66b34e76b042474fe91e7 Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Mon, 8 Jul 2024 15:47:51 -0700 Subject: [PATCH 03/15] Bump certifi from 2024.2.2 to 2024.7.4 in /docs (#529) Signed-off-by: dependabot[bot] Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> --- docs/requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/requirements.txt b/docs/requirements.txt index 8af51c7f6..8365d7241 100644 --- a/docs/requirements.txt +++ b/docs/requirements.txt @@ -2,7 +2,7 @@ accessible-pygments==0.0.5 alabaster==0.7.16 Babel==2.15.0 beautifulsoup4==4.12.3 -certifi==2024.2.2 +certifi==2024.7.4 charset-normalizer==3.3.2 docutils==0.21.2 idna==3.7 From 82566e5026444717991adffbd1c809989ae79aea Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Tue, 9 Jul 2024 15:12:22 +0200 Subject: [PATCH 04/15] Test symmetric with complex types and hermitian and conjtrans with real types --- .../sparse_blas/include/test_spmm.hpp | 7 +-- .../sparse_blas/include/test_spmv.hpp | 61 ++++++++++++++++--- .../sparse_blas/source/sparse_spmv_buffer.cpp | 20 ++---- .../sparse_blas/source/sparse_spmv_usm.cpp | 20 ++---- 4 files changed, 61 insertions(+), 47 deletions(-) diff --git a/tests/unit_tests/sparse_blas/include/test_spmm.hpp b/tests/unit_tests/sparse_blas/include/test_spmm.hpp index 049d58b88..6188d4268 100644 --- a/tests/unit_tests/sparse_blas/include/test_spmm.hpp +++ b/tests/unit_tests/sparse_blas/include/test_spmm.hpp @@ -205,7 +205,6 @@ void test_helper_with_format_with_transpose( /** * Helper function to test combination of transpose vals. - * Only test \p conjtrans if \p fpType is complex. * * @tparam fpType Complex or scalar, single or double precision type * @tparam testFunctorI32 Test functor for fpType and int32 @@ -223,10 +222,8 @@ void test_helper_with_format( const std::vector &non_default_algorithms, int &num_passed, int &num_skipped) { std::vector transpose_vals{ oneapi::mkl::transpose::nontrans, - oneapi::mkl::transpose::trans }; - if (complex_info::is_complex) { - transpose_vals.push_back(oneapi::mkl::transpose::conjtrans); - } + oneapi::mkl::transpose::trans, + oneapi::mkl::transpose::conjtrans }; for (auto transpose_A : transpose_vals) { for (auto transpose_B : transpose_vals) { test_helper_with_format_with_transpose( diff --git a/tests/unit_tests/sparse_blas/include/test_spmv.hpp b/tests/unit_tests/sparse_blas/include/test_spmv.hpp index 43599e9d3..6ee256adb 100644 --- a/tests/unit_tests/sparse_blas/include/test_spmv.hpp +++ b/tests/unit_tests/sparse_blas/include/test_spmv.hpp @@ -51,7 +51,7 @@ * The test functions will use different sizes if the configuration implies a symmetric matrix. */ template -void test_helper_with_format( +void test_helper_with_format_with_transpose( testFunctorI32 test_functor_i32, testFunctorI64 test_functor_i64, sycl::device *dev, sparse_matrix_format_t format, const std::vector &non_default_algorithms, @@ -153,22 +153,37 @@ void test_helper_with_format( no_reset_data, no_scalars_on_device), num_passed, num_skipped); if (transpose_val != oneapi::mkl::transpose::conjtrans) { - // Lower symmetric or hermitian + // Do not test conjtrans with symmetric or hermitian views as no backend supports it. + // Lower symmetric oneapi::mkl::sparse::matrix_view symmetric_view( - complex_info::is_complex ? oneapi::mkl::sparse::matrix_descr::hermitian - : oneapi::mkl::sparse::matrix_descr::symmetric); + oneapi::mkl::sparse::matrix_descr::symmetric); EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, fp_one, fp_zero, default_alg, symmetric_view, no_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); - // Upper symmetric or hermitian + // Upper symmetric symmetric_view.uplo_view = oneapi::mkl::uplo::upper; EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, fp_one, fp_zero, default_alg, symmetric_view, no_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); + // Lower hermitian + oneapi::mkl::sparse::matrix_view hermitian_view( + oneapi::mkl::sparse::matrix_descr::hermitian); + EXPECT_TRUE_OR_FUTURE_SKIP( + test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, + transpose_val, fp_one, fp_zero, default_alg, hermitian_view, + no_properties, no_reset_data, no_scalars_on_device), + num_passed, num_skipped); + // Upper hermitian + hermitian_view.uplo_view = oneapi::mkl::uplo::upper; + EXPECT_TRUE_OR_FUTURE_SKIP( + test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, + transpose_val, fp_one, fp_zero, default_alg, hermitian_view, + no_properties, no_reset_data, no_scalars_on_device), + num_passed, num_skipped); } // Test other algorithms for (auto alg : non_default_algorithms) { @@ -188,6 +203,34 @@ void test_helper_with_format( } } +/** + * Helper function to test combination of transpose vals. + * + * @tparam fpType Complex or scalar, single or double precision type + * @tparam testFunctorI32 Test functor for fpType and int32 + * @tparam testFunctorI64 Test functor for fpType and int64 + * @param dev Device to test + * @param format Sparse matrix format to use + * @param non_default_algorithms Algorithms compatible with the given format, other than default_alg + * @param num_passed Increase the number of configurations passed + * @param num_skipped Increase the number of configurations skipped + */ +template +void test_helper_with_format( + testFunctorI32 test_functor_i32, testFunctorI64 test_functor_i64, sycl::device *dev, + sparse_matrix_format_t format, + const std::vector &non_default_algorithms, int &num_passed, + int &num_skipped) { + std::vector transpose_vals{ oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::trans, + oneapi::mkl::transpose::conjtrans }; + for (auto transpose_A : transpose_vals) { + test_helper_with_format_with_transpose(test_functor_i32, test_functor_i64, dev, + format, non_default_algorithms, transpose_A, + num_passed, num_skipped); + } +} + /** * Helper function to test multiple sparse matrix format and choose valid algorithms. * @@ -195,24 +238,22 @@ void test_helper_with_format( * @tparam testFunctorI32 Test functor for fpType and int32 * @tparam testFunctorI64 Test functor for fpType and int64 * @param dev Device to test - * @param transpose_val Transpose value for the input matrix * @param num_passed Increase the number of configurations passed * @param num_skipped Increase the number of configurations skipped */ template void test_helper(testFunctorI32 test_functor_i32, testFunctorI64 test_functor_i64, - sycl::device *dev, oneapi::mkl::transpose transpose_val, int &num_passed, - int &num_skipped) { + sycl::device *dev, int &num_passed, int &num_skipped) { test_helper_with_format( test_functor_i32, test_functor_i64, dev, sparse_matrix_format_t::CSR, { oneapi::mkl::sparse::spmv_alg::no_optimize_alg, oneapi::mkl::sparse::spmv_alg::csr_alg1, oneapi::mkl::sparse::spmv_alg::csr_alg2, oneapi::mkl::sparse::spmv_alg::csr_alg3 }, - transpose_val, num_passed, num_skipped); + num_passed, num_skipped); test_helper_with_format( test_functor_i32, test_functor_i64, dev, sparse_matrix_format_t::COO, { oneapi::mkl::sparse::spmv_alg::no_optimize_alg, oneapi::mkl::sparse::spmv_alg::coo_alg1, oneapi::mkl::sparse::spmv_alg::coo_alg2 }, - transpose_val, num_passed, num_skipped); + num_passed, num_skipped); } /// Compute spmv reference as a dense operation diff --git a/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp index 12b449e61..0ba5afb9c 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp @@ -184,9 +184,7 @@ TEST_P(SparseSpmvBufferTests, RealSinglePrecision) { using fpType = float; int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -199,9 +197,7 @@ TEST_P(SparseSpmvBufferTests, RealDoublePrecision) { CHECK_DOUBLE_ON_DEVICE(GetParam()); int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -213,11 +209,7 @@ TEST_P(SparseSpmvBufferTests, ComplexSinglePrecision) { using fpType = std::complex; int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::conjtrans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -230,11 +222,7 @@ TEST_P(SparseSpmvBufferTests, ComplexDoublePrecision) { CHECK_DOUBLE_ON_DEVICE(GetParam()); int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::conjtrans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped diff --git a/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp b/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp index 85feacbda..fdeb57913 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp @@ -233,9 +233,7 @@ TEST_P(SparseSpmvUsmTests, RealSinglePrecision) { using fpType = float; int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -248,9 +246,7 @@ TEST_P(SparseSpmvUsmTests, RealDoublePrecision) { CHECK_DOUBLE_ON_DEVICE(GetParam()); int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -262,11 +258,7 @@ TEST_P(SparseSpmvUsmTests, ComplexSinglePrecision) { using fpType = std::complex; int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::conjtrans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped @@ -279,11 +271,7 @@ TEST_P(SparseSpmvUsmTests, ComplexDoublePrecision) { CHECK_DOUBLE_ON_DEVICE(GetParam()); int num_passed = 0, num_skipped = 0; test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::nontrans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::trans, num_passed, num_skipped); - test_helper(test_spmv, test_spmv, GetParam(), - oneapi::mkl::transpose::conjtrans, num_passed, num_skipped); + num_passed, num_skipped); if (num_skipped > 0) { // Mark that some tests were skipped GTEST_SKIP() << "Passed: " << num_passed << ", Skipped: " << num_skipped From ced6a4d4b4ca38009f7146f96060d5a6a9e7856b Mon Sep 17 00:00:00 2001 From: Andrei Fedorov Date: Tue, 9 Jul 2024 17:35:19 +0200 Subject: [PATCH 05/15] Update link to oneMKL Spec (#531) --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 386969e3c..e4c392799 100644 --- a/README.md +++ b/README.md @@ -2,7 +2,7 @@ # oneAPI Math Kernel Library (oneMKL) Interfaces -oneMKL Interfaces is an open-source implementation of the oneMKL Data Parallel C++ (DPC++) interface according to the [oneMKL specification](https://spec.oneapi.com/versions/latest/elements/oneMKL/source/index.html). It works with multiple devices (backends) using device-specific libraries underneath. +oneMKL Interfaces is an open-source implementation of the oneMKL Data Parallel C++ (DPC++) interface according to the [oneMKL specification](https://oneapi-spec.uxlfoundation.org/specifications/oneapi/latest/elements/onemkl/source/). It works with multiple devices (backends) using device-specific libraries underneath. oneMKL is part of the [UXL Foundation](http://www.uxlfoundation.org).

From c1cb8bd65c7c6c70d37f2776b1332d341cf97382 Mon Sep 17 00:00:00 2001 From: Nicolas Miller Date: Wed, 10 Jul 2024 09:08:01 +0100 Subject: [PATCH 06/15] [BLAS][HIP] Fix blas support for rocBLAS 4+ (#519) --- src/blas/backends/rocblas/rocblas_level3.cpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/src/blas/backends/rocblas/rocblas_level3.cpp b/src/blas/backends/rocblas/rocblas_level3.cpp index ec6dd220d..ef739a88b 100644 --- a/src/blas/backends/rocblas/rocblas_level3.cpp +++ b/src/blas/backends/rocblas/rocblas_level3.cpp @@ -381,10 +381,17 @@ inline void trmm(Func func, sycl::queue &queue, side left_right, uplo upper_lowe auto a_ = sc.get_mem(a_acc); auto b_ = sc.get_mem(b_acc); rocblas_status err; +#if ROCBLAS_VERSION_MAJOR >= 4 + ROCBLAS_ERROR_FUNC_SYNC(func, err, handle, get_rocblas_side_mode(left_right), + get_rocblas_fill_mode(upper_lower), + get_rocblas_operation(trans), get_rocblas_diag_type(unit_diag), + m, n, (rocDataType *)&alpha, a_, lda, b_, ldb, b_, ldb); +#else ROCBLAS_ERROR_FUNC_SYNC(func, err, handle, get_rocblas_side_mode(left_right), get_rocblas_fill_mode(upper_lower), get_rocblas_operation(trans), get_rocblas_diag_type(unit_diag), m, n, (rocDataType *)&alpha, a_, lda, b_, ldb); +#endif }); }); } @@ -805,10 +812,17 @@ inline sycl::event trmm(Func func, sycl::queue &queue, side left_right, uplo upp auto a_ = reinterpret_cast(a); auto b_ = reinterpret_cast(b); rocblas_status err; +#if ROCBLAS_VERSION_MAJOR >= 4 + ROCBLAS_ERROR_FUNC_SYNC(func, err, handle, get_rocblas_side_mode(left_right), + get_rocblas_fill_mode(upper_lower), + get_rocblas_operation(trans), get_rocblas_diag_type(unit_diag), + m, n, (rocDataType *)&alpha, a_, lda, b_, ldb, b_, ldb); +#else ROCBLAS_ERROR_FUNC_SYNC(func, err, handle, get_rocblas_side_mode(left_right), get_rocblas_fill_mode(upper_lower), get_rocblas_operation(trans), get_rocblas_diag_type(unit_diag), m, n, (rocDataType *)&alpha, a_, lda, b_, ldb); +#endif }); }); From bb1941d2c692f302765dcdb5c53c92cf9f2c7d68 Mon Sep 17 00:00:00 2001 From: Romain Biessy Date: Thu, 11 Jul 2024 10:29:24 +0200 Subject: [PATCH 07/15] Fix README documentation link (#530) --- README.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index e4c392799..e74e3b5ed 100644 --- a/README.md +++ b/README.md @@ -525,7 +525,8 @@ Product | Supported Version | License - [About](https://oneapi-src.github.io/oneMKL/introduction.html) - Get Started - [Selecting a Compiler](https://oneapi-src.github.io/oneMKL/selecting_a_compiler.html) - - [Building the Project](https://oneapi-src.github.io/oneMKL/building_the_project.html) + - [Building the Project with DPC++](https://oneapi-src.github.io/oneMKL/building_the_project_with_dpcpp.html) + - [Building the Project with AdaptiveCpp](https://oneapi-src.github.io/oneMKL/building_the_project_with_adaptivecpp.html) - Developer Reference - [oneMKL Defined Datatypes](https://oneapi-src.github.io/oneMKL/onemkl-datatypes.html) - [Dense Linear Algebra](https://oneapi-src.github.io/oneMKL/domains/dense_linear_algebra.html) From 0bac3d49c56c48499e661d0cd4199407cbaabcc9 Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Fri, 12 Jul 2024 11:36:52 +0200 Subject: [PATCH 08/15] Merge operations in one file --- .../backends/mklcpu/CMakeLists.txt | 4 +-- ...{mklcpu_spmm.cpp => mklcpu_operations.cpp} | 2 ++ .../backends/mklcpu/mklcpu_spmv.cpp | 30 ------------------- .../backends/mklcpu/mklcpu_spsv.cpp | 30 ------------------- .../backends/mklgpu/CMakeLists.txt | 4 +-- ...{mklgpu_spmv.cpp => mklgpu_operations.cpp} | 2 ++ .../backends/mklgpu/mklgpu_spmm.cpp | 30 ------------------- .../backends/mklgpu/mklgpu_spsv.cpp | 30 ------------------- 8 files changed, 6 insertions(+), 126 deletions(-) rename src/sparse_blas/backends/mklcpu/{mklcpu_spmm.cpp => mklcpu_operations.cpp} (91%) delete mode 100644 src/sparse_blas/backends/mklcpu/mklcpu_spmv.cpp delete mode 100644 src/sparse_blas/backends/mklcpu/mklcpu_spsv.cpp rename src/sparse_blas/backends/mklgpu/{mklgpu_spmv.cpp => mklgpu_operations.cpp} (91%) delete mode 100644 src/sparse_blas/backends/mklgpu/mklgpu_spmm.cpp delete mode 100644 src/sparse_blas/backends/mklgpu/mklgpu_spsv.cpp diff --git a/src/sparse_blas/backends/mklcpu/CMakeLists.txt b/src/sparse_blas/backends/mklcpu/CMakeLists.txt index c851db8bc..e41cae268 100644 --- a/src/sparse_blas/backends/mklcpu/CMakeLists.txt +++ b/src/sparse_blas/backends/mklcpu/CMakeLists.txt @@ -25,9 +25,7 @@ include(WarningsUtils) add_library(${LIB_NAME}) add_library(${LIB_OBJ} OBJECT mklcpu_handles.cpp - mklcpu_spmm.cpp - mklcpu_spmv.cpp - mklcpu_spsv.cpp + mklcpu_operations.cpp $<$: mklcpu_wrappers.cpp> ) add_dependencies(onemkl_backend_libs_sparse_blas ${LIB_NAME}) diff --git a/src/sparse_blas/backends/mklcpu/mklcpu_spmm.cpp b/src/sparse_blas/backends/mklcpu/mklcpu_operations.cpp similarity index 91% rename from src/sparse_blas/backends/mklcpu/mklcpu_spmm.cpp rename to src/sparse_blas/backends/mklcpu/mklcpu_operations.cpp index de34dcb4d..4e0242c2d 100644 --- a/src/sparse_blas/backends/mklcpu/mklcpu_spmm.cpp +++ b/src/sparse_blas/backends/mklcpu/mklcpu_operations.cpp @@ -26,5 +26,7 @@ namespace oneapi::mkl::sparse::mklcpu { #include "sparse_blas/backends/mkl_common/mkl_spmm.cxx" +#include "sparse_blas/backends/mkl_common/mkl_spmv.cxx" +#include "sparse_blas/backends/mkl_common/mkl_spsv.cxx" } // namespace oneapi::mkl::sparse::mklcpu diff --git a/src/sparse_blas/backends/mklcpu/mklcpu_spmv.cpp b/src/sparse_blas/backends/mklcpu/mklcpu_spmv.cpp deleted file mode 100644 index ffbba2f5e..000000000 --- a/src/sparse_blas/backends/mklcpu/mklcpu_spmv.cpp +++ /dev/null @@ -1,30 +0,0 @@ -/*************************************************************************** -* Copyright (C) Codeplay Software Limited -* 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 -* -* For your convenience, a copy of the License has been included in this -* repository. -* -* 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. -* -**************************************************************************/ - -#include "sparse_blas/backends/mkl_common/mkl_helper.hpp" -#include "sparse_blas/macros.hpp" -#include "sparse_blas/backends/mkl_common/mkl_handles.hpp" - -#include "oneapi/mkl/sparse_blas/detail/mklcpu/onemkl_sparse_blas_mklcpu.hpp" - -namespace oneapi::mkl::sparse::mklcpu { - -#include "sparse_blas/backends/mkl_common/mkl_spmv.cxx" - -} // namespace oneapi::mkl::sparse::mklcpu diff --git a/src/sparse_blas/backends/mklcpu/mklcpu_spsv.cpp b/src/sparse_blas/backends/mklcpu/mklcpu_spsv.cpp deleted file mode 100644 index aff9b4abc..000000000 --- a/src/sparse_blas/backends/mklcpu/mklcpu_spsv.cpp +++ /dev/null @@ -1,30 +0,0 @@ -/*************************************************************************** -* Copyright (C) Codeplay Software Limited -* 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 -* -* For your convenience, a copy of the License has been included in this -* repository. -* -* 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. -* -**************************************************************************/ - -#include "sparse_blas/backends/mkl_common/mkl_helper.hpp" -#include "sparse_blas/macros.hpp" -#include "sparse_blas/backends/mkl_common/mkl_handles.hpp" - -#include "oneapi/mkl/sparse_blas/detail/mklcpu/onemkl_sparse_blas_mklcpu.hpp" - -namespace oneapi::mkl::sparse::mklcpu { - -#include "sparse_blas/backends/mkl_common/mkl_spsv.cxx" - -} // namespace oneapi::mkl::sparse::mklcpu diff --git a/src/sparse_blas/backends/mklgpu/CMakeLists.txt b/src/sparse_blas/backends/mklgpu/CMakeLists.txt index b83a39297..cd25babc2 100644 --- a/src/sparse_blas/backends/mklgpu/CMakeLists.txt +++ b/src/sparse_blas/backends/mklgpu/CMakeLists.txt @@ -25,9 +25,7 @@ include(WarningsUtils) add_library(${LIB_NAME}) add_library(${LIB_OBJ} OBJECT mklgpu_handles.cpp - mklgpu_spmm.cpp - mklgpu_spmv.cpp - mklgpu_spsv.cpp + mklgpu_operations.cpp $<$: mklgpu_wrappers.cpp> ) add_dependencies(onemkl_backend_libs_sparse_blas ${LIB_NAME}) diff --git a/src/sparse_blas/backends/mklgpu/mklgpu_spmv.cpp b/src/sparse_blas/backends/mklgpu/mklgpu_operations.cpp similarity index 91% rename from src/sparse_blas/backends/mklgpu/mklgpu_spmv.cpp rename to src/sparse_blas/backends/mklgpu/mklgpu_operations.cpp index 6a7087a86..0c5a73fb0 100644 --- a/src/sparse_blas/backends/mklgpu/mklgpu_spmv.cpp +++ b/src/sparse_blas/backends/mklgpu/mklgpu_operations.cpp @@ -25,6 +25,8 @@ namespace oneapi::mkl::sparse::mklgpu { +#include "sparse_blas/backends/mkl_common/mkl_spmm.cxx" #include "sparse_blas/backends/mkl_common/mkl_spmv.cxx" +#include "sparse_blas/backends/mkl_common/mkl_spsv.cxx" } // namespace oneapi::mkl::sparse::mklgpu diff --git a/src/sparse_blas/backends/mklgpu/mklgpu_spmm.cpp b/src/sparse_blas/backends/mklgpu/mklgpu_spmm.cpp deleted file mode 100644 index ca2338787..000000000 --- a/src/sparse_blas/backends/mklgpu/mklgpu_spmm.cpp +++ /dev/null @@ -1,30 +0,0 @@ -/*************************************************************************** -* Copyright (C) Codeplay Software Limited -* 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 -* -* For your convenience, a copy of the License has been included in this -* repository. -* -* 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. -* -**************************************************************************/ - -#include "sparse_blas/backends/mkl_common/mkl_helper.hpp" -#include "sparse_blas/macros.hpp" -#include "sparse_blas/backends/mkl_common/mkl_handles.hpp" - -#include "oneapi/mkl/sparse_blas/detail/mklgpu/onemkl_sparse_blas_mklgpu.hpp" - -namespace oneapi::mkl::sparse::mklgpu { - -#include "sparse_blas/backends/mkl_common/mkl_spmm.cxx" - -} // namespace oneapi::mkl::sparse::mklgpu diff --git a/src/sparse_blas/backends/mklgpu/mklgpu_spsv.cpp b/src/sparse_blas/backends/mklgpu/mklgpu_spsv.cpp deleted file mode 100644 index b42d4539f..000000000 --- a/src/sparse_blas/backends/mklgpu/mklgpu_spsv.cpp +++ /dev/null @@ -1,30 +0,0 @@ -/*************************************************************************** -* Copyright (C) Codeplay Software Limited -* 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 -* -* For your convenience, a copy of the License has been included in this -* repository. -* -* 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. -* -**************************************************************************/ - -#include "sparse_blas/backends/mkl_common/mkl_helper.hpp" -#include "sparse_blas/macros.hpp" -#include "sparse_blas/backends/mkl_common/mkl_handles.hpp" - -#include "oneapi/mkl/sparse_blas/detail/mklgpu/onemkl_sparse_blas_mklgpu.hpp" - -namespace oneapi::mkl::sparse::mklgpu { - -#include "sparse_blas/backends/mkl_common/mkl_spsv.cxx" - -} // namespace oneapi::mkl::sparse::mklgpu From d04452ac2c020f5c0336d05290d402afeb123713 Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Fri, 12 Jul 2024 11:41:57 +0200 Subject: [PATCH 09/15] Make get_data_type constexpr --- src/sparse_blas/enum_data_types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/sparse_blas/enum_data_types.hpp b/src/sparse_blas/enum_data_types.hpp index 7114482ee..26946facb 100644 --- a/src/sparse_blas/enum_data_types.hpp +++ b/src/sparse_blas/enum_data_types.hpp @@ -40,7 +40,7 @@ inline std::string data_type_to_str(data_type data_type) { } template -data_type get_data_type() { +constexpr data_type get_data_type() { if constexpr (std::is_same_v) { return data_type::int32; } From 8abfdfcbb9fe7fe2e6eea1214478f62e30e4db18 Mon Sep 17 00:00:00 2001 From: Ivan Martianov Date: Fri, 12 Jul 2024 16:02:54 +0100 Subject: [PATCH 10/15] [RNG] Workaround for get_multi_ptr with AdaptiveCPP (#535) --- examples/rng/device/include/rng_example_helper.hpp | 10 ++++++++++ src/rng/backends/mklcpu/cpu_common.hpp | 9 +++++++-- .../rng/device/include/rng_device_test_common.hpp | 8 ++++++++ 3 files changed, 25 insertions(+), 2 deletions(-) diff --git a/examples/rng/device/include/rng_example_helper.hpp b/examples/rng/device/include/rng_example_helper.hpp index ad0fdea03..0bcf114b4 100644 --- a/examples/rng/device/include/rng_example_helper.hpp +++ b/examples/rng/device/include/rng_example_helper.hpp @@ -29,12 +29,22 @@ struct has_member_code_meta().get_multi_ template ::value>::type* = nullptr> auto get_multi_ptr(T acc) { +// Workaround for AdaptiveCPP, as they do not yet support the get_multi_ptr function +#ifndef __HIPSYCL__ return acc.get_multi_ptr(); +#else + return acc.get_pointer(); +#endif }; template ::value>::type* = nullptr> auto get_multi_ptr(T acc) { +// Workaround for AdaptiveCPP, as they do not yet support the get_multi_ptr function +#ifndef __HIPSYCL__ return acc.template get_multi_ptr(); +#else + return acc.get_pointer(); +#endif }; #endif // _RNG_EXAMPLE_HELPER_HPP__ diff --git a/src/rng/backends/mklcpu/cpu_common.hpp b/src/rng/backends/mklcpu/cpu_common.hpp index b5e10585c..cbd6cae59 100644 --- a/src/rng/backends/mklcpu/cpu_common.hpp +++ b/src/rng/backends/mklcpu/cpu_common.hpp @@ -56,9 +56,14 @@ class kernel_name {}; template class kernel_name_usm {}; -template -T *get_raw_ptr(sycl::accessor acc) { +template +typename Acc::value_type *get_raw_ptr(Acc acc) { +// Workaround for AdaptiveCPP, as they do not yet support the get_multi_ptr function +#ifndef __HIPSYCL__ return acc.template get_multi_ptr().get_raw(); +#else + return acc.get_pointer(); +#endif } } // namespace mklcpu diff --git a/tests/unit_tests/rng/device/include/rng_device_test_common.hpp b/tests/unit_tests/rng/device/include/rng_device_test_common.hpp index e6f2b5026..6b014f0ec 100644 --- a/tests/unit_tests/rng/device/include/rng_device_test_common.hpp +++ b/tests/unit_tests/rng/device/include/rng_device_test_common.hpp @@ -112,12 +112,20 @@ struct has_member_code_meta().get_multi_ template ::value>::type* = nullptr> auto get_multi_ptr(T acc) { +#ifndef __HIPSYCL__ return acc.get_multi_ptr(); +#else + return acc.get_pointer(); +#endif }; template ::value>::type* = nullptr> auto get_multi_ptr(T acc) { +#ifndef __HIPSYCL__ return acc.template get_multi_ptr(); +#else + return acc.get_pointer(); +#endif }; template From e8eac877710b5f802d6a9953914059417e074c2b Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Mon, 15 Jul 2024 18:20:03 +0200 Subject: [PATCH 11/15] Remove unused macro TEST_RUN_CT_SELECT --- tests/unit_tests/include/test_helper.hpp | 30 ++++++------------- .../sparse_blas/include/test_common.hpp | 2 +- 2 files changed, 10 insertions(+), 22 deletions(-) diff --git a/tests/unit_tests/include/test_helper.hpp b/tests/unit_tests/include/test_helper.hpp index 7e0024195..ad215761f 100644 --- a/tests/unit_tests/include/test_helper.hpp +++ b/tests/unit_tests/include/test_helper.hpp @@ -203,7 +203,7 @@ TEST_RUN_PORTFFT_SELECT_NO_ARGS(q, func); \ } while (0); -#define TEST_RUN_CT_SELECT(q, func, ...) \ +#define TEST_RUN_BLAS_CT_SELECT(q, func, ...) \ do { \ if (CHECK_HOST_OR_CPU(q)) \ TEST_RUN_INTELCPU_SELECT(q, func, __VA_ARGS__); \ @@ -214,21 +214,15 @@ TEST_RUN_INTELGPU_SELECT(q, func, __VA_ARGS__); \ else if (vendor_id == NVIDIA_ID) { \ TEST_RUN_NVIDIAGPU_CUBLAS_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_NVIDIAGPU_CUSOLVER_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_NVIDIAGPU_CURAND_SELECT(q, func, __VA_ARGS__); \ } \ else if (vendor_id == AMD_ID) { \ TEST_RUN_AMDGPU_ROCBLAS_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_AMDGPU_ROCRAND_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_AMDGPU_ROCSOLVER_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_AMDGPU_ROCFFT_SELECT(q, func, __VA_ARGS__); \ } \ } \ TEST_RUN_PORTBLAS_SELECT(q, func, __VA_ARGS__); \ - TEST_RUN_PORTFFT_SELECT(q, func, __VA_ARGS__); \ } while (0); -#define TEST_RUN_BLAS_CT_SELECT(q, func, ...) \ +#define TEST_RUN_RNG_CT_SELECT(q, func, ...) \ do { \ if (CHECK_HOST_OR_CPU(q)) \ TEST_RUN_INTELCPU_SELECT(q, func, __VA_ARGS__); \ @@ -238,16 +232,15 @@ if (vendor_id == INTEL_ID) \ TEST_RUN_INTELGPU_SELECT(q, func, __VA_ARGS__); \ else if (vendor_id == NVIDIA_ID) { \ - TEST_RUN_NVIDIAGPU_CUBLAS_SELECT(q, func, __VA_ARGS__); \ + TEST_RUN_NVIDIAGPU_CURAND_SELECT(q, func, __VA_ARGS__); \ } \ else if (vendor_id == AMD_ID) { \ - TEST_RUN_AMDGPU_ROCBLAS_SELECT(q, func, __VA_ARGS__); \ + TEST_RUN_AMDGPU_ROCRAND_SELECT(q, func, __VA_ARGS__); \ } \ } \ - TEST_RUN_PORTBLAS_SELECT(q, func, __VA_ARGS__); \ } while (0); -#define TEST_RUN_RNG_CT_SELECT(q, func, ...) \ +#define TEST_RUN_LAPACK_CT_SELECT(q, func, ...) \ do { \ if (CHECK_HOST_OR_CPU(q)) \ TEST_RUN_INTELCPU_SELECT(q, func, __VA_ARGS__); \ @@ -257,28 +250,23 @@ if (vendor_id == INTEL_ID) \ TEST_RUN_INTELGPU_SELECT(q, func, __VA_ARGS__); \ else if (vendor_id == NVIDIA_ID) { \ - TEST_RUN_NVIDIAGPU_CURAND_SELECT(q, func, __VA_ARGS__); \ + TEST_RUN_NVIDIAGPU_CUSOLVER_SELECT(q, func, __VA_ARGS__); \ } \ else if (vendor_id == AMD_ID) { \ - TEST_RUN_AMDGPU_ROCRAND_SELECT(q, func, __VA_ARGS__); \ + TEST_RUN_AMDGPU_ROCSOLVER_SELECT(q, func, __VA_ARGS__); \ } \ } \ } while (0); -#define TEST_RUN_LAPACK_CT_SELECT(q, func, ...) \ +#define TEST_RUN_SPARSE_CT_SELECT(q, func, ...) \ do { \ if (CHECK_HOST_OR_CPU(q)) \ TEST_RUN_INTELCPU_SELECT(q, func, __VA_ARGS__); \ else if (q.get_device().is_gpu()) { \ unsigned int vendor_id = static_cast( \ q.get_device().get_info()); \ - if (vendor_id == INTEL_ID) \ + if (vendor_id == INTEL_ID) { \ TEST_RUN_INTELGPU_SELECT(q, func, __VA_ARGS__); \ - else if (vendor_id == NVIDIA_ID) { \ - TEST_RUN_NVIDIAGPU_CUSOLVER_SELECT(q, func, __VA_ARGS__); \ - } \ - else if (vendor_id == AMD_ID) { \ - TEST_RUN_AMDGPU_ROCSOLVER_SELECT(q, func, __VA_ARGS__); \ } \ } \ } while (0); diff --git a/tests/unit_tests/sparse_blas/include/test_common.hpp b/tests/unit_tests/sparse_blas/include/test_common.hpp index 48d0b13ee..d579ff5fb 100644 --- a/tests/unit_tests/sparse_blas/include/test_common.hpp +++ b/tests/unit_tests/sparse_blas/include/test_common.hpp @@ -39,7 +39,7 @@ #ifdef CALL_RT_API #define CALL_RT_OR_CT(FUNC, QUEUE, ...) FUNC(QUEUE, __VA_ARGS__) #else -#define CALL_RT_OR_CT(FUNC, QUEUE, ...) TEST_RUN_CT_SELECT(QUEUE, FUNC, __VA_ARGS__); +#define CALL_RT_OR_CT(FUNC, QUEUE, ...) TEST_RUN_SPARSE_CT_SELECT(QUEUE, FUNC, __VA_ARGS__); #endif template From b4955cae8240c8ca6b088160abaadb7fd0a245b3 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Tue, 16 Jul 2024 09:11:36 +0100 Subject: [PATCH 12/15] [BLAS][DFT][SPARSE][TESTS] Remove .template from `get_host_access` to fix build (#534) Signed-off-by: JackAKirk --- tests/unit_tests/blas/batch/axpy_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/copy_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/dgmm_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/gemm_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/gemv_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/imatcopy_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/omatadd_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/omatcopy_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/syrk_batch_stride.cpp | 2 +- tests/unit_tests/blas/batch/trsm_batch_stride.cpp | 2 +- tests/unit_tests/blas/extensions/gemm_bias.cpp | 2 +- tests/unit_tests/blas/extensions/gemmt.cpp | 2 +- tests/unit_tests/blas/extensions/imatcopy.cpp | 2 +- tests/unit_tests/blas/extensions/omatadd.cpp | 2 +- tests/unit_tests/blas/extensions/omatcopy.cpp | 2 +- tests/unit_tests/blas/extensions/omatcopy2.cpp | 2 +- tests/unit_tests/blas/level1/asum.cpp | 2 +- tests/unit_tests/blas/level1/axpby.cpp | 2 +- tests/unit_tests/blas/level1/axpy.cpp | 2 +- tests/unit_tests/blas/level1/copy.cpp | 2 +- tests/unit_tests/blas/level1/dot.cpp | 2 +- tests/unit_tests/blas/level1/dotc.cpp | 2 +- tests/unit_tests/blas/level1/dotu.cpp | 2 +- tests/unit_tests/blas/level1/iamax.cpp | 2 +- tests/unit_tests/blas/level1/iamin.cpp | 2 +- tests/unit_tests/blas/level1/nrm2.cpp | 2 +- tests/unit_tests/blas/level1/rot.cpp | 4 ++-- tests/unit_tests/blas/level1/rotg.cpp | 8 ++++---- tests/unit_tests/blas/level1/rotm.cpp | 4 ++-- tests/unit_tests/blas/level1/rotmg.cpp | 8 ++++---- tests/unit_tests/blas/level1/scal.cpp | 2 +- tests/unit_tests/blas/level1/sdsdot.cpp | 2 +- tests/unit_tests/blas/level1/swap.cpp | 4 ++-- tests/unit_tests/blas/level2/gbmv.cpp | 2 +- tests/unit_tests/blas/level2/gemv.cpp | 2 +- tests/unit_tests/blas/level2/ger.cpp | 2 +- tests/unit_tests/blas/level2/gerc.cpp | 2 +- tests/unit_tests/blas/level2/geru.cpp | 2 +- tests/unit_tests/blas/level2/hbmv.cpp | 2 +- tests/unit_tests/blas/level2/hemv.cpp | 2 +- tests/unit_tests/blas/level2/her.cpp | 2 +- tests/unit_tests/blas/level2/her2.cpp | 2 +- tests/unit_tests/blas/level2/hpmv.cpp | 2 +- tests/unit_tests/blas/level2/hpr.cpp | 2 +- tests/unit_tests/blas/level2/hpr2.cpp | 2 +- tests/unit_tests/blas/level2/sbmv.cpp | 2 +- tests/unit_tests/blas/level2/spmv.cpp | 2 +- tests/unit_tests/blas/level2/spr.cpp | 2 +- tests/unit_tests/blas/level2/spr2.cpp | 2 +- tests/unit_tests/blas/level2/symv.cpp | 2 +- tests/unit_tests/blas/level2/syr.cpp | 2 +- tests/unit_tests/blas/level2/syr2.cpp | 2 +- tests/unit_tests/blas/level2/tbmv.cpp | 2 +- tests/unit_tests/blas/level2/tbsv.cpp | 2 +- tests/unit_tests/blas/level2/tpmv.cpp | 2 +- tests/unit_tests/blas/level2/tpsv.cpp | 2 +- tests/unit_tests/blas/level2/trmv.cpp | 2 +- tests/unit_tests/blas/level2/trsv.cpp | 2 +- tests/unit_tests/blas/level3/gemm.cpp | 2 +- tests/unit_tests/blas/level3/hemm.cpp | 2 +- tests/unit_tests/blas/level3/her2k.cpp | 2 +- tests/unit_tests/blas/level3/herk.cpp | 2 +- tests/unit_tests/blas/level3/symm.cpp | 2 +- tests/unit_tests/blas/level3/syr2k.cpp | 2 +- tests/unit_tests/blas/level3/syrk.cpp | 2 +- tests/unit_tests/blas/level3/trmm.cpp | 2 +- tests/unit_tests/blas/level3/trsm.cpp | 2 +- tests/unit_tests/dft/include/compute_inplace.hpp | 2 +- .../unit_tests/dft/include/compute_inplace_real_real.hpp | 8 ++++---- tests/unit_tests/dft/include/compute_out_of_place.hpp | 2 +- .../dft/include/compute_out_of_place_real_real.hpp | 8 ++++---- .../unit_tests/sparse_blas/source/sparse_gemm_buffer.cpp | 2 +- .../unit_tests/sparse_blas/source/sparse_gemv_buffer.cpp | 2 +- .../unit_tests/sparse_blas/source/sparse_trsv_buffer.cpp | 2 +- 74 files changed, 89 insertions(+), 89 deletions(-) diff --git a/tests/unit_tests/blas/batch/axpy_batch_stride.cpp b/tests/unit_tests/blas/batch/axpy_batch_stride.cpp index c082cc543..9bb1406ef 100644 --- a/tests/unit_tests/blas/batch/axpy_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/axpy_batch_stride.cpp @@ -150,7 +150,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t incx, int64_t incy, fp // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = true; for (i = 0; i < batch_size; i++) { good = good && check_equal_vector(y_accessor.get_pointer() + i * stride_y, diff --git a/tests/unit_tests/blas/batch/copy_batch_stride.cpp b/tests/unit_tests/blas/batch/copy_batch_stride.cpp index 8b8571cb8..a1da595f6 100644 --- a/tests/unit_tests/blas/batch/copy_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/copy_batch_stride.cpp @@ -147,7 +147,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t incx, int64_t incy, in // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = true; for (i = 0; i < batch_size; i++) { good = good && check_equal_vector(y_accessor.get_pointer() + i * stride_y, diff --git a/tests/unit_tests/blas/batch/dgmm_batch_stride.cpp b/tests/unit_tests/blas/batch/dgmm_batch_stride.cpp index 3016cd28b..bb642c3ee 100644 --- a/tests/unit_tests/blas/batch/dgmm_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/dgmm_batch_stride.cpp @@ -166,7 +166,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::side left_right, // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = true; for (i = 0; i < batch_size; i++) { good = good && diff --git a/tests/unit_tests/blas/batch/gemm_batch_stride.cpp b/tests/unit_tests/blas/batch/gemm_batch_stride.cpp index 12af18ec9..5241cb822 100644 --- a/tests/unit_tests/blas/batch/gemm_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/gemm_batch_stride.cpp @@ -223,7 +223,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t batch_size) { for (size_t i = 0; i < C_ref.size(); ++i) { C_cast_ref[i] = C_ref[i]; } - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_almost_equal_matrix(C_accessor, C_cast_ref, oneapi::mkl::layout::col_major, stride_c * batch_size, 1, stride_c * batch_size, error_mag, std::cout); diff --git a/tests/unit_tests/blas/batch/gemv_batch_stride.cpp b/tests/unit_tests/blas/batch/gemv_batch_stride.cpp index 0cf26d7da..bd92f70ca 100644 --- a/tests/unit_tests/blas/batch/gemv_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/gemv_batch_stride.cpp @@ -181,7 +181,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t incx, int64_t incy, in // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = true; for (i = 0; i < batch_size; i++) { good = good && check_equal_vector(y_accessor.get_pointer() + i * stride_y, diff --git a/tests/unit_tests/blas/batch/imatcopy_batch_stride.cpp b/tests/unit_tests/blas/batch/imatcopy_batch_stride.cpp index 0b7b1c316..ac8bbb2b4 100644 --- a/tests/unit_tests/blas/batch/imatcopy_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/imatcopy_batch_stride.cpp @@ -163,7 +163,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t batch_size) { // Compare the results of reference implementation and DPC++ implementation. - auto AB_accessor = AB_buffer.template get_host_access(read_only); + auto AB_accessor = AB_buffer.get_host_access(read_only); bool good = check_equal_matrix(AB_accessor, AB_ref, oneapi::mkl::layout::col_major, stride * batch_size, 1, stride * batch_size, 10, std::cout); diff --git a/tests/unit_tests/blas/batch/omatadd_batch_stride.cpp b/tests/unit_tests/blas/batch/omatadd_batch_stride.cpp index ed8928ac1..cc20d0e3b 100644 --- a/tests/unit_tests/blas/batch/omatadd_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/omatadd_batch_stride.cpp @@ -179,7 +179,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t batch_size) { // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, oneapi::mkl::layout::col_major, stride_c * batch_size, 1, stride_c * batch_size, 10, std::cout); diff --git a/tests/unit_tests/blas/batch/omatcopy_batch_stride.cpp b/tests/unit_tests/blas/batch/omatcopy_batch_stride.cpp index f5f4746f7..d08329fc6 100644 --- a/tests/unit_tests/blas/batch/omatcopy_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/omatcopy_batch_stride.cpp @@ -166,7 +166,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t batch_size) { // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_matrix(B_accessor, B_ref, oneapi::mkl::layout::col_major, stride_b * batch_size, 1, stride_b * batch_size, 10, std::cout); diff --git a/tests/unit_tests/blas/batch/syrk_batch_stride.cpp b/tests/unit_tests/blas/batch/syrk_batch_stride.cpp index d33ce7285..58dc4d7dc 100644 --- a/tests/unit_tests/blas/batch/syrk_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/syrk_batch_stride.cpp @@ -185,7 +185,7 @@ int test(device *dev, oneapi::mkl::layout layout, int64_t batch_size) { // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, oneapi::mkl::layout::col_major, stride_c * batch_size, 1, stride_c * batch_size, 10 * k, std::cout); diff --git a/tests/unit_tests/blas/batch/trsm_batch_stride.cpp b/tests/unit_tests/blas/batch/trsm_batch_stride.cpp index e7ca5c29c..cde6aa367 100644 --- a/tests/unit_tests/blas/batch/trsm_batch_stride.cpp +++ b/tests/unit_tests/blas/batch/trsm_batch_stride.cpp @@ -192,7 +192,7 @@ int test(device *dev, oneapi::mkl::layout layout) { } // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_trsm_matrix(B_accessor, B_ref, oneapi::mkl::layout::col_major, total_size_b, 1, total_size_b, 10 * std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/extensions/gemm_bias.cpp b/tests/unit_tests/blas/extensions/gemm_bias.cpp index 399514466..c6e99e829 100644 --- a/tests/unit_tests/blas/extensions/gemm_bias.cpp +++ b/tests/unit_tests/blas/extensions/gemm_bias.cpp @@ -155,7 +155,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::transpose transa, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, m, n, ldc, 10 * k, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/extensions/gemmt.cpp b/tests/unit_tests/blas/extensions/gemmt.cpp index c56beb778..228a85d33 100644 --- a/tests/unit_tests/blas/extensions/gemmt.cpp +++ b/tests/unit_tests/blas/extensions/gemmt.cpp @@ -136,7 +136,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, upper_lower, n, n, ldc, 10 * k, std::cout); diff --git a/tests/unit_tests/blas/extensions/imatcopy.cpp b/tests/unit_tests/blas/extensions/imatcopy.cpp index 773657b70..e21702775 100644 --- a/tests/unit_tests/blas/extensions/imatcopy.cpp +++ b/tests/unit_tests/blas/extensions/imatcopy.cpp @@ -154,7 +154,7 @@ int test(device *dev, oneapi::mkl::layout layout) { // Compare the results of reference implementation and DPC++ implementation. - auto AB_accessor = AB_buffer.template get_host_access(read_only); + auto AB_accessor = AB_buffer.get_host_access(read_only); bool good = check_equal_matrix(AB_accessor, AB_ref, oneapi::mkl::layout::col_major, size, 1, size, 10, std::cout); diff --git a/tests/unit_tests/blas/extensions/omatadd.cpp b/tests/unit_tests/blas/extensions/omatadd.cpp index 4f7e327ab..b2af98935 100644 --- a/tests/unit_tests/blas/extensions/omatadd.cpp +++ b/tests/unit_tests/blas/extensions/omatadd.cpp @@ -170,7 +170,7 @@ int test(device *dev, oneapi::mkl::layout layout) { // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, oneapi::mkl::layout::col_major, size_c, 1, size_c, 10, std::cout); diff --git a/tests/unit_tests/blas/extensions/omatcopy.cpp b/tests/unit_tests/blas/extensions/omatcopy.cpp index e14994aba..122ba2c79 100644 --- a/tests/unit_tests/blas/extensions/omatcopy.cpp +++ b/tests/unit_tests/blas/extensions/omatcopy.cpp @@ -163,7 +163,7 @@ int test(device *dev, oneapi::mkl::layout layout) { // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_matrix(B_accessor, B_ref, oneapi::mkl::layout::col_major, size_b, 1, size_b, 10, std::cout); diff --git a/tests/unit_tests/blas/extensions/omatcopy2.cpp b/tests/unit_tests/blas/extensions/omatcopy2.cpp index fbea9ab71..50dcd0f6b 100644 --- a/tests/unit_tests/blas/extensions/omatcopy2.cpp +++ b/tests/unit_tests/blas/extensions/omatcopy2.cpp @@ -162,7 +162,7 @@ int test(device *dev, oneapi::mkl::layout layout) { // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_matrix(B_accessor, B_ref, oneapi::mkl::layout::col_major, size_b, 1, size_b, 10, std::cout); diff --git a/tests/unit_tests/blas/level1/asum.cpp b/tests/unit_tests/blas/level1/asum.cpp index f0584369a..6969789e3 100644 --- a/tests/unit_tests/blas/level1/asum.cpp +++ b/tests/unit_tests/blas/level1/asum.cpp @@ -119,7 +119,7 @@ int test(device* dev, oneapi::mkl::layout layout, int64_t N, int64_t incx) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/axpby.cpp b/tests/unit_tests/blas/level1/axpby.cpp index 8924e09af..d43f9beda 100644 --- a/tests/unit_tests/blas/level1/axpby.cpp +++ b/tests/unit_tests/blas/level1/axpby.cpp @@ -124,7 +124,7 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy, fp // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/axpy.cpp b/tests/unit_tests/blas/level1/axpy.cpp index 8550f67c7..c81f2902d 100644 --- a/tests/unit_tests/blas/level1/axpy.cpp +++ b/tests/unit_tests/blas/level1/axpy.cpp @@ -124,7 +124,7 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy, fp // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/copy.cpp b/tests/unit_tests/blas/level1/copy.cpp index ecd286260..87a1c2f1b 100644 --- a/tests/unit_tests/blas/level1/copy.cpp +++ b/tests/unit_tests/blas/level1/copy.cpp @@ -122,7 +122,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx, int incy) { // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/dot.cpp b/tests/unit_tests/blas/level1/dot.cpp index 02b194bd5..11cb09bcc 100644 --- a/tests/unit_tests/blas/level1/dot.cpp +++ b/tests/unit_tests/blas/level1/dot.cpp @@ -123,7 +123,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx, int incy) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/dotc.cpp b/tests/unit_tests/blas/level1/dotc.cpp index 948dd0e60..cb8d0fc37 100644 --- a/tests/unit_tests/blas/level1/dotc.cpp +++ b/tests/unit_tests/blas/level1/dotc.cpp @@ -125,7 +125,7 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_reference, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/dotu.cpp b/tests/unit_tests/blas/level1/dotu.cpp index 9d8552359..bbef3ad8c 100644 --- a/tests/unit_tests/blas/level1/dotu.cpp +++ b/tests/unit_tests/blas/level1/dotu.cpp @@ -125,7 +125,7 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_reference, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/iamax.cpp b/tests/unit_tests/blas/level1/iamax.cpp index bf8cc7d96..977f12b5d 100644 --- a/tests/unit_tests/blas/level1/iamax.cpp +++ b/tests/unit_tests/blas/level1/iamax.cpp @@ -120,7 +120,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, 0, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/iamin.cpp b/tests/unit_tests/blas/level1/iamin.cpp index ea81b9048..a52862cb6 100644 --- a/tests/unit_tests/blas/level1/iamin.cpp +++ b/tests/unit_tests/blas/level1/iamin.cpp @@ -120,7 +120,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, 0, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/nrm2.cpp b/tests/unit_tests/blas/level1/nrm2.cpp index d5c8ab46b..423cecb59 100644 --- a/tests/unit_tests/blas/level1/nrm2.cpp +++ b/tests/unit_tests/blas/level1/nrm2.cpp @@ -120,7 +120,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx) { // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/rot.cpp b/tests/unit_tests/blas/level1/rot.cpp index 8cc3d6926..f65540182 100644 --- a/tests/unit_tests/blas/level1/rot.cpp +++ b/tests/unit_tests/blas/level1/rot.cpp @@ -125,9 +125,9 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy, fp_ // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good_x = check_equal_vector(x_accessor, x_ref, N, incx, N, std::cout); - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good_y = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); bool good = good_x && good_y; diff --git a/tests/unit_tests/blas/level1/rotg.cpp b/tests/unit_tests/blas/level1/rotg.cpp index dc9a7908f..1a0d569d8 100644 --- a/tests/unit_tests/blas/level1/rotg.cpp +++ b/tests/unit_tests/blas/level1/rotg.cpp @@ -130,13 +130,13 @@ int test(device *dev, oneapi::mkl::layout layout) { } // Compare the results of reference implementation and DPC++ implementation. - auto a_accessor = a_buffer.template get_host_access(read_only); + auto a_accessor = a_buffer.get_host_access(read_only); bool good_a = check_equal(a_accessor[0], a_ref, 4, std::cout); - auto b_accessor = b_buffer.template get_host_access(read_only); + auto b_accessor = b_buffer.get_host_access(read_only); bool good_b = check_equal(b_accessor[0], b_ref, 4, std::cout); - auto s_accessor = s_buffer.template get_host_access(read_only); + auto s_accessor = s_buffer.get_host_access(read_only); bool good_s = check_equal(s_accessor[0], s_ref, 4, std::cout); - auto c_accessor = c_buffer.template get_host_access(read_only); + auto c_accessor = c_buffer.get_host_access(read_only); bool good_c = check_equal(c_accessor[0], c_ref, 4, std::cout); bool good = good_a && good_b && good_c && good_s; diff --git a/tests/unit_tests/blas/level1/rotm.cpp b/tests/unit_tests/blas/level1/rotm.cpp index 93addfd9b..ab2c599bf 100644 --- a/tests/unit_tests/blas/level1/rotm.cpp +++ b/tests/unit_tests/blas/level1/rotm.cpp @@ -127,9 +127,9 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy, fp } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good_x = check_equal_vector(x_accessor, x_ref, N, incx, N, std::cout); - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good_y = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); bool good = good_x && good_y; diff --git a/tests/unit_tests/blas/level1/rotmg.cpp b/tests/unit_tests/blas/level1/rotmg.cpp index ff4ec64ee..f62bd1cf9 100644 --- a/tests/unit_tests/blas/level1/rotmg.cpp +++ b/tests/unit_tests/blas/level1/rotmg.cpp @@ -130,13 +130,13 @@ int test(device* dev, oneapi::mkl::layout layout) { int error_mag = 50; - auto d1_accessor = d1_buffer.template get_host_access(read_only); + auto d1_accessor = d1_buffer.get_host_access(read_only); bool good_d1 = check_equal(d1_accessor[0], d1_ref, error_mag, std::cout); - auto d2_accessor = d2_buffer.template get_host_access(read_only); + auto d2_accessor = d2_buffer.get_host_access(read_only); bool good_d2 = check_equal(d2_accessor[0], d2_ref, error_mag, std::cout); - auto x1_accessor = x1_buffer.template get_host_access(read_only); + auto x1_accessor = x1_buffer.get_host_access(read_only); bool good_x1 = check_equal(x1_accessor[0], x1_ref, error_mag, std::cout); - auto param_accessor = param_buffer.template get_host_access(read_only); + auto param_accessor = param_buffer.get_host_access(read_only); constexpr fp unit_matrix = -2; constexpr fp rescaled_matrix = -1; diff --git a/tests/unit_tests/blas/level1/scal.cpp b/tests/unit_tests/blas/level1/scal.cpp index 12a103635..8901bb424 100644 --- a/tests/unit_tests/blas/level1/scal.cpp +++ b/tests/unit_tests/blas/level1/scal.cpp @@ -120,7 +120,7 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx, fp_scalar alp } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_vector(x_accessor, x_ref, N, incx, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/sdsdot.cpp b/tests/unit_tests/blas/level1/sdsdot.cpp index e15c53402..7293a3699 100644 --- a/tests/unit_tests/blas/level1/sdsdot.cpp +++ b/tests/unit_tests/blas/level1/sdsdot.cpp @@ -123,7 +123,7 @@ int test(device *dev, oneapi::mkl::layout layout, int N, int incx, int incy, flo // Compare the results of reference implementation and DPC++ implementation. - auto result_accessor = result_buffer.template get_host_access(read_only); + auto result_accessor = result_buffer.get_host_access(read_only); bool good = check_equal(result_accessor[0], result_ref, N, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level1/swap.cpp b/tests/unit_tests/blas/level1/swap.cpp index 291001fc4..6c6721537 100644 --- a/tests/unit_tests/blas/level1/swap.cpp +++ b/tests/unit_tests/blas/level1/swap.cpp @@ -122,8 +122,8 @@ int test(device* dev, oneapi::mkl::layout layout, int N, int incx, int incy) { // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); - auto x_accessor = x_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good_y = check_equal_vector(y_accessor, y_ref, N, incy, N, std::cout); bool good_x = check_equal_vector(x_accessor, x_ref, N, incx, N, std::cout); bool good = good_x && good_y; diff --git a/tests/unit_tests/blas/level2/gbmv.cpp b/tests/unit_tests/blas/level2/gbmv.cpp index fe3956cc6..94fcbc906 100644 --- a/tests/unit_tests/blas/level2/gbmv.cpp +++ b/tests/unit_tests/blas/level2/gbmv.cpp @@ -136,7 +136,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::transpose transa, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, y_len, incy, std::max(m, n), std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/gemv.cpp b/tests/unit_tests/blas/level2/gemv.cpp index 72e901d5d..3bfff4324 100644 --- a/tests/unit_tests/blas/level2/gemv.cpp +++ b/tests/unit_tests/blas/level2/gemv.cpp @@ -133,7 +133,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::transpose transa, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, y_len, incy, std::max(m, n), std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/ger.cpp b/tests/unit_tests/blas/level2/ger.cpp index 93b968d5f..3b32d2827 100644 --- a/tests/unit_tests/blas/level2/ger.cpp +++ b/tests/unit_tests/blas/level2/ger.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, int m, int n, fp alpha, int in } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, m, n, lda, std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level2/gerc.cpp b/tests/unit_tests/blas/level2/gerc.cpp index 4204ccff7..c19c9f029 100644 --- a/tests/unit_tests/blas/level2/gerc.cpp +++ b/tests/unit_tests/blas/level2/gerc.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, int m, int n, fp alpha, int in } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, m, n, lda, std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level2/geru.cpp b/tests/unit_tests/blas/level2/geru.cpp index c71cecc7f..e0cb7c45d 100644 --- a/tests/unit_tests/blas/level2/geru.cpp +++ b/tests/unit_tests/blas/level2/geru.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, int m, int n, fp alpha, int in } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, m, n, lda, std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level2/hbmv.cpp b/tests/unit_tests/blas/level2/hbmv.cpp index b472ab2dc..119aef32a 100644 --- a/tests/unit_tests/blas/level2/hbmv.cpp +++ b/tests/unit_tests/blas/level2/hbmv.cpp @@ -132,7 +132,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/hemv.cpp b/tests/unit_tests/blas/level2/hemv.cpp index fd5b2093d..3636e3774 100644 --- a/tests/unit_tests/blas/level2/hemv.cpp +++ b/tests/unit_tests/blas/level2/hemv.cpp @@ -130,7 +130,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/her.cpp b/tests/unit_tests/blas/level2/her.cpp index c9c94ffb1..46ae9a879 100644 --- a/tests/unit_tests/blas/level2/her.cpp +++ b/tests/unit_tests/blas/level2/her.cpp @@ -125,7 +125,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, lda, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/her2.cpp b/tests/unit_tests/blas/level2/her2.cpp index 3d1709b69..e98c5cc8b 100644 --- a/tests/unit_tests/blas/level2/her2.cpp +++ b/tests/unit_tests/blas/level2/her2.cpp @@ -129,7 +129,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, lda, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/hpmv.cpp b/tests/unit_tests/blas/level2/hpmv.cpp index 69ac869ef..69e6ea9b2 100644 --- a/tests/unit_tests/blas/level2/hpmv.cpp +++ b/tests/unit_tests/blas/level2/hpmv.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/hpr.cpp b/tests/unit_tests/blas/level2/hpr.cpp index f06bec26a..b2e5548bd 100644 --- a/tests/unit_tests/blas/level2/hpr.cpp +++ b/tests/unit_tests/blas/level2/hpr.cpp @@ -125,7 +125,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, n, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/hpr2.cpp b/tests/unit_tests/blas/level2/hpr2.cpp index f6954d284..e2b19e2fd 100644 --- a/tests/unit_tests/blas/level2/hpr2.cpp +++ b/tests/unit_tests/blas/level2/hpr2.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, n, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/sbmv.cpp b/tests/unit_tests/blas/level2/sbmv.cpp index 3d1d79246..c0347dfda 100644 --- a/tests/unit_tests/blas/level2/sbmv.cpp +++ b/tests/unit_tests/blas/level2/sbmv.cpp @@ -130,7 +130,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/spmv.cpp b/tests/unit_tests/blas/level2/spmv.cpp index 06930e73d..799e7d775 100644 --- a/tests/unit_tests/blas/level2/spmv.cpp +++ b/tests/unit_tests/blas/level2/spmv.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/spr.cpp b/tests/unit_tests/blas/level2/spr.cpp index a224f9d31..4e4b5d8a9 100644 --- a/tests/unit_tests/blas/level2/spr.cpp +++ b/tests/unit_tests/blas/level2/spr.cpp @@ -124,7 +124,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, n, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/spr2.cpp b/tests/unit_tests/blas/level2/spr2.cpp index e192f766c..d9d00a4e8 100644 --- a/tests/unit_tests/blas/level2/spr2.cpp +++ b/tests/unit_tests/blas/level2/spr2.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, n, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/symv.cpp b/tests/unit_tests/blas/level2/symv.cpp index e02c68301..a22e48ff7 100644 --- a/tests/unit_tests/blas/level2/symv.cpp +++ b/tests/unit_tests/blas/level2/symv.cpp @@ -129,7 +129,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto y_accessor = y_buffer.template get_host_access(read_only); + auto y_accessor = y_buffer.get_host_access(read_only); bool good = check_equal_vector(y_accessor, y_ref, n, incy, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/syr.cpp b/tests/unit_tests/blas/level2/syr.cpp index 604d1e395..6b305582b 100644 --- a/tests/unit_tests/blas/level2/syr.cpp +++ b/tests/unit_tests/blas/level2/syr.cpp @@ -124,7 +124,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, lda, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/syr2.cpp b/tests/unit_tests/blas/level2/syr2.cpp index 0ec77d4b2..5da1e0106 100644 --- a/tests/unit_tests/blas/level2/syr2.cpp +++ b/tests/unit_tests/blas/level2/syr2.cpp @@ -128,7 +128,7 @@ int test(device *dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto A_accessor = A_buffer.template get_host_access(read_only); + auto A_accessor = A_buffer.get_host_access(read_only); bool good = check_equal_matrix(A_accessor, A_ref, layout, n, n, lda, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/tbmv.cpp b/tests/unit_tests/blas/level2/tbmv.cpp index 5fda0e767..554082a01 100644 --- a/tests/unit_tests/blas/level2/tbmv.cpp +++ b/tests/unit_tests/blas/level2/tbmv.cpp @@ -128,7 +128,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/tbsv.cpp b/tests/unit_tests/blas/level2/tbsv.cpp index c29945d8f..e653105e8 100644 --- a/tests/unit_tests/blas/level2/tbsv.cpp +++ b/tests/unit_tests/blas/level2/tbsv.cpp @@ -128,7 +128,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_trsv_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/tpmv.cpp b/tests/unit_tests/blas/level2/tpmv.cpp index f7c4c680d..ce45279bb 100644 --- a/tests/unit_tests/blas/level2/tpmv.cpp +++ b/tests/unit_tests/blas/level2/tpmv.cpp @@ -126,7 +126,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/tpsv.cpp b/tests/unit_tests/blas/level2/tpsv.cpp index 3fb1bcf71..2a12ab1da 100644 --- a/tests/unit_tests/blas/level2/tpsv.cpp +++ b/tests/unit_tests/blas/level2/tpsv.cpp @@ -126,7 +126,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_trsv_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/trmv.cpp b/tests/unit_tests/blas/level2/trmv.cpp index 616d652c7..8dfc517eb 100644 --- a/tests/unit_tests/blas/level2/trmv.cpp +++ b/tests/unit_tests/blas/level2/trmv.cpp @@ -126,7 +126,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level2/trsv.cpp b/tests/unit_tests/blas/level2/trsv.cpp index e4447c917..fb1e39e06 100644 --- a/tests/unit_tests/blas/level2/trsv.cpp +++ b/tests/unit_tests/blas/level2/trsv.cpp @@ -126,7 +126,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto x_accessor = x_buffer.template get_host_access(read_only); + auto x_accessor = x_buffer.get_host_access(read_only); bool good = check_equal_trsv_vector(x_accessor, x_ref, n, incx, n, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level3/gemm.cpp b/tests/unit_tests/blas/level3/gemm.cpp index 2d23a2bfd..564700b16 100644 --- a/tests/unit_tests/blas/level3/gemm.cpp +++ b/tests/unit_tests/blas/level3/gemm.cpp @@ -139,7 +139,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::transpose transa, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, m, n, ldc, 10 * k, std::cout); return (int)good; diff --git a/tests/unit_tests/blas/level3/hemm.cpp b/tests/unit_tests/blas/level3/hemm.cpp index b711849a5..ce050e97d 100644 --- a/tests/unit_tests/blas/level3/hemm.cpp +++ b/tests/unit_tests/blas/level3/hemm.cpp @@ -138,7 +138,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::side left_right, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, m, n, ldc, 10 * std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level3/her2k.cpp b/tests/unit_tests/blas/level3/her2k.cpp index 847454474..ce57041d9 100644 --- a/tests/unit_tests/blas/level3/her2k.cpp +++ b/tests/unit_tests/blas/level3/her2k.cpp @@ -139,7 +139,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, n, n, ldc, 10 * std::max(n, k), std::cout); diff --git a/tests/unit_tests/blas/level3/herk.cpp b/tests/unit_tests/blas/level3/herk.cpp index fb2acf220..f908a77b7 100644 --- a/tests/unit_tests/blas/level3/herk.cpp +++ b/tests/unit_tests/blas/level3/herk.cpp @@ -130,7 +130,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, n, n, ldc, 10 * std::max(n, k), std::cout); diff --git a/tests/unit_tests/blas/level3/symm.cpp b/tests/unit_tests/blas/level3/symm.cpp index 246c9b678..3f6920370 100644 --- a/tests/unit_tests/blas/level3/symm.cpp +++ b/tests/unit_tests/blas/level3/symm.cpp @@ -138,7 +138,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::side left_right, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, m, n, ldc, 10 * std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level3/syr2k.cpp b/tests/unit_tests/blas/level3/syr2k.cpp index 2b7611755..0153e9ec0 100644 --- a/tests/unit_tests/blas/level3/syr2k.cpp +++ b/tests/unit_tests/blas/level3/syr2k.cpp @@ -134,7 +134,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, n, n, ldc, 10 * std::max(n, k), std::cout); diff --git a/tests/unit_tests/blas/level3/syrk.cpp b/tests/unit_tests/blas/level3/syrk.cpp index 22b2d0916..a6b28735d 100644 --- a/tests/unit_tests/blas/level3/syrk.cpp +++ b/tests/unit_tests/blas/level3/syrk.cpp @@ -129,7 +129,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::uplo upper_lower, } // Compare the results of reference implementation and DPC++ implementation. - auto C_accessor = C_buffer.template get_host_access(read_only); + auto C_accessor = C_buffer.get_host_access(read_only); bool good = check_equal_matrix(C_accessor, C_ref, layout, n, n, ldc, 10 * std::max(n, k), std::cout); diff --git a/tests/unit_tests/blas/level3/trmm.cpp b/tests/unit_tests/blas/level3/trmm.cpp index 6415a91da..2a02aa0d1 100644 --- a/tests/unit_tests/blas/level3/trmm.cpp +++ b/tests/unit_tests/blas/level3/trmm.cpp @@ -138,7 +138,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::side left_right, } // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_matrix(B_accessor, B_ref, layout, m, n, ldb, 10 * std::max(m, n), std::cout); diff --git a/tests/unit_tests/blas/level3/trsm.cpp b/tests/unit_tests/blas/level3/trsm.cpp index d1ccf41f7..90b8d5c93 100644 --- a/tests/unit_tests/blas/level3/trsm.cpp +++ b/tests/unit_tests/blas/level3/trsm.cpp @@ -138,7 +138,7 @@ int test(device* dev, oneapi::mkl::layout layout, oneapi::mkl::side left_right, } // Compare the results of reference implementation and DPC++ implementation. - auto B_accessor = B_buffer.template get_host_access(read_only); + auto B_accessor = B_buffer.get_host_access(read_only); bool good = check_equal_trsm_matrix(B_accessor, B_ref, layout, m, n, ldb, 10 * std::max(m, n), std::cout); diff --git a/tests/unit_tests/dft/include/compute_inplace.hpp b/tests/unit_tests/dft/include/compute_inplace.hpp index ef4669e34..145106b7a 100644 --- a/tests/unit_tests/dft/include/compute_inplace.hpp +++ b/tests/unit_tests/dft/include/compute_inplace.hpp @@ -91,7 +91,7 @@ int DFT_Test::test_in_place_buffer() { oneapi::mkl::dft::compute_forward(descriptor, inout_buf); { - auto acc_host = inout_buf.template get_host_access(); + auto acc_host = inout_buf.get_host_access(); auto ptr_host = reinterpret_cast(acc_host.get_pointer()); for (std::int64_t i = 0; i < batches; i++) { EXPECT_TRUE(check_equal_strided( diff --git a/tests/unit_tests/dft/include/compute_inplace_real_real.hpp b/tests/unit_tests/dft/include/compute_inplace_real_real.hpp index dbc2100bf..d4af1a44a 100644 --- a/tests/unit_tests/dft/include/compute_inplace_real_real.hpp +++ b/tests/unit_tests/dft/include/compute_inplace_real_real.hpp @@ -121,8 +121,8 @@ int DFT_Test::test_in_place_real_real_buffer() { inout_im_buf); { - auto acc_inout_re = inout_re_buf.template get_host_access(); - auto acc_inout_im = inout_im_buf.template get_host_access(); + auto acc_inout_re = inout_re_buf.get_host_access(); + auto acc_inout_im = inout_im_buf.get_host_access(); std::vector output_data(size_total, static_cast(0)); for (std::size_t i = 0; i < output_data.size(); ++i) { output_data[i] = { acc_inout_re[i], acc_inout_im[i] }; @@ -136,8 +136,8 @@ int DFT_Test::test_in_place_real_real_buffer() { PrecisionType>(descriptor, inout_re_buf, inout_im_buf); { - auto acc_inout_re = inout_re_buf.template get_host_access(); - auto acc_inout_im = inout_im_buf.template get_host_access(); + auto acc_inout_re = inout_re_buf.get_host_access(); + auto acc_inout_im = inout_im_buf.get_host_access(); std::vector output_data(size_total, static_cast(0)); for (std::size_t i = 0; i < output_data.size(); ++i) { output_data[i] = { acc_inout_re[i], acc_inout_im[i] }; diff --git a/tests/unit_tests/dft/include/compute_out_of_place.hpp b/tests/unit_tests/dft/include/compute_out_of_place.hpp index 4ab42874f..069c2ce9e 100644 --- a/tests/unit_tests/dft/include/compute_out_of_place.hpp +++ b/tests/unit_tests/dft/include/compute_out_of_place.hpp @@ -70,7 +70,7 @@ int DFT_Test::test_out_of_place_buffer() { descriptor, fwd_buf, bwd_buf); { - auto acc_bwd = bwd_buf.template get_host_access(); + auto acc_bwd = bwd_buf.get_host_access(); auto bwd_ptr = acc_bwd.get_pointer(); for (std::int64_t i = 0; i < batches; i++) { EXPECT_TRUE(check_equal_strided( diff --git a/tests/unit_tests/dft/include/compute_out_of_place_real_real.hpp b/tests/unit_tests/dft/include/compute_out_of_place_real_real.hpp index a7b59edf6..fb3ecb4f2 100644 --- a/tests/unit_tests/dft/include/compute_out_of_place_real_real.hpp +++ b/tests/unit_tests/dft/include/compute_out_of_place_real_real.hpp @@ -127,8 +127,8 @@ int DFT_Test::test_out_of_place_real_real_buffer() { descriptor, in_dev_re, in_dev_im, out_dev_re, out_dev_im); { - auto acc_out_re = out_dev_re.template get_host_access(); - auto acc_out_im = out_dev_im.template get_host_access(); + auto acc_out_re = out_dev_re.get_host_access(); + auto acc_out_im = out_dev_im.get_host_access(); std::vector output_data(size_total, static_cast(0)); for (std::size_t i = 0; i < output_data.size(); ++i) { output_data[i] = { acc_out_re[i], acc_out_im[i] }; @@ -143,8 +143,8 @@ int DFT_Test::test_out_of_place_real_real_buffer() { descriptor, out_dev_re, out_dev_im, out_back_dev_re, out_back_dev_im); { - auto acc_back_out_re = out_back_dev_re.template get_host_access(); - auto acc_back_out_im = out_back_dev_im.template get_host_access(); + auto acc_back_out_re = out_back_dev_re.get_host_access(); + auto acc_back_out_im = out_back_dev_im.get_host_access(); std::vector output_data(size_total, static_cast(0)); for (std::size_t i = 0; i < output_data.size(); ++i) { output_data[i] = { acc_back_out_re[i], acc_back_out_im[i] }; diff --git a/tests/unit_tests/sparse_blas/source/sparse_gemm_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_gemm_buffer.cpp index cc6fae6db..1c9549fcc 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_gemm_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_gemm_buffer.cpp @@ -120,7 +120,7 @@ int test(sycl::device *dev, intType nrows_A, intType ncols_A, intType ncols_C, c_ref_host.data()); // Compare the results of reference implementation and DPC++ implementation. - auto c_acc = c_buf.template get_host_access(sycl::read_only); + auto c_acc = c_buf.get_host_access(sycl::read_only); bool valid = check_equal_vector(c_acc, c_ref_host); ev_release.wait_and_throw(); diff --git a/tests/unit_tests/sparse_blas/source/sparse_gemv_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_gemv_buffer.cpp index b6506ed5f..b95636831 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_gemv_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_gemv_buffer.cpp @@ -111,7 +111,7 @@ int test(sycl::device *dev, intType nrows, intType ncols, double density_A_matri y_ref_host.data()); // Compare the results of reference implementation and DPC++ implementation. - auto y_acc = y_buf.template get_host_access(sycl::read_only); + auto y_acc = y_buf.get_host_access(sycl::read_only); bool valid = check_equal_vector(y_acc, y_ref_host); ev_release.wait_and_throw(); diff --git a/tests/unit_tests/sparse_blas/source/sparse_trsv_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_trsv_buffer.cpp index 00ec6e5ed..4e82ae1f0 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_trsv_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_trsv_buffer.cpp @@ -119,7 +119,7 @@ int test(sycl::device *dev, intType m, double density_A_matrix, oneapi::mkl::ind y_ref_host.data()); // Compare the results of reference implementation and DPC++ implementation. - auto y_acc = y_buf.template get_host_access(sycl::read_only); + auto y_acc = y_buf.get_host_access(sycl::read_only); bool valid = check_equal_vector(y_acc, y_ref_host); ev_release.wait_and_throw(); From 2a59f228356de672ca30bc8ef00659d4459050d8 Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Wed, 17 Jul 2024 12:08:35 +0200 Subject: [PATCH 13/15] clang-format --- .../run_time_dispatching/sparse_blas_spmv_usm.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/examples/sparse_blas/run_time_dispatching/sparse_blas_spmv_usm.cpp b/examples/sparse_blas/run_time_dispatching/sparse_blas_spmv_usm.cpp index 4a8a4280b..f93569a92 100644 --- a/examples/sparse_blas/run_time_dispatching/sparse_blas_spmv_usm.cpp +++ b/examples/sparse_blas/run_time_dispatching/sparse_blas_spmv_usm.cpp @@ -148,8 +148,10 @@ int run_sparse_matrix_vector_multiply_example(const sycl::device &dev) { // Create and initialize dense vector handles oneapi::mkl::sparse::dense_vector_handle_t x_handle = nullptr; oneapi::mkl::sparse::dense_vector_handle_t y_handle = nullptr; - oneapi::mkl::sparse::init_dense_vector(main_queue, &x_handle, static_cast(sizevec), x); - oneapi::mkl::sparse::init_dense_vector(main_queue, &y_handle, static_cast(sizevec), y); + oneapi::mkl::sparse::init_dense_vector(main_queue, &x_handle, + static_cast(sizevec), x); + oneapi::mkl::sparse::init_dense_vector(main_queue, &y_handle, + static_cast(sizevec), y); // Create operation descriptor oneapi::mkl::sparse::spmv_descr_t descr = nullptr; From 43f4669048634c40cef3dcd8c7e0c0615a71e8c0 Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Thu, 18 Jul 2024 19:58:14 +0200 Subject: [PATCH 14/15] Take string as reference --- src/sparse_blas/backends/mkl_common/mkl_spmm.cxx | 2 +- src/sparse_blas/backends/mkl_common/mkl_spmv.cxx | 2 +- src/sparse_blas/backends/mkl_common/mkl_spsv.cxx | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx index aa292921e..604db11a7 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx @@ -29,7 +29,7 @@ sycl::event release_spmm_descr(sycl::queue &queue, oneapi::mkl::sparse::spmm_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spmm(const std::string function_name, sycl::queue &queue, +void check_valid_spmm(const std::string &function_name, sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_matrix_handle_t B_handle, diff --git a/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx b/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx index 6a954feba..b35ad0847 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx @@ -29,7 +29,7 @@ sycl::event release_spmv_descr(sycl::queue &queue, oneapi::mkl::sparse::spmv_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spmv(const std::string function_name, sycl::queue &queue, +void check_valid_spmv(const std::string &function_name, sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, diff --git a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx index ea8410a88..4ca4ee9d8 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx @@ -29,7 +29,7 @@ sycl::event release_spsv_descr(sycl::queue &queue, oneapi::mkl::sparse::spsv_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spsv(const std::string function_name, sycl::queue &queue, +void check_valid_spsv(const std::string &function_name, sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, From 2f59edce37db07bd9eb897892572b136dded53cf Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Fri, 19 Jul 2024 11:20:46 +0200 Subject: [PATCH 15/15] Reduce number of calls to get_pointer_type --- .../backends/mkl_common/mkl_helper.hpp | 13 ++-- .../backends/mkl_common/mkl_spmm.cxx | 60 +++++++++++-------- .../backends/mkl_common/mkl_spmv.cxx | 47 ++++++++++----- .../backends/mkl_common/mkl_spsv.cxx | 35 +++++++---- 4 files changed, 97 insertions(+), 58 deletions(-) diff --git a/src/sparse_blas/backends/mkl_common/mkl_helper.hpp b/src/sparse_blas/backends/mkl_common/mkl_helper.hpp index d1303d949..c76af5cb6 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_helper.hpp +++ b/src/sparse_blas/backends/mkl_common/mkl_helper.hpp @@ -43,10 +43,10 @@ inline bool is_ptr_accessible_on_host(sycl::queue &queue, const T *host_or_devic } /// Throw an exception if the scalar is not accessible in the host -template -void check_ptr_is_host_accessible(const std::string &function_name, const std::string &scalar_name, - sycl::queue &queue, const T *host_or_device_ptr) { - if (!is_ptr_accessible_on_host(queue, host_or_device_ptr)) { +inline void check_ptr_is_host_accessible(const std::string &function_name, + const std::string &scalar_name, + bool is_ptr_accessible_on_host) { + if (!is_ptr_accessible_on_host) { throw mkl::invalid_argument( "sparse_blas", function_name, "Scalar " + scalar_name + " must be accessible on the host for buffer functions."); @@ -56,8 +56,9 @@ void check_ptr_is_host_accessible(const std::string &function_name, const std::s /// Return a scalar on the host from a pointer to host or device memory /// Used for USM functions template -inline T get_scalar_on_host(sycl::queue &queue, const T *host_or_device_ptr) { - if (is_ptr_accessible_on_host(queue, host_or_device_ptr)) { +inline T get_scalar_on_host(sycl::queue &queue, const T *host_or_device_ptr, + bool is_ptr_accessible_on_host) { + if (is_ptr_accessible_on_host) { return *host_or_device_ptr; } T scalar; diff --git a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx index 604db11a7..3c2a9f161 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spmm.cxx @@ -29,12 +29,12 @@ sycl::event release_spmm_descr(sycl::queue &queue, oneapi::mkl::sparse::spmm_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spmm(const std::string &function_name, sycl::queue &queue, - oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, +void check_valid_spmm(const std::string &function_name, oneapi::mkl::transpose opA, + oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_matrix_handle_t B_handle, - oneapi::mkl::sparse::dense_matrix_handle_t C_handle, const void *alpha, - const void *beta) { + oneapi::mkl::sparse::dense_matrix_handle_t C_handle, + bool is_alpha_host_accessible, bool is_beta_host_accessible) { THROW_IF_NULLPTR(function_name, A_handle); THROW_IF_NULLPTR(function_name, B_handle); THROW_IF_NULLPTR(function_name, C_handle); @@ -42,11 +42,10 @@ void check_valid_spmm(const std::string &function_name, sycl::queue &queue, auto internal_A_handle = detail::get_internal_handle(A_handle); detail::check_all_containers_compatible(function_name, internal_A_handle, B_handle, C_handle); if (internal_A_handle->all_use_buffer()) { - detail::check_ptr_is_host_accessible("spmm", "alpha", queue, alpha); - detail::check_ptr_is_host_accessible("spmm", "beta", queue, beta); + detail::check_ptr_is_host_accessible("spmm", "alpha", is_alpha_host_accessible); + detail::check_ptr_is_host_accessible("spmm", "beta", is_beta_host_accessible); } - if (detail::is_ptr_accessible_on_host(queue, alpha) != - detail::is_ptr_accessible_on_host(queue, beta)) { + if (is_alpha_host_accessible != is_beta_host_accessible) { throw mkl::invalid_argument( "sparse_blas", function_name, "Alpha and beta must both be placed on host memory or device memory."); @@ -91,7 +90,10 @@ void spmm_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::sparse::spmm_descr_t /*spmm_descr*/, std::size_t &temp_buffer_size) { // TODO: Add support for external workspace once the close-source oneMKL backend supports it. - check_valid_spmm(__func__, queue, opA, A_view, A_handle, B_handle, C_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmm(__func__, opA, A_view, A_handle, B_handle, C_handle, is_alpha_host_accessible, + is_beta_host_accessible); temp_buffer_size = 0; } @@ -103,7 +105,10 @@ void spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl:: oneapi::mkl::sparse::spmm_alg alg, oneapi::mkl::sparse::spmm_descr_t /*spmm_descr*/, sycl::buffer /*workspace*/) { - check_valid_spmm(__func__, queue, opA, A_view, A_handle, B_handle, C_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmm(__func__, opA, A_view, A_handle, B_handle, C_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); if (!internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -124,7 +129,10 @@ sycl::event spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::sparse::spmm_alg alg, oneapi::mkl::sparse::spmm_descr_t /*spmm_descr*/, void * /*workspace*/, const std::vector &dependencies) { - check_valid_spmm(__func__, queue, opA, A_view, A_handle, B_handle, C_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmm(__func__, opA, A_view, A_handle, B_handle, C_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); if (internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -138,17 +146,17 @@ sycl::event spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, } template -sycl::event internal_spmm(sycl::queue &queue, oneapi::mkl::transpose opA, - oneapi::mkl::transpose opB, const void *alpha, - oneapi::mkl::sparse::matrix_view /*A_view*/, - oneapi::mkl::sparse::matrix_handle_t A_handle, - oneapi::mkl::sparse::dense_matrix_handle_t B_handle, const void *beta, - oneapi::mkl::sparse::dense_matrix_handle_t C_handle, - oneapi::mkl::sparse::spmm_alg /*alg*/, - oneapi::mkl::sparse::spmm_descr_t /*spmm_descr*/, - const std::vector &dependencies) { - T host_alpha = detail::get_scalar_on_host(queue, static_cast(alpha)); - T host_beta = detail::get_scalar_on_host(queue, static_cast(beta)); +sycl::event internal_spmm( + sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::transpose opB, const void *alpha, + oneapi::mkl::sparse::matrix_view /*A_view*/, oneapi::mkl::sparse::matrix_handle_t A_handle, + oneapi::mkl::sparse::dense_matrix_handle_t B_handle, const void *beta, + oneapi::mkl::sparse::dense_matrix_handle_t C_handle, oneapi::mkl::sparse::spmm_alg /*alg*/, + oneapi::mkl::sparse::spmm_descr_t /*spmm_descr*/, const std::vector &dependencies, + bool is_alpha_host_accessible, bool is_beta_host_accessible) { + T host_alpha = + detail::get_scalar_on_host(queue, static_cast(alpha), is_alpha_host_accessible); + T host_beta = + detail::get_scalar_on_host(queue, static_cast(beta), is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); internal_A_handle->can_be_reset = false; auto layout = B_handle->dense_layout; @@ -177,8 +185,12 @@ sycl::event spmm(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::tr oneapi::mkl::sparse::dense_matrix_handle_t C_handle, oneapi::mkl::sparse::spmm_alg alg, oneapi::mkl::sparse::spmm_descr_t spmm_descr, const std::vector &dependencies) { - check_valid_spmm(__func__, queue, opA, A_view, A_handle, B_handle, C_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmm(__func__, opA, A_view, A_handle, B_handle, C_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto value_type = detail::get_internal_handle(A_handle)->get_value_type(); DISPATCH_MKL_OPERATION("spmm", value_type, internal_spmm, queue, opA, opB, alpha, A_view, - A_handle, B_handle, beta, C_handle, alg, spmm_descr, dependencies); + A_handle, B_handle, beta, C_handle, alg, spmm_descr, dependencies, + is_alpha_host_accessible, is_beta_host_accessible); } diff --git a/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx b/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx index b35ad0847..930e1ec87 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spmv.cxx @@ -29,12 +29,12 @@ sycl::event release_spmv_descr(sycl::queue &queue, oneapi::mkl::sparse::spmv_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spmv(const std::string &function_name, sycl::queue &queue, - oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, +void check_valid_spmv(const std::string &function_name, oneapi::mkl::transpose opA, + oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, - oneapi::mkl::sparse::dense_vector_handle_t y_handle, const void *alpha, - const void *beta) { + oneapi::mkl::sparse::dense_vector_handle_t y_handle, + bool is_alpha_host_accessible, bool is_beta_host_accessible) { THROW_IF_NULLPTR(function_name, A_handle); THROW_IF_NULLPTR(function_name, x_handle); THROW_IF_NULLPTR(function_name, y_handle); @@ -42,11 +42,10 @@ void check_valid_spmv(const std::string &function_name, sycl::queue &queue, auto internal_A_handle = detail::get_internal_handle(A_handle); detail::check_all_containers_compatible(function_name, internal_A_handle, x_handle, y_handle); if (internal_A_handle->all_use_buffer()) { - detail::check_ptr_is_host_accessible("spmv", "alpha", queue, alpha); - detail::check_ptr_is_host_accessible("spmv", "beta", queue, beta); + detail::check_ptr_is_host_accessible("spmv", "alpha", is_alpha_host_accessible); + detail::check_ptr_is_host_accessible("spmv", "beta", is_beta_host_accessible); } - if (detail::is_ptr_accessible_on_host(queue, alpha) != - detail::is_ptr_accessible_on_host(queue, beta)) { + if (is_alpha_host_accessible != is_beta_host_accessible) { throw mkl::invalid_argument( "sparse_blas", function_name, "Alpha and beta must both be placed on host memory or device memory."); @@ -81,7 +80,10 @@ void spmv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void oneapi::mkl::sparse::spmv_descr_t /*spmv_descr*/, std::size_t &temp_buffer_size) { // TODO: Add support for external workspace once the close-source oneMKL backend supports it. - check_valid_spmv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + is_beta_host_accessible); temp_buffer_size = 0; } @@ -93,7 +95,10 @@ void spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a oneapi::mkl::sparse::spmv_alg alg, oneapi::mkl::sparse::spmv_descr_t /*spmv_descr*/, sycl::buffer /*workspace*/) { - check_valid_spmv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); if (!internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -127,7 +132,10 @@ sycl::event spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const oneapi::mkl::sparse::spmv_alg alg, oneapi::mkl::sparse::spmv_descr_t /*spmv_descr*/, void * /*workspace*/, const std::vector &dependencies) { - check_valid_spmv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); if (internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -158,9 +166,12 @@ sycl::event internal_spmv(sycl::queue &queue, oneapi::mkl::transpose opA, const oneapi::mkl::sparse::dense_vector_handle_t y_handle, oneapi::mkl::sparse::spmv_alg /*alg*/, oneapi::mkl::sparse::spmv_descr_t /*spmv_descr*/, - const std::vector &dependencies) { - T host_alpha = detail::get_scalar_on_host(queue, static_cast(alpha)); - T host_beta = detail::get_scalar_on_host(queue, static_cast(beta)); + const std::vector &dependencies, + bool is_alpha_host_accessible, bool is_beta_host_accessible) { + T host_alpha = + detail::get_scalar_on_host(queue, static_cast(alpha), is_alpha_host_accessible); + T host_beta = + detail::get_scalar_on_host(queue, static_cast(beta), is_beta_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); internal_A_handle->can_be_reset = false; auto backend_handle = internal_A_handle->backend_handle; @@ -210,8 +221,12 @@ sycl::event spmv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp oneapi::mkl::sparse::dense_vector_handle_t y_handle, oneapi::mkl::sparse::spmv_alg alg, oneapi::mkl::sparse::spmv_descr_t spmv_descr, const std::vector &dependencies) { - check_valid_spmv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, beta); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + bool is_beta_host_accessible = detail::is_ptr_accessible_on_host(queue, beta); + check_valid_spmv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + is_beta_host_accessible); auto value_type = detail::get_internal_handle(A_handle)->get_value_type(); DISPATCH_MKL_OPERATION("spmv", value_type, internal_spmv, queue, opA, alpha, A_view, A_handle, - x_handle, beta, y_handle, alg, spmv_descr, dependencies); + x_handle, beta, y_handle, alg, spmv_descr, dependencies, + is_alpha_host_accessible, is_beta_host_accessible); } diff --git a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx index 4ca4ee9d8..849919f12 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_spsv.cxx @@ -29,12 +29,12 @@ sycl::event release_spsv_descr(sycl::queue &queue, oneapi::mkl::sparse::spsv_des return detail::collapse_dependencies(queue, dependencies); } -void check_valid_spsv(const std::string &function_name, sycl::queue &queue, - oneapi::mkl::transpose opA, oneapi::mkl::sparse::matrix_view A_view, +void check_valid_spsv(const std::string &function_name, oneapi::mkl::transpose opA, + oneapi::mkl::sparse::matrix_view A_view, oneapi::mkl::sparse::matrix_handle_t A_handle, oneapi::mkl::sparse::dense_vector_handle_t x_handle, - oneapi::mkl::sparse::dense_vector_handle_t y_handle, const void *alpha, - oneapi::mkl::sparse::spsv_alg alg) { + oneapi::mkl::sparse::dense_vector_handle_t y_handle, + bool is_alpha_host_accessible, oneapi::mkl::sparse::spsv_alg alg) { THROW_IF_NULLPTR(function_name, A_handle); THROW_IF_NULLPTR(function_name, x_handle); THROW_IF_NULLPTR(function_name, y_handle); @@ -67,7 +67,7 @@ void check_valid_spsv(const std::string &function_name, sycl::queue &queue, } if (internal_A_handle->all_use_buffer()) { - detail::check_ptr_is_host_accessible("spsv", "alpha", queue, alpha); + detail::check_ptr_is_host_accessible("spsv", "alpha", is_alpha_host_accessible); } } @@ -80,7 +80,9 @@ void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, std::size_t &temp_buffer_size) { // TODO: Add support for external workspace once the close-source oneMKL backend supports it. - check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + check_valid_spsv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + alg); temp_buffer_size = 0; } @@ -92,7 +94,9 @@ void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, sycl::buffer /*workspace*/) { - check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + check_valid_spsv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + alg); auto internal_A_handle = detail::get_internal_handle(A_handle); if (!internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -115,7 +119,9 @@ sycl::event spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, void * /*workspace*/, const std::vector &dependencies) { - check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + check_valid_spsv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + alg); auto internal_A_handle = detail::get_internal_handle(A_handle); if (internal_A_handle->all_use_buffer()) { detail::throw_incompatible_container(__func__); @@ -136,8 +142,10 @@ sycl::event internal_spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const oneapi::mkl::sparse::dense_vector_handle_t y_handle, oneapi::mkl::sparse::spsv_alg /*alg*/, oneapi::mkl::sparse::spsv_descr_t /*spsv_descr*/, - const std::vector &dependencies) { - T host_alpha = detail::get_scalar_on_host(queue, static_cast(alpha)); + const std::vector &dependencies, + bool is_alpha_host_accessible) { + T host_alpha = + detail::get_scalar_on_host(queue, static_cast(alpha), is_alpha_host_accessible); auto internal_A_handle = detail::get_internal_handle(A_handle); internal_A_handle->can_be_reset = false; if (internal_A_handle->all_use_buffer()) { @@ -162,8 +170,11 @@ sycl::event spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp oneapi::mkl::sparse::dense_vector_handle_t y_handle, oneapi::mkl::sparse::spsv_alg alg, oneapi::mkl::sparse::spsv_descr_t spsv_descr, const std::vector &dependencies) { - check_valid_spsv(__func__, queue, opA, A_view, A_handle, x_handle, y_handle, alpha, alg); + bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); + check_valid_spsv(__func__, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, + alg); auto value_type = detail::get_internal_handle(A_handle)->get_value_type(); DISPATCH_MKL_OPERATION("spsv", value_type, internal_spsv, queue, opA, alpha, A_view, A_handle, - x_handle, y_handle, alg, spsv_descr, dependencies); + x_handle, y_handle, alg, spsv_descr, dependencies, + is_alpha_host_accessible); }