diff --git a/src/memory/notifier.hpp b/src/memory/notifier.hpp index f01f7f8301..bcc030ce0a 100644 --- a/src/memory/notifier.hpp +++ b/src/memory/notifier.hpp @@ -59,25 +59,25 @@ class Notifier { return; } - uint32_t done = signal_ + 1; + uint32_t done {signal_ + 1}; __syncthreads(); - uint32_t retval = 0; + uint32_t retval {0}; bool executor {!threadIdx.x && !threadIdx.y && !threadIdx.z}; if (executor) { retval = detail::atomic::fetch_add(&count_, 1, orders_); - detail::atomic::threadfence(); + fence(); } __syncthreads(); if (retval == ((gridDim.x * gridDim.y * gridDim.z) - 1)) { if (executor) { detail::atomic::store(&count_, 0, orders_); - detail::atomic::threadfence(); - auto x = detail::atomic::fetch_add(&signal_, 1, orders_); - detail::atomic::threadfence(); + fence(); + detail::atomic::fetch_add(&signal_, 1, orders_); } } + while (detail::atomic::load(&signal_, orders_) != done) { ; }