commit 222ab3410ac5ef4f146e3a16d636e812e9aadacc
author Kyle Roarty <kyleroarty1716@gmail.com> Fri Jul 23 13:34:19 2021 -0500
committer Matt Sinclair <mattdsinclair@gmail.com> Mon Jul 26 18:17:29 2021 +0000
tree f816f00130396a724e020dec38d6ce3ee7242f7d
parent 9366e6d0857a77211ae15be475a0099163898f12

resources: Update heterosync to use s_sleep

Previously gem5 didn't support the s_sleep instruction used by
heterosync. This patch updates heterosync to use the instruction instead
of a spin loop

Change-Id: Ib333d9c361985b863fa3e57464c90e4c548d8c27
Reviewed-on: https://gem5-review.googlesource.com/c/public/gem5-resources/+/48483
Reviewed-by: Matt Sinclair <mattdsinclair@gmail.com>
Reviewed-by: Bobby R. Bruce <bbruce@ucdavis.edu>
Maintainer: Matt Sinclair <mattdsinclair@gmail.com>
Maintainer: Bobby R. Bruce <bbruce@ucdavis.edu>
Tested-by: Bobby R. Bruce <bbruce@ucdavis.edu>

src/gpu/heterosync/src/hipLocks.h

diff --git a/src/gpu/heterosync/src/hipLocks.h b/src/gpu/heterosync/src/hipLocks.h
index 690ce4f..2a8dafd 100644
--- a/src/gpu/heterosync/src/hipLocks.h
+++ b/src/gpu/heterosync/src/hipLocks.h
@@ -7,9 +7,6 @@
 /*
   Shared sleep function.  Since s_sleep only takes in consstants (between 1 and 128),
   need code to handle long tail.
-
-  Currently s_sleep is unsupported in gem5, so sleepFunc is commented out and
-  replaced with a spin in the lock implementations
  */
 inline __device__ void sleepFunc(int backoff) {
   int backoffCopy = backoff;

src/gpu/heterosync/src/hipLocksBarrierAtomic.h

diff --git a/src/gpu/heterosync/src/hipLocksBarrierAtomic.h b/src/gpu/heterosync/src/hipLocksBarrierAtomic.h
index a51a77f..38fbc9d 100644
--- a/src/gpu/heterosync/src/hipLocksBarrierAtomic.h
+++ b/src/gpu/heterosync/src/hipLocksBarrierAtomic.h
@@ -56,8 +56,7 @@
     // do exponential backoff to reduce the number of times we pound the global
     // barrier
     if (!*done) {
-      //sleepFunc(backoff);
-      for (int j = 0; j < backoff; ++j) { ; }
+      sleepFunc(backoff);
       __syncthreads();
     }
   }

src/gpu/heterosync/src/hipLocksMutexEBO.h

diff --git a/src/gpu/heterosync/src/hipLocksMutexEBO.h b/src/gpu/heterosync/src/hipLocksMutexEBO.h
index 0adaac0..69ab38d 100644
--- a/src/gpu/heterosync/src/hipLocksMutexEBO.h
+++ b/src/gpu/heterosync/src/hipLocksMutexEBO.h
@@ -43,8 +43,7 @@
       {
         // if we failed in acquiring the lock, wait for a little while before
         // trying again
-        //sleepFunc(backoff);
-        for (int j = 0; j < backoff; ++j) { ; }
+        sleepFunc(backoff);
         // (capped) exponential backoff
         backoff = (((backoff << 1) + 1) & (MAX_BACKOFF-1));
       }
@@ -100,8 +99,7 @@
       {
         // if we failed in acquiring the lock, wait for a little while before
         // trying again
-        //sleepFunc(backoff);
-        for (int j = 0; j < backoff; ++j) { ; }
+        sleepFunc(backoff);
         // (capped) exponential backoff
         backoff = (((backoff << 1) + 1) & (MAX_BACKOFF-1));
       }

src/gpu/heterosync/src/hipLocksMutexSleep.h

diff --git a/src/gpu/heterosync/src/hipLocksMutexSleep.h b/src/gpu/heterosync/src/hipLocksMutexSleep.h
index b9a1461..c49d401 100644
--- a/src/gpu/heterosync/src/hipLocksMutexSleep.h
+++ b/src/gpu/heterosync/src/hipLocksMutexSleep.h
@@ -79,8 +79,7 @@
       {
         // if we failed in acquiring the lock, wait for a little while before
         // trying again
-        //sleepFunc(backoff);
-        for (int j = 0; j < backoff; ++j) { ; }
+        sleepFunc(backoff);
         // (capped) exponential backoff
         backoff = (((backoff << 1) + 1) & (MAX_BACKOFF-1));
       }

src/gpu/heterosync/src/hipLocksSemaphoreEBO.h

diff --git a/src/gpu/heterosync/src/hipLocksSemaphoreEBO.h b/src/gpu/heterosync/src/hipLocksSemaphoreEBO.h
index 0128de3..69520be 100644
--- a/src/gpu/heterosync/src/hipLocksSemaphoreEBO.h
+++ b/src/gpu/heterosync/src/hipLocksSemaphoreEBO.h
@@ -162,8 +162,7 @@
     {
       // if we failed to enter the semaphore, wait for a little while before
       // trying again
-      //sleepFunc(backoff);
-      for (int j = 0; j < backoff; ++j) { ; }
+      sleepFunc(backoff);
       /*
         for writers increse backoff a lot because failing means readers are in
         the CS currently -- most important for non-unique because all WGs on
@@ -385,8 +384,7 @@
         if we failed to enter the semaphore, wait for a little while before
         trying again
       */
-      //sleepFunc(backoff);
-      for (int j = 0; j < backoff; ++j) { ; }
+      sleepFunc(backoff);
       // (capped) exponential backoff
       backoff = (((backoff << 1) + 1) & (MAX_BACKOFF-1));
     }