[Openmp-commits] [openmp] r363807 - [OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Wed Jun 19 07:20:34 2019
New Revision: 363807

URL: http://llvm.org/viewvc/llvm-project?rev=363807&view=rev
Log:
[OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0.

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

Reviewers: grokos

Subscribers: guansong, jdoerfert, openmp-commits, kkwli0, caomhin

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D63515

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

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h?rev=363807&r1=363806&r2=363807&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h Wed Jun 19 07:20:34 2019
@@ -61,7 +61,12 @@
 #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 {