[CUDA] Conservatively mark inline asm as convergent.

[jlebar]

Summary

This is particularly important because a some convergent CUDA intrinsics
(e.g. __shfl_down) are implemented in terms of inline asm.

[tra]

I guess we would not be able to remove convergent from inline asm automatically. Do we need a way to explicitly remove convergent from inline asm?

[jlebar]

In D20836#444911, @tra wrote:

I guess we would not be able to remove convergent from inline asm automatically. Do we need a way to explicitly remove convergent from inline asm?

We can think about it. I'm not sure it will make a big difference, frankly. Like, if this encourages people to write less inline asm, I'm onboard with that. :)

cfe/trunk/lib/CodeGen/CGStmt.cpp

void CodeGenFunction::EmitAsmStmt(const AsmStmt &S) {
   } else {
     // At least put the line number on MS inline asm blobs.
     auto Loc = llvm::ConstantInt::get(Int32Ty, S.getAsmLoc().getRawEncoding());
     Result->setMetadata("srcloc",
                         llvm::MDNode::get(getLLVMContext(),
                                           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) {
     RegResults.push_back(Result);
   } else {
     for (unsigned i = 0, e = ResultRegTypes.size(); i != e; ++i) {
       llvm::Value *Tmp = Builder.CreateExtractValue(Result, i, "asmresult");
       RegResults.push_back(Tmp);

cfe/trunk/test/CodeGenCUDA/convergent.cu

 // HOST-NEXT: define void @_Z3barv
 // DEVICE: Function Attrs:
 // DEVICE-SAME: convergent
 // DEVICE-NEXT: define void @_Z3barv
 __host__ __device__ void baz();
 __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]] = {
 // HOST-NOT: convergent
 // NOST-SAME: }