[PATCH] D17056: Mark all CUDA device-side function defs and decls as convergent.
David Majnemer via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 17 16:42:52 PST 2016
majnemer added inline comments.
================
Comment at: lib/CodeGen/CodeGenModule.cpp:817-823
@@ -816,1 +816,9 @@
+
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+ // Conservatively, mark all functions in CUDA as convergent (meaning, they
+ // may call an intrinsically convergent op, such as __syncthreads(), and so
+ // can't have certain optimizations applied around them). LLVM will remove
+ // this attribute where it safely can.
+ F->addFnAttr(llvm::Attribute::Convergent);
+ }
}
----------------
Sorry for the goose chase...
`ConstructAttributeList` seems better.
http://reviews.llvm.org/D17056
More information about the cfe-commits
mailing list