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

Alexey Bataev via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Wed Jun 19 07:18:45 PDT 2019


This revision was automatically updated to reflect the committed changes.
Closed by commit rL363807: [OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0. (authored by ABataev, committed by ).
Herald added a project: LLVM.
Herald added a subscriber: llvm-commits.

Changed prior to commit:
  https://reviews.llvm.org/D63515?vs=205427&id=205585#toc

Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D63515/new/

https://reviews.llvm.org/D63515

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


Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ openmp/trunk/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.205585.patch
Type: text/x-patch
Size: 688 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190619/038f0dbe/attachment.bin>


More information about the Openmp-commits mailing list