From 197cb705f92944923da4b0ddbcb88a360525bedf Mon Sep 17 00:00:00 2001 From: nscipione Date: Tue, 8 Oct 2024 10:46:25 +0100 Subject: [PATCH] Update sparse_blas function_tables operator[] calls Now function_tables operator[] requires a tuple, not a single value. Update call in sparse blas --- src/sparse_blas/sparse_blas_loader.cpp | 211 +++++++++++++------------ 1 file changed, 109 insertions(+), 102 deletions(-) diff --git a/src/sparse_blas/sparse_blas_loader.cpp b/src/sparse_blas/sparse_blas_loader.cpp index cdc3ae6b2..4304479d1 100644 --- a/src/sparse_blas/sparse_blas_loader.cpp +++ b/src/sparse_blas/sparse_blas_loader.cpp @@ -31,31 +31,34 @@ static oneapi::mkl::detail::table_initializer \ - void init_dense_vector(sycl::queue &queue, dense_vector_handle_t *p_dvhandle, \ - std::int64_t size, sycl::buffer val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].init_dense_vector_buffer##FP_SUFFIX(queue, p_dvhandle, size, val); \ - } \ - template <> \ - void init_dense_vector(sycl::queue &queue, dense_vector_handle_t *p_dvhandle, \ - std::int64_t size, FP_TYPE *val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].init_dense_vector_usm##FP_SUFFIX(queue, p_dvhandle, size, val); \ - } \ - template <> \ - void set_dense_vector_data(sycl::queue &queue, dense_vector_handle_t dvhandle, \ - std::int64_t size, sycl::buffer val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].set_dense_vector_data_buffer##FP_SUFFIX(queue, dvhandle, size, \ - val); \ - } \ - template <> \ - void set_dense_vector_data(sycl::queue &queue, dense_vector_handle_t dvhandle, \ - std::int64_t size, FP_TYPE *val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].set_dense_vector_data_usm##FP_SUFFIX(queue, dvhandle, size, val); \ +#define DEFINE_DENSE_VECTOR_FUNCS(FP_TYPE, FP_SUFFIX) \ + template <> \ + void init_dense_vector(sycl::queue &queue, dense_vector_handle_t *p_dvhandle, \ + std::int64_t size, sycl::buffer val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].init_dense_vector_buffer##FP_SUFFIX(queue, p_dvhandle, \ + size, val); \ + } \ + template <> \ + void init_dense_vector(sycl::queue &queue, dense_vector_handle_t *p_dvhandle, \ + std::int64_t size, FP_TYPE *val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].init_dense_vector_usm##FP_SUFFIX(queue, p_dvhandle, \ + size, val); \ + } \ + template <> \ + void set_dense_vector_data(sycl::queue &queue, dense_vector_handle_t dvhandle, \ + std::int64_t size, sycl::buffer val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].set_dense_vector_data_buffer##FP_SUFFIX( \ + queue, dvhandle, size, val); \ + } \ + template <> \ + void set_dense_vector_data(sycl::queue &queue, dense_vector_handle_t dvhandle, \ + std::int64_t size, FP_TYPE *val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].set_dense_vector_data_usm##FP_SUFFIX(queue, dvhandle, \ + size, val); \ } FOR_EACH_FP_TYPE(DEFINE_DENSE_VECTOR_FUNCS); #undef DEFINE_DENSE_VECTOR_FUNCS @@ -63,42 +66,42 @@ FOR_EACH_FP_TYPE(DEFINE_DENSE_VECTOR_FUNCS); sycl::event release_dense_vector(sycl::queue &queue, dense_vector_handle_t dvhandle, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_dense_vector(queue, dvhandle, dependencies); + return function_tables[{ libkey, queue }].release_dense_vector(queue, dvhandle, dependencies); } // Dense matrix -#define DEFINE_DENSE_MATRIX_FUNCS(FP_TYPE, FP_SUFFIX) \ - template <> \ - void init_dense_matrix(sycl::queue &queue, dense_matrix_handle_t *p_dmhandle, \ - std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ - layout dense_layout, sycl::buffer val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].init_dense_matrix_buffer##FP_SUFFIX( \ - queue, p_dmhandle, num_rows, num_cols, ld, dense_layout, val); \ - } \ - template <> \ - void init_dense_matrix(sycl::queue &queue, dense_matrix_handle_t *p_dmhandle, \ - std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ - layout dense_layout, FP_TYPE *val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].init_dense_matrix_usm##FP_SUFFIX(queue, p_dmhandle, num_rows, \ - num_cols, ld, dense_layout, val); \ - } \ - template <> \ - void set_dense_matrix_data(sycl::queue &queue, dense_matrix_handle_t dmhandle, \ - std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ - layout dense_layout, sycl::buffer val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].set_dense_matrix_data_buffer##FP_SUFFIX( \ - queue, dmhandle, num_rows, num_cols, ld, dense_layout, val); \ - } \ - template <> \ - void set_dense_matrix_data(sycl::queue &queue, dense_matrix_handle_t dmhandle, \ - std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ - layout dense_layout, FP_TYPE *val) { \ - auto libkey = get_device_id(queue); \ - function_tables[libkey].set_dense_matrix_data_usm##FP_SUFFIX( \ - queue, dmhandle, num_rows, num_cols, ld, dense_layout, val); \ +#define DEFINE_DENSE_MATRIX_FUNCS(FP_TYPE, FP_SUFFIX) \ + template <> \ + void init_dense_matrix(sycl::queue &queue, dense_matrix_handle_t *p_dmhandle, \ + std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ + layout dense_layout, sycl::buffer val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].init_dense_matrix_buffer##FP_SUFFIX( \ + queue, p_dmhandle, num_rows, num_cols, ld, dense_layout, val); \ + } \ + template <> \ + void init_dense_matrix(sycl::queue &queue, dense_matrix_handle_t *p_dmhandle, \ + std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ + layout dense_layout, FP_TYPE *val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].init_dense_matrix_usm##FP_SUFFIX( \ + queue, p_dmhandle, num_rows, num_cols, ld, dense_layout, val); \ + } \ + template <> \ + void set_dense_matrix_data(sycl::queue &queue, dense_matrix_handle_t dmhandle, \ + std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ + layout dense_layout, sycl::buffer val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].set_dense_matrix_data_buffer##FP_SUFFIX( \ + queue, dmhandle, num_rows, num_cols, ld, dense_layout, val); \ + } \ + template <> \ + void set_dense_matrix_data(sycl::queue &queue, dense_matrix_handle_t dmhandle, \ + std::int64_t num_rows, std::int64_t num_cols, std::int64_t ld, \ + layout dense_layout, FP_TYPE *val) { \ + auto libkey = get_device_id(queue); \ + function_tables[{ libkey, queue }].set_dense_matrix_data_usm##FP_SUFFIX( \ + queue, dmhandle, num_rows, num_cols, ld, dense_layout, val); \ } FOR_EACH_FP_TYPE(DEFINE_DENSE_MATRIX_FUNCS); #undef DEFINE_DENSE_MATRIX_FUNCS @@ -106,7 +109,7 @@ FOR_EACH_FP_TYPE(DEFINE_DENSE_MATRIX_FUNCS); sycl::event release_dense_matrix(sycl::queue &queue, dense_matrix_handle_t dmhandle, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_dense_matrix(queue, dmhandle, dependencies); + return function_tables[{ libkey, queue }].release_dense_matrix(queue, dmhandle, dependencies); } // COO matrix @@ -117,7 +120,7 @@ sycl::event release_dense_matrix(sycl::queue &queue, dense_matrix_handle_t dmhan sycl::buffer row_ind, sycl::buffer col_ind, \ sycl::buffer val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].init_coo_matrix_buffer##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].init_coo_matrix_buffer##FP_SUFFIX##INT_SUFFIX( \ queue, p_smhandle, num_rows, num_cols, nnz, index, row_ind, col_ind, val); \ } \ template <> \ @@ -125,7 +128,7 @@ sycl::event release_dense_matrix(sycl::queue &queue, dense_matrix_handle_t dmhan std::int64_t num_cols, std::int64_t nnz, index_base index, \ INT_TYPE *row_ind, INT_TYPE *col_ind, FP_TYPE *val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].init_coo_matrix_usm##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].init_coo_matrix_usm##FP_SUFFIX##INT_SUFFIX( \ queue, p_smhandle, num_rows, num_cols, nnz, index, row_ind, col_ind, val); \ } \ template <> \ @@ -134,7 +137,7 @@ sycl::event release_dense_matrix(sycl::queue &queue, dense_matrix_handle_t dmhan sycl::buffer row_ind, sycl::buffer col_ind, \ sycl::buffer val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].set_coo_matrix_data_buffer##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].set_coo_matrix_data_buffer##FP_SUFFIX##INT_SUFFIX( \ queue, smhandle, num_rows, num_cols, nnz, index, row_ind, col_ind, val); \ } \ template <> \ @@ -142,7 +145,7 @@ sycl::event release_dense_matrix(sycl::queue &queue, dense_matrix_handle_t dmhan std::int64_t num_cols, std::int64_t nnz, index_base index, \ INT_TYPE *row_ind, INT_TYPE *col_ind, FP_TYPE *val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].set_coo_matrix_data_usm##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].set_coo_matrix_data_usm##FP_SUFFIX##INT_SUFFIX( \ queue, smhandle, num_rows, num_cols, nnz, index, row_ind, col_ind, val); \ } FOR_EACH_FP_AND_INT_TYPE(DEFINE_COO_MATRIX_FUNCS); @@ -156,7 +159,7 @@ FOR_EACH_FP_AND_INT_TYPE(DEFINE_COO_MATRIX_FUNCS); sycl::buffer row_ptr, sycl::buffer col_ind, \ sycl::buffer val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].init_csr_matrix_buffer##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].init_csr_matrix_buffer##FP_SUFFIX##INT_SUFFIX( \ queue, p_smhandle, num_rows, num_cols, nnz, index, row_ptr, col_ind, val); \ } \ template <> \ @@ -164,7 +167,7 @@ FOR_EACH_FP_AND_INT_TYPE(DEFINE_COO_MATRIX_FUNCS); std::int64_t num_cols, std::int64_t nnz, index_base index, \ INT_TYPE *row_ptr, INT_TYPE *col_ind, FP_TYPE *val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].init_csr_matrix_usm##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].init_csr_matrix_usm##FP_SUFFIX##INT_SUFFIX( \ queue, p_smhandle, num_rows, num_cols, nnz, index, row_ptr, col_ind, val); \ } \ template <> \ @@ -173,7 +176,7 @@ FOR_EACH_FP_AND_INT_TYPE(DEFINE_COO_MATRIX_FUNCS); sycl::buffer row_ptr, sycl::buffer col_ind, \ sycl::buffer val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].set_csr_matrix_data_buffer##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].set_csr_matrix_data_buffer##FP_SUFFIX##INT_SUFFIX( \ queue, smhandle, num_rows, num_cols, nnz, index, row_ptr, col_ind, val); \ } \ template <> \ @@ -181,7 +184,7 @@ FOR_EACH_FP_AND_INT_TYPE(DEFINE_COO_MATRIX_FUNCS); std::int64_t num_cols, std::int64_t nnz, index_base index, \ INT_TYPE *row_ptr, INT_TYPE *col_ind, FP_TYPE *val) { \ auto libkey = get_device_id(queue); \ - function_tables[libkey].set_csr_matrix_data_usm##FP_SUFFIX##INT_SUFFIX( \ + function_tables[{ libkey, queue }].set_csr_matrix_data_usm##FP_SUFFIX##INT_SUFFIX( \ queue, smhandle, num_rows, num_cols, nnz, index, row_ptr, col_ind, val); \ } FOR_EACH_FP_AND_INT_TYPE(DEFINE_INIT_CSR_MATRIX_FUNCS); @@ -191,24 +194,24 @@ FOR_EACH_FP_AND_INT_TYPE(DEFINE_INIT_CSR_MATRIX_FUNCS); sycl::event release_sparse_matrix(sycl::queue &queue, matrix_handle_t smhandle, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_sparse_matrix(queue, smhandle, dependencies); + return function_tables[{ libkey, queue }].release_sparse_matrix(queue, smhandle, dependencies); } bool set_matrix_property(sycl::queue &queue, matrix_handle_t smhandle, matrix_property property) { auto libkey = get_device_id(queue); - return function_tables[libkey].set_matrix_property(queue, smhandle, property); + return function_tables[{ libkey, queue }].set_matrix_property(queue, smhandle, property); } // SPMM void init_spmm_descr(sycl::queue &queue, spmm_descr_t *p_spmm_descr) { auto libkey = get_device_id(queue); - function_tables[libkey].init_spmm_descr(queue, p_spmm_descr); + function_tables[{ libkey, queue }].init_spmm_descr(queue, p_spmm_descr); } sycl::event release_spmm_descr(sycl::queue &queue, spmm_descr_t spmm_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_spmm_descr(queue, spmm_descr, dependencies); + return function_tables[{ libkey, queue }].release_spmm_descr(queue, spmm_descr, dependencies); } void spmm_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::transpose opB, @@ -217,8 +220,9 @@ void spmm_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mk dense_matrix_handle_t C_handle, spmm_alg alg, spmm_descr_t spmm_descr, std::size_t &temp_buffer_size) { auto libkey = get_device_id(queue); - function_tables[libkey].spmm_buffer_size(queue, opA, opB, alpha, A_view, A_handle, B_handle, - beta, C_handle, alg, spmm_descr, temp_buffer_size); + function_tables[{ libkey, queue }].spmm_buffer_size(queue, opA, opB, alpha, A_view, A_handle, + B_handle, beta, C_handle, alg, spmm_descr, + temp_buffer_size); } void spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::transpose opB, @@ -226,8 +230,9 @@ void spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl:: dense_matrix_handle_t B_handle, const void *beta, dense_matrix_handle_t C_handle, spmm_alg alg, spmm_descr_t spmm_descr, sycl::buffer workspace) { auto libkey = get_device_id(queue); - function_tables[libkey].spmm_optimize_buffer(queue, opA, opB, alpha, A_view, A_handle, B_handle, - beta, C_handle, alg, spmm_descr, workspace); + function_tables[{ libkey, queue }].spmm_optimize_buffer(queue, opA, opB, alpha, A_view, + A_handle, B_handle, beta, C_handle, alg, + spmm_descr, workspace); } sycl::event spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, @@ -237,9 +242,9 @@ sycl::event spmm_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, spmm_descr_t spmm_descr, void *workspace, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spmm_optimize_usm(queue, opA, opB, alpha, A_view, A_handle, - B_handle, beta, C_handle, alg, spmm_descr, - workspace, dependencies); + return function_tables[{ libkey, queue }].spmm_optimize_usm( + queue, opA, opB, alpha, A_view, A_handle, B_handle, beta, C_handle, alg, spmm_descr, + workspace, dependencies); } sycl::event spmm(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::transpose opB, @@ -248,20 +253,21 @@ sycl::event spmm(sycl::queue &queue, oneapi::mkl::transpose opA, oneapi::mkl::tr spmm_alg alg, spmm_descr_t spmm_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spmm(queue, opA, opB, alpha, A_view, A_handle, B_handle, beta, - C_handle, alg, spmm_descr, dependencies); + return function_tables[{ libkey, queue }].spmm(queue, opA, opB, alpha, A_view, A_handle, + B_handle, beta, C_handle, alg, spmm_descr, + dependencies); } // SPMV void init_spmv_descr(sycl::queue &queue, spmv_descr_t *p_spmv_descr) { auto libkey = get_device_id(queue); - function_tables[libkey].init_spmv_descr(queue, p_spmv_descr); + function_tables[{ libkey, queue }].init_spmv_descr(queue, p_spmv_descr); } sycl::event release_spmv_descr(sycl::queue &queue, spmv_descr_t spmv_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_spmv_descr(queue, spmv_descr, dependencies); + return function_tables[{ libkey, queue }].release_spmv_descr(queue, spmv_descr, dependencies); } void spmv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -269,8 +275,9 @@ void spmv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void const void *beta, dense_vector_handle_t y_handle, spmv_alg alg, spmv_descr_t spmv_descr, std::size_t &temp_buffer_size) { auto libkey = get_device_id(queue); - function_tables[libkey].spmv_buffer_size(queue, opA, alpha, A_view, A_handle, x_handle, beta, - y_handle, alg, spmv_descr, temp_buffer_size); + function_tables[{ libkey, queue }].spmv_buffer_size(queue, opA, alpha, A_view, A_handle, + x_handle, beta, y_handle, alg, spmv_descr, + temp_buffer_size); } void spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -278,8 +285,8 @@ void spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a const void *beta, dense_vector_handle_t y_handle, spmv_alg alg, spmv_descr_t spmv_descr, sycl::buffer workspace) { auto libkey = get_device_id(queue); - function_tables[libkey].spmv_optimize_buffer(queue, opA, alpha, A_view, A_handle, x_handle, - beta, y_handle, alg, spmv_descr, workspace); + function_tables[{ libkey, queue }].spmv_optimize_buffer( + queue, opA, alpha, A_view, A_handle, x_handle, beta, y_handle, alg, spmv_descr, workspace); } sycl::event spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -288,9 +295,9 @@ sycl::event spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const dense_vector_handle_t y_handle, spmv_alg alg, spmv_descr_t spmv_descr, void *workspace, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spmv_optimize_usm(queue, opA, alpha, A_view, A_handle, x_handle, - beta, y_handle, alg, spmv_descr, workspace, - dependencies); + return function_tables[{ libkey, queue }].spmv_optimize_usm( + queue, opA, alpha, A_view, A_handle, x_handle, beta, y_handle, alg, spmv_descr, workspace, + dependencies); } sycl::event spmv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -298,20 +305,20 @@ sycl::event spmv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp const void *beta, dense_vector_handle_t y_handle, spmv_alg alg, spmv_descr_t spmv_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spmv(queue, opA, alpha, A_view, A_handle, x_handle, beta, - y_handle, alg, spmv_descr, dependencies); + return function_tables[{ libkey, queue }].spmv(queue, opA, alpha, A_view, A_handle, x_handle, + beta, y_handle, alg, spmv_descr, dependencies); } // SPSV void init_spsv_descr(sycl::queue &queue, spsv_descr_t *p_spsv_descr) { auto libkey = get_device_id(queue); - function_tables[libkey].init_spsv_descr(queue, p_spsv_descr); + function_tables[{ libkey, queue }].init_spsv_descr(queue, p_spsv_descr); } sycl::event release_spsv_descr(sycl::queue &queue, spsv_descr_t spsv_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].release_spsv_descr(queue, spsv_descr, dependencies); + return function_tables[{ libkey, queue }].release_spsv_descr(queue, spsv_descr, dependencies); } void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -319,8 +326,8 @@ void spsv_buffer_size(sycl::queue &queue, oneapi::mkl::transpose opA, const void dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr, std::size_t &temp_buffer_size) { auto libkey = get_device_id(queue); - function_tables[libkey].spsv_buffer_size(queue, opA, alpha, A_view, A_handle, x_handle, - y_handle, alg, spsv_descr, temp_buffer_size); + function_tables[{ libkey, queue }].spsv_buffer_size( + queue, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, spsv_descr, temp_buffer_size); } void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -328,8 +335,8 @@ void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr, sycl::buffer workspace) { auto libkey = get_device_id(queue); - function_tables[libkey].spsv_optimize_buffer(queue, opA, alpha, A_view, A_handle, x_handle, - y_handle, alg, spsv_descr, workspace); + function_tables[{ libkey, queue }].spsv_optimize_buffer( + queue, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, spsv_descr, workspace); } sycl::event spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -338,9 +345,9 @@ sycl::event spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const spsv_alg alg, spsv_descr_t spsv_descr, void *workspace, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spsv_optimize_usm(queue, opA, alpha, A_view, A_handle, x_handle, - y_handle, alg, spsv_descr, workspace, - dependencies); + return function_tables[{ libkey, queue }].spsv_optimize_usm(queue, opA, alpha, A_view, A_handle, + x_handle, y_handle, alg, spsv_descr, + workspace, dependencies); } sycl::event spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha, @@ -348,8 +355,8 @@ sycl::event spsv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp dense_vector_handle_t y_handle, spsv_alg alg, spsv_descr_t spsv_descr, const std::vector &dependencies) { auto libkey = get_device_id(queue); - return function_tables[libkey].spsv(queue, opA, alpha, A_view, A_handle, x_handle, y_handle, - alg, spsv_descr, dependencies); + return function_tables[{ libkey, queue }].spsv(queue, opA, alpha, A_view, A_handle, x_handle, + y_handle, alg, spsv_descr, dependencies); } } // namespace oneapi::mkl::sparse