diff --git a/CHANGELOG.md b/CHANGELOG.md index 692f40d1a..a1d899dec 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -31,6 +31,7 @@ Versioning](http://semver.org/spec/v2.0.0.html). - Celerity warns on excessive calls to `queue::wait()` or `distr_queue::slow_full_sync()` in a long running program. This operation has a much more pronounced performance penalty than its SYCL counterpart (#283) - On systems that do not support device-to-device copies, data is now staged in linearized buffers for better performance (#287) +- Removed the flush_async workaround for newer ACPP versions, keeping compatibility with older versions (#333) - The `access::neighborhood` built-in range mapper now receives a `range` instead of a coordinate list (#292) - Overhauled the [installation](docs/installation.md) and [configuration](docs/configuration.md) documentation (#309) - Celerity will now queue up several command groups in order to combine allocations and elide resize operations. diff --git a/src/backend/sycl_backend.cc b/src/backend/sycl_backend.cc index 40ccc9f6b..4631a96da 100644 --- a/src/backend/sycl_backend.cc +++ b/src/backend/sycl_backend.cc @@ -65,12 +65,17 @@ std::optional delayed_async_event::get_native_executio return m_state->m_event.get_native_execution_time(); } +template +static void try_flush_async(DagManager& dag) { + // AdaptiveCpp (prior to https://github.com/AdaptiveCpp/AdaptiveCpp/pull/1798) does not guarantee that command groups are actually scheduled until an + // explicit await operation, which we cannot insert without blocking the executor loop (see https://github.com/AdaptiveCpp/AdaptiveCpp/issues/599). Instead, + // we explicitly flush the queue to be able to continue using our polling-based approach. + if constexpr(requires { dag.flush_async(); }) { dag.flush_async(); } +} + void flush(sycl::queue& queue) { #if CELERITY_WORKAROUND(ACPP) - // AdaptiveCpp does not guarantee that command groups are actually scheduled until an explicit await operation, which we cannot insert without - // blocking the executor loop (see https://github.com/AdaptiveCpp/AdaptiveCpp/issues/599). Instead, we explicitly flush the queue to be able to continue - // using our polling-based approach. - queue.get_context().AdaptiveCpp_runtime()->dag().flush_async(); + try_flush_async(queue.get_context().AdaptiveCpp_runtime()->dag()); #else (void)queue; #endif