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

Reduce s_sleep cycles

Browse Source
This commit is contained in:
Wenkai Du
2022-09-19 19:39:04 +00:00
committed by Wen-Heng (Jack) Chung
parent 49c811ecf9
commit a3c8ef8c03
4 changed files with 4 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
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"); \
} \
} \
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;
}
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;
}
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));