summaryrefslogtreecommitdiffstats
path: root/clang/lib
diff options
context:
space:
mode:
authorMatt Arsenault <Matthew.Arsenault@amd.com>2017-11-13 22:40:55 +0000
committerMatt Arsenault <Matthew.Arsenault@amd.com>2017-11-13 22:40:55 +0000
commita5888a730dfb4c11e4386ee56c65df65d706cf7e (patch)
treeca093ce39ed6e67063d1af4e90884bf4dda55fb4 /clang/lib
parentcfa8aa7edb708155b15c657129f76af9cdb581ba (diff)
downloadbcm5719-llvm-a5888a730dfb4c11e4386ee56c65df65d706cf7e.tar.gz
bcm5719-llvm-a5888a730dfb4c11e4386ee56c65df65d706cf7e.zip
OpenCL: Assume inline asm is convergent
Already done for CUDA. llvm-svn: 318098
Diffstat (limited to 'clang/lib')
-rw-r--r--clang/lib/CodeGen/CGStmt.cpp9
1 files changed, 5 insertions, 4 deletions
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 6a78865fc94..7b1afab5f2c 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -2149,10 +2149,11 @@ void CodeGenFunction::EmitAsmStmt(const AsmStmt &S) {
llvm::ConstantAsMetadata::get(Loc)));
}
- if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
- // Conservatively, mark all inline asm blocks in CUDA as convergent
- // (meaning, they may call an intrinsically convergent op, such as bar.sync,
- // and so can't have certain optimizations applied around them).
+ if (getLangOpts().assumeFunctionsAreConvergent()) {
+ // Conservatively, mark all inline asm blocks in CUDA or OpenCL as
+ // convergent (meaning, they may call an intrinsically convergent op, such
+ // as bar.sync, and so can't have certain optimizations applied around
+ // them).
Result->addAttribute(llvm::AttributeList::FunctionIndex,
llvm::Attribute::Convergent);
}
OpenPOWER on IntegriCloud