[Openmp-commits] [PATCH] D63515: [OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0.

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added a reviewer: grokos.
Herald added subscribers: openmp-commits, jdoerfert, guansong.
Herald added a project: OpenMP.

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


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D63515

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


Index: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -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 {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D63515.205427.patch
Type: text/x-patch
Size: 648 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190618/5bfc8225/attachment.bin>