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

Remove extra L1 cache invalidate and restore __ATOMIC_SEQ_CST atomics (#426)

Browse Source 
[ROCm/rccl commit: 5ae3f3f954]
This commit is contained in:
Wenkai Du
2021-09-14 18:30:16 -07:00
committed by GitHub
parent f4387b2954
commit 1eaf495391
1 changed files with 6 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/collectives/device/primitives.h
+6 -2
View File
@@ -136,7 +136,7 @@ class ncclPrimitives {
spins = 0;
while (connHeadCache + NCCL_STEPS < step + SLICESTEPS) {
__builtin_amdgcn_s_sleep(8);
connHeadCache = atomicAdd_system((unsigned long long *)connHeadPtr, 0);
connHeadCache = LOAD(connHeadPtr);
if (checkAbort()) break;
}
__asm__ __volatile__("s_wakeup");
@@ -158,7 +158,7 @@ class ncclPrimitives {
#endif
while (connTailCache < step + SLICESTEPS) {
__builtin_amdgcn_s_sleep(8);
connTailCache = atomicAdd_system((unsigned long long *)connTailPtr, 0);
connTailCache =LOAD(connTailPtr);
if (checkAbort()) break;
}
__asm__ __volatile__("s_wakeup");
@@ -207,7 +207,11 @@ class ncclPrimitives {
}
}
barrier();
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
// L1 cache invalidation by __threadfence_system() is already done in above barrier()
#else
if (SEND && (role & ROLE_POST_SEND) && realSize > 0 && index == 0) __threadfence_system();
#endif
__syncwarp();
if (SEND && (role & ROLE_POST_SEND)) postSend();
if (RECV && (role & ROLE_POST_RECV)) postRecv();