From 0a12c0d90c68d17e1093dba47058f0949f6d4db2 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 27 Jun 2024 09:50:32 +0800 Subject: [PATCH] Add helper test Signed-off-by: Jiang, Zhiwei --- help_function/src/sparse_utils_2_buffer.cpp | 199 ++++++++++++++++++++ help_function/src/sparse_utils_2_usm.cpp | 199 ++++++++++++++++++++ 2 files changed, 398 insertions(+) diff --git a/help_function/src/sparse_utils_2_buffer.cpp b/help_function/src/sparse_utils_2_buffer.cpp index eb2ef4e0c..1809c365c 100644 --- a/help_function/src/sparse_utils_2_buffer.cpp +++ b/help_function/src/sparse_utils_2_buffer.cpp @@ -1967,6 +1967,204 @@ void test_cusparseSpSV() { } } +// A * B = C +// +// | 0 1 2 | | 1 0 0 0 | | 2 3 10 12 | +// | 0 0 3 | * | 2 3 0 0 | = | 0 0 15 18 | +// | 4 0 0 | | 0 0 5 6 | | 4 0 0 0 | +void test_cusparseTcsrgemm() { + dpct::device_ext &dev_ct1 = dpct::get_current_device(); + sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); + std::vector a_val_vec = {1, 2, 3, 4}; + Data a_s_val(a_val_vec.data(), 4); + Data a_d_val(a_val_vec.data(), 4); + Data a_c_val(a_val_vec.data(), 4); + Data a_z_val(a_val_vec.data(), 4); + std::vector a_row_ptr_vec = {0, 2, 3, 4}; + Data a_row_ptr(a_row_ptr_vec.data(), 4); + std::vector a_col_ind_vec = {1, 2, 2, 0}; + Data a_col_ind(a_col_ind_vec.data(), 4); + + std::vector b_val_vec = {1, 2, 3, 5, 6}; + Data b_s_val(b_val_vec.data(), 5); + Data b_d_val(b_val_vec.data(), 5); + Data b_c_val(b_val_vec.data(), 5); + Data b_z_val(b_val_vec.data(), 5); + std::vector b_row_ptr_vec = {0, 1, 3, 5}; + Data b_row_ptr(b_row_ptr_vec.data(), 4); + std::vector b_col_ind_vec = {0, 0, 1, 2, 3}; + Data b_col_ind(b_col_ind_vec.data(), 5); + + float alpha = 1; + Data alpha_s(&alpha, 1); + Data alpha_d(&alpha, 1); + Data alpha_c(&alpha, 1); + Data alpha_z(&alpha, 1); + + float beta = 0; + Data beta_s(&beta, 1); + Data beta_d(&beta, 1); + Data beta_c(&beta, 1); + Data beta_z(&beta, 1); + + dpct::sparse::descriptor_ptr handle; + handle = new dpct::sparse::descriptor(); + + /* + DPCT1026:38: The call to cusparseSetPointerMode was removed because this + functionality is redundant in SYCL. + */ + + a_s_val.H2D(); + a_d_val.H2D(); + a_c_val.H2D(); + a_z_val.H2D(); + a_row_ptr.H2D(); + a_col_ind.H2D(); + b_s_val.H2D(); + b_d_val.H2D(); + b_c_val.H2D(); + b_z_val.H2D(); + b_row_ptr.H2D(); + b_col_ind.H2D(); + alpha_s.H2D(); + alpha_d.H2D(); + alpha_c.H2D(); + alpha_z.H2D(); + beta_s.H2D(); + beta_d.H2D(); + beta_c.H2D(); + beta_z.H2D(); + + Data c_s_row_ptr(4); + Data c_d_row_ptr(4); + Data c_c_row_ptr(4); + Data c_z_row_ptr(4); + + std::shared_ptr descrA; + std::shared_ptr descrB; + std::shared_ptr descrC; + descrA = std::make_shared(); + descrB = std::make_shared(); + descrC = std::make_shared(); + descrA->set_index_base(oneapi::mkl::index_base::zero); + descrB->set_index_base(oneapi::mkl::index_base::zero); + descrC->set_index_base(oneapi::mkl::index_base::zero); + descrA->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + descrB->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + descrC->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + + int c_nnz_s; + int c_nnz_d; + int c_nnz_c; + int c_nnz_z; + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_s_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_s_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_s_row_ptr.d_data, &c_nnz_s); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_d_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_d_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_d_row_ptr.d_data, &c_nnz_d); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_c_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_c_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_c_row_ptr.d_data, &c_nnz_c); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_z_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_z_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_z_row_ptr.d_data, &c_nnz_z); + + Data c_s_val(c_nnz_s); + Data c_d_val(c_nnz_d); + Data c_c_val(c_nnz_c); + Data c_z_val(c_nnz_z); + Data c_s_col_ind(c_nnz_s); + Data c_d_col_ind(c_nnz_d); + Data c_c_col_ind(c_nnz_c); + Data c_z_col_ind(c_nnz_z); + + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_s_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_s_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_s_val.d_data, + c_s_row_ptr.d_data, c_s_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_d_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_d_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_d_val.d_data, + c_d_row_ptr.d_data, c_d_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_c_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_c_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_c_val.d_data, + c_c_row_ptr.d_data, c_c_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_z_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_z_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_z_val.d_data, + c_z_row_ptr.d_data, c_z_col_ind.d_data); + + q_ct1.wait(); + + /* + DPCT1026:39: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + /* + DPCT1026:40: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + /* + DPCT1026:41: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + delete (handle); + + c_s_val.D2H(); + c_d_val.D2H(); + c_c_val.D2H(); + c_z_val.D2H(); + c_s_row_ptr.D2H(); + c_d_row_ptr.D2H(); + c_c_row_ptr.D2H(); + c_z_row_ptr.D2H(); + c_s_col_ind.D2H(); + c_d_col_ind.D2H(); + c_c_col_ind.D2H(); + c_z_col_ind.D2H(); + + float expect_c_val[7] = {2.000000, 3.000000, 10.000000, 12.000000, 15.000000, 18.000000, 4.000000}; + float expect_c_row_ptr[4] = {0.000000, 4.000000, 6.000000, 7.000000}; + float expect_c_col_ind[7] = {0.000000, 1.000000, 2.000000, 3.000000, 2.000000, 3.000000, 0.000000}; + if (compare_result(expect_c_val, c_s_val.h_data, 7) && + compare_result(expect_c_val, c_d_val.h_data, 7) && + compare_result(expect_c_val, c_c_val.h_data, 7) && + compare_result(expect_c_val, c_z_val.h_data, 7) && + compare_result(expect_c_row_ptr, c_s_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_d_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_c_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_z_row_ptr.h_data, 4) && + compare_result(expect_c_col_ind, c_s_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_d_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_c_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_z_col_ind.h_data, 7) + ) + printf("Tcsrgemm pass\n"); + else { + printf("Tcsrgemm fail\n"); + test_passed = false; + } +} + int main() { test_cusparseSetGetStream(); test_cusparseTcsrmv_ge(); @@ -1979,6 +2177,7 @@ int main() { test_cusparseCsrmvEx(); test_cusparseSpGEMM(); test_cusparseSpSV(); + test_cusparseTcsrgemm(); if (test_passed) return 0; diff --git a/help_function/src/sparse_utils_2_usm.cpp b/help_function/src/sparse_utils_2_usm.cpp index 1cd1c99d0..aea44a6af 100644 --- a/help_function/src/sparse_utils_2_usm.cpp +++ b/help_function/src/sparse_utils_2_usm.cpp @@ -2016,6 +2016,204 @@ void test_cusparseSpSV() { } } +// A * B = C +// +// | 0 1 2 | | 1 0 0 0 | | 2 3 10 12 | +// | 0 0 3 | * | 2 3 0 0 | = | 0 0 15 18 | +// | 4 0 0 | | 0 0 5 6 | | 4 0 0 0 | +void test_cusparseTcsrgemm() { + dpct::device_ext &dev_ct1 = dpct::get_current_device(); + sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + std::vector a_val_vec = {1, 2, 3, 4}; + Data a_s_val(a_val_vec.data(), 4); + Data a_d_val(a_val_vec.data(), 4); + Data a_c_val(a_val_vec.data(), 4); + Data a_z_val(a_val_vec.data(), 4); + std::vector a_row_ptr_vec = {0, 2, 3, 4}; + Data a_row_ptr(a_row_ptr_vec.data(), 4); + std::vector a_col_ind_vec = {1, 2, 2, 0}; + Data a_col_ind(a_col_ind_vec.data(), 4); + + std::vector b_val_vec = {1, 2, 3, 5, 6}; + Data b_s_val(b_val_vec.data(), 5); + Data b_d_val(b_val_vec.data(), 5); + Data b_c_val(b_val_vec.data(), 5); + Data b_z_val(b_val_vec.data(), 5); + std::vector b_row_ptr_vec = {0, 1, 3, 5}; + Data b_row_ptr(b_row_ptr_vec.data(), 4); + std::vector b_col_ind_vec = {0, 0, 1, 2, 3}; + Data b_col_ind(b_col_ind_vec.data(), 5); + + float alpha = 1; + Data alpha_s(&alpha, 1); + Data alpha_d(&alpha, 1); + Data alpha_c(&alpha, 1); + Data alpha_z(&alpha, 1); + + float beta = 0; + Data beta_s(&beta, 1); + Data beta_d(&beta, 1); + Data beta_c(&beta, 1); + Data beta_z(&beta, 1); + + dpct::sparse::descriptor_ptr handle; + handle = new dpct::sparse::descriptor(); + + /* + DPCT1026:38: The call to cusparseSetPointerMode was removed because this + functionality is redundant in SYCL. + */ + + a_s_val.H2D(); + a_d_val.H2D(); + a_c_val.H2D(); + a_z_val.H2D(); + a_row_ptr.H2D(); + a_col_ind.H2D(); + b_s_val.H2D(); + b_d_val.H2D(); + b_c_val.H2D(); + b_z_val.H2D(); + b_row_ptr.H2D(); + b_col_ind.H2D(); + alpha_s.H2D(); + alpha_d.H2D(); + alpha_c.H2D(); + alpha_z.H2D(); + beta_s.H2D(); + beta_d.H2D(); + beta_c.H2D(); + beta_z.H2D(); + + Data c_s_row_ptr(4); + Data c_d_row_ptr(4); + Data c_c_row_ptr(4); + Data c_z_row_ptr(4); + + std::shared_ptr descrA; + std::shared_ptr descrB; + std::shared_ptr descrC; + descrA = std::make_shared(); + descrB = std::make_shared(); + descrC = std::make_shared(); + descrA->set_index_base(oneapi::mkl::index_base::zero); + descrB->set_index_base(oneapi::mkl::index_base::zero); + descrC->set_index_base(oneapi::mkl::index_base::zero); + descrA->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + descrB->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + descrC->set_matrix_type(dpct::sparse::matrix_info::matrix_type::ge); + + int c_nnz_s; + int c_nnz_d; + int c_nnz_c; + int c_nnz_z; + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_s_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_s_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_s_row_ptr.d_data, &c_nnz_s); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_d_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_d_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_d_row_ptr.d_data, &c_nnz_d); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_c_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_c_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_c_row_ptr.d_data, &c_nnz_c); + dpct::sparse::csrgemm_nnz( + handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, 4, a_z_val.d_data, + a_row_ptr.d_data, a_col_ind.d_data, descrB, 5, b_z_val.d_data, + b_row_ptr.d_data, b_col_ind.d_data, descrC, c_z_row_ptr.d_data, &c_nnz_z); + + Data c_s_val(c_nnz_s); + Data c_d_val(c_nnz_d); + Data c_c_val(c_nnz_c); + Data c_z_val(c_nnz_z); + Data c_s_col_ind(c_nnz_s); + Data c_d_col_ind(c_nnz_d); + Data c_c_col_ind(c_nnz_c); + Data c_z_col_ind(c_nnz_z); + + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_s_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_s_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_s_val.d_data, + c_s_row_ptr.d_data, c_s_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_d_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_d_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_d_val.d_data, + c_d_row_ptr.d_data, c_d_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_c_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_c_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_c_val.d_data, + c_c_row_ptr.d_data, c_c_col_ind.d_data); + dpct::sparse::csrgemm(handle, oneapi::mkl::transpose::nontrans, + oneapi::mkl::transpose::nontrans, 3, 3, 4, descrA, + a_z_val.d_data, a_row_ptr.d_data, a_col_ind.d_data, + descrB, b_z_val.d_data, b_row_ptr.d_data, + b_col_ind.d_data, descrC, c_z_val.d_data, + c_z_row_ptr.d_data, c_z_col_ind.d_data); + + q_ct1.wait(); + + /* + DPCT1026:39: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + /* + DPCT1026:40: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + /* + DPCT1026:41: The call to cusparseDestroyMatDescr was removed because this + functionality is redundant in SYCL. + */ + delete (handle); + + c_s_val.D2H(); + c_d_val.D2H(); + c_c_val.D2H(); + c_z_val.D2H(); + c_s_row_ptr.D2H(); + c_d_row_ptr.D2H(); + c_c_row_ptr.D2H(); + c_z_row_ptr.D2H(); + c_s_col_ind.D2H(); + c_d_col_ind.D2H(); + c_c_col_ind.D2H(); + c_z_col_ind.D2H(); + + float expect_c_val[7] = {2.000000, 3.000000, 10.000000, 12.000000, 15.000000, 18.000000, 4.000000}; + float expect_c_row_ptr[4] = {0.000000, 4.000000, 6.000000, 7.000000}; + float expect_c_col_ind[7] = {0.000000, 1.000000, 2.000000, 3.000000, 2.000000, 3.000000, 0.000000}; + if (compare_result(expect_c_val, c_s_val.h_data, 7) && + compare_result(expect_c_val, c_d_val.h_data, 7) && + compare_result(expect_c_val, c_c_val.h_data, 7) && + compare_result(expect_c_val, c_z_val.h_data, 7) && + compare_result(expect_c_row_ptr, c_s_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_d_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_c_row_ptr.h_data, 4) && + compare_result(expect_c_row_ptr, c_z_row_ptr.h_data, 4) && + compare_result(expect_c_col_ind, c_s_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_d_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_c_col_ind.h_data, 7) && + compare_result(expect_c_col_ind, c_z_col_ind.h_data, 7) + ) + printf("Tcsrgemm pass\n"); + else { + printf("Tcsrgemm fail\n"); + test_passed = false; + } +} + int main() { test_cusparseSetGetStream(); test_cub_spmv_csrmv(); @@ -2029,6 +2227,7 @@ int main() { test_cusparseCsrmvEx(); test_cusparseSpGEMM(); test_cusparseSpSV(); + test_cusparseTcsrgemm(); if (test_passed) return 0;