AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #621 from whchung/exp_reduce_sleep_cycles

Browse Source 
Reduce s_sleep cycles

[ROCm/rccl commit: 35f1fe3434]
This commit is contained in:
Wen-Heng (Jack) Chung
2022-09-23 15:31:16 -05:00
committed by GitHub
parent 183f1e6b32 0b56e397cc
commit bd19566413
4 changed files with 4 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/collectives/device/primitives.h
+1 -1
View File
@@ -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"); \
} \
} \
projects/rccl/src/collectives/device/prims_ll.h
+1 -1
View File
@@ -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;
}
projects/rccl/src/collectives/device/prims_ll128.h
+1 -1
View File
@@ -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;
}
projects/rccl/src/collectives/device/prims_simple.h
+1 -1
View File
@@ -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));