[OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0.

[ABataev]

Summary

The problems with syncthreads() were fixed in clang >= 9.0 and the
original syncthreads() can be used instead of the ptx instruction.

[hfinkel]

LGTM. Thanks!

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

64 ↗

(On Diff #205427)

Please end comments with a period.

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

 #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)
 #endif
 
 #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)
+#endif
 
 // arguments needed for L0 parallelism only.
 class omptarget_nvptx_SharedArgs {
 public:
   // All these methods must be called by the master thread only.
   INLINE void Init() {
     args  = buffer;
     nArgs = MAX_SHARED_ARGS;