From a798ac7bbd8c780ceb334d9f5ef3c91fa0e55230 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Wed, 22 Mar 2023 11:34:09 -0400 Subject: [PATCH] Explain acquire_team_scratch_space --- core/src/SYCL/Kokkos_SYCL_Instance.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/core/src/SYCL/Kokkos_SYCL_Instance.cpp b/core/src/SYCL/Kokkos_SYCL_Instance.cpp index f60b198578..24f38d7c06 100644 --- a/core/src/SYCL/Kokkos_SYCL_Instance.cpp +++ b/core/src/SYCL/Kokkos_SYCL_Instance.cpp @@ -137,6 +137,9 @@ void SYCLInternal::initialize(const sycl::queue& q) { } int SYCLInternal::acquire_team_scratch_space() { + // Grab the next scratch memory allocation. We must make sure that the last + // kernel using the allocation has completed, so we wait for the event that + // was registered with that kernel. int current_team_scratch = desul::atomic_fetch_inc_mod( &m_current_team_scratch, m_n_team_scratch - 1, desul::MemoryOrderRelaxed(), desul::MemoryScopeDevice());