diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc index f4d6eea59645..6885185efda1 100644 --- a/clang/include/clang/DPCT/DPCTOptions.inc +++ b/clang/include/clang/DPCT/DPCTOptions.inc @@ -632,7 +632,7 @@ DPCT_ENUM_OPTION( "syclcompat", int(clang::dpct::ExplicitNamespace::EN_SYCLCompat), "DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Generate code with syclcompat:: namespace.", false)), llvm::cl::desc("Define the namespaces to use explicitly in generated " - "code. The is a comma\n" + "code. The is a comma\n" "separated list. Default: dpct/syclcompat, sycl.\n" "Possible values are:"), llvm::cl::CommaSeparated, llvm::cl::value_desc("value"), @@ -1019,7 +1019,7 @@ DPCT_FLAG_OPTION( "use-syclcompat", llvm::cl::desc( "DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Use SYCLcompat header-only library (syclcompat:: namespace) to assist " - "the migration of input source code. \nDefault: off.\n"), + "the migration of input source code.\nDefault: off.\n"), llvm::cl::cat(CtHelpCatCodeGen), llvm::cl::cat(CtHelpCatAll)) DPCT_FLAG_OPTION( diff --git a/clang/lib/DPCT/RulesLangLib/CUB/RewriterDeviceSpmv.cpp b/clang/lib/DPCT/RulesLangLib/CUB/RewriterDeviceSpmv.cpp index 0101fa3e1f3e..c0714fb9aa5e 100644 --- a/clang/lib/DPCT/RulesLangLib/CUB/RewriterDeviceSpmv.cpp +++ b/clang/lib/DPCT/RulesLangLib/CUB/RewriterDeviceSpmv.cpp @@ -28,10 +28,10 @@ RewriterMap dpct::createDeviceSpmvRewriterMap() { "cub::DeviceSpmv::CsrMV", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", STREAM(10), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), - ARG(7), ARG(8))), + ARG(7), ARG(8), ARG(9))), CALL_FACTORY_ENTRY( "cub::DeviceSpmv::CsrMV", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", QUEUESTR, ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), - ARG(7), ARG(8)))))))}; + ARG(7), ARG(8), ARG(9)))))))}; } diff --git a/clang/lib/DPCT/RulesMathLib/APINamesCUSPARSE.inc b/clang/lib/DPCT/RulesMathLib/APINamesCUSPARSE.inc index 4928f1f51dea..bc9f0eef91b5 100644 --- a/clang/lib/DPCT/RulesMathLib/APINamesCUSPARSE.inc +++ b/clang/lib/DPCT/RulesMathLib/APINamesCUSPARSE.inc @@ -74,28 +74,28 @@ FEATURE_REQUEST_FACTORY( ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseScsrsv_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8))))) FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseDcsrsv_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8))))) FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCcsrsv_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8))))) FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseZcsrsv_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -103,28 +103,28 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseScsrmv", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseDcsrmv", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCcsrmv", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseZcsrmv", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -132,28 +132,28 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseScsrmv_mp", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseDcsrmv_mp", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCcsrmv_mp", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseZcsrmv_mp", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(5), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12))))) ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cusparseCsrmvEx_bufferSize", DEREF(20), @@ -163,7 +163,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseCsrmvEx", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmv", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(2), - ARG(3), ARG(4), ARG(6), ARG(7), ARG(8), ARG(9), + ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15), ARG(16), ARG(17), ARG(18))))) @@ -172,7 +172,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseScsrmm", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -180,7 +180,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseDcsrmm", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -188,7 +188,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseCcsrmm", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -196,7 +196,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseZcsrmm", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(6), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15))))) @@ -205,7 +205,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseScsrmm2", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(5), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15), ARG(16))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -213,7 +213,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseDcsrmm2", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(5), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15), ARG(16))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -221,7 +221,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseCcsrmm2", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(5), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15), ARG(16))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, @@ -229,7 +229,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "cusparseZcsrmm2", CALL(MapNames::getLibraryHelperNamespace() + "sparse::csrmm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), - ARG(2), ARG(3), ARG(4), ARG(5), ARG(7), ARG(8), + ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9), ARG(10), ARG(11), ARG(12), ARG(13), ARG(14), ARG(15), ARG(16))))) @@ -451,22 +451,22 @@ ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cusparseZcsrsv2_bufferSize", DEREF(9), ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseScsrsv2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), ARG(5), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseDcsrsv2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), ARG(5), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCcsrsv2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), ARG(5), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseZcsrsv2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), ARG(5), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( @@ -502,7 +502,7 @@ ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cusparseZcsrsv2_bufferSizeExt", ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCsrsv_analysisEx", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsv", - MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(4), ARG(5), + MEMBER_CALL(ARG(0), true, "get_queue"), ARG(1), ARG(2), ARG(3), ARG(4), ARG(5), ARG(6), ARG(7), ARG(8), ARG(9)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCsrsv_solveEx", @@ -537,22 +537,22 @@ ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseScsrsm2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(2), ARG(3), ARG(4), ARG(5), - ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) + ARG(6), ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseDcsrsm2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(2), ARG(3), ARG(4), ARG(5), - ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) + ARG(6), ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseCcsrsm2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(2), ARG(3), ARG(4), ARG(5), - ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) + ARG(6), ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseZcsrsm2_analysis", CALL(MapNames::getLibraryHelperNamespace() + "sparse::optimize_csrsm", MEMBER_CALL(ARG(0), true, "get_queue"), ARG(2), ARG(3), ARG(4), ARG(5), - ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) + ARG(6), ARG(8), ARG(9), ARG(10), ARG(11), ARG(14)))) ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY( "cusparseScsrsm2_solve", diff --git a/clang/runtime/dpct-rt/include/dpct/detail/sparse_utils_detail.hpp b/clang/runtime/dpct-rt/include/dpct/detail/sparse_utils_detail.hpp index b341e7878d19..698b07bcdbed 100644 --- a/clang/runtime/dpct-rt/include/dpct/detail/sparse_utils_detail.hpp +++ b/clang/runtime/dpct-rt/include/dpct/detail/sparse_utils_detail.hpp @@ -128,18 +128,18 @@ class matrix_handle_manager public: using handle_manager::handle_manager; template - void set_matrix_data(const int rows, const int cols, + void set_matrix_data(const int rows, const int cols, const int nnz, oneapi::mkl::index_base base, const void *row_ptr, const void *col_ind, const void *val) { #ifdef DPCT_USM_LEVEL_NONE _row_ptr_buf = dpct::detail::get_memory((int *)row_ptr); _col_ind_buf = dpct::detail::get_memory((int *)col_ind); _val_buf = dpct::detail::get_memory((Ty *)val); - oneapi::mkl::sparse::set_csr_data(*_q, get_handle(), rows, cols, base, + oneapi::mkl::sparse::set_csr_data(*_q, get_handle(), rows, cols, nnz, base, _row_ptr_buf, _col_ind_buf, std::get>(_val_buf)); #else - oneapi::mkl::sparse::set_csr_data(*_q, get_handle(), rows, cols, base, + oneapi::mkl::sparse::set_csr_data(*_q, get_handle(), rows, cols, nnz, base, (int *)row_ptr, (int *)col_ind, (Ty *)val); #endif @@ -171,6 +171,41 @@ class matrix_handle_manager #endif template struct optimize_csrsv_impl { + void operator()(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, int nnz, + const std::shared_ptr info, const void *val, + const int *row_ptr, const int *col_ind, + std::shared_ptr optimize_info) { + using Ty = typename ::dpct::detail::lib_data_traits_t; + auto temp_row_ptr = dpct::detail::get_memory(row_ptr); + auto temp_col_ind = dpct::detail::get_memory(col_ind); + auto temp_val = dpct::detail::get_memory(val); +#ifdef DPCT_USM_LEVEL_NONE + optimize_info->_row_ptr_buf = temp_row_ptr; + optimize_info->_col_ind_buf = temp_col_ind; + optimize_info->_val_buf = temp_val; + auto &data_row_ptr = optimize_info->_row_ptr_buf; + auto &data_col_ind = optimize_info->_col_ind_buf; + auto &data_val = std::get>(optimize_info->_val_buf); +#else + auto data_row_ptr = temp_row_ptr; + auto data_col_ind = temp_col_ind; + auto data_val = temp_val; +#endif + oneapi::mkl::sparse::set_csr_data(queue, optimize_info->get_matrix_handle(), + row_col, row_col, nnz, info->get_index_base(), + data_row_ptr, data_col_ind, data_val); + if (info->get_matrix_type() != matrix_info::matrix_type::tr) + throw std::runtime_error("dpct::sparse::optimize_csrsv_impl()(): " + "oneapi::mkl::sparse::optimize_trsv " + "only accept triangular matrix."); + SPARSE_CALL(oneapi::mkl::sparse::optimize_trsv( + queue, info->get_uplo(), trans, info->get_diag(), + optimize_info->get_matrix_handle()), + optimize_info); + } +}; + +template struct optimize_csrsv_impl_deprecated { void operator()(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, const std::shared_ptr info, const void *val, const int *row_ptr, const int *col_ind, @@ -204,6 +239,7 @@ template struct optimize_csrsv_impl { optimize_info); } }; + template struct csrsv_impl { void operator()(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, const void *alpha, const std::shared_ptr info, @@ -224,6 +260,46 @@ template struct csrsv_impl { }; template struct optimize_csrsm_impl { + void operator()(sycl::queue &queue, oneapi::mkl::transpose transa, + oneapi::mkl::transpose transb, int row_col, int nrhs, int nnz, + const std::shared_ptr info, const void *val, + const int *row_ptr, const int *col_ind, + std::shared_ptr optimize_info) { + using Ty = typename ::dpct::detail::lib_data_traits_t; + auto temp_row_ptr = dpct::detail::get_memory(row_ptr); + auto temp_col_ind = dpct::detail::get_memory(col_ind); + auto temp_val = dpct::detail::get_memory(val); +#ifdef DPCT_USM_LEVEL_NONE + optimize_info->_row_ptr_buf = temp_row_ptr; + optimize_info->_col_ind_buf = temp_col_ind; + optimize_info->_val_buf = temp_val; + auto &data_row_ptr = optimize_info->_row_ptr_buf; + auto &data_col_ind = optimize_info->_col_ind_buf; + auto &data_val = std::get>(optimize_info->_val_buf); +#else + auto data_row_ptr = temp_row_ptr; + auto data_col_ind = temp_col_ind; + auto data_val = temp_val; +#endif + oneapi::mkl::sparse::set_csr_data(queue, optimize_info->get_matrix_handle(), + row_col, row_col, nnz, info->get_index_base(), + data_row_ptr, data_col_ind, data_val); + if (info->get_matrix_type() != matrix_info::matrix_type::tr) + throw std::runtime_error("dpct::sparse::optimize_csrsv_impl()(): " + "oneapi::mkl::sparse::optimize_trsm " + "only accept triangular matrix."); + SPARSE_CALL(oneapi::mkl::sparse::optimize_trsm( + queue, + transb == oneapi::mkl::transpose::nontrans + ? oneapi::mkl::layout::col_major + : oneapi::mkl::layout::row_major, + info->get_uplo(), transa, info->get_diag(), + optimize_info->get_matrix_handle(), nrhs), + optimize_info); + } +}; + +template struct optimize_csrsm_impl_deprecated { void operator()(sycl::queue &queue, oneapi::mkl::transpose transa, oneapi::mkl::transpose transb, int row_col, int nrhs, const std::shared_ptr info, const void *val, @@ -262,6 +338,7 @@ template struct optimize_csrsm_impl { optimize_info); } }; + template struct csrsm_impl { void operator()(sycl::queue &queue, oneapi::mkl::transpose transa, oneapi::mkl::transpose transb, int row_col, int nrhs, @@ -470,10 +547,10 @@ template struct csr2csc_impl { ::dpct::cs::malloc(sizeof(Ty) * nnz, ::dpct::cs::get_default_queue()); } auto data_to_val = dpct::detail::get_memory(new_to_value); - oneapi::mkl::sparse::set_csr_data(queue, from_handle, m, n, base, + oneapi::mkl::sparse::set_csr_data(queue, from_handle, m, n, nnz, base, data_from_row_ptr, data_from_col_ind, data_from_val); - oneapi::mkl::sparse::set_csr_data(queue, to_handle, n, m, base, + oneapi::mkl::sparse::set_csr_data(queue, to_handle, n, m, nnz, base, data_to_col_ptr, data_to_row_ind, data_to_val); sycl::event e1 = oneapi::mkl::sparse::omatcopy( @@ -514,6 +591,73 @@ inline void spblas_shim(library_data_t type, args_t &&...args) { } template struct csrmv_impl { + void operator()(sycl::queue &queue, oneapi::mkl::transpose trans, + int num_rows, int num_cols, int nnz, const void *alpha, + const std::shared_ptr info, const void *val, + const int *row_ptr, const int *col_ind, const void *x, + const void *beta, void *y) { +#ifndef __INTEL_MKL__ + throw std::runtime_error( + "The oneAPI Math Kernel Library (oneMKL) Interfaces " + "Project does not support this API."); +#else + using Ty = typename ::dpct::detail::lib_data_traits_t; + auto alpha_value = + dpct::detail::get_value(reinterpret_cast(alpha), queue); + auto beta_value = + dpct::detail::get_value(reinterpret_cast(beta), queue); + + oneapi::mkl::sparse::matrix_handle_t *sparse_matrix_handle = + new oneapi::mkl::sparse::matrix_handle_t; + oneapi::mkl::sparse::init_matrix_handle(sparse_matrix_handle); + auto data_row_ptr = dpct::detail::get_memory(row_ptr); + auto data_col_ind = dpct::detail::get_memory(col_ind); + auto data_val = dpct::detail::get_memory(val); + oneapi::mkl::sparse::set_csr_data(queue, *sparse_matrix_handle, num_rows, + num_cols, nnz, info->get_index_base(), + data_row_ptr, data_col_ind, data_val); + + auto data_x = dpct::detail::get_memory(x); + auto data_y = dpct::detail::get_memory(y); + switch (info->get_matrix_type()) { + case matrix_info::matrix_type::ge: { + oneapi::mkl::sparse::optimize_gemv(queue, trans, *sparse_matrix_handle); + oneapi::mkl::sparse::gemv(queue, trans, alpha_value, + *sparse_matrix_handle, data_x, beta_value, + data_y); + break; + } + case matrix_info::matrix_type::sy: { + oneapi::mkl::sparse::symv(queue, info->get_uplo(), alpha_value, + *sparse_matrix_handle, data_x, beta_value, + data_y); + break; + } + case matrix_info::matrix_type::tr: { + oneapi::mkl::sparse::optimize_trmv(queue, info->get_uplo(), trans, + info->get_diag(), + *sparse_matrix_handle); + oneapi::mkl::sparse::trmv( + queue, info->get_uplo(), trans, info->get_diag(), alpha_value, + *sparse_matrix_handle, data_x, beta_value, data_y); + break; + } + default: + throw std::runtime_error( + "the spmv does not support matrix_info::matrix_type::he"); + } + + sycl::event e = + oneapi::mkl::sparse::release_matrix_handle(queue, sparse_matrix_handle); + queue.submit([&](sycl::handler &cgh) { + cgh.depends_on(e); + cgh.host_task([=] { delete sparse_matrix_handle; }); + }); +#endif + } +}; + +template struct csrmv_impl_deprecated { void operator()(sycl::queue &queue, oneapi::mkl::transpose trans, int num_rows, int num_cols, const void *alpha, const std::shared_ptr info, const void *val, diff --git a/clang/runtime/dpct-rt/include/dpct/sparse_utils.hpp b/clang/runtime/dpct-rt/include/dpct/sparse_utils.hpp index 4833ac92e2cd..585a10abe600 100644 --- a/clang/runtime/dpct-rt/include/dpct/sparse_utils.hpp +++ b/clang/runtime/dpct-rt/include/dpct/sparse_utils.hpp @@ -26,16 +26,21 @@ class matrix_info { auto get_diag() const { return _diag; } auto get_uplo() const { return _uplo; } auto get_index_base() const { return _index_base; } + /// Only used for csrgemm_nnz and csrgemm + auto get_nnz() const { return _nnz; } void set_matrix_type(matrix_type mt) { _matrix_type = mt; } void set_diag(oneapi::mkl::diag d) { _diag = d; } void set_uplo(oneapi::mkl::uplo u) { _uplo = u; } void set_index_base(oneapi::mkl::index_base ib) { _index_base = ib; } + /// Only used for csrgemm_nnz and csrgemm + auto set_nnz(int nnz) { _nnz = nnz; } private: matrix_type _matrix_type = matrix_type::ge; oneapi::mkl::diag _diag = oneapi::mkl::diag::nonunit; oneapi::mkl::uplo _uplo = oneapi::mkl::uplo::upper; oneapi::mkl::index_base _index_base = oneapi::mkl::index_base::zero; + int _nnz = 0; // Only used for csrgemm_nnz and csrgemm }; /// Sparse matrix data format @@ -50,7 +55,9 @@ enum class conversion_scope : int { index = 0, index_and_value }; // Forward declaration namespace detail { template struct optimize_csrsv_impl; +template struct optimize_csrsv_impl_deprecated; template struct optimize_csrsm_impl; +template struct optimize_csrsm_impl_deprecated; } /// Saving the optimization information for solving a system of linear @@ -75,7 +82,9 @@ class optimize_info { } #ifdef DPCT_USM_LEVEL_NONE template friend struct detail::optimize_csrsv_impl; + template friend struct detail::optimize_csrsv_impl_deprecated; template friend struct detail::optimize_csrsm_impl; + template friend struct detail::optimize_csrsm_impl_deprecated; #endif private: @@ -353,7 +362,7 @@ class sparse_matrix_desc { if (_data_format == matrix_format::csr) oneapi::mkl::sparse::set_csr_data( ::dpct::cs::get_default_queue(), _matrix_handle, _row_num, _col_num, - _base, std::get(_data_row_ptr), + _nnz, _base, std::get(_data_row_ptr), std::get(_data_col_ind), std::get(_data_value)); else @@ -593,6 +602,7 @@ using descriptor_ptr = descriptor *; /// \param [in] trans The operation applied to the matrix A. /// \param [in] num_rows Number of rows of the matrix A. /// \param [in] num_cols Number of columns of the matrix A. +/// \param [in] nnz NNZ of the matrix A. /// \param [in] alpha Scaling factor for the matrix A. /// \param [in] info Matrix info of the matrix A. /// \param [in] val An array containing the non-zero elements of the matrix A. @@ -604,14 +614,42 @@ using descriptor_ptr = descriptor *; /// \param [in, out] y Data of the vector y. template void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, int num_rows, - int num_cols, const T *alpha, + int num_cols, int nnz, const T *alpha, const std::shared_ptr info, const T *val, const int *row_ptr, const int *col_ind, const T *x, const T *beta, T *y) { - detail::csrmv_impl()(queue, trans, num_rows, num_cols, alpha, info, val, + detail::csrmv_impl()(queue, trans, num_rows, num_cols, nnz, alpha, info, val, row_ptr, col_ind, x, beta, y); } +/// Computes a CSR format sparse matrix-dense vector product. +/// y = alpha * op(A) * x + beta * y +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] trans The operation applied to the matrix A. +/// \param [in] num_rows Number of rows of the matrix A. +/// \param [in] num_cols Number of columns of the matrix A. +/// \param [in] alpha Scaling factor for the matrix A. +/// \param [in] info Matrix info of the matrix A. +/// \param [in] val An array containing the non-zero elements of the matrix A. +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [in] x Data of the vector x. +/// \param [in] beta Scaling factor for the vector x. +/// \param [in, out] y Data of the vector y. +template +[[deprecated("Use dpct::sparse::csrmv(queue, trans, num_rows, " + "num_cols, nnz, ...) instead.")]] +void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, int num_rows, + int num_cols, const T *alpha, + const std::shared_ptr info, const T *val, + const int *row_ptr, const int *col_ind, const T *x, const T *beta, + T *y) { + detail::csrmv_impl_deprecated()(queue, trans, num_rows, num_cols, alpha, + info, val, row_ptr, col_ind, x, beta, y); +} + /// Computes a CSR format sparse matrix-dense vector product. y = A * x /// /// \param [in] queue The queue where the routine should be executed. It must @@ -624,17 +662,47 @@ void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, int num_rows, /// \param [in, out] vector_y Data of the vector y. /// \param [in] num_rows Number of rows of the matrix A. /// \param [in] num_cols Number of columns of the matrix A. +/// \param [in] nnz NNZ of the matrix A. template void csrmv(sycl::queue &queue, const T *values, const int *row_offsets, const int *column_indices, const T *vector_x, T *vector_y, - int num_rows, int num_cols) { + int num_rows, int num_cols, int nnz) { T alpha{1}, beta{0}; auto matrix_info = std::make_shared(); matrix_info->set_index_base(oneapi::mkl::index_base::zero); matrix_info->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); detail::csrmv_impl()(queue, oneapi::mkl::transpose::nontrans, num_rows, num_cols, &alpha, matrix_info, values, row_offsets, - column_indices, vector_x, &beta, vector_y); + column_indices, vector_x, &beta, vector_y, nnz); +} + +/// Computes a CSR format sparse matrix-dense vector product. y = A * x +/// +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] values An array containing the non-zero elements of the matrix. +/// \param [in] row_offsets An array of length \p num_rows + 1. +/// \param [in] column_indices An array containing the column indices in +/// index-based numbering. +/// \param [in] vector_x Data of the vector x. +/// \param [in, out] vector_y Data of the vector y. +/// \param [in] num_rows Number of rows of the matrix A. +/// \param [in] num_cols Number of columns of the matrix A. +template +[[deprecated( + "Use dpct::sparse::csrmv(queue, values, row_offsets, " + "column_indices, vector_x, vector_y, num_rows, num_cols, nnz) instead.")]] +void csrmv(sycl::queue &queue, const T *values, const int *row_offsets, + const int *column_indices, const T *vector_x, T *vector_y, + int num_rows, int num_cols) { + T alpha{1}, beta{0}; + auto matrix_info = std::make_shared(); + matrix_info->set_index_base(oneapi::mkl::index_base::zero); + matrix_info->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + detail::csrmv_impl_deprecated()(queue, oneapi::mkl::transpose::nontrans, + num_rows, num_cols, &alpha, matrix_info, + values, row_offsets, column_indices, + vector_x, &beta, vector_y); } /// Computes a CSR format sparse matrix-dense vector product. @@ -644,6 +712,7 @@ void csrmv(sycl::queue &queue, const T *values, const int *row_offsets, /// \param [in] trans The operation applied to the matrix A. /// \param [in] num_rows Number of rows of the matrix A. /// \param [in] num_cols Number of columns of the matrix A. +/// \param [in] nnz NNZ of the matrix A. /// \param [in] alpha Scaling factor for the matrix A. /// \param [in] alpha_type Data type of \p alpha . /// \param [in] info Matrix info of the matrix A. @@ -659,7 +728,7 @@ void csrmv(sycl::queue &queue, const T *values, const int *row_offsets, /// \param [in, out] y Data of the vector y. /// \param [in] y_type Data type of \p y . inline void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, - int num_rows, int num_cols, const void *alpha, + int num_rows, int num_cols, int nnz, const void *alpha, library_data_t alpha_type, const std::shared_ptr info, const void *val, library_data_t val_type, const int *row_ptr, @@ -671,6 +740,120 @@ inline void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, col_ind, x, beta, y); } +/// Computes a CSR format sparse matrix-dense vector product. +/// y = alpha * op(A) * x + beta * y +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] trans The operation applied to the matrix A. +/// \param [in] num_rows Number of rows of the matrix A. +/// \param [in] num_cols Number of columns of the matrix A. +/// \param [in] alpha Scaling factor for the matrix A. +/// \param [in] alpha_type Data type of \p alpha . +/// \param [in] info Matrix info of the matrix A. +/// \param [in] val An array containing the non-zero elements of the matrix A. +/// \param [in] val_type Data type of \p val . +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [in] x Data of the vector x. +/// \param [in] x_type Data type of \p x . +/// \param [in] beta Scaling factor for the vector x. +/// \param [in] beta_type Data type of \p beta . +/// \param [in, out] y Data of the vector y. +/// \param [in] y_type Data type of \p y . +[[deprecated("Use dpct::sparse::csrmv(queue, trans, num_rows, " + "num_cols, nnz, ...) instead.")]] +inline void +csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, int num_rows, + int num_cols, const void *alpha, library_data_t alpha_type, + const std::shared_ptr info, const void *val, + library_data_t val_type, const int *row_ptr, const int *col_ind, + const void *x, library_data_t x_type, const void *beta, + library_data_t beta_type, void *y, library_data_t y_type) { + detail::spblas_shim( + val_type, queue, trans, num_rows, num_cols, alpha, info, val, row_ptr, + col_ind, x, beta, y); +} + +/// Computes a CSR format sparse matrix-dense matrix product. +/// C = alpha * op(A) * op(B) + beta * C +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] trans_a The operation applied to the matrix A. +/// \param [in] trans_b The operation applied to the matrix B. +/// \param [in] sparse_rows Number of rows of the matrix A. +/// \param [in] dense_cols Number of columns of the matrix B or C. +/// \param [in] sparse_cols Number of columns of the matrix A. +/// \param [in] nnz NNZ of the matrix A. +/// \param [in] alpha Scaling factor for the matrix A. +/// \param [in] info Matrix info of the matrix A. +/// \param [in] val An array containing the non-zero elements of the matrix A. +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [in] b Data of the matrix B. +/// \param [in] ldb Leading dimension of the matrix B. +/// \param [in] beta Scaling factor for the matrix B. +/// \param [in, out] c Data of the matrix C. +/// \param [in] ldc Leading dimension of the matrix C. +template +void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans_a, + oneapi::mkl::transpose trans_b, int sparse_rows, int dense_cols, + int sparse_cols, int nnz, const T *alpha, + const std::shared_ptr info, const T *val, + const int *row_ptr, const int *col_ind, const T *b, int ldb, + const T *beta, T *c, int ldc) { +#ifndef __INTEL_MKL__ + throw std::runtime_error("The oneAPI Math Kernel Library (oneMKL) Interfaces " + "Project does not support this API."); +#else + using Ty = typename ::dpct::detail::lib_data_traits_t; + auto alpha_value = + dpct::detail::get_value(reinterpret_cast(alpha), queue); + auto beta_value = + dpct::detail::get_value(reinterpret_cast(beta), queue); + + oneapi::mkl::sparse::matrix_handle_t *sparse_matrix_handle = + new oneapi::mkl::sparse::matrix_handle_t; + oneapi::mkl::sparse::init_matrix_handle(sparse_matrix_handle); + auto data_row_ptr = dpct::detail::get_memory(row_ptr); + auto data_col_ind = dpct::detail::get_memory(col_ind); + auto data_val = dpct::detail::get_memory(val); + oneapi::mkl::sparse::set_csr_data(queue, *sparse_matrix_handle, sparse_rows, + sparse_cols, nnz, info->get_index_base(), + data_row_ptr, data_col_ind, data_val); + + auto data_b = dpct::detail::get_memory(b); + auto data_c = dpct::detail::get_memory(c); + sycl::event gemm_event; + switch (info->get_matrix_type()) { + case matrix_info::matrix_type::ge: { +#ifndef DPCT_USM_LEVEL_NONE + gemm_event = +#endif + oneapi::mkl::sparse::gemm(queue, oneapi::mkl::layout::col_major, + trans_a, trans_b, alpha_value, + *sparse_matrix_handle, data_b, dense_cols, + ldb, beta_value, data_c, ldc); + break; + } + default: + throw std::runtime_error( + "the csrmm does not support matrix_info::matrix_type::sy, " + "matrix_info::matrix_type::tr and matrix_info::matrix_type::he"); + } +#ifdef DPCT_USM_LEVEL_NONE + queue.wait(); +#endif + sycl::event e = oneapi::mkl::sparse::release_matrix_handle( + queue, sparse_matrix_handle, {gemm_event}); + queue.submit([&](sycl::handler &cgh) { + cgh.depends_on(e); + cgh.host_task([=] { delete sparse_matrix_handle; }); + }); +#endif +} + /// Computes a CSR format sparse matrix-dense matrix product. /// C = alpha * op(A) * op(B) + beta * C /// \param [in] queue The queue where the routine should be executed. It must @@ -692,6 +875,8 @@ inline void csrmv(sycl::queue &queue, oneapi::mkl::transpose trans, /// \param [in, out] c Data of the matrix C. /// \param [in] ldc Leading dimension of the matrix C. template +[[deprecated("Use dpct::sparse::csrmm(queue, trans_a, trans_b, sparse_rows, " + "dense_cols, sparse_cols, nnz, ...) instead.")]] void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans_a, oneapi::mkl::transpose trans_b, int sparse_rows, int dense_cols, int sparse_cols, const T *alpha, @@ -749,6 +934,37 @@ void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans_a, #endif } +/// Computes a CSR format sparse matrix-dense matrix product. +/// C = alpha * op(A) * B + beta * C +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] trans The operation applied to the matrix A. +/// \param [in] sparse_rows Number of rows of the matrix A. +/// \param [in] dense_cols Number of columns of the matrix op(B) or C. +/// \param [in] sparse_cols Number of columns of the matrix A. +/// \param [in] nnz NNZ of the matrix A. +/// \param [in] alpha Scaling factor for the matrix A. +/// \param [in] info Matrix info of the matrix A. +/// \param [in] val An array containing the non-zero elements of the matrix A. +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [in] b Data of the matrix B. +/// \param [in] ldb Leading dimension of the matrix B. +/// \param [in] beta Scaling factor for the matrix B. +/// \param [in, out] c Data of the matrix C. +/// \param [in] ldc Leading dimension of the matrix C. +template +void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans, int sparse_rows, + int dense_cols, int sparse_cols, int nnz, const T *alpha, + const std::shared_ptr info, const T *val, + const int *row_ptr, const int *col_ind, const T *b, int ldb, + const T *beta, T *c, int ldc) { + csrmm(queue, trans, oneapi::mkl::transpose::nontrans, sparse_rows, + dense_cols, sparse_cols, nnz, alpha, info, val, row_ptr, col_ind, b, + ldb, beta, c, ldc); +} + /// Computes a CSR format sparse matrix-dense matrix product. /// C = alpha * op(A) * B + beta * C /// \param [in] queue The queue where the routine should be executed. It must @@ -769,6 +985,8 @@ void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans_a, /// \param [in, out] c Data of the matrix C. /// \param [in] ldc Leading dimension of the matrix C. template +[[deprecated("Use dpct::sparse::csrmm(queue, trans, sparse_rows, " + "dense_cols, sparse_cols, nnz, ...) instead.")]] void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans, int sparse_rows, int dense_cols, int sparse_cols, const T *alpha, const std::shared_ptr info, const T *val, @@ -786,6 +1004,7 @@ void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans, int sparse_rows, /// have the in_order property when using the USM mode. /// \param [in] trans The operation applied to the sparse matrix. /// \param [in] row_col Number of rows of the sparse matrix. +/// \param [in] nnz NNZ of the sparse matrix. /// \param [in] info Matrix info of the sparse matrix. /// \param [in] val An array containing the non-zero elements of the sparse matrix. /// \param [in] row_ptr An array of length \p num_rows + 1. @@ -793,20 +1012,56 @@ void csrmm(sycl::queue &queue, oneapi::mkl::transpose trans, int sparse_rows, /// numbering. /// \param [out] optimize_info The result of the optimizations. template +void optimize_csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, + int row_col, int nnz, + const std::shared_ptr info, const T *val, + const int *row_ptr, const int *col_ind, + std::shared_ptr optimize_info) { + detail::optimize_csrsv_impl()(queue, trans, row_col, nnz, info, val, + row_ptr, col_ind, optimize_info); +} + +/// Performs internal optimizations for solving a system of linear equations for +/// a CSR format sparse matrix. +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] trans The operation applied to the sparse matrix. +/// \param [in] row_col Number of rows of the sparse matrix. +/// \param [in] info Matrix info of the sparse matrix. +/// \param [in] val An array containing the non-zero elements of the sparse matrix. +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [out] optimize_info The result of the optimizations. +template +[[deprecated("Use oneapi::mkl::sparse::optimize_csrsv(queue, trans, " + "row_col, nnz, ...) instead.")]] void optimize_csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, const std::shared_ptr info, const T *val, const int *row_ptr, const int *col_ind, std::shared_ptr optimize_info) { - detail::optimize_csrsv_impl()(queue, trans, row_col, info, val, row_ptr, - col_ind, optimize_info); + detail::optimize_csrsv_impl_deprecated()(queue, trans, row_col, info, val, + row_ptr, col_ind, optimize_info); } inline void optimize_csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, - int row_col, const std::shared_ptr info, + int row_col, int nnz, const std::shared_ptr info, const void *val, library_data_t val_type, const int *row_ptr, const int *col_ind, std::shared_ptr optimize_info) { detail::spblas_shim( + val_type, queue, trans, row_col, nnz, info, val, row_ptr, col_ind, + optimize_info); +} + +[[deprecated("Use oneapi::mkl::sparse::optimize_csrsv(queue, trans, " + "row_col, nnz, ...) instead.")]] +inline void optimize_csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, + int row_col, const std::shared_ptr info, + const void *val, library_data_t val_type, + const int *row_ptr, const int *col_ind, + std::shared_ptr optimize_info) { + detail::spblas_shim( val_type, queue, trans, row_col, info, val, row_ptr, col_ind, optimize_info); } @@ -841,6 +1096,7 @@ inline void csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, /// \param [in] transb The operation applied to B and X. /// \param [in] row_col Number of rows and columns of A. /// \param [in] nrhs Number of columns op_b(B). +/// \param [in] nnz NNZ of the matrix A. /// \param [in] info Matrix info of A. /// \param [in] val An array containing the non-zero elements of A. /// \param [in] row_ptr An array of length \p num_rows + 1. @@ -849,14 +1105,42 @@ inline void csrsv(sycl::queue &queue, oneapi::mkl::transpose trans, int row_col, /// \param [out] optimize_info The result of the optimizations. template void optimize_csrsm(sycl::queue &queue, oneapi::mkl::transpose transa, - oneapi::mkl::transpose transb, int row_col, int nrhs, + oneapi::mkl::transpose transb, int row_col, int nrhs, int nnz, const std::shared_ptr info, const T *val, const int *row_ptr, const int *col_ind, std::shared_ptr optimize_info) { - detail::optimize_csrsm_impl()(queue, transa, transb, row_col, nrhs, info, + detail::optimize_csrsm_impl()(queue, transa, transb, row_col, nrhs, nnz, info, val, row_ptr, col_ind, optimize_info); } +/// Performs internal optimizations for dpct::sparse::csrsm by analyzing +/// the provided matrix structure and operation parameters. The matrix A must be +/// a triangular sparse matrix with the CSR format. +/// \param [in] queue The queue where the routine should be executed. It must +/// have the in_order property when using the USM mode. +/// \param [in] transa The operation applied to A. +/// \param [in] transb The operation applied to B and X. +/// \param [in] row_col Number of rows and columns of A. +/// \param [in] nrhs Number of columns op_b(B). +/// \param [in] info Matrix info of A. +/// \param [in] val An array containing the non-zero elements of A. +/// \param [in] row_ptr An array of length \p num_rows + 1. +/// \param [in] col_ind An array containing the column indices in index-based +/// numbering. +/// \param [out] optimize_info The result of the optimizations. +template +[[deprecated("Use oneapi::mkl::sparse::optimize_csrsm(queue, transa, transb, " + "row_col, nrhs, nnz, ...) instead.")]] +void optimize_csrsm(sycl::queue &queue, oneapi::mkl::transpose transa, + oneapi::mkl::transpose transb, int row_col, int nrhs, + const std::shared_ptr info, const T *val, + const int *row_ptr, const int *col_ind, + std::shared_ptr optimize_info) { + detail::optimize_csrsm_impl_deprecated()(queue, transa, transb, row_col, + nrhs, info, val, row_ptr, col_ind, + optimize_info); +} + /// Solves the sparse triangular system op_a(A) * op_b(X) = alpha * op_b(B). /// A is a sparse triangular matrix with the CSR format of size \p row_col /// by \p row_col . @@ -1298,14 +1582,14 @@ void csrgemm_nnz(descriptor_ptr desc, oneapi::mkl::transpose trans_a, int cols_b = (trans_b == oneapi::mkl::transpose::nontrans) ? n : k; info.matrix_handle_a.set_matrix_data( - rows_a, cols_a, info_a->get_index_base(), row_ptr_a, col_ind_a, val_a); + rows_a, cols_a, nnz_a, info_a->get_index_base(), row_ptr_a, col_ind_a, val_a); info.matrix_handle_b.set_matrix_data( - rows_b, cols_b, info_b->get_index_base(), row_ptr_b, col_ind_b, val_b); + rows_b, cols_b, nnz_b, info_b->get_index_base(), row_ptr_b, col_ind_b, val_b); // In the future, oneMKL will allow nullptr to be passed in for row_ptr_c in // the initial calls before matmat. But currently, it needs an array of // length row_number + 1. info.matrix_handle_c.set_matrix_data( - rows_c, cols_c, info_c->get_index_base(), row_ptr_c, nullptr, nullptr); + rows_c, cols_c, 0, info_c->get_index_base(), row_ptr_c, nullptr, nullptr); oneapi::mkl::sparse::set_matmat_data( info.matmat_desc.get_handle(), @@ -1345,6 +1629,7 @@ void csrgemm_nnz(descriptor_ptr desc, oneapi::mkl::transpose trans_a, nnz_c_int = *nnz_c; #endif + info_c->set_nnz(nnz_c_int); if (nnz_host_ptr) { *nnz_host_ptr = nnz_c_int; } @@ -1396,7 +1681,7 @@ void csrgemm(descriptor_ptr desc, oneapi::mkl::transpose trans_a, int rows_c = m; int cols_c = n; info.matrix_handle_c.set_matrix_data( - rows_c, cols_c, info_c->get_index_base(), row_ptr_c, col_ind_c, val_c); + rows_c, cols_c, info_c->get_nnz(), info_c->get_index_base(), row_ptr_c, col_ind_c, val_c); sycl::event e; #ifndef DPCT_USM_LEVEL_NONE e = @@ -1424,6 +1709,7 @@ class csrgemm2_info { void *row_ptr_axb = nullptr; void *col_ind_axb = nullptr; void *val_axb = nullptr; + int c_nnz = 0; #ifdef DPCT_USM_LEVEL_NONE sycl::buffer *temp_buffer_1_size; sycl::buffer *temp_buffer_2_size; @@ -1555,15 +1841,15 @@ void csrgemm2_get_buffer_size( info->row_ptr_axb = (int *)::dpct::cs::malloc((m + 1) * sizeof(int), queue); info->init(&queue); - info->matrix_handle_a.set_matrix_data(m, k, info_a->get_index_base(), + info->matrix_handle_a.set_matrix_data(m, k, nnz_a, info_a->get_index_base(), row_ptr_a, col_ind_a, nullptr); - info->matrix_handle_b.set_matrix_data(k, n, info_b->get_index_base(), + info->matrix_handle_b.set_matrix_data(k, n, nnz_b, info_b->get_index_base(), row_ptr_b, col_ind_b, nullptr); info->matrix_handle_axb.set_matrix_data( - m, n, oneapi::mkl::index_base::zero, info->row_ptr_axb, nullptr, nullptr); - info->matrix_handle_d.set_matrix_data(m, n, info_d->get_index_base(), + m, n, 0, oneapi::mkl::index_base::zero, info->row_ptr_axb, nullptr, nullptr); + info->matrix_handle_d.set_matrix_data(m, n, nnz_d, info_d->get_index_base(), row_ptr_d, col_ind_d, nullptr); - info->matrix_handle_c.set_matrix_data(m, n, oneapi::mkl::index_base::zero, + info->matrix_handle_c.set_matrix_data(m, n, 0, oneapi::mkl::index_base::zero, nullptr, nullptr, nullptr); oneapi::mkl::sparse::set_matmat_data( @@ -1628,7 +1914,7 @@ void csrgemm2_get_buffer_size( info->col_ind_axb = (int *)::dpct::cs::malloc(nnz_axb * sizeof(int), queue); info->val_axb = (Ty *)::dpct::cs::malloc(nnz_axb * sizeof(Ty), queue); info->matrix_handle_axb.set_matrix_data( - m, n, oneapi::mkl::index_base::zero, info->row_ptr_axb, info->col_ind_axb, + m, n, nnz_axb, oneapi::mkl::index_base::zero, info->row_ptr_axb, info->col_ind_axb, info->val_axb); __MATMAT(finalize_structure, nullptr, nullptr); @@ -1682,19 +1968,19 @@ inline void csrgemm2_nnz(descriptor_ptr desc, int m, int n, int k, sycl::queue &queue = desc->get_queue(); if (info->matrix_c_datatype == csrgemm2_info::matrix_c_datatype_t::mcd_float) { - info->matrix_handle_c.set_matrix_data(m, n, info_c->get_index_base(), + info->matrix_handle_c.set_matrix_data(m, n, 0, info_c->get_index_base(), nullptr, nullptr, nullptr); } else if (info->matrix_c_datatype == csrgemm2_info::matrix_c_datatype_t::mcd_double) { info->matrix_handle_c.set_matrix_data( - m, n, info_c->get_index_base(), nullptr, nullptr, nullptr); + m, n, 0, info_c->get_index_base(), nullptr, nullptr, nullptr); } else if (info->matrix_c_datatype == csrgemm2_info::matrix_c_datatype_t::mcd_float2) { info->matrix_handle_c.set_matrix_data>( - m, n, info_c->get_index_base(), nullptr, nullptr, nullptr); + m, n, 0, info_c->get_index_base(), nullptr, nullptr, nullptr); } else { info->matrix_handle_c.set_matrix_data>( - m, n, info_c->get_index_base(), nullptr, nullptr, nullptr); + m, n, 0, info_c->get_index_base(), nullptr, nullptr, nullptr); } auto data_buffer = dpct::detail::get_memory(buffer); oneapi::mkl::sparse::omatadd_analyze( @@ -1719,6 +2005,7 @@ inline void csrgemm2_nnz(descriptor_ptr desc, int m, int n, int k, assert((nnz_c >= INT_MIN && nnz_c <= INT_MAX) && "nnz_c is out of range."); int nnz_c_int = nnz_c; + info->nnz_c = nnz_c_int; if (nnz_ptr) ::dpct::cs::memcpy(queue, nnz_ptr, &nnz_c_int, sizeof(int)).wait(); int row_ptr_c_0 = @@ -1780,16 +2067,16 @@ void csrgemm2(descriptor_ptr desc, int m, int n, int k, const T *alpha, auto beta_value = dpct::detail::get_value(reinterpret_cast(beta), queue); - info->matrix_handle_a.set_matrix_data(m, k, info_a->get_index_base(), + info->matrix_handle_a.set_matrix_data(m, k, nnz_a, info_a->get_index_base(), row_ptr_a, col_ind_a, val_a); - info->matrix_handle_b.set_matrix_data(k, n, info_b->get_index_base(), + info->matrix_handle_b.set_matrix_data(k, n, nnz_b, info_b->get_index_base(), row_ptr_b, col_ind_b, val_b); __MATMAT(compute, nullptr, nullptr); __MATMAT(finalize, nullptr, nullptr); - info->matrix_handle_d.set_matrix_data(m, n, info_d->get_index_base(), + info->matrix_handle_d.set_matrix_data(m, n, nnz_d, info_d->get_index_base(), row_ptr_d, col_ind_d, val_d); - info->matrix_handle_c.set_matrix_data(m, n, info_c->get_index_base(), + info->matrix_handle_c.set_matrix_data(m, n, info_c->nnz_c, info_c->get_index_base(), row_ptr_c, col_ind_c, val_c); sycl::event e; diff --git a/clang/test/dpct/asm/ld.cu b/clang/test/dpct/asm/ld.cu index d8d94365d138..86a0fbfc2774 100644 --- a/clang/test/dpct/asm/ld.cu +++ b/clang/test/dpct/asm/ld.cu @@ -7,6 +7,7 @@ // clang-format off #include #include +#include #include #include diff --git a/clang/test/dpct/cub/devicelevel/device_spmv.cu b/clang/test/dpct/cub/devicelevel/device_spmv.cu index 5d192a5756f5..3061f743574d 100644 --- a/clang/test/dpct/cub/devicelevel/device_spmv.cu +++ b/clang/test/dpct/cub/devicelevel/device_spmv.cu @@ -41,7 +41,7 @@ int main() { d_vector_y, num_rows, num_cols, num_nonzeros); cudaMalloc(&d_temp_storage, temp_storage_bytes); - // CHECK: dpct::sparse::csrmv(q_ct1, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols); + // CHECK: dpct::sparse::csrmv(q_ct1, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols, num_nonzeros); cub::DeviceSpmv::CsrMV(d_temp_storage, temp_storage_bytes, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols, num_nonzeros); @@ -60,7 +60,7 @@ int main() { d_vector_y, num_rows, num_cols, num_nonzeros, S); cudaMalloc(&d_temp_storage, temp_storage_bytes); - // CHECK: dpct::sparse::csrmv(*S, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols); + // CHECK: dpct::sparse::csrmv(*S, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols, num_nonzeros); cub::DeviceSpmv::CsrMV(d_temp_storage, temp_storage_bytes, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols, num_nonzeros, S); diff --git a/clang/test/dpct/cusparse-usm.cu b/clang/test/dpct/cusparse-usm.cu index 3ca5ffade2ff..b2c87eb8e97f 100644 --- a/clang/test/dpct/cusparse-usm.cu +++ b/clang/test/dpct/cusparse-usm.cu @@ -71,7 +71,7 @@ int foo(int aaaaa){ //CHECK-NEXT: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), (oneapi::mkl::transpose)aaaaa, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), (oneapi::mkl::transpose)aaaaa, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); cusparseCreate(&handle); cusparseCreateMatDescr(&descrA); cusparseSetMatType(descrA, (cusparseMatrixType_t)aaaaa); @@ -83,44 +83,44 @@ int foo(int aaaaa){ //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); cusparseZcsrmv(handle, transA, m, n, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); cusparseDcsrmv_mp(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); cusparseZcsrmv_mp(handle, transA, m, n, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, &beta_Z, y_Z); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); cusparseDcsrmm(handle, transA, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); cusparseZcsrmm(handle, transA, m, n, k, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); cusparseOperation_t transB = CUSPARSE_OPERATION_NON_TRANSPOSE; //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); cusparseDcsrmm2(handle, transA, transB, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); cusparseZcsrmm2(handle, transA, transB, m, n, k, nnz, &alpha_Z, descrA, csrValA_Z, csrRowPtrA, csrColIndA, x_Z, ldb, &beta_Z, y_Z, ldc); size_t ws_size; @@ -128,7 +128,7 @@ int foo(int aaaaa){ //CHECK: int alg = 0; //CHECK-NEXT: alg = 1; //CHECK-NEXT: ws_size = 0; - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, dpct::library_data_t::real_double, descrA, csrValA, dpct::library_data_t::real_double, csrRowPtrA, csrColIndA, x, dpct::library_data_t::real_double, &beta_value, dpct::library_data_t::real_double, y, dpct::library_data_t::real_double); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, dpct::library_data_t::real_double, descrA, csrValA, dpct::library_data_t::real_double, csrRowPtrA, csrColIndA, x, dpct::library_data_t::real_double, &beta_value, dpct::library_data_t::real_double, y, dpct::library_data_t::real_double); cusparseAlgMode_t alg = CUSPARSE_ALG0; alg = CUSPARSE_ALG1; cusparseCsrmvEx_bufferSize(handle, alg, transA, m, n, nnz, &alpha, CUDA_R_64F, descrA, csrValA, CUDA_R_64F, csrRowPtrA, csrColIndA, x, CUDA_R_64F, &beta_value, CUDA_R_64F, y, CUDA_R_64F, CUDA_R_64F, &ws_size); @@ -140,31 +140,31 @@ int foo(int aaaaa){ //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: if(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} + //CHECK-NEXT: if(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} if(status = cusparseDcsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)){} //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: for(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y));;){} + //CHECK-NEXT: for(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y));;){} for(status = cusparseDcsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y);;){} //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: switch(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} + //CHECK-NEXT: switch(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} switch(status = cusparseDcsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)){} //CHECK: std::shared_ptr info; //CHECK-NEXT: info = std::make_shared(); - //CHECK-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, descrA, csrValA, csrRowPtrA, csrColIndA, info); + //CHECK-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, nnz, descrA, csrValA, csrRowPtrA, csrColIndA, info); //CHECK-NEXT: info.reset(); cusparseSolveAnalysisInfo_t info; cusparseCreateSolveAnalysisInfo(&info); cusparseDcsrsv_analysis(handle, transA, m, nnz, descrA, csrValA, csrRowPtrA, csrColIndA, info); cusparseDestroySolveAnalysisInfo(info); - //CHECK: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, descrA, csrValA_Z, csrRowPtrA, csrColIndA, info); + //CHECK: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, nnz, descrA, csrValA_Z, csrRowPtrA, csrColIndA, info); cusparseZcsrsv_analysis(handle, transA, m, nnz, descrA, csrValA_Z, csrRowPtrA, csrColIndA, info); //CHECK: /* @@ -179,7 +179,7 @@ int foo(int aaaaa){ //CHECK-NEXT: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ -//CHECK-NEXT: return DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)); +//CHECK-NEXT: return DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)); //CHECK-NEXT: } int foo(cusparseMatDescr_t descrB){ return cusparseDcsrmv(handle, transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); diff --git a/clang/test/dpct/cusparse.cu b/clang/test/dpct/cusparse.cu index ebf9a2f83a50..3e9e86d3bcd4 100644 --- a/clang/test/dpct/cusparse.cu +++ b/clang/test/dpct/cusparse.cu @@ -80,7 +80,7 @@ int main(){ //CHECK-NEXT: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), (oneapi::mkl::transpose)zero, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), (oneapi::mkl::transpose)zero, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); cusparseCreateMatDescr(&descrA); cusparseSetMatType(descrA, CUSPARSE_MATRIX_TYPE_GENERAL); cusparseSetMatIndexBase(descrA, CUSPARSE_INDEX_BASE_ZERO); @@ -91,44 +91,44 @@ int main(){ //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); cusparseCcsrmv(handle, transA, m, n, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); cusparseScsrmv_mp(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); + //CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); cusparseCcsrmv_mp(handle, transA, m, n, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, &beta_C, y_C); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); cusparseScsrmm(handle, transA, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, m, n, k, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); cusparseCcsrmm(handle, transA, m, n, k, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); cusparseOperation_t transB = CUSPARSE_OPERATION_NON_TRANSPOSE; //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); cusparseScsrmm2(handle, transA, transB, m, n, k, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, ldb, &beta_value, y, ldc); //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); + //CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), transA, transB, m, n, k, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); cusparseCcsrmm2(handle, transA, transB, m, n, k, nnz, &alpha_C, descrA, csrValA_C, csrRowPtrA, csrColIndA, x_C, ldb, &beta_C, y_C, ldc); //CHECK:int status; @@ -137,31 +137,31 @@ int main(){ //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: if(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} + //CHECK-NEXT: if(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} if(status = cusparseScsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)){} //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: for(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y));;){} + //CHECK-NEXT: for(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y));;){} for(status = cusparseScsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y);;){} //CHECK: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ - //CHECK-NEXT: switch(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} + //CHECK-NEXT: switch(status = DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y))){} switch(status = cusparseScsrmv(handle, transA, m, n, nnz, &alpha, descrA, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)){} //CHECK: std::shared_ptr info; //CHECK-NEXT: info = std::make_shared(); - //CHECK-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, descrA, csrValA, csrRowPtrA, csrColIndA, info); + //CHECK-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, nnz, descrA, csrValA, csrRowPtrA, csrColIndA, info); //CHECK-NEXT: info.reset(); cusparseSolveAnalysisInfo_t info; cusparseCreateSolveAnalysisInfo(&info); cusparseScsrsv_analysis(handle, transA, m, nnz, descrA, csrValA, csrRowPtrA, csrColIndA, info); cusparseDestroySolveAnalysisInfo(info); - //CHECK: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, descrA, csrValA_C, csrRowPtrA, csrColIndA, info); + //CHECK: dpct::sparse::optimize_csrsv(handle->get_queue(), transA, m, nnz, descrA, csrValA_C, csrRowPtrA, csrColIndA, info); cusparseCcsrsv_analysis(handle, transA, m, nnz, descrA, csrValA_C, csrRowPtrA, csrColIndA, info); //CHECK: /* @@ -176,7 +176,7 @@ int main(){ //CHECK-NEXT: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ -//CHECK-NEXT: return DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)); +//CHECK-NEXT: return DPCT_CHECK_ERROR(dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y)); //CHECK-NEXT:} int foo(cusparseMatDescr_t descrB){ return cusparseScsrmv(handle, transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); @@ -186,7 +186,7 @@ int foo(cusparseMatDescr_t descrB){ //CHECK-NEXT: /* //CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general/symmetric/triangular sparse matrix type. You may need to adjust the code. //CHECK-NEXT: */ -//CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); +//CHECK-NEXT: dpct::sparse::csrmv(handle->get_queue(), transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); //CHECK-NEXT:} void foo2(cusparseMatDescr_t descrB){ cusparseScsrmv(handle, transA, m, n, nnz, &alpha, descrB, csrValA, csrRowPtrA, csrColIndA, x, &beta_value, y); @@ -227,19 +227,19 @@ void foo3(){ // CHECK: /* // CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. // CHECK-NEXT: */ - // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, alpha_s, descrA, val_a_s, row_ptr_a, col_ind_a, val_b_s, 5, beta_s, val_c_s, 4); + // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, 9, alpha_s, descrA, val_a_s, row_ptr_a, col_ind_a, val_b_s, 5, beta_s, val_c_s, 4); // CHECK-NEXT: /* // CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. // CHECK-NEXT: */ - // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, alpha_d, descrA, val_a_d, row_ptr_a, col_ind_a, val_b_d, 5, beta_d, val_c_d, 4); + // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, 9, alpha_d, descrA, val_a_d, row_ptr_a, col_ind_a, val_b_d, 5, beta_d, val_c_d, 4); // CHECK-NEXT: /* // CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. // CHECK-NEXT: */ - // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, alpha_c, descrA, val_a_c, row_ptr_a, col_ind_a, val_b_c, 5, beta_c, val_c_c, 4); + // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, 9, alpha_c, descrA, val_a_c, row_ptr_a, col_ind_a, val_b_c, 5, beta_c, val_c_c, 4); // CHECK-NEXT: /* // CHECK-NEXT: DPCT1045:{{[0-9]+}}: Migration is only supported for this API for the general sparse matrix type. You may need to adjust the code. // CHECK-NEXT: */ - // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, alpha_z, descrA, val_a_z, row_ptr_a, col_ind_a, val_b_z, 5, beta_z, val_c_z, 4); + // CHECK-NEXT: dpct::sparse::csrmm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::nontrans, 4, 2, 5, 9, alpha_z, descrA, val_a_z, row_ptr_a, col_ind_a, val_b_z, 5, beta_z, val_c_z, 4); cusparseScsrmm2(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, 4, 2, 5, 9, alpha_s, descrA, val_a_s, row_ptr_a, col_ind_a, val_b_s, 5, beta_s, val_c_s, 4); cusparseDcsrmm2(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, 4, 2, 5, 9, alpha_d, descrA, val_a_d, row_ptr_a, col_ind_a, val_b_d, 5, beta_d, val_c_d, 4); cusparseCcsrmm2(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, 4, 2, 5, 9, alpha_c, descrA, val_a_c, row_ptr_a, col_ind_a, val_b_c, 5, beta_c, val_c_c, 4); @@ -331,10 +331,10 @@ void foo4() { cudaMalloc(&buffer_c, buffer_size_c); cudaMalloc(&buffer_z, buffer_size_z); - //CHECK:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, descrA, a_s_val, a_row_ptr, a_col_ind, info2); - //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, descrA, a_d_val, a_row_ptr, a_col_ind, info2); - //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, descrA, a_c_val, a_row_ptr, a_col_ind, info2); - //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, descrA, a_z_val, a_row_ptr, a_col_ind, info2); + //CHECK:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, 6, descrA, a_s_val, a_row_ptr, a_col_ind, info2); + //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, 6, descrA, a_d_val, a_row_ptr, a_col_ind, info2); + //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, 6, descrA, a_c_val, a_row_ptr, a_col_ind, info2); + //CHECK-NEXT:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, 6, descrA, a_z_val, a_row_ptr, a_col_ind, info2); //CHECK-NEXT:dpct::sparse::csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, &alpha_s, descrA, a_s_val, a_row_ptr, a_col_ind, info2, f_s, x_s); //CHECK-NEXT:dpct::sparse::csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, &alpha_d, descrA, a_d_val, a_row_ptr, a_col_ind, info2, f_d, x_d); //CHECK-NEXT:dpct::sparse::csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, &alpha_c, descrA, a_c_val, a_row_ptr, a_col_ind, info2, f_c, x_c); @@ -364,7 +364,7 @@ void foo5() { float *x_s; float alpha_s = 1; - //CHECK:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, descrA, a_s_val, dpct::library_data_t::real_float, a_row_ptr, a_col_ind, info); + //CHECK:dpct::sparse::optimize_csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, 6, descrA, a_s_val, dpct::library_data_t::real_float, a_row_ptr, a_col_ind, info); //CHECK-NEXT:dpct::sparse::csrsv(handle->get_queue(), oneapi::mkl::transpose::nontrans, 3, &alpha_s, dpct::library_data_t::real_float, descrA, a_s_val, dpct::library_data_t::real_float, a_row_ptr, a_col_ind, info, f_s, dpct::library_data_t::real_float, x_s, dpct::library_data_t::real_float); cusparseCsrsv_analysisEx(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, 3, 6, descrA, a_s_val, CUDA_R_32F, a_row_ptr, a_col_ind, info, CUDA_R_32F); cusparseCsrsv_solveEx(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, 3, &alpha_s, CUDA_R_32F, descrA, a_s_val, CUDA_R_32F, a_row_ptr, a_col_ind, info, f_s, CUDA_R_32F, x_s, CUDA_R_32F, CUDA_R_32F); diff --git a/clang/test/dpct/cusparse_9_to_11.cu b/clang/test/dpct/cusparse_9_to_11.cu index 21e75b5b1632..b2a7db4d669f 100644 --- a/clang/test/dpct/cusparse_9_to_11.cu +++ b/clang/test/dpct/cusparse_9_to_11.cu @@ -45,10 +45,10 @@ void foo() { cusparseCcsrsm2_bufferSizeExt(handle, 0, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_TRANSPOSE, m, nrhs, nnz, &alpha_c, descrA, val_c, row_ptr, col_ind, b_c, nrhs, info, policy, &buffer_size); cusparseZcsrsm2_bufferSizeExt(handle, 0, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_TRANSPOSE, m, nrhs, nnz, &alpha_z, descrA, val_z, row_ptr, col_ind, b_z, nrhs, info, policy, &buffer_size); - // CHECK: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, descrA, val_s, row_ptr, col_ind, info); - // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, descrA, val_d, row_ptr, col_ind, info); - // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, descrA, val_c, row_ptr, col_ind, info); - // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, descrA, val_z, row_ptr, col_ind, info); + // CHECK: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, nnz, descrA, val_s, row_ptr, col_ind, info); + // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, nnz, descrA, val_d, row_ptr, col_ind, info); + // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, nnz, descrA, val_c, row_ptr, col_ind, info); + // CHECK-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), oneapi::mkl::transpose::nontrans, oneapi::mkl::transpose::trans, m, nrhs, nnz, descrA, val_z, row_ptr, col_ind, info); cusparseScsrsm2_analysis(handle, 0, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_TRANSPOSE, m, nrhs, nnz, &alpha_s, descrA, val_s, row_ptr, col_ind, b_s, nrhs, info, policy, buffer); cusparseDcsrsm2_analysis(handle, 0, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_TRANSPOSE, m, nrhs, nnz, &alpha_d, descrA, val_d, row_ptr, col_ind, b_d, nrhs, info, policy, buffer); cusparseCcsrsm2_analysis(handle, 0, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_TRANSPOSE, m, nrhs, nnz, &alpha_c, descrA, val_c, row_ptr, col_ind, b_c, nrhs, info, policy, buffer); diff --git a/clang/test/dpct/help_option_check/lin/help_all.txt b/clang/test/dpct/help_option_check/lin/help_all.txt index 22186a896daa..e76c51ba0366 100644 --- a/clang/test/dpct/help_option_check/lin/help_all.txt +++ b/clang/test/dpct/help_option_check/lin/help_all.txt @@ -180,8 +180,8 @@ All DPCT options =sycl - Generate code with sycl:: namespace. Cannot be used with cl or sycl-math values. =sycl-math - Generate code with sycl:: namespace, applied only for SYCL math functions. Cannot be used with cl or sycl values. - =syclcompat - Generate code with syclcompat:: namespace. - --use-syclcompat - Use SYCLcompat header-only library (syclcompat:: namespace) to assist the migration of input source code. + =syclcompat - DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Generate code with syclcompat:: namespace. + --use-syclcompat - DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Use SYCLcompat header-only library (syclcompat:: namespace) to assist the migration of input source code. Default: off. --usm-level= - Set the Unified Shared Memory (USM) level to use in source code generation. =none - Uses helper functions from DPCT header files for memory management migration. diff --git a/clang/test/dpct/help_option_check/win/help_all.txt b/clang/test/dpct/help_option_check/win/help_all.txt index 11be80121350..27cd182e751a 100644 --- a/clang/test/dpct/help_option_check/win/help_all.txt +++ b/clang/test/dpct/help_option_check/win/help_all.txt @@ -179,8 +179,8 @@ All DPCT options =sycl - Generate code with sycl:: namespace. Cannot be used with cl or sycl-math values. =sycl-math - Generate code with sycl:: namespace, applied only for SYCL math functions. Cannot be used with cl or sycl values. - =syclcompat - Generate code with syclcompat:: namespace. - --use-syclcompat - Use SYCLcompat header-only library (syclcompat:: namespace) to assist the migration of input source code. + =syclcompat - DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Generate code with syclcompat:: namespace. + --use-syclcompat - DEPRECATED (Intel(R) DPC++ Compiler 2025.3 is deprecating SYCLCompat and will remove it in next release). Use SYCLcompat header-only library (syclcompat:: namespace) to assist the migration of input source code. Default: off. --usm-level= - Set the Unified Shared Memory (USM) level to use in source code generation. =none - Uses helper functions from DPCT header files for memory management migration. diff --git a/clang/test/dpct/query_api_mapping/CUB/cub_device.cu b/clang/test/dpct/query_api_mapping/CUB/cub_device.cu index 0c475ccf238b..032689079d34 100644 --- a/clang/test/dpct/query_api_mapping/CUB/cub_device.cu +++ b/clang/test/dpct/query_api_mapping/CUB/cub_device.cu @@ -526,4 +526,4 @@ // CHECK_DEVICESPMV_CSRMV: cudaMalloc(&d_temp_storage, temp_storage_bytes); // CHECK_DEVICESPMV_CSRMV: cub::DeviceSpmv::CsrMV(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_values/*float **/, d_row_offsets/*int **/, d_column_indices/*int **/, d_vector_x/*float **/, d_vector_y/*float **/, num_rows/*int*/, num_cols/*int*/, num_nonzeros/*int*/); // CHECK_DEVICESPMV_CSRMV: Is migrated to: -// CHECK_DEVICESPMV_CSRMV: dpct::sparse::csrmv(q_ct1, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols); +// CHECK_DEVICESPMV_CSRMV: dpct::sparse::csrmv(q_ct1, d_values, d_row_offsets, d_column_indices, d_vector_x, d_vector_y, num_rows, num_cols, num_nonzeros); diff --git a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_112-118.cu b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_112-118.cu index e234dd0fe80c..0d2ee3a582ce 100644 --- a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_112-118.cu +++ b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_112-118.cu @@ -75,7 +75,7 @@ // cusparseCsrmvEx-NEXT: beta /*const void **/, beta_type /*cudaDataType*/, y /*void **/, // cusparseCsrmvEx-NEXT: y_type /*cudaDataType*/, exec_type /*cudaDataType*/, buffer /*void **/); // cusparseCsrmvEx-NEXT: Is migrated to: -// cusparseCsrmvEx-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, alpha_type, desc, value, value_type, row_ptr, col_idx, x, x_type, beta, beta_type, y, y_type); +// cusparseCsrmvEx-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, alpha_type, desc, value, value_type, row_ptr, col_idx, x, x_type, beta, beta_type, y, y_type); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCsrmvEx_bufferSize | FileCheck %s -check-prefix=cusparseCsrmvEx_bufferSize // cusparseCsrmvEx_bufferSize: CUDA API: @@ -150,7 +150,7 @@ // cusparseScsrsv2_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, info /*csrsv2Info_t*/, // cusparseScsrsv2_analysis-NEXT: policy /*cusparseSolvePolicy_t*/, buffer /*void **/); // cusparseScsrsv2_analysis-NEXT: Is migrated to: -// cusparseScsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseScsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrsv2_analysis | FileCheck %s -check-prefix=cusparseDcsrsv2_analysis // cusparseDcsrsv2_analysis: CUDA API: @@ -160,7 +160,7 @@ // cusparseDcsrsv2_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, info /*csrsv2Info_t*/, // cusparseDcsrsv2_analysis-NEXT: policy /*cusparseSolvePolicy_t*/, buffer /*void **/); // cusparseDcsrsv2_analysis-NEXT: Is migrated to: -// cusparseDcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseDcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrsv2_analysis | FileCheck %s -check-prefix=cusparseCcsrsv2_analysis // cusparseCcsrsv2_analysis: CUDA API: @@ -170,7 +170,7 @@ // cusparseCcsrsv2_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, info /*csrsv2Info_t*/, // cusparseCcsrsv2_analysis-NEXT: policy /*cusparseSolvePolicy_t*/, buffer /*void **/); // cusparseCcsrsv2_analysis-NEXT: Is migrated to: -// cusparseCcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseCcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrsv2_analysis | FileCheck %s -check-prefix=cusparseZcsrsv2_analysis // cusparseZcsrsv2_analysis: CUDA API: @@ -181,4 +181,4 @@ // cusparseZcsrsv2_analysis-NEXT: col_idx /*const int **/, info /*csrsv2Info_t*/, // cusparseZcsrsv2_analysis-NEXT: policy /*cusparseSolvePolicy_t*/, buffer /*void **/); // cusparseZcsrsv2_analysis-NEXT: Is migrated to: -// cusparseZcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseZcsrsv2_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); diff --git a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-102.cu b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-102.cu index 2cd75f2317d5..4740ff8c830f 100644 --- a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-102.cu +++ b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-102.cu @@ -23,7 +23,7 @@ // cusparseScsrmv-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseScsrmv-NEXT: x /*const float **/, beta /*const float **/, y /*float **/); // cusparseScsrmv-NEXT: Is migrated to: -// cusparseScsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseScsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrmv | FileCheck %s -check-prefix=cusparseDcsrmv // cusparseDcsrmv: CUDA API: @@ -33,7 +33,7 @@ // cusparseDcsrmv-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseDcsrmv-NEXT: x /*const double **/, beta /*const double **/, y /*double **/); // cusparseDcsrmv-NEXT: Is migrated to: -// cusparseDcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseDcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrmv | FileCheck %s -check-prefix=cusparseCcsrmv // cusparseCcsrmv: CUDA API: @@ -44,7 +44,7 @@ // cusparseCcsrmv-NEXT: x /*const cuComplex **/, beta /*const cuComplex **/, // cusparseCcsrmv-NEXT: y /*cuComplex **/); // cusparseCcsrmv-NEXT: Is migrated to: -// cusparseCcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseCcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrmv | FileCheck %s -check-prefix=cusparseZcsrmv // cusparseZcsrmv: CUDA API: @@ -55,7 +55,7 @@ // cusparseZcsrmv-NEXT: col_idx /*const int **/, x /*const cuDoubleComplex **/, // cusparseZcsrmv-NEXT: beta /*const cuDoubleComplex **/, y /*cuDoubleComplex **/); // cusparseZcsrmv-NEXT: Is migrated to: -// cusparseZcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseZcsrmv-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsrmv_mp | FileCheck %s -check-prefix=cusparseScsrmv_mp // cusparseScsrmv_mp: CUDA API: @@ -65,7 +65,7 @@ // cusparseScsrmv_mp-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseScsrmv_mp-NEXT: x /*const float **/, beta /*const float **/, y /*float **/); // cusparseScsrmv_mp-NEXT: Is migrated to: -// cusparseScsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseScsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrmv_mp | FileCheck %s -check-prefix=cusparseDcsrmv_mp // cusparseDcsrmv_mp: CUDA API: @@ -76,7 +76,7 @@ // cusparseDcsrmv_mp-NEXT: x /*const double **/, beta /*const double **/, // cusparseDcsrmv_mp-NEXT: y /*double **/); // cusparseDcsrmv_mp-NEXT: Is migrated to: -// cusparseDcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseDcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrmv_mp | FileCheck %s -check-prefix=cusparseCcsrmv_mp // cusparseCcsrmv_mp: CUDA API: @@ -87,7 +87,7 @@ // cusparseCcsrmv_mp-NEXT: col_idx /*const int **/, x /*const cuComplex **/, // cusparseCcsrmv_mp-NEXT: beta /*const cuComplex **/, y /*cuComplex **/); // cusparseCcsrmv_mp-NEXT: Is migrated to: -// cusparseCcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseCcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrmv_mp | FileCheck %s -check-prefix=cusparseZcsrmv_mp // cusparseZcsrmv_mp: CUDA API: @@ -99,7 +99,7 @@ // cusparseZcsrmv_mp-NEXT: col_idx /*const int **/, x /*const cuDoubleComplex **/, // cusparseZcsrmv_mp-NEXT: beta /*const cuDoubleComplex **/, y /*cuDoubleComplex **/); // cusparseZcsrmv_mp-NEXT: Is migrated to: -// cusparseZcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, alpha, desc, value, row_ptr, col_idx, x, beta, y); +// cusparseZcsrmv_mp-NEXT: dpct::sparse::csrmv(handle->get_queue(), trans, m, n, nnz, alpha, desc, value, row_ptr, col_idx, x, beta, y); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsrsv_analysis | FileCheck %s -check-prefix=cusparseScsrsv_analysis // cusparseScsrsv_analysis: CUDA API: @@ -109,7 +109,7 @@ // cusparseScsrsv_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseScsrsv_analysis-NEXT: info /*cusparseSolveAnalysisInfo_t*/); // cusparseScsrsv_analysis-NEXT: Is migrated to: -// cusparseScsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseScsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrsv_analysis | FileCheck %s -check-prefix=cusparseDcsrsv_analysis // cusparseDcsrsv_analysis: CUDA API: @@ -119,7 +119,7 @@ // cusparseDcsrsv_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseDcsrsv_analysis-NEXT: info /*cusparseSolveAnalysisInfo_t*/); // cusparseDcsrsv_analysis-NEXT: Is migrated to: -// cusparseDcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseDcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrsv_analysis | FileCheck %s -check-prefix=cusparseCcsrsv_analysis // cusparseCcsrsv_analysis: CUDA API: @@ -129,7 +129,7 @@ // cusparseCcsrsv_analysis-NEXT: row_ptr /*const int **/, col_idx /*const int **/, // cusparseCcsrsv_analysis-NEXT: info /*cusparseSolveAnalysisInfo_t*/); // cusparseCcsrsv_analysis-NEXT: Is migrated to: -// cusparseCcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseCcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrsv_analysis | FileCheck %s -check-prefix=cusparseZcsrsv_analysis // cusparseZcsrsv_analysis: CUDA API: @@ -139,7 +139,7 @@ // cusparseZcsrsv_analysis-NEXT: value /*const cuDoubleComplex **/, row_ptr /*const int **/, // cusparseZcsrsv_analysis-NEXT: col_idx /*const int **/, info /*cusparseSolveAnalysisInfo_t*/); // cusparseZcsrsv_analysis-NEXT: Is migrated to: -// cusparseZcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, row_ptr, col_idx, info); +// cusparseZcsrsv_analysis-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCsrsv_analysisEx | FileCheck %s -check-prefix=cusparseCsrsv_analysisEx // cusparseCsrsv_analysisEx: CUDA API: @@ -150,7 +150,7 @@ // cusparseCsrsv_analysisEx-NEXT: col_idx /*const int **/, info /*cusparseSolveAnalysisInfo_t*/, // cusparseCsrsv_analysisEx-NEXT: exec_type /*cudaDataType*/); // cusparseCsrsv_analysisEx-NEXT: Is migrated to: -// cusparseCsrsv_analysisEx-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, desc, value, value_type, row_ptr, col_idx, info); +// cusparseCsrsv_analysisEx-NEXT: dpct::sparse::optimize_csrsv(handle->get_queue(), trans, m, nnz, desc, value, value_type, row_ptr, col_idx, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsrsv_solve | FileCheck %s -check-prefix=cusparseScsrsv_solve // cusparseScsrsv_solve: CUDA API: @@ -217,7 +217,7 @@ // cusparseScsrmm-NEXT: col_idx /*const int **/, B /*const float **/, ldb /*int*/, // cusparseScsrmm-NEXT: beta /*const float **/, C /*float **/, ldc /*int*/); // cusparseScsrmm-NEXT: Is migrated to: -// cusparseScsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseScsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrmm | FileCheck %s -check-prefix=cusparseDcsrmm // cusparseDcsrmm: CUDA API: @@ -228,7 +228,7 @@ // cusparseDcsrmm-NEXT: col_idx /*const int **/, B /*const double **/, ldb /*int*/, // cusparseDcsrmm-NEXT: beta /*const double **/, C /*double **/, ldc /*int*/); // cusparseDcsrmm-NEXT: Is migrated to: -// cusparseDcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseDcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrmm | FileCheck %s -check-prefix=cusparseCcsrmm // cusparseCcsrmm: CUDA API: @@ -239,7 +239,7 @@ // cusparseCcsrmm-NEXT: col_idx /*const int **/, B /*const cuComplex **/, ldb /*int*/, // cusparseCcsrmm-NEXT: beta /*const cuComplex **/, C /*cuComplex **/, ldc /*int*/); // cusparseCcsrmm-NEXT: Is migrated to: -// cusparseCcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseCcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrmm | FileCheck %s -check-prefix=cusparseZcsrmm // cusparseZcsrmm: CUDA API: @@ -251,7 +251,7 @@ // cusparseZcsrmm-NEXT: ldb /*int*/, beta /*const cuDoubleComplex **/, // cusparseZcsrmm-NEXT: C /*cuDoubleComplex **/, ldc /*int*/); // cusparseZcsrmm-NEXT: Is migrated to: -// cusparseZcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseZcsrmm-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsrmm2 | FileCheck %s -check-prefix=cusparseScsrmm2 // cusparseScsrmm2: CUDA API: @@ -263,7 +263,7 @@ // cusparseScsrmm2-NEXT: B /*const float **/, ldb /*int*/, beta /*const float **/, // cusparseScsrmm2-NEXT: C /*float **/, ldc /*int*/); // cusparseScsrmm2-NEXT: Is migrated to: -// cusparseScsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseScsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrmm2 | FileCheck %s -check-prefix=cusparseDcsrmm2 // cusparseDcsrmm2: CUDA API: @@ -275,7 +275,7 @@ // cusparseDcsrmm2-NEXT: B /*const double **/, ldb /*int*/, beta /*const double **/, // cusparseDcsrmm2-NEXT: C /*double **/, ldc /*int*/); // cusparseDcsrmm2-NEXT: Is migrated to: -// cusparseDcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseDcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrmm2 | FileCheck %s -check-prefix=cusparseCcsrmm2 // cusparseCcsrmm2: CUDA API: @@ -287,7 +287,7 @@ // cusparseCcsrmm2-NEXT: B /*const cuComplex **/, ldb /*int*/, // cusparseCcsrmm2-NEXT: beta /*const cuComplex **/, C /*cuComplex **/, ldc /*int*/); // cusparseCcsrmm2-NEXT: Is migrated to: -// cusparseCcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseCcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrmm2 | FileCheck %s -check-prefix=cusparseZcsrmm2 // cusparseZcsrmm2: CUDA API: @@ -300,7 +300,7 @@ // cusparseZcsrmm2-NEXT: ldb /*int*/, beta /*const cuDoubleComplex **/, // cusparseZcsrmm2-NEXT: C /*cuDoubleComplex **/, ldc /*int*/); // cusparseZcsrmm2-NEXT: Is migrated to: -// cusparseZcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); +// cusparseZcsrmm2-NEXT: dpct::sparse::csrmm(handle->get_queue(), trans_a, trans_b, m, n, k, nnz, alpha, desc, value, row_ptr, col_idx, B, ldb, beta, C, ldc); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsr2csc | FileCheck %s -check-prefix=cusparseScsr2csc // cusparseScsr2csc: CUDA API: diff --git a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-118.cu b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-118.cu index 3a81f4fecc4e..e292caf5d633 100644 --- a/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-118.cu +++ b/clang/test/dpct/query_api_mapping/cuSPARSE/lit_cuSPARSE_92-118.cu @@ -78,7 +78,7 @@ // cusparseScsrsm2_analysis-NEXT: ldb /*int*/, info /*csrsm2Info_t*/, policy /*cusparseSolvePolicy_t*/, // cusparseScsrsm2_analysis-NEXT: buffer /*void **/); // cusparseScsrsm2_analysis-NEXT: Is migrated to: -// cusparseScsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, descr, value, row_ptr, col_ind, info); +// cusparseScsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, nnz, descr, value, row_ptr, col_ind, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseDcsrsm2_analysis | FileCheck %s -check-prefix=cusparseDcsrsm2_analysis // cusparseDcsrsm2_analysis: CUDA API: @@ -91,7 +91,7 @@ // cusparseDcsrsm2_analysis-NEXT: ldb /*int*/, info /*csrsm2Info_t*/, policy /*cusparseSolvePolicy_t*/, // cusparseDcsrsm2_analysis-NEXT: buffer /*void **/); // cusparseDcsrsm2_analysis-NEXT: Is migrated to: -// cusparseDcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, descr, value, row_ptr, col_ind, info); +// cusparseDcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, nnz, descr, value, row_ptr, col_ind, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseCcsrsm2_analysis | FileCheck %s -check-prefix=cusparseCcsrsm2_analysis // cusparseCcsrsm2_analysis: CUDA API: @@ -104,7 +104,7 @@ // cusparseCcsrsm2_analysis-NEXT: ldb /*int*/, info /*csrsm2Info_t*/, policy /*cusparseSolvePolicy_t*/, // cusparseCcsrsm2_analysis-NEXT: buffer /*void **/); // cusparseCcsrsm2_analysis-NEXT: Is migrated to: -// cusparseCcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, descr, value, row_ptr, col_ind, info); +// cusparseCcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, nnz, descr, value, row_ptr, col_ind, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseZcsrsm2_analysis | FileCheck %s -check-prefix=cusparseZcsrsm2_analysis // cusparseZcsrsm2_analysis: CUDA API: @@ -117,7 +117,7 @@ // cusparseZcsrsm2_analysis-NEXT: b /*const cuDoubleComplex **/, ldb /*int*/, info /*csrsm2Info_t*/, // cusparseZcsrsm2_analysis-NEXT: policy /*cusparseSolvePolicy_t*/, buffer /*void **/); // cusparseZcsrsm2_analysis-NEXT: Is migrated to: -// cusparseZcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, descr, value, row_ptr, col_ind, info); +// cusparseZcsrsm2_analysis-NEXT: dpct::sparse::optimize_csrsm(handle->get_queue(), trans_a, trans_b, m, nrhs, nnz, descr, value, row_ptr, col_ind, info); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cusparseScsrsm2_solve | FileCheck %s -check-prefix=cusparseScsrsm2_solve // cusparseScsrsm2_solve: CUDA API: