[OPENMP][NVPTX]Use __syncwarp() to reconverge the threads.

[ABataev]

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.

[grokos]

Does this behavior of CUDA >= 9.0 affect only the parallel level counter? Do we need to propagate these changes to other functions?

[ABataev]

In D65013#1595977, @grokos wrote:

Does this behavior of CUDA >= 9.0 affect only the parallel level counter? Do we need to propagate these changes to other functions?

Seems to me, it affects only parallelLevel counter because currently only parallelLevel counter is used on per-warp level.

[grokos]

OK, looks good.

[jdoerfert]

I'm confused, partly about the "convergent" part.

The code looks vastly different but no tests are affected?
Could you please point out how to reproduce the problem?
Where did the shuffles go?
Why is there a threadfence and syncwrap now?
Which old accesses were problematic and why?

[ABataev]

In D65013#1596821, @jdoerfert wrote:

I'm confused, partly about the "convergent" part.

The code looks vastly different but no tests are affected?
Could you please point out how to reproduce the problem?
Where did the shuffles go?
Why is there a threadfence and syncwrap now?
Which old accesses were problematic and why?

There is a problem with at least 1 test in Cuda 9+: spmd_parallel_regions.cpp. To fix this problem we need 3 things: fix the test itself (see D65112), fix the runtime part (this patch) and fix the handling of critical sections in compiler (the 3rd patch that depends on this one).

[jdoerfert]

In D65013#1596990, @ABataev wrote:

In D65013#1596821, @jdoerfert wrote:

I'm confused, partly about the "convergent" part.

The code looks vastly different but no tests are affected?
Could you please point out how to reproduce the problem?
Where did the shuffles go?
Why is there a threadfence and syncwrap now?
Which old accesses were problematic and why?

There is a problem with at least 1 test in Cuda 9+: spmd_parallel_regions.cpp. To fix this problem we need 3 things: fix the test itself (see D65112), fix the runtime part (this patch) and fix the handling of critical sections in compiler (the 3rd patch that depends on this one).

There seems to be a problem with this "fix", not the test. At least so far, the argument was CUDA 9 semantics which should be irrelevant to the test. If there is a problem, than that the runtime doesn't implement OpenMP semantics properly for that test. Modifying the test will only hide that problem.

[ABataev]

Reworked to fix the test spmd_parallel_regions.cpp and fix problems with SPMD mode in CUDA9+ found during testing.

[ABataev]

Rebase.

[ABataev]

Ping

[ABataev]

Ping

[jdoerfert]

Generally, this seems fine but I was hoping we could say what configuration and test file can be used to reproduce this error.

[ABataev]

In D65013#1643020, @jdoerfert wrote:

Generally, this seems fine but I was hoping we could say what configuration and test file can be used to reproduce this error.

Generally speaking, we just switched to __syncwarp function instead of shfl_sync in this patch (In cuda <= 8 we just don't need to reconverge the threads because of the architecture, in Cuda 9+ there is a special function __syncwarp for this). It will just improve performance in Cuda 8 and won't affect Cuda9+ at all.
The problem is in threads divergence within the warp. We use parallel level counter on per-warp basis (because we're very limited in shared memory). Previously, it was even worse, we had parallel level counter on per-block basis.

I think, it would be better to rename the patch. Just like I said, this is not a fix, just a small improvement in the way we reconverge the threads.

[ABataev]

Renamed.

[jdoerfert]

thx for the explanation. LGTM.

[jdoerfert]

Make sure to update the commit message as well.

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 // Also, __ballot(1) in Cuda 8.0 is replaced with __activemask().
 #ifndef CUDA_VERSION
 #error CUDA_VERSION macro is undefined, something wrong with cuda.
 #elif CUDA_VERSION >= 9000
 #define __SHFL_SYNC(mask, var, srcLane) __shfl_sync((mask), (var), (srcLane))
 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down_sync((mask), (var), (delta), (width))
 #define __ACTIVEMASK() __activemask()
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
+#define __SYNCWARP() __syncwarp()
+#else
+// For .target sm_6x or below, all threads in mask must execute the same
+// __syncwarp() in convergence, and the union of all values in mask must be
+// equal to the active mask. Otherwise, the behavior is undefined.
+// (https://docs.nvidia.com/cuda/archive/9.0/cuda-c-programming-guide/index.html#synchronization-functions)
+#define __SYNCWARP() __syncwarp(__activemask())
+#endif // __CUDA_ARCH__
 #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()
 #endif // CUDA_VERSION
 
 #define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
 // Use original __syncthreads if compiled by nvcc or clang >= 9.0.
 #if !defined(__clang__) || __clang_major__ >= 9
 #define __SYNCTHREADS() __syncthreads()
 #else
 #define __SYNCTHREADS() __SYNCTHREADS_N(0)

libomptarget/deviceRTLs/nvptx/src/supporti.h

 // Masters
 
 INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Parallel level
 
 INLINE void IncParallelLevel(bool ActiveParallel) {
-  unsigned tnum = __ACTIVEMASK();
-  int leader = __ffs(tnum) - 1;
-  __SHFL_SYNC(tnum, leader, leader);
-  if (GetLaneId() == leader) {
+  __SYNCWARP();
+  unsigned Active = __ACTIVEMASK();
+  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();
 }
 
 INLINE void DecParallelLevel(bool ActiveParallel) {
-  unsigned tnum = __ACTIVEMASK();
-  int leader = __ffs(tnum) - 1;
-  __SHFL_SYNC(tnum, leader, leader);
-  if (GetLaneId() == leader) {
+  __SYNCWARP();
+  unsigned Active = __ACTIVEMASK();
+  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();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // get OpenMP number of procs
 
 // Get the number of processors in the device.
 INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
   if (!isSPMDExecutionMode)