[PATCH] D20836: [CUDA] Conservatively mark inline asm as convergent.

[Justin Lebar via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL271336: [CUDA] Conservatively mark inline asm as convergent. (authored by jlebar).

Changed prior to commit:
  http://reviews.llvm.org/D20836?vs=59130&id=59133#toc

Repository:
  rL LLVM

http://reviews.llvm.org/D20836

Files:
  cfe/trunk/lib/CodeGen/CGStmt.cpp
  cfe/trunk/test/CodeGenCUDA/convergent.cu

Index: cfe/trunk/test/CodeGenCUDA/convergent.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/convergent.cu
+++ cfe/trunk/test/CodeGenCUDA/convergent.cu
@@ -25,13 +25,19 @@
 __host__ __device__ void bar() {
   // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
   baz();
+  // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]]
+  int x;
+  asm ("trap;" : "=l"(x));
+  // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]]
+  asm volatile ("trap;");
 }
 
 // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // DEVICE: attributes [[BAZ_ATTR]] = {
 // DEVICE-SAME: convergent
 // DEVICE-SAME: }
 // DEVICE: attributes [[CALL_ATTR]] = { convergent }
+// DEVICE: attributes [[ASM_ATTR]] = { convergent
 
 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // HOST: attributes [[BAZ_ATTR]] = {
Index: cfe/trunk/lib/CodeGen/CGStmt.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGStmt.cpp
+++ cfe/trunk/lib/CodeGen/CGStmt.cpp
@@ -2054,6 +2054,14 @@
                                           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).
+    Result->addAttribute(llvm::AttributeSet::FunctionIndex,
+                         llvm::Attribute::Convergent);
+  }
+
   // Extract all of the register value results from the asm.
   std::vector<llvm::Value*> RegResults;
   if (ResultRegTypes.size() == 1) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D20836.59133.patch
Type: text/x-patch
Size: 1720 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160531/f2aa6569/attachment.bin>