From 45c29e7734385ceb5de6ae26db197e208c543752 Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Fri, 16 Aug 2024 12:33:13 -0700 Subject: [PATCH] Minor updates to Nofifier sync method --- src/memory/notifier.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) 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) { ; }