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

Improve barrier implementation

Browse Source 
[ROCm/rccl commit: adb8d63352]
This commit is contained in:
Wenkai Du
2021-09-08 16:14:32 -05:00
parent d75504e9dc
commit d2580c8cf5
2 changed files with 12 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/collectives/device/primitives.h
+10 -6
View File
@@ -33,12 +33,16 @@
} while (0)
#define barrier_by_group() do { \
const int w = threadIdx.x/WARP_SIZE; \
const int wid = threadIdx.x%WARP_SIZE; \
if (wid == 0) { \
barrier_next[w] += nthreads/WARP_SIZE; \
__atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \
while (LOAD(barriers) < barrier_next[w]) /* spin */; \
if (nthreads == NCCL_MAX_NTHREADS) \
__syncthreads(); \
else { \
const int w = threadIdx.x/WARP_SIZE; \
const int wid = threadIdx.x%WARP_SIZE; \
if (wid == 0) { \
barrier_next[w] += nthreads/WARP_SIZE; \
__atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \
while (LOAD(barriers) < barrier_next[w]) /* spin */; \
} \
} \
} while (0)
projects/rccl/src/include/devcomm.h
+2 -2
View File
@@ -18,8 +18,8 @@
// Convert volatile access to atomic
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST)
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST)
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_ACQUIRE)
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_RELEASE)
#else
#define LOAD(VAR) *(VAR)
#define STORE(DST, SRC) *(DST) = (SRC)