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;
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();
     }
   }
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));
       }
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));
       }
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));
     }