[OpenMP] Apply OpenMP assumptions to applicable call sites

[jhuber6]

Summary

This patch adds OpenMP assumption attributes to call sites in applicable
regions. Currently this applies the caller's assumption attributes to
any calls contained within it. So, if a call occurs inside an OpenMP
assumes region to a function outside that region, we will assume that
call respects the assumptions. This is primarily useful for inline
assembly calls used heavily in the OpenMP GPU device runtime, which
allows us to then make judgements about what the ASM will do.

[jhuber6]

Fix test.

[jdoerfert]

LG, once we can propagate assumptions we can remove this again.

clang/lib/CodeGen/CGCall.cpp

static void AddAttributesFromFunctionProtoType(ASTContext &Ctx,
   if (!FPT)
     return;
 
   if (!isUnresolvedExceptionSpec(FPT->getExceptionSpecType()) &&
       FPT->isNothrow())
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
 }
 
+static void AddAttributesFromAssumes(llvm::AttrBuilder &FuncAttrs,
+                                     const Decl *Callee, const Decl *Caller,
+                                     bool AssumptionOnCallSite) {
+  if (!Callee)
+    return;
+
+  SmallVector<StringRef, 4> Attrs;
+
+  for (const AssumptionAttr *AA : Callee->specific_attrs<AssumptionAttr>())
+    AA->getAssumption().split(Attrs, ",");
+
+  if (Caller && Caller->hasAttrs() && AssumptionOnCallSite)
+    for (const AssumptionAttr *AA : Caller->specific_attrs<AssumptionAttr>())
+      AA->getAssumption().split(Attrs, ",");
+
+  if (!Attrs.empty())
+    FuncAttrs.addAttribute(llvm::AssumptionAttrKey,
+                           llvm::join(Attrs.begin(), Attrs.end(), ","));
+}
+
 bool CodeGenModule::MayDropFunctionReturn(const ASTContext &Context,
                                           QualType ReturnType) {
   // We can't just discard the return value for a record type with a
   // complex destructor or a non-trivially copyable type.
   if (const RecordType *RT =
           ReturnType.getCanonicalType()->getAs<RecordType>()) {
     if (const auto *ClassDecl = dyn_cast<CXXRecordDecl>(RT->getDecl()))
       return ClassDecl->hasTrivialDestructor();

 ///     modules such as CUDA's libdevice.
 ///
 ///   - ConstructAttributeList builds on top of getDefaultFunctionAttributes
 ///     and adds declaration-specific, convention-specific, and
 ///     frontend-specific logic.  The last is of particular importance:
 ///     attributes that restrict how the frontend generates code must be
 ///     added here rather than getDefaultFunctionAttributes.
 ///
-void CodeGenModule::ConstructAttributeList(StringRef Name,
-                                           const CGFunctionInfo &FI,
-                                           CGCalleeInfo CalleeInfo,
-                                           llvm::AttributeList &AttrList,
-                                           unsigned &CallingConv,
-                                           bool AttrOnCallSite, bool IsThunk) {
+void CodeGenModule::ConstructAttributeList(
+    StringRef Name, const CGFunctionInfo &FI, CGCalleeInfo CalleeInfo,
+    llvm::AttributeList &AttrList, unsigned &CallingConv, bool AttrOnCallSite,
+    bool IsThunk, const Decl *Caller) {
   llvm::AttrBuilder FuncAttrs;
   llvm::AttrBuilder RetAttrs;
 
   // Collect function IR attributes from the CC lowering.
   // We'll collect the paramete and result attributes later.
   CallingConv = FI.getEffectiveCallingConvention();
   if (FI.isNoReturn())
     FuncAttrs.addAttribute(llvm::Attribute::NoReturn);
   if (FI.isCmseNSCall())
     FuncAttrs.addAttribute("cmse_nonsecure_call");
 
   // Collect function IR attributes from the callee prototype if we have one.
   AddAttributesFromFunctionProtoType(getContext(), FuncAttrs,
                                      CalleeInfo.getCalleeFunctionProtoType());
 
   const Decl *TargetDecl = CalleeInfo.getCalleeDecl().getDecl();
 
+  // Only attach assumptions to call sites in OpenMP mode.
+  bool AssumptionOnCallSite = getLangOpts().OpenMP && AttrOnCallSite;
+
+  // Attach assumption attributes to the declaration. If this is a call
+  // site, attach assumptions from the caller to the call as well.
+  AddAttributesFromAssumes(FuncAttrs, TargetDecl, Caller, AssumptionOnCallSite);
+
   bool HasOptnone = false;
   // The NoBuiltinAttr attached to the target FunctionDecl.
   const NoBuiltinAttr *NBA = nullptr;
 
   // Collect function IR attributes based on declaration-specific
   // information.
   // FIXME: handle sseregparm someday...
   if (TargetDecl) {
if (TargetDecl->hasAttr<OpenCLKernelAttr>()) {
         // '-cl-uniform-work-group-size' compile option gets a hint
         // to the compiler that the global work-size be a multiple of
         // the work-group size specified to clEnqueueNDRangeKernel
         // (i.e. work groups are uniform).
         FuncAttrs.addAttribute("uniform-work-group-size",
                                llvm::toStringRef(CodeGenOpts.UniformWGSize));
       }
     }
-
-    std::string AssumptionValueStr;
-    for (AssumptionAttr *AssumptionA :
-         TargetDecl->specific_attrs<AssumptionAttr>()) {
-      std::string AS = AssumptionA->getAssumption().str();
-      if (!AS.empty() && !AssumptionValueStr.empty())
-        AssumptionValueStr += ",";
-      AssumptionValueStr += AS;
-    }
-
-    if (!AssumptionValueStr.empty())
-      FuncAttrs.addAttribute(llvm::AssumptionAttrKey, AssumptionValueStr);
   }
 
   // Attach "no-builtins" attributes to:
   // * call sites: both `nobuiltin` and "no-builtins" or "no-builtin-<name>".
   // * definitions: "no-builtins" or "no-builtin-<name>" only.
   // The attributes can come from:
   // * LangOpts: -ffreestanding, -fno-builtin, -fno-builtin-<name>
   // * FunctionDecl attributes: __attribute__((no_builtin(...)))
#endif
   }
 
   // Compute the calling convention and attributes.
   unsigned CallingConv;
   llvm::AttributeList Attrs;
   CGM.ConstructAttributeList(CalleePtr->getName(), CallInfo,
                              Callee.getAbstractInfo(), Attrs, CallingConv,
                              /*AttrOnCallSite=*/true,
-                             /*IsThunk=*/false);
+                             /*IsThunk=*/false, CurFuncDecl);
 
   if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(CurFuncDecl))
     if (FD->hasAttr<StrictFPAttr>())
       // All calls within a strictfp function are marked strictfp
       Attrs = Attrs.addFnAttribute(getLLVMContext(), llvm::Attribute::StrictFP);
 
   // Add call-site nomerge attribute if exists.
   if (InNoMergeAttributedStmt)

clang/lib/CodeGen/CGStmt.cpp

 #include "clang/AST/StmtVisitor.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/DiagnosticSema.h"
 #include "clang/Basic/PrettyStackTrace.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetInfo.h"
 #include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/StringExtras.h"
+#include "llvm/IR/Assumptions.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/MDBuilder.h"
 #include "llvm/Support/SaveAndRestore.h"
 
 using namespace clang;
 using namespace CodeGen;
static void UpdateAsmCallInst(llvm::CallBase &Result, bool HasSideEffect,
   // Attach readnone and readonly attributes.
   if (!HasSideEffect) {
     if (ReadNone)
       Result.addFnAttr(llvm::Attribute::ReadNone);
     else if (ReadOnly)
       Result.addFnAttr(llvm::Attribute::ReadOnly);
   }
 
+  // Attach OpenMP assumption attributes from the caller, if they exist.
+  if (CGF.CGM.getLangOpts().OpenMP) {
+    SmallVector<StringRef, 4> Attrs;
+
+    for (const AssumptionAttr *AA :
+         CGF.CurFuncDecl->specific_attrs<AssumptionAttr>())
+      AA->getAssumption().split(Attrs, ",");
+
+    if (!Attrs.empty())
+      Result.addFnAttr(
+          llvm::Attribute::get(CGF.getLLVMContext(), llvm::AssumptionAttrKey,
+                               llvm::join(Attrs.begin(), Attrs.end(), ",")));
+  }
+
   // Slap the source location of the inline asm into a !srcloc metadata on the
   // call.
   if (const auto *gccAsmStmt = dyn_cast<GCCAsmStmt>(&S))
     Result.setMetadata("srcloc",
                        getAsmSrcLocInfo(gccAsmStmt->getAsmString(), CGF));
   else {
     // At least put the line number on MS inline asm blobs.
     llvm::Constant *Loc =

clang/lib/CodeGen/CGVTables.cpp

void CodeGenFunction::EmitMustTailThunk(GlobalDecl GD,
   llvm::CallInst *Call = Builder.CreateCall(Callee, Args);
   Call->setTailCallKind(llvm::CallInst::TCK_MustTail);
 
   // Apply the standard set of call attributes.
   unsigned CallingConv;
   llvm::AttributeList Attrs;
   CGM.ConstructAttributeList(Callee.getCallee()->getName(), *CurFnInfo, GD,
                              Attrs, CallingConv, /*AttrOnCallSite=*/true,
-                             /*IsThunk=*/false);
+                             /*IsThunk=*/false, CurFuncDecl);
   Call->setAttributes(Attrs);
   Call->setCallingConv(static_cast<llvm::CallingConv::ID>(CallingConv));
 
   if (Call->getType()->isVoidTy())
     Builder.CreateRetVoid();
   else
     Builder.CreateRet(Call);
 

clang/lib/CodeGen/CodeGenModule.h

public:
   /// \param CalleeInfo - The callee information these attributes are being
   /// constructed for. If valid, the attributes applied to this decl may
   /// contribute to the function attributes and calling convention.
   /// \param Attrs [out] - On return, the attribute list to use.
   /// \param CallingConv [out] - On return, the LLVM calling convention to use.
   void ConstructAttributeList(StringRef Name, const CGFunctionInfo &Info,
                               CGCalleeInfo CalleeInfo,
                               llvm::AttributeList &Attrs, unsigned &CallingConv,
-                              bool AttrOnCallSite, bool IsThunk);
+                              bool AttrOnCallSite, bool IsThunk,
+                              const Decl *Caller = nullptr);
 
   /// Adds attributes to F according to our CodeGenOptions and LangOptions, as
   /// though we had emitted it ourselves.  We remove any attributes on F that
   /// conflict with the attributes we add here.
   ///
   /// This is useful for adding attrs to bitcode modules that you want to link
   /// with but don't control, such as CUDA's libdevice.  When linking with such
   /// a bitcode library, you might want to set e.g. its functions'

clang/test/OpenMP/assumes_codegen.cpp

 
 #pragma omp begin assumes ext_lambda_assumption
 int lambda_outer() {
   auto lambda_inner = []() { return 42; };
   return lambda_inner();
 }
 #pragma omp end assumes
 
+void no_assume() {
+  foo();
+}
+
+#pragma omp begin assumes ext_call_site
+void assume() {
+  foo();
+}
+
+void assembly() {
+  asm ("nop");
+}
+#pragma omp end assumes ext_call_site
+
 // AST:      void foo() __attribute__((assume("omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses"))) __attribute__((assume("omp_no_openmp"))) {
 // AST-NEXT: }
 // AST-NEXT: class BAR {
 // AST-NEXT: public:
 // AST-NEXT:     BAR() __attribute__((assume("ompx_range_bar_only"))) __attribute__((assume("ompx_range_bar_only_2"))) __attribute__((assume("omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses"))) __attribute__((assume("omp_no_openmp")))     {
 // AST-NEXT:     }
 // AST-NEXT:     void bar1() __attribute__((assume("ompx_range_bar_only"))) __attribute__((assume("ompx_range_bar_only_2"))) __attribute__((assume("omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses"))) __attribute__((assume("omp_no_openmp")))     {
 // AST-NEXT:     }

 
 #endif
 
 // CHECK: define{{.*}} void @_Z3foov()
 // CHECK-SAME: [[attr0:#[0-9]]]
 // CHECK: define{{.*}} void @_Z3barv()
 // CHECK-SAME: [[attr1:#[0-9]]]
 // CHECK:   call{{.*}} @_ZN3BARC1Ev(%class.BAR*{{.*}} %b)
-// CHECK-SAME: [[attr9:#[0-9]]]
+// CHECK-SAME: [[attr10:#[0-9]]]
 // CHECK: define{{.*}} void @_ZN3BARC1Ev(%class.BAR*{{.*}} %this)
 // CHECK-SAME: [[attr2:#[0-9]]]
 // CHECK:   call{{.*}} @_ZN3BARC2Ev(%class.BAR*{{.*}} %this1)
-// CHECK-SAME: [[attr9]]
+// CHECK-SAME: [[attr10]]
 // CHECK: define{{.*}} void @_ZN3BARC2Ev(%class.BAR*{{.*}} %this)
 // CHECK-SAME: [[attr3:#[0-9]]]
 // CHECK: define{{.*}} void @_Z3bazv()
 // CHECK-SAME: [[attr4:#[0-9]]]
 // CHECK:   call{{.*}} @_ZN3BAZIfEC1Ev(%class.BAZ*{{.*}} %b)
-// CHECK-SAME: [[attr10:#[0-9]]]
+// CHECK-SAME: [[attr11:#[0-9]]]
 // CHECK: define{{.*}} void @_ZN3BAZIfEC1Ev(%class.BAZ*{{.*}} %this)
 // CHECK-SAME: [[attr5:#[0-9]]]
 // CHECK:   call{{.*}} @_ZN3BAZIfEC2Ev(%class.BAZ*{{.*}} %this1)
-// CHECK-SAME: [[attr10]]
+// CHECK-SAME: [[attr12:#[0-9]]]
 // CHECK: define{{.*}} void @_ZN3BAZIfEC2Ev(%class.BAZ*{{.*}} %this)
 // CHECK-SAME: [[attr6:#[0-9]]]
 // CHECK: define{{.*}} i32 @_Z12lambda_outerv()
 // CHECK-SAME: [[attr7:#[0-9]]]
 // CHECK: call{{.*}} @"_ZZ12lambda_outervENK3$_0clEv"
-// CHECK-SAME: [[attr11:#[0-9]]]
+// CHECK-SAME: [[attr13:#[0-9]]]
 // CHECK: define{{.*}} i32 @"_ZZ12lambda_outervENK3$_0clEv"(%class.anon*{{.*}} %this)
 // CHECK-SAME: [[attr8:#[0-9]]]
+// CHECK: define{{.*}} void @_Z9no_assumev()
+// CHECK-SAME: [[attr0:#[0-9]]]
+// CHECK: call{{.*}} @_Z3foov()
+// CHECK-SAME: [[attr14:#[0-9]]]
+// CHECK: define{{.*}} void @_Z6assumev()
+// CHECK-SAME: [[attr9:#[0-9]]]
+// CHECK: call{{.*}} @_Z3foov()
+// CHECK-SAME: [[attr15:#[0-9]]]
+// CHECK: define{{.*}} void @_Z8assemblyv()
+// CHECK-SAME: [[attr9:#[0-9]]]
+// CHECK: call{{.*}} void asm sideeffect "nop", "~{dirflag},~{fpsr},~{flags}"()
+// CHECK-SAME: [[attr16:#[0-9]]]
 
 // CHECK:     attributes [[attr0]]
 // CHECK-SAME:  "llvm.assume"="omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr1]]
 // CHECK-SAME:  "llvm.assume"="ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr2]]
 // CHECK-SAME:  "llvm.assume"="ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr3]]
 // CHECK-SAME:  "llvm.assume"="ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr4]]
 // CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr5]]
 // CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr6]]
 // CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr7]]
 // CHECK-SAME:  "llvm.assume"="ompx_lambda_assumption,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr8]]
 // CHECK-SAME:  "llvm.assume"="ompx_lambda_assumption,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr9]]
-// CHECK-SAME:  "llvm.assume"="ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK-SAME:  "llvm.assume"="ompx_call_site,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr10]]
-// CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK-SAME:  "llvm.assume"="ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_range_bar_only,ompx_range_bar_only_2,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
 // CHECK:     attributes [[attr11]]
-// CHECK-SAME:  "llvm.assume"="ompx_lambda_assumption,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK:     attributes [[attr12]]
+// CHECK-SAME:  "llvm.assume"="ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_1234,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK:     attributes [[attr13]]
+// CHECK-SAME:  "llvm.assume"="ompx_lambda_assumption,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_lambda_assumption,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK:     attributes [[attr14]]
+// CHECK-SAME:  "llvm.assume"="omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK:     attributes [[attr15]]
+// CHECK-SAME:  "llvm.assume"="omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp,ompx_call_site,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"
+// CHECK:     attributes [[attr16]]
+// CHECK-SAME:  "llvm.assume"="ompx_call_site,omp_no_openmp_routines,ompx_another_warning,ompx_after_invalid_clauses,omp_no_openmp"

clang/test/OpenMP/assumes_include_nvptx.cpp

 // CHECK: %call = call double @_Z3sind(double 0.000000e+00) [[attr2:#[0-9]]]
 // CHECK: declare double @_Z3sind(double) [[attr1]]
 
 // CHECK:       attributes [[attr0]]
 // CHECK-NOT:  "llvm.assume"
 // CHECK:       attributes [[attr1]]
 // CHECK-SAME:  "llvm.assume"="ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations"
 // CHECK:       attributes [[attr2]]
-// CHECK-SAME:  "llvm.assume"="ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations"
+// CHECK-SAME:  "llvm.assume"="ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations,ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations"
 
 
 template <typename T>
 void foo() {
   cos(T(0));
 }
 
 template <typename T>