diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp index 8a0fb7fe1..230687174 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spmv.cpp @@ -151,6 +151,9 @@ void spmv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a spmv_optimize_impl(cu_handle, opA, alpha, A_handle, x_handle, beta, y_handle, alg, workspace_ptr, is_alpha_host_accessible); }; + + // The accessor can only be bound to the cgh if the buffer size is + // greater than 0 sycl::accessor workspace_placeholder_acc(workspace); event = dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, x_handle, y_handle); diff --git a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp index 7f2d6287f..027d51e08 100644 --- a/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp +++ b/src/sparse_blas/backends/cusparse/operations/cusparse_spsv.cpp @@ -128,18 +128,33 @@ void spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *a // Ignore spsv_alg::no_optimize_alg as this step is mandatory for cuSPARSE // Copy the buffer to extend its lifetime until the descriptor is free'd. spsv_descr->workspace.set_buffer_untyped(workspace); - auto functor = [=](CusparseScopedContextHandler &sc, - sycl::accessor workspace_acc) { - auto cu_handle = sc.get_handle(queue); - auto workspace_ptr = sc.get_mem(workspace_acc); - spsv_optimize_impl(cu_handle, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, - spsv_descr, workspace_ptr, is_alpha_host_accessible); - }; - sycl::accessor workspace_placeholder_acc(workspace); - auto event = dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, - x_handle, y_handle); - event.wait_and_throw(); + if (workspace.size() > 0) { + auto functor = [=](CusparseScopedContextHandler &sc, + sycl::accessor workspace_acc) { + auto cu_handle = sc.get_handle(queue); + auto workspace_ptr = sc.get_mem(workspace_acc); + spsv_optimize_impl(cu_handle, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, + spsv_descr, workspace_ptr, is_alpha_host_accessible); + }; + + // The accessor can only be bound to the cgh if the buffer size is + // greater than 0 + sycl::accessor workspace_placeholder_acc(workspace); + auto event = dispatch_submit(__func__, queue, functor, A_handle, workspace_placeholder_acc, + x_handle, y_handle); + event.wait_and_throw(); + } + else { + auto functor = [=](CusparseScopedContextHandler &sc) { + auto cu_handle = sc.get_handle(queue); + spsv_optimize_impl(cu_handle, opA, alpha, A_view, A_handle, x_handle, y_handle, alg, + spsv_descr, nullptr, is_alpha_host_accessible); + }; + + auto event = dispatch_submit(__func__, queue, functor, A_handle, x_handle, y_handle); + event.wait_and_throw(); + } } sycl::event spsv_optimize(sycl::queue &queue, oneapi::mkl::transpose opA, const void *alpha,