diff options
| author | Alexey Bataev <a.bataev@hotmail.com> | 2019-01-03 17:43:46 +0000 |
|---|---|---|
| committer | Alexey Bataev <a.bataev@hotmail.com> | 2019-01-03 17:43:46 +0000 |
| commit | c0c737ac56d8031e83118d231bb273a54493a489 (patch) | |
| tree | 04af855e820f8941f345569523861664e46624d3 /libomptarget/deviceRTLs/nvptx/src/sync.cu | |
| parent | 270adfa61aae21ac9c1a643f6f4a39a50601e6e2 (diff) | |
[OPENMP][NVPTX]Fix incompatibility of __syncthreads with LLVM, NFC.
Summary:
One of the LLVM optimizations, split critical edges, also clones tail
instructions. This is a dangerous operation for __syncthreads()
functions and this transformation leads to undefined behavior or
incorrect results. Patch fixes this problem by replacing __syncthreads()
function with the assembler instruction, which cost is too high and
wich cannot be copied.
Reviewers: grokos, gtbercea, kkwli0
Subscribers: guansong, openmp-commits, caomhin
Differential Revision: https://reviews.llvm.org/D56274
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@350333 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'libomptarget/deviceRTLs/nvptx/src/sync.cu')
| -rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/sync.cu | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/libomptarget/deviceRTLs/nvptx/src/sync.cu b/libomptarget/deviceRTLs/nvptx/src/sync.cu index 7cdb7ff..5f6aef9 100644 --- a/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ b/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -74,7 +74,8 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) { // parallel region and that all worker threads participate. EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) { PRINT0(LD_SYNC, "call kmpc_barrier_simple_spmd\n"); - __syncthreads(); + // FIXME: use __syncthreads instead when the function copy is fixed in LLVM. + __SYNCTHREADS(); PRINT0(LD_SYNC, "completed kmpc_barrier_simple_spmd\n"); } |