From 615cfb857ec84e5e65fc4d4d20464628e059daa6 Mon Sep 17 00:00:00 2001
From: "romain.biessy" <romain.biessy@codeplay.com>
Date: Fri, 19 Jul 2024 16:34:31 +0200
Subject: [PATCH] Ensure empty placeholder accessors are avoided
---
.../cusparse/operations/cusparse_spmv.cpp | 3 ++
.../cusparse/operations/cusparse_spsv.cpp | 37 +++++++++++++------
2 files changed, 29 insertions(+), 11 deletions(-)
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<std::uint8_t, 1> 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<std::uint8_t> 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<std::uint8_t, 1> 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<std::uint8_t> 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<std::uint8_t, 1> 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,