tree 04af855e820f8941f345569523861664e46624d3
parent 270adfa61aae21ac9c1a643f6f4a39a50601e6e2
author Alexey Bataev <a.bataev@hotmail.com> 1546537426 +0000
committer Alexey Bataev <a.bataev@hotmail.com> 1546537426 +0000

[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