Mark all CUDA device-side function defs and decls as convergent.

[jlebar]

Summary

This is important for e.g. the following case:

void sync() { __syncthreads(); }
void foo() {
  do_something();
  sync();
  do_something_else():
}

Without this change, if the optimizer does not inline sync() (which it
won't because syncthreads is also marked as noduplicate, for now
anyway), it is free to perform optimizations on sync() that it would not
be able to perform on syncthreads(), because sync() is not marked as
convergent.

This chagne is conservative; the optimizer will remove these attrs where
it can, see r260318, r260319.

[tra]

lib/CodeGen/CodeGenModule.cpp

1880

intrinsically?

test/CodeGenCUDA/convergent.cu

4

You may want to add -disable-llvm-passes to make sure you're checking results of front-end actions only.

[jlebar]

Update per tra's review.

[majnemer]

SetLLVMFunctionAttributesForDefinition seems like the more obvious place
for this logic.

[jlebar]

Move code into SetLLVMFunctionAttributesForDefinition.

[jlebar]

In D17056#355198, @jlebar wrote:

Move code into SetLLVMFunctionAttributesForDefinition.

Actually, this doesn't work -- we don't annotate

__host__ __device__ void baz();

as convergent. (I ran the tests, but of course I didn't notice it failing...)

[jlebar]

Move coded into SetLLVMFunctionAttributes (not ForDefinition).

[jlebar]

In D17056#355228, @jlebar wrote:

Move coded into SetLLVMFunctionAttributes (not ForDefinition).

Much better. :)

[majnemer]

lib/CodeGen/CodeGenModule.cpp

817–823

Sorry for the goose chase...

ConstructAttributeList seems better.

[jlebar]

Move code into ConstructAttributeList. Now it applies to both functions and calls.

lib/CodeGen/CodeGenModule.cpp

817–823

Actually, this is much better; now I don't need D17313.

[jlebar]

Friendly ping -- are we happy with this?

[majnemer]

LGTM

lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::SetLLVMFunctionAttributes(const Decl *D,
                                               llvm::Function *F) {
   unsigned CallingConv;
   AttributeListType AttributeList;
   ConstructAttributeList(F->getName(), Info, D, AttributeList, CallingConv,
                          false);
   F->setAttributes(llvm::AttributeSet::get(getLLVMContext(), AttributeList));
   F->setCallingConv(static_cast<llvm::CallingConv::ID>(CallingConv));
 }
 
 /// Determines whether the language options require us to model
 /// unwind exceptions.  We treat -fexceptions as mandating this
 /// except under the fragile ObjC ABI with only ObjC exceptions
 /// enabled.  This means, for example, that C with -fexceptions
 /// enables this.
 static bool hasUnwindExceptions(const LangOptions &LangOpts) {
   // If exceptions are completely disabled, obviously this is false.
   if (!LangOpts.Exceptions) return false;
 
   // If C++ exceptions are enabled, this is true.
   if (LangOpts.CXXExceptions) return true;
 
   // If ObjC exceptions are enabled, this depends on the ABI.
   if (LangOpts.ObjCExceptions) {
CodeGenModule::GetOrCreateLLVMFunction(StringRef MangledName,
   if (ExtraAttrs.hasAttributes(llvm::AttributeSet::FunctionIndex)) {
     llvm::AttrBuilder B(ExtraAttrs, llvm::AttributeSet::FunctionIndex);
     F->addAttributes(llvm::AttributeSet::FunctionIndex,
                      llvm::AttributeSet::get(VMContext,
                                              llvm::AttributeSet::FunctionIndex,
                                              B));
   }
 
+  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+    // Conservatively, mark all functions in CUDA as convergent (meaning, they
+    // may call an intrinsicly 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);
+  }
+
   if (!DontDefer) {
     // All MSVC dtors other than the base dtor are linkonce_odr and delegate to
     // each other bottoming out with the base dtor.  Therefore we emit non-base
     // dtors on usage, even if there is no dtor definition in the TU.
     if (D && isa<CXXDestructorDecl>(D) &&
         getCXXABI().useThunkForDtorVariant(cast<CXXDestructorDecl>(D),
                                            GD.getDtorType()))
       addDeferredDeclToEmit(F, GD);

test/CodeGenCUDA/convergent.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
+// RUN:   -o - %s | FileCheck -check-prefix DEVICE %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | \
+// RUN:  FileCheck -check-prefix HOST %s
+
+#include "Inputs/cuda.h"
+
+// DEVICE: Function Attrs:
+// DEVICE-SAME: convergent
+// DEVICE-NEXT: define void @_Z3foov
+__device__ void foo() {}
+
+// HOST: Function Attrs:
+// HOST-NOT: convergent
+// HOST-NEXT: define void @_Z3barv
+// DEVICE: Function Attrs:
+// DEVICE-SAME: convergent
+// DEVICE-NEXT: define void @_Z3barv
+__host__ __device__ void baz();
+__host__ __device__ void bar() { baz(); }
+
+// DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// DEVICE: attributes [[BAZ_ATTR]] = {
+// DEVICE-SAME: convergent
+// DEVICE-SAME: }
+
+// HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// HOST: attributes [[BAZ_ATTR]] = {
+// HOST-NOT: convergent
+// NOST-SAME: }