diff --git a/docs/domains/sparse_linear_algebra.rst b/docs/domains/sparse_linear_algebra.rst index 11c59407b..80b9b4419 100644 --- a/docs/domains/sparse_linear_algebra.rst +++ b/docs/domains/sparse_linear_algebra.rst @@ -45,14 +45,17 @@ cuSPARSE backend Currently known limitations: +- The COO format requires the indices to be sorted by row. See the `cuSPARSE + documentation + `_. Sparse + operations using matrices with the COO format without the property + ``matrix_property::sorted_by_rows`` or ``matrix_property::sorted`` will throw + an ``oneapi::mkl::unimplemented`` exception. - Using ``spmm`` with the algorithm ``spmm_alg::csr_alg3`` and an ``opA`` other than ``transpose::nontrans`` or an ``opB`` ``transpose::conjtrans`` will throw an ``oneapi::mkl::unimplemented`` exception. - Using ``spmv`` with a ``type_view`` other than ``matrix_descr::general`` will throw an ``oneapi::mkl::unimplemented`` exception. -- The COO format requires the indices to be sorted by row. See the `cuSPARSE - documentation - `_. Operation algorithms mapping diff --git a/include/oneapi/mkl/sparse_blas/types.hpp b/include/oneapi/mkl/sparse_blas/types.hpp index d4aea3e88..39256f3c8 100644 --- a/include/oneapi/mkl/sparse_blas/types.hpp +++ b/include/oneapi/mkl/sparse_blas/types.hpp @@ -36,6 +36,7 @@ namespace sparse { enum class matrix_property { symmetric, sorted, + sorted_by_rows, }; enum class spmm_alg { diff --git a/src/sparse_blas/backends/cusparse/cusparse_handles.cpp b/src/sparse_blas/backends/cusparse/cusparse_handles.cpp index 1909c8d3c..09bade903 100644 --- a/src/sparse_blas/backends/cusparse/cusparse_handles.cpp +++ b/src/sparse_blas/backends/cusparse/cusparse_handles.cpp @@ -263,8 +263,8 @@ void init_coo_matrix(sycl::queue &queue, matrix_handle_t *p_smhandle, std::int64 CUSPARSE_ERR_FUNC(cusparseCreateCoo, &cu_smhandle, num_rows, num_cols, nnz, sc.get_mem(row_acc), sc.get_mem(col_acc), sc.get_mem(val_acc), cuda_index_type, cuda_index_base, cuda_value_type); - *p_smhandle = new matrix_handle(cu_smhandle, row_ind, col_ind, val, num_rows, num_cols, - nnz, index); + *p_smhandle = new matrix_handle(cu_smhandle, row_ind, col_ind, val, detail::sparse_format::COO, + num_rows, num_cols, nnz, index); }); }); event.wait_and_throw(); @@ -284,8 +284,8 @@ void init_coo_matrix(sycl::queue &queue, matrix_handle_t *p_smhandle, std::int64 cusparseSpMatDescr_t cu_smhandle; CUSPARSE_ERR_FUNC(cusparseCreateCoo, &cu_smhandle, num_rows, num_cols, nnz, row_ind, col_ind, val, cuda_index_type, cuda_index_base, cuda_value_type); - *p_smhandle = new matrix_handle(cu_smhandle, row_ind, col_ind, val, num_rows, num_cols, - nnz, index); + *p_smhandle = new matrix_handle(cu_smhandle, row_ind, col_ind, val, detail::sparse_format::COO, + num_rows, num_cols, nnz, index); }); }); event.wait_and_throw(); @@ -388,8 +388,8 @@ void init_csr_matrix(sycl::queue &queue, matrix_handle_t *p_smhandle, std::int64 CUSPARSE_ERR_FUNC(cusparseCreateCsr, &cu_smhandle, num_rows, num_cols, nnz, sc.get_mem(row_acc), sc.get_mem(col_acc), sc.get_mem(val_acc), cuda_index_type, cuda_index_type, cuda_index_base, cuda_value_type); - *p_smhandle = new matrix_handle(cu_smhandle, row_ptr, col_ind, val, num_rows, num_cols, - nnz, index); + *p_smhandle = new matrix_handle(cu_smhandle, row_ptr, col_ind, val, detail::sparse_format::CSR, + num_rows, num_cols, nnz, index); }); }); event.wait_and_throw(); @@ -410,8 +410,8 @@ void init_csr_matrix(sycl::queue &queue, matrix_handle_t *p_smhandle, std::int64 CUSPARSE_ERR_FUNC(cusparseCreateCsr, &cu_smhandle, num_rows, num_cols, nnz, row_ptr, col_ind, val, cuda_index_type, cuda_index_type, cuda_index_base, cuda_value_type); - *p_smhandle = new matrix_handle(cu_smhandle, row_ptr, col_ind, val, num_rows, num_cols, - nnz, index); + *p_smhandle = new matrix_handle(cu_smhandle, row_ptr, col_ind, val, detail::sparse_format::CSR, + num_rows, num_cols, nnz, index); }); }); event.wait_and_throw(); diff --git a/src/sparse_blas/backends/cusparse/cusparse_handles.hpp b/src/sparse_blas/backends/cusparse/cusparse_handles.hpp index ac22d33ae..2653d84c1 100644 --- a/src/sparse_blas/backends/cusparse/cusparse_handles.hpp +++ b/src/sparse_blas/backends/cusparse/cusparse_handles.hpp @@ -59,20 +59,33 @@ struct dense_matrix_handle : public detail::generic_dense_matrix_handle { template matrix_handle(cusparseSpMatDescr_t cu_descr, intType* row_ptr, intType* col_ptr, - fpType* value_ptr, std::int64_t num_rows, std::int64_t num_cols, std::int64_t nnz, - oneapi::mkl::index_base index) + fpType* value_ptr, detail::sparse_format format, std::int64_t num_rows, + std::int64_t num_cols, std::int64_t nnz, oneapi::mkl::index_base index) : detail::generic_sparse_handle( - cu_descr, row_ptr, col_ptr, value_ptr, num_rows, num_cols, nnz, index) {} + cu_descr, row_ptr, col_ptr, value_ptr, format, num_rows, num_cols, nnz, index) {} template matrix_handle(cusparseSpMatDescr_t cu_descr, const sycl::buffer row_buffer, const sycl::buffer col_buffer, - const sycl::buffer value_buffer, std::int64_t num_rows, - std::int64_t num_cols, std::int64_t nnz, oneapi::mkl::index_base index) - : detail::generic_sparse_handle( - cu_descr, row_buffer, col_buffer, value_buffer, num_rows, num_cols, nnz, index) {} + const sycl::buffer value_buffer, detail::sparse_format format, + std::int64_t num_rows, std::int64_t num_cols, std::int64_t nnz, + oneapi::mkl::index_base index) + : detail::generic_sparse_handle(cu_descr, row_buffer, col_buffer, + value_buffer, format, num_rows, + num_cols, nnz, index) {} }; +inline void check_valid_matrix_properties(const std::string& function_name, + matrix_handle_t sm_handle) { + if (sm_handle->format == detail::sparse_format::COO && + !(sm_handle->has_matrix_property(matrix_property::sorted_by_rows) || + sm_handle->has_matrix_property(matrix_property::sorted))) { + throw mkl::unimplemented( + "sparse_blas", function_name, + "The backend does not support unsorted COO format. Use `set_matrix_property` to set the property `matrix_property::sorted_by_rows` or `matrix_property::sorted`"); + } +} + } // namespace oneapi::mkl::sparse #endif // _ONEMKL_SRC_SPARSE_BLAS_BACKENDS_CUSPARSE_HANDLES_HPP_ diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp index 8c8db04eb..116bbffe9 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp @@ -77,6 +77,7 @@ void check_valid_spmm(const std::string& function_name, oneapi::mkl::transpose o bool is_alpha_host_accessible, bool is_beta_host_accessible, spmm_alg alg) { detail::check_valid_spmm_common(function_name, A_view, A_handle, B_handle, C_handle, is_alpha_host_accessible, is_beta_host_accessible); + check_valid_matrix_properties(function_name, A_handle); if (alg == spmm_alg::csr_alg3 && opA != oneapi::mkl::transpose::nontrans) { throw mkl::unimplemented( "sparse_blas", function_name, diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp index f75464e91..03cdd15e0 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp @@ -73,6 +73,7 @@ void check_valid_spmv(const std::string &function_name, oneapi::mkl::transpose o bool is_beta_host_accessible) { detail::check_valid_spmv_common(function_name, opA, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible, is_beta_host_accessible); + check_valid_matrix_properties(function_name, A_handle); if (A_view.type_view != matrix_descr::general) { throw mkl::unimplemented( "sparse_blas", function_name, diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp index 5eedeca70..c06335100 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp @@ -69,13 +69,20 @@ inline auto get_cuda_spsv_alg(spsv_alg /*alg*/) { return CUSPARSE_SPSV_ALG_DEFAULT; } +void check_valid_spsv(const std::string &function_name, matrix_view A_view, + matrix_handle_t A_handle, dense_vector_handle_t x_handle, + dense_vector_handle_t y_handle, bool is_alpha_host_accessible) { + detail::check_valid_spsv_common(function_name, A_view, A_handle, x_handle, y_handle, + is_alpha_host_accessible); + check_valid_matrix_properties(function_name, A_handle); +} + void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, matrix_view A_view, matrix_handle_t A_handle, dense_vector_handle_t x_handle, dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr, std::size_t &temp_buffer_size) { bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); - detail::check_valid_spsv_common(__func__, A_view, A_handle, x_handle, y_handle, - is_alpha_host_accessible); + check_valid_spsv(__func__, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible); auto functor = [=, &temp_buffer_size](CusparseScopedContextHandler &sc) { auto cu_handle = sc.get_handle(queue); auto cu_a = A_handle->backend_handle; @@ -101,8 +108,8 @@ inline void common_spsv_optimize(oneapi::mkl::transpose opA, bool is_alpha_host_ matrix_view A_view, matrix_handle_t A_handle, dense_vector_handle_t x_handle, dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr) { - detail::check_valid_spsv_common("spsv_optimize", A_view, A_handle, x_handle, y_handle, - is_alpha_host_accessible); + check_valid_spsv("spsv_optimize", A_view, A_handle, x_handle, y_handle, + is_alpha_host_accessible); if (!spsv_descr->buffer_size_called) { throw mkl::uninitialized("sparse_blas", "spsv_optimize", "spsv_buffer_size must be called before spsv_optimize."); @@ -202,8 +209,7 @@ sycl::event spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr, const std::vector &dependencies) { bool is_alpha_host_accessible = detail::is_ptr_accessible_on_host(queue, alpha); - detail::check_valid_spsv_common(__func__, A_view, A_handle, x_handle, y_handle, - is_alpha_host_accessible); + check_valid_spsv(__func__, A_view, A_handle, x_handle, y_handle, is_alpha_host_accessible); if (A_handle->all_use_buffer() != spsv_descr->workspace.use_buffer()) { detail::throw_incompatible_container(__func__); } diff --git a/src/sparse_blas/backends/mkl_common/mkl_handles.cxx b/src/sparse_blas/backends/mkl_common/mkl_handles.cxx index 8d1a923b6..6ddbd43ef 100644 --- a/src/sparse_blas/backends/mkl_common/mkl_handles.cxx +++ b/src/sparse_blas/backends/mkl_common/mkl_handles.cxx @@ -109,8 +109,8 @@ void init_coo_matrix(sycl::queue &queue, oneapi::mkl::sparse::matrix_handle_t *p sycl::buffer col_ind, sycl::buffer val) { oneapi::mkl::sparse::matrix_handle_t mkl_handle; oneapi::mkl::sparse::init_matrix_handle(&mkl_handle); - auto internal_smhandle = new detail::sparse_matrix_handle(mkl_handle, row_ind, col_ind, val, - num_rows, num_cols, nnz, index); + auto internal_smhandle = new detail::sparse_matrix_handle( + mkl_handle, row_ind, col_ind, val, detail::sparse_format::COO, num_rows, num_cols, nnz, index); // The backend handle must use the buffers from the internal handle as they will be kept alive until the handle is released. oneapi::mkl::sparse::set_coo_data(queue, mkl_handle, static_cast(num_rows), static_cast(num_cols), static_cast(nnz), @@ -127,8 +127,8 @@ void init_coo_matrix(sycl::queue &queue, oneapi::mkl::sparse::matrix_handle_t *p fpType *val) { oneapi::mkl::sparse::matrix_handle_t mkl_handle; oneapi::mkl::sparse::init_matrix_handle(&mkl_handle); - auto internal_smhandle = new detail::sparse_matrix_handle(mkl_handle, row_ind, col_ind, val, - num_rows, num_cols, nnz, index); + auto internal_smhandle = new detail::sparse_matrix_handle( + mkl_handle, row_ind, col_ind, val, detail::sparse_format::COO, num_rows, num_cols, nnz, index); auto event = oneapi::mkl::sparse::set_coo_data( queue, mkl_handle, static_cast(num_rows), static_cast(num_cols), static_cast(nnz), index, row_ind, col_ind, val); @@ -189,8 +189,8 @@ void init_csr_matrix(sycl::queue &queue, oneapi::mkl::sparse::matrix_handle_t *p sycl::buffer col_ind, sycl::buffer val) { oneapi::mkl::sparse::matrix_handle_t mkl_handle; oneapi::mkl::sparse::init_matrix_handle(&mkl_handle); - auto internal_smhandle = new detail::sparse_matrix_handle(mkl_handle, row_ptr, col_ind, val, - num_rows, num_cols, nnz, index); + auto internal_smhandle = new detail::sparse_matrix_handle( + mkl_handle, row_ptr, col_ind, val, detail::sparse_format::CSR, num_rows, num_cols, nnz, index); // The backend deduces nnz from row_ptr. // The backend handle must use the buffers from the internal handle as they will be kept alive until the handle is released. oneapi::mkl::sparse::set_csr_data(queue, mkl_handle, static_cast(num_rows), @@ -208,8 +208,8 @@ void init_csr_matrix(sycl::queue &queue, oneapi::mkl::sparse::matrix_handle_t *p fpType *val) { oneapi::mkl::sparse::matrix_handle_t mkl_handle; oneapi::mkl::sparse::init_matrix_handle(&mkl_handle); - auto internal_smhandle = new detail::sparse_matrix_handle(mkl_handle, row_ptr, col_ind, val, - num_rows, num_cols, nnz, index); + auto internal_smhandle = new detail::sparse_matrix_handle( + mkl_handle, row_ptr, col_ind, val, detail::sparse_format::CSR, num_rows, num_cols, nnz, index); // The backend deduces nnz from row_ptr. auto event = oneapi::mkl::sparse::set_csr_data( queue, mkl_handle, static_cast(num_rows), static_cast(num_cols), index, diff --git a/src/sparse_blas/generic_container.hpp b/src/sparse_blas/generic_container.hpp index 33adf3abb..09d408a77 100644 --- a/src/sparse_blas/generic_container.hpp +++ b/src/sparse_blas/generic_container.hpp @@ -205,6 +205,8 @@ struct generic_dense_matrix_handle : public detail::generic_dense_handle struct generic_sparse_handle { @@ -214,6 +216,7 @@ struct generic_sparse_handle { generic_container col_container; generic_container value_container; + sparse_format format; std::int64_t num_rows; std::int64_t num_cols; std::int64_t nnz; @@ -223,12 +226,13 @@ struct generic_sparse_handle { template generic_sparse_handle(BackendHandleT backend_handle, intType* row_ptr, intType* col_ptr, - fpType* value_ptr, std::int64_t num_rows, std::int64_t num_cols, - std::int64_t nnz, oneapi::mkl::index_base index) + fpType* value_ptr, sparse_format format, std::int64_t num_rows, + std::int64_t num_cols, std::int64_t nnz, oneapi::mkl::index_base index) : backend_handle(backend_handle), row_container(generic_container(row_ptr)), col_container(generic_container(col_ptr)), value_container(generic_container(value_ptr)), + format(format), num_rows(num_rows), num_cols(num_cols), nnz(nnz), @@ -239,12 +243,14 @@ struct generic_sparse_handle { template generic_sparse_handle(BackendHandleT backend_handle, const sycl::buffer row_buffer, const sycl::buffer col_buffer, - const sycl::buffer value_buffer, std::int64_t num_rows, - std::int64_t num_cols, std::int64_t nnz, oneapi::mkl::index_base index) + const sycl::buffer value_buffer, sparse_format format, + std::int64_t num_rows, std::int64_t num_cols, std::int64_t nnz, + oneapi::mkl::index_base index) : backend_handle(backend_handle), row_container(row_buffer), col_container(col_buffer), value_container(value_buffer), + format(format), num_rows(num_rows), num_cols(num_cols), nnz(nnz), @@ -266,6 +272,11 @@ struct generic_sparse_handle { } void set_matrix_property(matrix_property property) { + if (format == sparse_format::CSR && property == matrix_property::sorted_by_rows) { + throw mkl::invalid_argument( + "sparse_blas", "set_matrix_property", + "Property `matrix_property::sorted_by_rows` is not compatible with CSR format."); + } properties_mask |= matrix_property_to_mask(property); } @@ -278,6 +289,7 @@ struct generic_sparse_handle { switch (property) { case matrix_property::symmetric: return 1 << 0; case matrix_property::sorted: return 1 << 1; + case matrix_property::sorted_by_rows: return 1 << 2; default: throw oneapi::mkl::invalid_argument( "sparse_blas", "set_matrix_property", diff --git a/tests/unit_tests/sparse_blas/include/test_common.hpp b/tests/unit_tests/sparse_blas/include/test_common.hpp index 892e0969c..3848c5bf1 100644 --- a/tests/unit_tests/sparse_blas/include/test_common.hpp +++ b/tests/unit_tests/sparse_blas/include/test_common.hpp @@ -59,12 +59,39 @@ enum sparse_matrix_format_t { COO, }; -static std::vector> test_matrix_properties{ - { oneapi::mkl::sparse::matrix_property::sorted }, - { oneapi::mkl::sparse::matrix_property::symmetric }, - { oneapi::mkl::sparse::matrix_property::sorted, - oneapi::mkl::sparse::matrix_property::symmetric } -}; +inline std::set get_default_matrix_properties( + sycl::queue queue, sparse_matrix_format_t format) { + auto vendor_id = oneapi::mkl::get_device_id(queue); + if (vendor_id == oneapi::mkl::device::nvidiagpu && format == sparse_matrix_format_t::COO) { + return { oneapi::mkl::sparse::matrix_property::sorted_by_rows }; + } + return {}; +} + +/// Return the combinations of matrix_properties to test other than the default +inline std::vector> +get_all_matrix_properties_combinations(sycl::queue queue, sparse_matrix_format_t format) { + auto vendor_id = oneapi::mkl::get_device_id(queue); + if (vendor_id == oneapi::mkl::device::nvidiagpu && format == sparse_matrix_format_t::COO) { + // Ensure all the sets have the sorted or sorted_by_rows properties + return { { oneapi::mkl::sparse::matrix_property::sorted }, + { oneapi::mkl::sparse::matrix_property::sorted_by_rows, + oneapi::mkl::sparse::matrix_property::symmetric }, + { oneapi::mkl::sparse::matrix_property::sorted, + oneapi::mkl::sparse::matrix_property::symmetric } }; + } + + std::vector> properties_combinations{ + { oneapi::mkl::sparse::matrix_property::sorted }, + { oneapi::mkl::sparse::matrix_property::symmetric }, + { oneapi::mkl::sparse::matrix_property::sorted, + oneapi::mkl::sparse::matrix_property::symmetric } + }; + if (format == sparse_matrix_format_t::COO) { + properties_combinations.push_back({ oneapi::mkl::sparse::matrix_property::sorted_by_rows }); + } + return properties_combinations; +} void print_error_code(sycl::exception const &e); @@ -332,18 +359,23 @@ intType generate_random_matrix(sparse_matrix_format_t format, const intType nrow throw std::runtime_error("Unsupported sparse format"); } -inline bool require_coo_sorted_by_row(sycl::queue queue) { - auto vendor_id = oneapi::mkl::get_device_id(queue); - return vendor_id == oneapi::mkl::device::nvidiagpu; -} - /// Shuffle the 3arrays CSR or COO representation (ia, ja, values) /// of any sparse matrix. /// In CSR format, the elements within a row are shuffled without changing ia. /// In COO format, all the elements are shuffled. template -void shuffle_sparse_matrix(sycl::queue queue, sparse_matrix_format_t format, intType indexing, - intType *ia, intType *ja, fpType *a, intType nnz, std::size_t nrows) { +void shuffle_sparse_matrix_if_needed( + sparse_matrix_format_t format, + const std::set &matrix_properties, intType indexing, + intType *ia, intType *ja, fpType *a, intType nnz, std::size_t nrows) { + const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != + matrix_properties.cend(); + if (is_sorted) { + return; + } + const bool is_sorted_by_rows = + matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted_by_rows) != + matrix_properties.cend(); if (format == sparse_matrix_format_t::CSR) { for (std::size_t i = 0; i < nrows; ++i) { intType nnz_row = ia[i + 1] - ia[i]; @@ -354,9 +386,10 @@ void shuffle_sparse_matrix(sycl::queue queue, sparse_matrix_format_t format, int std::swap(a[q], a[j]); } } + // sorted_by_rows does not impact CSR } else if (format == sparse_matrix_format_t::COO) { - if (require_coo_sorted_by_row(queue)) { + if (is_sorted_by_rows) { std::size_t linear_idx = 0; for (std::size_t i = 0; i < nrows; ++i) { // Count the number of non-zero elements for the given row @@ -386,7 +419,7 @@ void shuffle_sparse_matrix(sycl::queue queue, sparse_matrix_format_t format, int } } else { - throw oneapi::mkl::exception("sparse_blas", "shuffle_sparse_matrix", + throw oneapi::mkl::exception("sparse_blas", "shuffle_sparse_matrix_if_needed", "Internal error: unsupported format"); } } diff --git a/tests/unit_tests/sparse_blas/include/test_spmm.hpp b/tests/unit_tests/sparse_blas/include/test_spmm.hpp index 6188d4268..d47b1732c 100644 --- a/tests/unit_tests/sparse_blas/include/test_spmm.hpp +++ b/tests/unit_tests/sparse_blas/include/test_spmm.hpp @@ -65,10 +65,13 @@ void test_helper_with_format_with_transpose( oneapi::mkl::layout col_major = oneapi::mkl::layout::col_major; oneapi::mkl::sparse::spmm_alg default_alg = oneapi::mkl::sparse::spmm_alg::default_alg; oneapi::mkl::sparse::matrix_view default_A_view; - std::set no_properties; bool no_reset_data = false; bool no_scalars_on_device = false; + // Queue is only used to get which matrix_property should be used for the tests. + sycl::queue properties_queue(*dev); + auto default_properties = get_default_matrix_properties(properties_queue, format); + { int m = 4, k = 6, n = 5; int nrows_A = (transpose_A != oneapi::mkl::transpose::nontrans) ? k : m; @@ -84,34 +87,34 @@ void test_helper_with_format_with_transpose( EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb, ldc, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Reset data EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb, ldc, - default_alg, default_A_view, no_properties, true, + default_alg, default_A_view, default_properties, true, no_scalars_on_device), num_passed, num_skipped); // Test alpha and beta on the device EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb, ldc, - default_alg, default_A_view, no_properties, no_reset_data, true), + default_alg, default_A_view, default_properties, no_reset_data, true), num_passed, num_skipped); // Test index_base 1 EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, oneapi::mkl::index_base::one, col_major, transpose_A, transpose_B, - fp_one, fp_zero, ldb, ldc, default_alg, default_A_view, no_properties, - no_reset_data, no_scalars_on_device), + fp_one, fp_zero, ldb, ldc, default_alg, default_A_view, + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default alpha EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, set_fp_value()(2.f, 1.5f), - fp_zero, ldb, ldc, default_alg, default_A_view, no_properties, + fp_zero, ldb, ldc, default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default beta @@ -119,42 +122,43 @@ void test_helper_with_format_with_transpose( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, set_fp_value()(3.2f, 1.f), ldb, ldc, default_alg, - default_A_view, no_properties, no_reset_data, no_scalars_on_device), + default_A_view, default_properties, no_reset_data, + no_scalars_on_device), num_passed, num_skipped); // Test 0 alpha EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_zero, fp_one, ldb, ldc, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test 0 alpha and beta EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_zero, fp_zero, ldb, ldc, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default ldb EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb + 5, ldc, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default ldc EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb, ldc + 6, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test row major layout EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, oneapi::mkl::layout::row_major, transpose_A, transpose_B, fp_one, - fp_zero, ncols_B, ncols_C, default_alg, default_A_view, no_properties, - no_reset_data, no_scalars_on_device), + fp_zero, ncols_B, ncols_C, default_alg, default_A_view, + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test int64 indices long long_nrows_A = 27, long_ncols_A = 13, long_ncols_C = 6; @@ -163,19 +167,19 @@ void test_helper_with_format_with_transpose( test_functor_i64(dev, format, long_nrows_A, long_ncols_A, long_ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, long_ldb, long_ldc, default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test other algorithms for (auto alg : non_default_algorithms) { EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, - ldb, ldc, alg, default_A_view, no_properties, no_reset_data, + ldb, ldc, alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); } // Test matrix properties - for (auto properties : test_matrix_properties) { + for (auto properties : get_all_matrix_properties_combinations(properties_queue, format)) { EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, @@ -197,7 +201,7 @@ void test_helper_with_format_with_transpose( EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, ncols_C, density_A_matrix, index_zero, col_major, transpose_A, transpose_B, fp_one, fp_zero, ldb, ldc, - default_alg, default_A_view, no_properties, no_reset_data, + default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); } diff --git a/tests/unit_tests/sparse_blas/include/test_spmv.hpp b/tests/unit_tests/sparse_blas/include/test_spmv.hpp index f141db893..66af38a7c 100644 --- a/tests/unit_tests/sparse_blas/include/test_spmv.hpp +++ b/tests/unit_tests/sparse_blas/include/test_spmv.hpp @@ -63,62 +63,65 @@ void test_helper_with_format_with_transpose( oneapi::mkl::index_base index_zero = oneapi::mkl::index_base::zero; oneapi::mkl::sparse::spmv_alg default_alg = oneapi::mkl::sparse::spmv_alg::default_alg; oneapi::mkl::sparse::matrix_view default_A_view; - std::set no_properties; bool no_reset_data = false; bool no_scalars_on_device = false; + // Queue is only used to get which matrix_property should be used for the tests. + sycl::queue properties_queue(*dev); + auto default_properties = get_default_matrix_properties(properties_queue, format); + // Basic test 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, default_A_view, no_properties, no_reset_data, - no_scalars_on_device), + fp_one, fp_zero, default_alg, default_A_view, default_properties, + no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Reset data 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, default_A_view, no_properties, true, + fp_one, fp_zero, default_alg, default_A_view, default_properties, true, no_scalars_on_device), num_passed, num_skipped); // Test alpha and beta on the device 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, default_A_view, no_properties, no_reset_data, - true), + fp_one, fp_zero, default_alg, default_A_view, default_properties, + no_reset_data, true), num_passed, num_skipped); // Test index_base 1 EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, oneapi::mkl::index_base::one, transpose_val, fp_one, fp_zero, default_alg, - default_A_view, no_properties, no_reset_data, no_scalars_on_device), + default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default alpha EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, set_fp_value()(2.f, 1.5f), fp_zero, default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default beta EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, fp_one, set_fp_value()(3.2f, 1.f), default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test 0 alpha EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, - fp_zero, fp_one, default_alg, default_A_view, no_properties, no_reset_data, - no_scalars_on_device), + fp_zero, fp_one, default_alg, default_A_view, default_properties, + no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test 0 alpha and beta EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, nrows_A, ncols_A, density_A_matrix, index_zero, transpose_val, - fp_zero, fp_zero, default_alg, default_A_view, no_properties, + fp_zero, fp_zero, default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test int64 indices EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i64(dev, format, 27L, 13L, density_A_matrix, index_zero, transpose_val, fp_one, - fp_zero, default_alg, default_A_view, no_properties, no_reset_data, + fp_zero, default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Lower triangular @@ -126,14 +129,14 @@ void test_helper_with_format_with_transpose( oneapi::mkl::sparse::matrix_descr::triangular); 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, triangular_A_view, no_properties, + fp_one, fp_zero, default_alg, triangular_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Upper triangular triangular_A_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, triangular_A_view, no_properties, + fp_one, fp_zero, default_alg, triangular_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Lower triangular unit diagonal @@ -142,54 +145,54 @@ void test_helper_with_format_with_transpose( triangular_unit_A_view.diag_view = oneapi::mkl::diag::unit; 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, triangular_unit_A_view, no_properties, + fp_one, fp_zero, default_alg, triangular_unit_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Upper triangular unit diagonal triangular_A_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, triangular_unit_A_view, no_properties, + fp_one, fp_zero, default_alg, triangular_unit_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Lower symmetric oneapi::mkl::sparse::matrix_view symmetric_view(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), + fp_one, fp_zero, default_alg, symmetric_view, default_properties, + no_reset_data, no_scalars_on_device), num_passed, num_skipped); // 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), + fp_one, fp_zero, default_alg, symmetric_view, default_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), + fp_one, fp_zero, default_alg, hermitian_view, default_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), + fp_one, fp_zero, default_alg, hermitian_view, default_properties, + no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test other algorithms for (auto alg : non_default_algorithms) { 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, alg, default_A_view, no_properties, - no_reset_data, no_scalars_on_device), + transpose_val, fp_one, fp_zero, alg, default_A_view, + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); } // Test matrix properties - for (auto properties : test_matrix_properties) { + for (auto properties : get_all_matrix_properties_combinations(properties_queue, format)) { 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, default_A_view, diff --git a/tests/unit_tests/sparse_blas/include/test_spsv.hpp b/tests/unit_tests/sparse_blas/include/test_spsv.hpp index bdf9210f8..ca58dfd7a 100644 --- a/tests/unit_tests/sparse_blas/include/test_spsv.hpp +++ b/tests/unit_tests/sparse_blas/include/test_spsv.hpp @@ -60,76 +60,83 @@ void test_helper_with_format(testFunctorI32 test_functor_i32, testFunctorI64 tes oneapi::mkl::sparse::matrix_view default_A_view(oneapi::mkl::sparse::matrix_descr::triangular); oneapi::mkl::sparse::matrix_view upper_A_view(oneapi::mkl::sparse::matrix_descr::triangular); upper_A_view.uplo_view = oneapi::mkl::uplo::upper; - std::set no_properties; bool no_reset_data = false; bool no_scalars_on_device = false; + // Queue is only used to get which matrix_property should be used for the tests. + sycl::queue properties_queue(*dev); + auto default_properties = get_default_matrix_properties(properties_queue, format); + // Basic test - EXPECT_TRUE_OR_FUTURE_SKIP(test_functor_i32(dev, format, m, density_A_matrix, index_zero, - transpose_val, alpha, default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), - num_passed, num_skipped); - // Reset data EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, - default_alg, default_A_view, no_properties, true, no_scalars_on_device), + default_alg, default_A_view, default_properties, no_reset_data, + no_scalars_on_device), num_passed, num_skipped); + // Reset data + EXPECT_TRUE_OR_FUTURE_SKIP(test_functor_i32(dev, format, m, density_A_matrix, index_zero, + transpose_val, alpha, default_alg, default_A_view, + default_properties, true, no_scalars_on_device), + num_passed, num_skipped); // Test alpha on the device EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, - default_alg, default_A_view, no_properties, no_reset_data, true), + default_alg, default_A_view, default_properties, no_reset_data, true), num_passed, num_skipped); // Test index_base 1 EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, oneapi::mkl::index_base::one, - transpose_val, alpha, default_alg, default_A_view, no_properties, + transpose_val, alpha, default_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test upper triangular matrix - EXPECT_TRUE_OR_FUTURE_SKIP(test_functor_i32(dev, format, m, density_A_matrix, index_zero, - transpose_val, alpha, default_alg, upper_A_view, - no_properties, no_reset_data, no_scalars_on_device), - num_passed, num_skipped); + EXPECT_TRUE_OR_FUTURE_SKIP( + test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, + default_alg, upper_A_view, default_properties, no_reset_data, + no_scalars_on_device), + num_passed, num_skipped); // Test lower triangular unit diagonal matrix oneapi::mkl::sparse::matrix_view triangular_unit_A_view( oneapi::mkl::sparse::matrix_descr::triangular); triangular_unit_A_view.diag_view = oneapi::mkl::diag::unit; EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, - default_alg, triangular_unit_A_view, no_properties, no_reset_data, + default_alg, triangular_unit_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test upper triangular unit diagonal matrix triangular_unit_A_view.uplo_view = oneapi::mkl::uplo::upper; EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, - default_alg, triangular_unit_A_view, no_properties, no_reset_data, + default_alg, triangular_unit_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test non-default alpha EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, set_fp_value()(2.f, 1.5f), default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), + default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test int64 indices - EXPECT_TRUE_OR_FUTURE_SKIP(test_functor_i64(dev, format, 15L, density_A_matrix, index_zero, - transpose_val, alpha, default_alg, default_A_view, - no_properties, no_reset_data, no_scalars_on_device), - num_passed, num_skipped); + EXPECT_TRUE_OR_FUTURE_SKIP( + test_functor_i64(dev, format, 15L, density_A_matrix, index_zero, transpose_val, alpha, + default_alg, default_A_view, default_properties, no_reset_data, + no_scalars_on_device), + num_passed, num_skipped); // Test lower no_optimize_alg EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, - no_optimize_alg, default_A_view, no_properties, no_reset_data, + no_optimize_alg, default_A_view, default_properties, no_reset_data, no_scalars_on_device), num_passed, num_skipped); // Test upper no_optimize_alg - EXPECT_TRUE_OR_FUTURE_SKIP(test_functor_i32(dev, format, m, density_A_matrix, index_zero, - transpose_val, alpha, no_optimize_alg, upper_A_view, - no_properties, no_reset_data, no_scalars_on_device), - num_passed, num_skipped); + EXPECT_TRUE_OR_FUTURE_SKIP( + test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, + no_optimize_alg, upper_A_view, default_properties, no_reset_data, + no_scalars_on_device), + num_passed, num_skipped); // Test matrix properties - for (auto properties : test_matrix_properties) { + for (auto properties : get_all_matrix_properties_combinations(properties_queue, format)) { // Basic test with matrix properties EXPECT_TRUE_OR_FUTURE_SKIP( test_functor_i32(dev, format, m, density_A_matrix, index_zero, transpose_val, alpha, diff --git a/tests/unit_tests/sparse_blas/source/sparse_spmm_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_spmm_buffer.cpp index df6fb850b..4a37e8c7c 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmm_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmm_buffer.cpp @@ -51,8 +51,6 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, auto [opa_nrows, opa_ncols] = swap_if_transposed(transpose_A, nrows_A, ncols_A); auto [opb_nrows, opb_ncols] = swap_if_transposed(transpose_B, opa_ncols, ncols_C); intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -73,10 +71,9 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, std::vector c_ref_host(c_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, + static_cast(nrows_A)); auto ia_buf = make_buffer(ia_host); auto ja_buf = make_buffer(ja_host); @@ -119,10 +116,9 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, intType reset_nnz = generate_random_matrix( format, nrows_A, ncols_A, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, + static_cast(nrows_A)); if (reset_nnz > nnz) { ia_buf = make_buffer(ia_host); ja_buf = make_buffer(ja_host); diff --git a/tests/unit_tests/sparse_blas/source/sparse_spmm_usm.cpp b/tests/unit_tests/sparse_blas/source/sparse_spmm_usm.cpp index 7d30426c4..8070633fc 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmm_usm.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmm_usm.cpp @@ -47,8 +47,6 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, auto [opa_nrows, opa_ncols] = swap_if_transposed(transpose_A, nrows_A, ncols_A); auto [opb_nrows, opb_ncols] = swap_if_transposed(transpose_B, opa_ncols, ncols_C); intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -69,10 +67,9 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, std::vector c_ref_host(c_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, + static_cast(nrows_A)); auto ia_usm_uptr = malloc_device_uptr(main_queue, ia_host.size()); auto ja_usm_uptr = malloc_device_uptr(main_queue, ja_host.size()); @@ -152,10 +149,9 @@ int test_spmm(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, intType reset_nnz = generate_random_matrix( format, nrows_A, ncols_A, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, + static_cast(nrows_A)); if (reset_nnz > nnz) { // Wait before freeing usm pointers ev_spmm.wait_and_throw(); 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 e03c09ebe..f56deaf91 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmv_buffer.cpp @@ -45,8 +45,6 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, } auto [opa_nrows, opa_ncols] = swap_if_transposed(transpose_val, nrows_A, ncols_A); intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -66,10 +64,9 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, std::vector y_ref_host(y_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, + static_cast(nrows_A)); auto ia_buf = make_buffer(ia_host); auto ja_buf = make_buffer(ja_host); @@ -109,10 +106,9 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, intType reset_nnz = generate_random_matrix( format, nrows_A, ncols_A, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, + static_cast(nrows_A)); if (reset_nnz > nnz) { ia_buf = make_buffer(ia_host); ja_buf = make_buffer(ja_host); 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 eb54f6a5d..2852a2495 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spmv_usm.cpp @@ -41,8 +41,6 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, } auto [opa_nrows, opa_ncols] = swap_if_transposed(transpose_val, nrows_A, ncols_A); intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -62,10 +60,9 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, std::vector y_ref_host(y_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, + static_cast(nrows_A)); auto ia_usm_uptr = malloc_device_uptr(main_queue, ia_host.size()); auto ja_usm_uptr = malloc_device_uptr(main_queue, ja_host.size()); @@ -144,10 +141,9 @@ int test_spmv(sycl::device *dev, sparse_matrix_format_t format, intType nrows_A, intType reset_nnz = generate_random_matrix( format, nrows_A, ncols_A, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, static_cast(nrows_A)); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, + static_cast(nrows_A)); if (reset_nnz > nnz) { // Wait before freeing usm pointers ev_spmv.wait_and_throw(); diff --git a/tests/unit_tests/sparse_blas/source/sparse_spsv_buffer.cpp b/tests/unit_tests/sparse_blas/source/sparse_spsv_buffer.cpp index 3a9d153d6..ebf47fd5e 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spsv_buffer.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spsv_buffer.cpp @@ -41,8 +41,6 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; const std::size_t mu = static_cast(m); - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -72,10 +70,8 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl std::vector y_ref_host(y_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, mu); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, mu); auto ia_buf = make_buffer(ia_host); auto ja_buf = make_buffer(ja_host); @@ -112,10 +108,8 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl intType reset_nnz = generate_random_matrix( format, m, m, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric, require_diagonal); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, mu); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, mu); if (reset_nnz > nnz) { ia_buf = make_buffer(ia_host); ja_buf = make_buffer(ja_host); diff --git a/tests/unit_tests/sparse_blas/source/sparse_spsv_usm.cpp b/tests/unit_tests/sparse_blas/source/sparse_spsv_usm.cpp index 6529069f9..03edf8d37 100644 --- a/tests/unit_tests/sparse_blas/source/sparse_spsv_usm.cpp +++ b/tests/unit_tests/sparse_blas/source/sparse_spsv_usm.cpp @@ -37,8 +37,6 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl intType indexing = (index == oneapi::mkl::index_base::zero) ? 0 : 1; const std::size_t mu = static_cast(m); - const bool is_sorted = matrix_properties.find(oneapi::mkl::sparse::matrix_property::sorted) != - matrix_properties.cend(); const bool is_symmetric = matrix_properties.find(oneapi::mkl::sparse::matrix_property::symmetric) != matrix_properties.cend(); @@ -68,10 +66,8 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl std::vector y_ref_host(y_host); // Shuffle ordering of column indices/values to test sortedness - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), nnz, mu); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), nnz, mu); auto ia_usm_uptr = malloc_device_uptr(main_queue, ia_host.size()); auto ja_usm_uptr = malloc_device_uptr(main_queue, ja_host.size()); @@ -143,10 +139,8 @@ int test_spsv(sycl::device *dev, sparse_matrix_format_t format, intType m, doubl intType reset_nnz = generate_random_matrix( format, m, m, density_A_matrix, indexing, ia_host, ja_host, a_host, is_symmetric, require_diagonal); - if (!is_sorted) { - shuffle_sparse_matrix(main_queue, format, indexing, ia_host.data(), ja_host.data(), - a_host.data(), reset_nnz, mu); - } + shuffle_sparse_matrix_if_needed(format, matrix_properties, indexing, ia_host.data(), + ja_host.data(), a_host.data(), reset_nnz, mu); if (reset_nnz > nnz) { // Wait before freeing usm pointers ev_spsv.wait_and_throw();