Merge pull request #621 from whchung/exp_reduce_sleep_cycles

Reduce s_sleep cycles

src/collectives/device/primitives.h

@@ -23,7 +23,7 @@
     if (wid == 0) { \
       barrier_next[w] += nthreads/WARP_SIZE; \
       atomicAdd((unsigned long long *)barriers, 1); \
-      while (atomicAdd((unsigned long long *)barriers, 0) < barrier_next[w]) __builtin_amdgcn_s_sleep(8); \
+      while (atomicAdd((unsigned long long *)barriers, 0) < barrier_next[w]) __builtin_amdgcn_s_sleep(1); \
       __asm__ __volatile__("s_wakeup"); \
     } \
   } \

src/collectives/device/prims_ll.h

@@ -99,7 +99,7 @@ private:
     if (sendConnHeadPtr) {
       int spins = 0;
       while (sendConnHeadCache + NCCL_STEPS < sendConnHead + 1) {
-        __builtin_amdgcn_s_sleep(8);
+        __builtin_amdgcn_s_sleep(1);
         sendConnHeadCache = atomicAdd_system((unsigned long long *)sendConnHeadPtr, 0);
         if (checkAbort(spins, 1)) break;
       }

src/collectives/device/prims_ll128.h

@@ -88,7 +88,7 @@ private:
     if (sendConnHeadPtr) {
       int spins = 0;
       while (sendConnHeadCache + NCCL_STEPS < sendConnHead + 1) {
-        __builtin_amdgcn_s_sleep(8);
+        __builtin_amdgcn_s_sleep(1);
         sendConnHeadCache = atomicAdd_system((unsigned long long *)sendConnHeadPtr, 0);
         if (checkAbort(spins, wid, 1)) break;
       }

src/collectives/device/prims_simple.h

@@ -105,7 +105,7 @@ private:
         ((flags & (Send*RoleWaitSend)) && !noSendWait)) {
       int spins = 0;
       while (connStepCache + (isSendNotRecv ? NCCL_STEPS : 0) < step + StepPerSlice) {
-        __builtin_amdgcn_s_sleep(8);
+        __builtin_amdgcn_s_sleep(1);
         connStepCache = atomicAdd_system((unsigned long long *)connStepPtr, 0);
         if (checkAbort(spins)) break;
         //if (spins == 0) printf("r=%d b=%d t=%d SPUN OUT got=%d want=%d\n", ncclShmem.comm.rank, blockIdx.x, threadIdx.x, int(connStepCache + (isSendNotRecv ? NCCL_STEPS : 0)), int(step+StepPerSlice));