diff options
author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-11-13 22:40:55 +0000 |
---|---|---|
committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-11-13 22:40:55 +0000 |
commit | a5888a730dfb4c11e4386ee56c65df65d706cf7e (patch) | |
tree | ca093ce39ed6e67063d1af4e90884bf4dda55fb4 /clang/lib | |
parent | cfa8aa7edb708155b15c657129f76af9cdb581ba (diff) | |
download | bcm5719-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.cpp | 9 |
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); } |