[OPENMP][NVPTX]Use __syncwarp() to reconverge the threads.
Summary:
In Cuda 9.0 it is not guaranteed that threads in the warps are
convergent. We need to use __syncwarp() function to reconverge
the threads and to guarantee the memory ordering among threads in the
warps.
This is the first patch to fix the problem with the test
libomptarget/deviceRTLs/nvptx/src/sync.cu on Cuda9+.
This patch just replaces calls to __shfl_sync() function with the call
of __syncwarp() function where we need to reconverge the threads when we
try to modify the value of the parallel level counter.
Reviewers: grokos
Subscribers: guansong, jfb, jdoerfert, caomhin, kkwli0, openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D65013
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@369796 91177308-0d34-0410-b5e6-96231b3b80d8
diff --git a/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
index f28284d..a5e4a71 100644
--- a/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ b/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -55,11 +55,14 @@
#define __SHFL_DOWN_SYNC(mask, var, delta, width) \
__shfl_down_sync((mask), (var), (delta), (width))
#define __ACTIVEMASK() __activemask()
+#define __SYNCWARP(Mask) __syncwarp(Mask)
#else
#define __SHFL_SYNC(mask, var, srcLane) __shfl((var), (srcLane))
#define __SHFL_DOWN_SYNC(mask, var, delta, width) \
__shfl_down((var), (delta), (width))
#define __ACTIVEMASK() __ballot(1)
+// In Cuda < 9.0 no need to sync threads in warps.
+#define __SYNCWARP(Mask)
#endif // CUDA_VERSION
#define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
diff --git a/libomptarget/deviceRTLs/nvptx/src/supporti.h b/libomptarget/deviceRTLs/nvptx/src/supporti.h
index ceb3951..ceed7d3 100644
--- a/libomptarget/deviceRTLs/nvptx/src/supporti.h
+++ b/libomptarget/deviceRTLs/nvptx/src/supporti.h
@@ -202,25 +202,31 @@
// Parallel level
INLINE void IncParallelLevel(bool ActiveParallel) {
- unsigned tnum = __ACTIVEMASK();
- int leader = __ffs(tnum) - 1;
- __SHFL_SYNC(tnum, leader, leader);
- if (GetLaneId() == leader) {
+ unsigned Active = __ACTIVEMASK();
+ __SYNCWARP(Active);
+ unsigned LaneMaskLt;
+ asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
+ unsigned Rank = __popc(Active & LaneMaskLt);
+ if (Rank == 0) {
parallelLevel[GetWarpId()] +=
(1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
+ __threadfence();
}
- __SHFL_SYNC(tnum, leader, leader);
+ __SYNCWARP(Active);
}
INLINE void DecParallelLevel(bool ActiveParallel) {
- unsigned tnum = __ACTIVEMASK();
- int leader = __ffs(tnum) - 1;
- __SHFL_SYNC(tnum, leader, leader);
- if (GetLaneId() == leader) {
+ unsigned Active = __ACTIVEMASK();
+ __SYNCWARP(Active);
+ unsigned LaneMaskLt;
+ asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
+ unsigned Rank = __popc(Active & LaneMaskLt);
+ if (Rank == 0) {
parallelLevel[GetWarpId()] -=
(1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
+ __threadfence();
}
- __SHFL_SYNC(tnum, leader, leader);
+ __SYNCWARP(Active);
}
////////////////////////////////////////////////////////////////////////////////
