[OPENMP 4.0] Codegen for 'declare reduction' construct.

Emit function for 'combiner' part of 'declare reduction' construct and
'initialilzer' part, if any.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@262699 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/AST/DeclOpenMP.h

@@ -99,7 +99,7 @@ public:
 /// \endcode
 ///
 /// Here 'omp_out += omp_in' is a combiner and 'omp_priv = 0' is an initializer.
-class OMPDeclareReductionDecl final : public NamedDecl, public DeclContext {
+class OMPDeclareReductionDecl final : public ValueDecl, public DeclContext {
 private:
   friend class ASTDeclReader;
   /// \brief Combiner for declare reduction construct.
@@ -110,21 +110,18 @@ private:
   /// scope with the same name. Required for proper templates instantiation if
   /// the declare reduction construct is declared inside compound statement.
   LazyDeclPtr PrevDeclInScope;
-  /// \brief Type of declare reduction construct.
-  QualType Ty;
 
   virtual void anchor();
 
   OMPDeclareReductionDecl(Kind DK, DeclContext *DC, SourceLocation L,
                           DeclarationName Name, QualType Ty,
                           OMPDeclareReductionDecl *PrevDeclInScope)
-      : NamedDecl(DK, DC, L, Name), DeclContext(DK), Combiner(nullptr),
-        Initializer(nullptr), PrevDeclInScope(PrevDeclInScope), Ty(Ty) {}
+      : ValueDecl(DK, DC, L, Name, Ty), DeclContext(DK), Combiner(nullptr),
+        Initializer(nullptr), PrevDeclInScope(PrevDeclInScope) {}
 
   void setPrevDeclInScope(OMPDeclareReductionDecl *Prev) {
     PrevDeclInScope = Prev;
   }
-  void setType(QualType T) { Ty = T; }
 
 public:
   /// \brief Create declare reduction node.
@@ -153,8 +150,6 @@ public:
   OMPDeclareReductionDecl *getPrevDeclInScope();
   const OMPDeclareReductionDecl *getPrevDeclInScope() const;
 
-  QualType getType() const { return Ty; }
-
   static bool classof(const Decl *D) { return classofKind(D->getKind()); }
   static bool classofKind(Kind K) { return K == OMPDeclareReduction; }
   static DeclContext *castToDeclContext(const OMPDeclareReductionDecl *D) {

include/clang/AST/GlobalDecl.h

@@ -17,6 +17,7 @@
 
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenMP.h"
 #include "clang/Basic/ABI.h"
 
 namespace clang {
@@ -43,6 +44,7 @@ public:
   GlobalDecl(const BlockDecl *D) { Init(D); }
   GlobalDecl(const CapturedDecl *D) { Init(D); }
   GlobalDecl(const ObjCMethodDecl *D) { Init(D); }
+  GlobalDecl(const OMPDeclareReductionDecl *D) { Init(D); }
 
   GlobalDecl(const CXXConstructorDecl *D, CXXCtorType Type)
   : Value(D, Type) {}

include/clang/Basic/DeclNodes.td

@@ -37,6 +37,7 @@ def Named : Decl<1>;
     def EnumConstant : DDecl<Value>;
     def UnresolvedUsingValue : DDecl<Value>;
     def IndirectField : DDecl<Value>;
+    def OMPDeclareReduction : DDecl<Value>, DeclContext;
     def Declarator : DDecl<Value, 1>;
       def Field : DDecl<Declarator>;
         def ObjCIvar : DDecl<Field>;
@@ -75,7 +76,6 @@ def Named : Decl<1>;
       def ObjCImplementation : DDecl<ObjCImpl>;
   def ObjCProperty : DDecl<Named>;
   def ObjCCompatibleAlias : DDecl<Named>;
-  def OMPDeclareReduction : DDecl<Named>, DeclContext;
 def LinkageSpec : Decl, DeclContext;
 def ObjCPropertyImpl : Decl;
 def FileScopeAsm : Decl;

lib/AST/ASTContext.cpp

@@ -8508,7 +8508,9 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   else if (isa<PragmaDetectMismatchDecl>(D))
     return true;
   else if (isa<OMPThreadPrivateDecl>(D))
-    return true;
+    return !D->getDeclContext()->isDependentContext();
+  else if (isa<OMPDeclareReductionDecl>(D))
+    return !D->getDeclContext()->isDependentContext();
   else
     return false;
 

lib/CodeGen/CGDecl.cpp

@@ -16,6 +16,7 @@
 #include "CGCleanup.h"
 #include "CGDebugInfo.h"
 #include "CGOpenCLRuntime.h"
+#include "CGOpenMPRuntime.h"
 #include "CodeGenModule.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/CharUnits.h"
@@ -120,7 +121,7 @@ void CodeGenFunction::EmitDecl(const Decl &D) {
   }
 
   case Decl::OMPDeclareReduction:
-    return CGM.EmitOMPDeclareReduction(cast<OMPDeclareReductionDecl>(&D));
+    return CGM.EmitOMPDeclareReduction(cast<OMPDeclareReductionDecl>(&D), this);
 
   case Decl::Typedef:      // typedef int X;
   case Decl::TypeAlias: {  // using X = int; [C++0x]
@@ -1867,6 +1868,10 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
     EmitVarAnnotations(&D, DeclPtr.getPointer());
 }
 
-void CodeGenModule::EmitOMPDeclareReduction(
-    const OMPDeclareReductionDecl * /*D*/) {}
+void CodeGenModule::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D,
+                                            CodeGenFunction *CGF) {
+  if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed()))
+    return;
+  getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
+}
 

lib/CodeGen/CGOpenMPRuntime.cpp

@@ -593,8 +593,7 @@ LValue CGOpenMPTaskOutlinedRegionInfo::getThreadIDVariableLValue(
 }
 
 CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
-    : CGM(CGM), DefaultOpenMPPSource(nullptr), KmpRoutineEntryPtrTy(nullptr),
-      OffloadEntriesInfoManager(CGM) {
+    : CGM(CGM), OffloadEntriesInfoManager(CGM) {
   IdentTy = llvm::StructType::create(
       "ident_t", CGM.Int32Ty /* reserved_1 */, CGM.Int32Ty /* flags */,
       CGM.Int32Ty /* reserved_2 */, CGM.Int32Ty /* reserved_3 */,
@@ -612,6 +611,82 @@ void CGOpenMPRuntime::clear() {
   InternalVars.clear();
 }
 
+static llvm::Function *
+emitCombinerOrInitializer(CodeGenModule &CGM, QualType Ty,
+                          const Expr *CombinerInitializer, const VarDecl *In,
+                          const VarDecl *Out, bool IsCombiner) {
+  // void .omp_combiner.(Ty *in, Ty *out);
+  auto &C = CGM.getContext();
+  QualType PtrTy = C.getPointerType(Ty).withRestrict();
+  FunctionArgList Args;
+  ImplicitParamDecl OmpInParm(C, /*DC=*/nullptr, In->getLocation(),
+                              /*Id=*/nullptr, PtrTy);
+  ImplicitParamDecl OmpOutParm(C, /*DC=*/nullptr, Out->getLocation(),
+                               /*Id=*/nullptr, PtrTy);
+  Args.push_back(&OmpInParm);
+  Args.push_back(&OmpOutParm);
+  FunctionType::ExtInfo Info;
+  auto &FnInfo =
+      CGM.getTypes().arrangeFreeFunctionDeclaration(C.VoidTy, Args, Info,
+                                                    /*isVariadic=*/false);
+  auto *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+  auto *Fn = llvm::Function::Create(
+      FnTy, llvm::GlobalValue::InternalLinkage,
+      IsCombiner ? ".omp_combiner." : ".omp_initializer.", &CGM.getModule());
+  CGM.SetInternalFunctionAttributes(/*D=*/nullptr, Fn, FnInfo);
+  CodeGenFunction CGF(CGM);
+  // Map "T omp_in;" variable to "*omp_in_parm" value in all expressions.
+  // Map "T omp_out;" variable to "*omp_out_parm" value in all expressions.
+  CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FnInfo, Args);
+  CodeGenFunction::OMPPrivateScope Scope(CGF);
+  Address AddrIn = CGF.GetAddrOfLocalVar(&OmpInParm);
+  Scope.addPrivate(In, [&CGF, AddrIn, PtrTy]() -> Address {
+    return CGF.EmitLoadOfPointerLValue(AddrIn, PtrTy->castAs<PointerType>())
+        .getAddress();
+  });
+  Address AddrOut = CGF.GetAddrOfLocalVar(&OmpOutParm);
+  Scope.addPrivate(Out, [&CGF, AddrOut, PtrTy]() -> Address {
+    return CGF.EmitLoadOfPointerLValue(AddrOut, PtrTy->castAs<PointerType>())
+        .getAddress();
+  });
+  (void)Scope.Privatize();
+  CGF.EmitIgnoredExpr(CombinerInitializer);
+  Scope.ForceCleanup();
+  CGF.FinishFunction();
+  return Fn;
+}
+
+void CGOpenMPRuntime::emitUserDefinedReduction(
+    CodeGenFunction *CGF, const OMPDeclareReductionDecl *D) {
+  if (UDRMap.count(D) > 0)
+    return;
+  auto &C = CGM.getContext();
+  if (!In || !Out) {
+    In = &C.Idents.get("omp_in");
+    Out = &C.Idents.get("omp_out");
+  }
+  llvm::Function *Combiner = emitCombinerOrInitializer(
+      CGM, D->getType(), D->getCombiner(), cast<VarDecl>(D->lookup(In).front()),
+      cast<VarDecl>(D->lookup(Out).front()),
+      /*IsCombiner=*/true);
+  llvm::Function *Initializer = nullptr;
+  if (auto *Init = D->getInitializer()) {
+    if (!Priv || !Orig) {
+      Priv = &C.Idents.get("omp_priv");
+      Orig = &C.Idents.get("omp_orig");
+    }
+    Initializer = emitCombinerOrInitializer(
+        CGM, D->getType(), Init, cast<VarDecl>(D->lookup(Orig).front()),
+        cast<VarDecl>(D->lookup(Priv).front()),
+        /*IsCombiner=*/false);
+  }
+  UDRMap.insert(std::make_pair(D, std::make_pair(Combiner, Initializer)));
+  if (CGF) {
+    auto &Decls = FunctionUDRMap.FindAndConstruct(CGF->CurFn);
+    Decls.second.push_back(D);
+  }
+}
+
 // Layout information for ident_t.
 static CharUnits getIdentAlign(CodeGenModule &CGM) {
   return CGM.getPointerAlign();
@@ -801,6 +876,12 @@ void CGOpenMPRuntime::functionFinished(CodeGenFunction &CGF) {
   assert(CGF.CurFn && "No function in current CodeGenFunction.");
   if (OpenMPLocThreadIDMap.count(CGF.CurFn))
     OpenMPLocThreadIDMap.erase(CGF.CurFn);
+  if (FunctionUDRMap.count(CGF.CurFn) > 0) {
+    for(auto *D : FunctionUDRMap[CGF.CurFn]) {
+      UDRMap.erase(D);
+    }
+    FunctionUDRMap.erase(CGF.CurFn);
+  }
 }
 
 llvm::Type *CGOpenMPRuntime::getIdentTyPointerTy() {

lib/CodeGen/CGOpenMPRuntime.h

@@ -38,6 +38,8 @@ class Expr;
 class GlobalDecl;
 class OMPExecutableDirective;
 class VarDecl;
+class OMPDeclareReductionDecl;
+class IdentifierInfo;
 
 namespace CodeGen {
 class Address;
@@ -50,7 +52,7 @@ class CGOpenMPRuntime {
   CodeGenModule &CGM;
   /// \brief Default const ident_t object used for initialization of all other
   /// ident_t objects.
-  llvm::Constant *DefaultOpenMPPSource;
+  llvm::Constant *DefaultOpenMPPSource = nullptr;
   /// \brief Map of flags and corresponding default locations.
   typedef llvm::DenseMap<unsigned, llvm::Value *> OpenMPDefaultLocMapTy;
   OpenMPDefaultLocMapTy OpenMPDefaultLocMap;
@@ -73,6 +75,20 @@ class CGOpenMPRuntime {
   typedef llvm::DenseMap<llvm::Function *, DebugLocThreadIdTy>
       OpenMPLocThreadIDMapTy;
   OpenMPLocThreadIDMapTy OpenMPLocThreadIDMap;
+  /// Map of UDRs and corresponding combiner/initializer.
+  typedef llvm::DenseMap<const OMPDeclareReductionDecl *,
+                         std::pair<llvm::Function *, llvm::Function *>>
+      UDRMapTy;
+  UDRMapTy UDRMap;
+  /// Map of functions and locally defined UDRs.
+  typedef llvm::DenseMap<llvm::Function *,
+                         SmallVector<const OMPDeclareReductionDecl *, 4>>
+      FunctionUDRMapTy;
+  FunctionUDRMapTy FunctionUDRMap;
+  IdentifierInfo *In = nullptr;
+  IdentifierInfo *Out = nullptr;
+  IdentifierInfo *Priv = nullptr;
+  IdentifierInfo *Orig = nullptr;
   /// \brief Type kmp_critical_name, originally defined as typedef kmp_int32
   /// kmp_critical_name[8];
   llvm::ArrayType *KmpCriticalNameTy;
@@ -84,7 +100,7 @@ class CGOpenMPRuntime {
   llvm::StringMap<llvm::AssertingVH<llvm::Constant>, llvm::BumpPtrAllocator>
       InternalVars;
   /// \brief Type typedef kmp_int32 (* kmp_routine_entry_t)(kmp_int32, void *);
-  llvm::Type *KmpRoutineEntryPtrTy;
+  llvm::Type *KmpRoutineEntryPtrTy = nullptr;
   QualType KmpRoutineEntryPtrQTy;
   /// \brief Type typedef struct kmp_task {
   ///    void *              shareds; /**< pointer to block of pointers to
@@ -364,6 +380,9 @@ public:
   virtual ~CGOpenMPRuntime() {}
   virtual void clear();
 
+  /// Emit code for the specified user defined reduction construct.
+  virtual void emitUserDefinedReduction(CodeGenFunction *CGF,
+                                        const OMPDeclareReductionDecl *D);
   /// \brief Emits outlined function for the specified OpenMP parallel directive
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

@@ -13,6 +13,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "CGOpenMPRuntimeNVPTX.h"
+#include "clang/AST/DeclOpenMP.h"
 
 using namespace clang;
 using namespace CodeGen;

lib/CodeGen/CodeGenModule.cpp

@@ -1543,10 +1543,17 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
     }
   }
 
-  // If this is OpenMP device, check if it is legal to emit this global
-  // normally.
-  if (OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(GD))
-    return;
+  if (LangOpts.OpenMP) {
+    // If this is OpenMP device, check if it is legal to emit this global
+    // normally.
+    if (OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(GD))
+      return;
+    if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Global)) {
+      if (MustBeEmitted(Global))
+        EmitOMPDeclareReduction(DRD);
+      return;
+    }
+  }
 
   // Ignore declarations, they will be emitted on their first use.
   if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {

lib/CodeGen/CodeGenModule.h

@@ -1112,7 +1112,8 @@ public:
   void EmitOMPThreadPrivateDecl(const OMPThreadPrivateDecl *D);
 
   /// \brief Emit a code for declare reduction construct.
-  void EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D);
+  void EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D,
+                               CodeGenFunction *CGF = nullptr);
 
   /// Returns whether we need bit sets attached to vtables.
   bool NeedVTableBitSets();

lib/CodeGen/ModuleBuilder.cpp

@@ -187,6 +187,15 @@ namespace {
           }
         }
       }
+      // For OpenMP emit declare reduction functions, if required.
+      if (Ctx->getLangOpts().OpenMP) {
+        for (Decl *Member : D->decls()) {
+          if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Member)) {
+            if (Ctx->DeclMustBeEmitted(DRD))
+              Builder->EmitGlobal(DRD);
+          }
+        }
+      }
     }
 
     void HandleTagDeclRequiredDefinition(const TagDecl *D) override {

lib/Sema/SemaExpr.cpp

@@ -2861,6 +2861,7 @@ ExprResult Sema::BuildDeclarationNameExpr(
     // Unresolved using declarations are dependent.
     case Decl::EnumConstant:
     case Decl::UnresolvedUsingValue:
+    case Decl::OMPDeclareReduction:
       valueKind = VK_RValue;
       break;
 

lib/Serialization/ASTReaderDecl.cpp

@@ -2396,12 +2396,11 @@ void ASTDeclReader::VisitOMPThreadPrivateDecl(OMPThreadPrivateDecl *D) {
 }
 
 void ASTDeclReader::VisitOMPDeclareReductionDecl(OMPDeclareReductionDecl *D) {
-  VisitNamedDecl(D);
+  VisitValueDecl(D);
   D->setLocation(Reader.ReadSourceLocation(F, Record, Idx));
   D->setCombiner(Reader.ReadExpr(F));
   D->setInitializer(Reader.ReadExpr(F));
   D->PrevDeclInScope = Reader.ReadDeclID(F, Record, Idx);
-  D->setType(Reader.readType(F, Record, Idx));
 }
 
 void ASTDeclReader::VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D) {
@@ -2458,10 +2457,10 @@ static bool isConsumerInterestedIn(Decl *D, bool HasBody) {
       isa<ObjCImplDecl>(D) ||
       isa<ImportDecl>(D) ||
       isa<PragmaCommentDecl>(D) ||
-      isa<PragmaDetectMismatchDecl>(D) ||
-      isa<OMPThreadPrivateDecl>(D) ||
-      isa<OMPDeclareReductionDecl>(D))
+      isa<PragmaDetectMismatchDecl>(D))
     return true;
+  if (isa<OMPThreadPrivateDecl>(D) || isa<OMPDeclareReductionDecl>(D))
+    return !D->getDeclContext()->isFunctionOrMethod();
   if (VarDecl *Var = dyn_cast<VarDecl>(D))
     return Var->isFileVarDecl() &&
            Var->isThisDeclarationADefinition() == VarDecl::Definition;

lib/Serialization/ASTWriterDecl.cpp

@@ -1655,12 +1655,11 @@ void ASTDeclWriter::VisitOMPThreadPrivateDecl(OMPThreadPrivateDecl *D) {
 }
 
 void ASTDeclWriter::VisitOMPDeclareReductionDecl(OMPDeclareReductionDecl *D) {
-  VisitNamedDecl(D);
+  VisitValueDecl(D);
   Writer.AddSourceLocation(D->getLocStart(), Record);
   Writer.AddStmt(D->getCombiner());
   Writer.AddStmt(D->getInitializer());
   Writer.AddDeclRef(D->getPrevDeclInScope(), Record);
-  Writer.AddTypeRef(D->getType(), Record);
   Code = serialization::DECL_OMP_DECLARE_REDUCTION;
 }
 

test/OpenMP/declare_reduction_codegen.c

@@ -0,0 +1,158 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -emit-llvm %s -triple %itanium_abi_triple -o - -femit-all-decls | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c -triple %itanium_abi_triple -emit-pch -o %t %s -femit-all-decls
+// RUN: %clang_cc1 -fopenmp -x c -triple %itanium_abi_triple -include-pch %t -verify %s -emit-llvm -o - -femit-all-decls | FileCheck --check-prefix=CHECK-LOAD %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// CHECK: [[SSS_INT:.+]] = type { i32 }
+// CHECK-LOAD: [[SSS_INT:.+]] = type { i32 }
+
+#pragma omp declare reduction(+ : int, char : omp_out *= omp_in)
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[MUL:%.+]] = mul nsw i32
+// CHECK-NEXT: store i32 [[MUL]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: store i32 [[MUL]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+// CHECK: sext i8
+// CHECK: sext i8
+// CHECK: [[MUL:%.+]] = mul nsw i32
+// CHECK-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+// CHECK-NEXT: store i8 [[TRUNC]], i8*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+// CHECK-LOAD-NEXT: store i8 [[TRUNC]], i8*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+#pragma omp declare reduction(fun : float : omp_out += omp_in) initializer(omp_priv = 15 + omp_orig)
+// CHECK: define internal {{.*}}void @{{[^(]+}}(float* noalias, float* noalias)
+// CHECK: [[ADD:%.+]] = fadd float
+// CHECK-NEXT: store float [[ADD]], float*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK: define internal {{.*}}void @{{[^(]+}}(float* noalias, float* noalias)
+// CHECK: [[ADD:%.+]] = fadd float 1.5
+// CHECK-NEXT: store float [[ADD]], float*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(float* noalias, float* noalias)
+// CHECK-LOAD: [[ADD:%.+]] = fadd float
+// CHECK-LOAD-NEXT: store float [[ADD]], float*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(float* noalias, float* noalias)
+// CHECK-LOAD: [[ADD:%.+]] = fadd float 1.5
+// CHECK-LOAD-NEXT: store float [[ADD]], float*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+struct SSS {
+  int field;
+#pragma omp declare reduction(+ : int, char : omp_out *= omp_in)
+  // CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+  // CHECK: [[MUL:%.+]] = mul nsw i32
+  // CHECK-NEXT: store i32 [[MUL]], i32*
+  // CHECK-NEXT: ret void
+  // CHECK-NEXT: }
+
+  // CHECK: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+  // CHECK: sext i8
+  // CHECK: sext i8
+  // CHECK: [[MUL:%.+]] = mul nsw i32
+  // CHECK-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+  // CHECK-NEXT: store i8 [[TRUNC]], i8*
+  // CHECK-NEXT: ret void
+  // CHECK-NEXT: }
+};
+
+void init(struct SSS *priv, struct SSS orig);
+
+#pragma omp declare reduction(fun : struct SSS : omp_out = omp_in) initializer(init(&omp_priv, omp_orig))
+// CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK: call void @llvm.memcpy
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK: call void @init(
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK-LOAD: call void @llvm.memcpy
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK-LOAD: call void @init(
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK-LABEL: @main
+// CHECK-LOAD-LABEL: @main
+int main() {
+#pragma omp declare reduction(fun : struct SSS : omp_out = omp_in) initializer(init(&omp_priv, omp_orig))
+  // CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+  // CHECK: call void @llvm.memcpy
+  // CHECK-NEXT: ret void
+  // CHECK-NEXT: }
+  // CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+  // CHECK: call void @init(
+  // CHECK-NEXT: ret void
+  // CHECK-NEXT: }
+  // CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+  // CHECK-LOAD: call void @llvm.memcpy
+  // CHECK-LOAD-NEXT: ret void
+  // CHECK-LOAD-NEXT: }
+  // CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+  // CHECK-LOAD: call void @init(
+  // CHECK-LOAD-NEXT: ret void
+  // CHECK-LOAD-NEXT: }
+  {
+#pragma omp declare reduction(fun : struct SSS : omp_out = omp_in) initializer(init(&omp_priv, omp_orig))
+    // CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+    // CHECK: call void @llvm.memcpy
+    // CHECK-NEXT: ret void
+    // CHECK-NEXT: }
+    // CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+    // CHECK: call void @init(
+    // CHECK-NEXT: ret void
+    // CHECK-NEXT: }
+    // CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+    // CHECK-LOAD: call void @llvm.memcpy
+    // CHECK-LOAD-NEXT: ret void
+    // CHECK-LOAD-NEXT: }
+    // CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+    // CHECK-LOAD: call void @init(
+    // CHECK-LOAD-NEXT: ret void
+    // CHECK-LOAD-NEXT: }
+  }
+  return 0;
+}
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: store i32 [[MUL]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+// CHECK-LOAD-NEXT: store i8 [[TRUNC]], i8*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+#endif

test/OpenMP/declare_reduction_codegen.cpp

@@ -0,0 +1,170 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -emit-llvm %s -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -o - -femit-all-decls | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s -femit-all-decls
+// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - -femit-all-decls | FileCheck --check-prefix=CHECK-LOAD %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// CHECK: [[SSS_INT:.+]] = type { i32 }
+// CHECK-LOAD: [[SSS_INT:.+]] = type { i32 }
+
+#pragma omp declare reduction(+ : int, char : omp_out *= omp_in)
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[MUL:%.+]] = mul nsw i32
+// CHECK-NEXT: store i32 [[MUL]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: store i32 [[MUL]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+// CHECK: sext i8
+// CHECK: sext i8
+// CHECK: [[MUL:%.+]] = mul nsw i32
+// CHECK-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+// CHECK-NEXT: store i8 [[TRUNC]], i8*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i8* noalias, i8* noalias)
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: sext i8
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32
+// CHECK-LOAD-NEXT: [[TRUNC:%.+]] = trunc i32 [[MUL]] to i8
+// CHECK-LOAD-NEXT: store i8 [[TRUNC]], i8*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+template <class T>
+struct SSS {
+  T a;
+#pragma omp declare reduction(fun : T : omp_out ^= omp_in) initializer(omp_priv = 24 + omp_orig)
+};
+
+SSS<int> d;
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[XOR:%.+]] = xor i32
+// CHECK-NEXT: store i32 [[XOR]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[ADD:%.+]] = add nsw i32 24,
+// CHECK-NEXT: store i32 [[ADD]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+// CHECK: define void [[INIT:@[^(]+]]([[SSS_INT]]*
+// CHECK-LOAD: define void [[INIT:@[^(]+]]([[SSS_INT]]*
+void init(SSS<int> &lhs, SSS<int> &rhs) {}
+
+#pragma omp declare reduction(fun : SSS < int > : omp_out = omp_in) initializer(init(omp_priv, omp_orig))
+// CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK: call void @llvm.memcpy
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK: call void [[INIT]](
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK-LOAD: call void @llvm.memcpy
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}([[SSS_INT]]* noalias, [[SSS_INT]]* noalias)
+// CHECK-LOAD: call void [[INIT]](
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+template <typename T>
+T foo(T a) {
+#pragma omp declare reduction(fun : T : omp_out += omp_in) initializer(omp_priv = 15 * omp_orig)
+  {
+#pragma omp declare reduction(fun : T : omp_out /= omp_in) initializer(omp_priv = 11 - omp_orig)
+  }
+  return a;
+}
+
+// CHECK-LABEL: @main
+int main() {
+  int i = 0;
+  SSS<int> sss;
+  // TODO: Add support for scoped reduction identifiers
+  //  #pragma omp parallel reduction(SSS<int>::fun : i)
+  // TODO-CHECK: #pragma omp parallel reduction(SSS<int>::fun: i)
+  {
+    i += 1;
+  }
+  // #pragma omp parallel reduction(::fun:sss)
+  // TODO-CHECK: #pragma omp parallel reduction(::fun: sss)
+  {
+  }
+  return foo(15);
+}
+
+// CHECK-LABEL: i32 @{{.+}}foo{{[^(].+}}(i32
+// CHECK-LOAD-LABEL: i32 @{{.+}}foo{{[^(].+}}(i32
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[XOR:%.+]] = xor i32
+// CHECK-LOAD-NEXT: store i32 [[XOR]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[ADD:%.+]] = add nsw i32 24,
+// CHECK-LOAD-NEXT: store i32 [[ADD]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[ADD:%.+]] = add nsw i32
+// CHECK-NEXT: store i32 [[ADD]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[ADD:%.+]] = add nsw i32
+// CHECK-LOAD-NEXT: store i32 [[ADD]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[MUL:%.+]] = mul nsw i32 15,
+// CHECK-NEXT: store i32 [[MUL]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[MUL:%.+]] = mul nsw i32 15,
+// CHECK-LOAD-NEXT: store i32 [[MUL]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[DIV:%.+]] = sdiv i32
+// CHECK-NEXT: store i32 [[DIV]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[DIV:%.+]] = sdiv i32
+// CHECK-LOAD-NEXT: store i32 [[DIV]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+// CHECK: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK: [[SUB:%.+]] = sub nsw i32 11,
+// CHECK-NEXT: store i32 [[SUB]], i32*
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+// CHECK-LOAD: define internal {{.*}}void @{{[^(]+}}(i32* noalias, i32* noalias)
+// CHECK-LOAD: [[SUB:%.+]] = sub nsw i32 11,
+// CHECK-LOAD-NEXT: store i32 [[SUB]], i32*
+// CHECK-LOAD-NEXT: ret void
+// CHECK-LOAD-NEXT: }
+
+#endif

test/OpenMP/threadprivate_codegen.cpp

@@ -948,5 +948,5 @@ int foobar() {
 // CHECK-TLS:      call void [[ST_S4_ST_CXX_INIT]]
 // CHECK-TLS:      [[DONE_LABEL]]
 
-// CHECK-TLS:      declare {{.*}} void [[GS3_TLS_INIT]]
-// CHECK-TLS:      declare {{.*}} void [[STATIC_S_TLS_INIT]]
+// CHECK-TLS-DAG:      declare {{.*}} void [[GS3_TLS_INIT]]
+// CHECK-TLS-DAG:      declare {{.*}} void [[STATIC_S_TLS_INIT]]