From 4db2187e21ca99000cbab08cafa768e94e5df14a Mon Sep 17 00:00:00 2001 From: "romain.biessy" Date: Tue, 24 Sep 2024 15:08:23 +0200 Subject: [PATCH] Avoid placeholder accessor --- .../backends/cusparse/cusparse_task.hpp | 75 +++++++++---------- .../cusparse/operations/cusparse_spmm.cpp | 9 +-- .../cusparse/operations/cusparse_spmv.cpp | 9 +-- .../cusparse/operations/cusparse_spsv.cpp | 4 +- src/sparse_blas/generic_container.hpp | 5 -- 5 files changed, 44 insertions(+), 58 deletions(-) diff --git a/src/sparse_blas/backends/cusparse/cusparse_task.hpp b/src/sparse_blas/backends/cusparse/cusparse_task.hpp index 0b1deb3f9..c6f34d49a 100644 --- a/src/sparse_blas/backends/cusparse/cusparse_task.hpp +++ b/src/sparse_blas/backends/cusparse/cusparse_task.hpp @@ -89,7 +89,6 @@ void submit_host_task_with_acc(sycl::handler &cgh, sycl::queue &queue, Functor f // specification but should be true for all the implementations. This // assumption avoids the overhead of resetting the pointer of all data // handles for each enqueued command. - cgh.require(workspace_placeholder_acc); cgh.host_task([functor, queue, workspace_placeholder_acc, capture_only_accessors...](sycl::interop_handle ih) { auto unused = std::make_tuple(capture_only_accessors...); @@ -151,7 +150,6 @@ void submit_native_command_ext_with_acc(sycl::handler &cgh, sycl::queue &queue, // assumption avoids the overhead of resetting the pointer of all data // handles for each enqueued command. #ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND - cgh.require(workspace_placeholder_acc); cgh.ext_codeplay_enqueue_native_command([functor, queue, dependencies, workspace_placeholder_acc, capture_only_accessors...](sycl::interop_handle ih) { @@ -196,36 +194,36 @@ template &dependencies, Functor functor, matrix_handle_t sm_handle, - sycl::accessor workspace_placeholder_acc, + sycl::buffer workspace_buffer, Ts... other_containers) { if (sm_handle->all_use_buffer()) { detail::data_type value_type = sm_handle->get_value_type(); detail::data_type int_type = sm_handle->get_int_type(); -#define ONEMKL_CUSPARSE_SUBMIT(FP_TYPE, INT_TYPE) \ - return queue.submit([&](sycl::handler &cgh) { \ - cgh.depends_on(dependencies); \ - auto fp_accs = get_fp_accessors(cgh, sm_handle, other_containers...); \ - auto int_accs = get_int_accessors(cgh, sm_handle); \ - if constexpr (UseWorkspace) { \ - if constexpr (UseEnqueueNativeCommandExt) { \ - submit_native_command_ext_with_acc(cgh, queue, functor, dependencies, \ - workspace_placeholder_acc, fp_accs, int_accs); \ - } \ - else { \ - submit_host_task_with_acc(cgh, queue, functor, workspace_placeholder_acc, fp_accs, \ - int_accs); \ - } \ - } \ - else { \ - (void)workspace_placeholder_acc; \ - if constexpr (UseEnqueueNativeCommandExt) { \ - submit_native_command_ext(cgh, queue, functor, dependencies, fp_accs, int_accs); \ - } \ - else { \ - submit_host_task(cgh, queue, functor, fp_accs, int_accs); \ - } \ - } \ +#define ONEMKL_CUSPARSE_SUBMIT(FP_TYPE, INT_TYPE) \ + return queue.submit([&](sycl::handler &cgh) { \ + cgh.depends_on(dependencies); \ + auto fp_accs = get_fp_accessors(cgh, sm_handle, other_containers...); \ + auto int_accs = get_int_accessors(cgh, sm_handle); \ + auto workspace_acc = workspace_buffer.get_access(cgh); \ + if constexpr (UseWorkspace) { \ + if constexpr (UseEnqueueNativeCommandExt) { \ + submit_native_command_ext_with_acc(cgh, queue, functor, dependencies, \ + workspace_acc, fp_accs, int_accs); \ + } \ + else { \ + submit_host_task_with_acc(cgh, queue, functor, workspace_acc, fp_accs, int_accs); \ + } \ + } \ + else { \ + (void)workspace_buffer; \ + if constexpr (UseEnqueueNativeCommandExt) { \ + submit_native_command_ext(cgh, queue, functor, dependencies, fp_accs, int_accs); \ + } \ + else { \ + submit_host_task(cgh, queue, functor, fp_accs, int_accs); \ + } \ + } \ }) #define ONEMKL_CUSPARSE_SUBMIT_INT(FP_TYPE) \ if (int_type == detail::data_type::int32) { \ @@ -318,14 +316,12 @@ sycl::event dispatch_submit_impl_fp(const std::string &function_name, sycl::queu /// Helper function for dispatch_submit_impl_fp_int template sycl::event dispatch_submit(const std::string &function_name, sycl::queue queue, Functor functor, - matrix_handle_t sm_handle, - sycl::accessor workspace_placeholder_acc, + matrix_handle_t sm_handle, sycl::buffer workspace_buffer, Ts... other_containers) { constexpr bool UseWorkspace = true; constexpr bool UseEnqueueNativeCommandExt = false; return dispatch_submit_impl_fp_int( - function_name, queue, {}, functor, sm_handle, workspace_placeholder_acc, - other_containers...); + function_name, queue, {}, functor, sm_handle, workspace_buffer, other_containers...); } /// Helper function for dispatch_submit_impl_fp_int @@ -335,8 +331,9 @@ sycl::event dispatch_submit(const std::string &function_name, sycl::queue queue, matrix_handle_t sm_handle, Ts... other_containers) { constexpr bool UseWorkspace = false; constexpr bool UseEnqueueNativeCommandExt = false; + sycl::buffer no_workspace(sycl::range<1>(0)); return dispatch_submit_impl_fp_int( - function_name, queue, dependencies, functor, sm_handle, {}, other_containers...); + function_name, queue, dependencies, functor, sm_handle, no_workspace, other_containers...); } /// Helper function for dispatch_submit_impl_fp_int @@ -345,15 +342,16 @@ sycl::event dispatch_submit(const std::string &function_name, sycl::queue queue, matrix_handle_t sm_handle, Ts... other_containers) { constexpr bool UseWorkspace = false; constexpr bool UseEnqueueNativeCommandExt = false; + sycl::buffer no_workspace(sycl::range<1>(0)); return dispatch_submit_impl_fp_int( - function_name, queue, {}, functor, sm_handle, {}, other_containers...); + function_name, queue, {}, functor, sm_handle, no_workspace, other_containers...); } /// Helper function for dispatch_submit_impl_fp_int template sycl::event dispatch_submit_native_ext(const std::string &function_name, sycl::queue queue, Functor functor, matrix_handle_t sm_handle, - sycl::accessor workspace_placeholder_acc, + sycl::buffer workspace_buffer, Ts... other_containers) { constexpr bool UseWorkspace = true; #ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND @@ -362,8 +360,7 @@ sycl::event dispatch_submit_native_ext(const std::string &function_name, sycl::q constexpr bool UseEnqueueNativeCommandExt = false; #endif return dispatch_submit_impl_fp_int( - function_name, queue, {}, functor, sm_handle, workspace_placeholder_acc, - other_containers...); + function_name, queue, {}, functor, sm_handle, workspace_buffer, other_containers...); } /// Helper function for dispatch_submit_impl_fp_int @@ -378,8 +375,9 @@ sycl::event dispatch_submit_native_ext(const std::string &function_name, sycl::q #else constexpr bool UseEnqueueNativeCommandExt = false; #endif + sycl::buffer no_workspace(sycl::range<1>(0)); return dispatch_submit_impl_fp_int( - function_name, queue, dependencies, functor, sm_handle, {}, other_containers...); + function_name, queue, dependencies, functor, sm_handle, no_workspace, other_containers...); } /// Helper function for dispatch_submit_impl_fp_int @@ -393,8 +391,9 @@ sycl::event dispatch_submit_native_ext(const std::string &function_name, sycl::q #else constexpr bool UseEnqueueNativeCommandExt = false; #endif + sycl::buffer no_workspace(sycl::range<1>(0)); return dispatch_submit_impl_fp_int( - function_name, queue, {}, functor, sm_handle, {}, other_containers...); + function_name, queue, {}, functor, sm_handle, no_workspace, other_containers...); } } // namespace oneapi::mkl::sparse::cusparse diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp index 116bbffe9..ca301a1e2 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spmm.cpp @@ -185,9 +185,7 @@ void spmm_optimize(sycl::queue& queue, oneapi::mkl::transpose opA, oneapi::mkl:: workspace_ptr, is_alpha_host_accessible); }; - sycl::accessor workspace_placeholder_acc(workspace); - dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, B_handle, - C_handle); + dispatch_submit(__func__, queue, functor, A_handle, workspace, B_handle, C_handle); } sycl::event spmm_optimize(sycl::queue& queue, oneapi::mkl::transpose opA, @@ -268,10 +266,9 @@ sycl::event spmm(sycl::queue& queue, oneapi::mkl::transpose opA, oneapi::mkl::tr auto workspace_ptr = sc.get_mem(workspace_acc); compute_functor(sc, workspace_ptr); }; - sycl::accessor workspace_placeholder_acc( - spmm_descr->workspace.get_buffer()); return dispatch_submit_native_ext(__func__, queue, functor_buffer, A_handle, - workspace_placeholder_acc, B_handle, C_handle); + spmm_descr->workspace.get_buffer(), + B_handle, C_handle); } else { // The same dispatch_submit can be used for USM or buffers if no diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp index 03cdd15e0..a0db00d8a 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp @@ -181,9 +181,7 @@ void spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a // The accessor can only be bound to the cgh if the buffer size is // greater than 0 - sycl::accessor workspace_placeholder_acc(workspace); - dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, x_handle, - y_handle); + dispatch_submit(__func__, queue, functor, A_handle, workspace, x_handle, y_handle); } else { auto functor = [=](CusparseScopedContextHandler &sc) { @@ -284,10 +282,9 @@ sycl::event spmv(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alp auto workspace_ptr = sc.get_mem(workspace_acc); compute_functor(sc, workspace_ptr); }; - sycl::accessor workspace_placeholder_acc( - spmv_descr->workspace.get_buffer()); return dispatch_submit_native_ext(__func__, queue, functor_buffer, A_handle, - workspace_placeholder_acc, x_handle, y_handle); + spmv_descr->workspace.get_buffer(), + x_handle, y_handle); } else { // The same dispatch_submit can be used for USM or buffers if no diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp index c06335100..4488d1d02 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp @@ -168,9 +168,7 @@ void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a // The accessor can only be bound to the cgh if the buffer size is // greater than 0 - sycl::accessor workspace_placeholder_acc(workspace); - dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, x_handle, - y_handle); + dispatch_submit(__func__, queue, functor, A_handle, workspace, x_handle, y_handle); } else { auto functor = [=](CusparseScopedContextHandler &sc) { diff --git a/src/sparse_blas/generic_container.hpp b/src/sparse_blas/generic_container.hpp index 09d408a77..5fa278497 100644 --- a/src/sparse_blas/generic_container.hpp +++ b/src/sparse_blas/generic_container.hpp @@ -272,11 +272,6 @@ struct generic_sparse_handle { } void set_matrix_property(matrix_property property) { - if (format == sparse_format::CSR && property == matrix_property::sorted_by_rows) { - throw mkl::invalid_argument( - "sparse_blas", "set_matrix_property", - "Property `matrix_property::sorted_by_rows` is not compatible with CSR format."); - } properties_mask |= matrix_property_to_mask(property); }