[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