[OpenMP] Add initial support for `omp [begin/end] assumes`

[jdoerfert]

Summary

The assumes directive is an OpenMP 5.1 feature that allows the user to
provide assumptions to the optimizer. Assumptions can refer to
directives (absent and contains clauses), expressions (holds
clause), or generic properties (no_openmp_routines, ext_ABCD, ...).

The assumes spelling is used for assumptions in the global scope while
assume is used for executable contexts with an associated structured
block.

This patch only implements the global spellings. While clauses with
arguments are "accepted" by the parser, they will simply be ignored for
now. The implementation lowers the assumptions directly to the
AssumptionAttr.

[ABataev]

clang/include/clang/Sema/Sema.h

10335

It may have side-effects with template instantiations. What if we have something like this:

#pragma omp begin assumes ...
template <typename T>
struct S {
...
}
#pragma omp end assumes

void bar() {
#pragma omp assumes ...
  {
    S<int> s;
    s.foo();
  }
}

?

struct S<int> will be instantiated in the second assume region context though I doubt this is the user intention.

10335–10338

const member functions.

[jdoerfert]

clang/include/clang/Sema/Sema.h

10335

I'm not sure what problem you expect here. Could you elaborate what you think should happen or should not?

[ABataev]

clang/include/clang/Sema/Sema.h

10335

May it lead to the wrong compiler assumptions and result in the incorrect codegen?

[jdoerfert]

clang/include/clang/Sema/Sema.h

10335

I'm still not following your example. What wrong assumptions and/or incorrect code generation do you expect? Could you provide an example with actual assumptions and then describe what the problem is please.

[ABataev]

clang/include/clang/Sema/Sema.h

10335

Say, you have something like this:

template <typename T>
struct S {
  int a;
  void foo() {
    #pragma omp parallel
    ...
  }
};

void bar() {
#pragma omp assumems no_openmp
  S<int> s; // S<int> is instantiated here and function S<int>::foo() gets no_openmp attribute.
  s.a = 0; // S<int>::foo is instantiated here but not used and the attribute no_openmp should not be propagated for this function.
#pragma omp end assumes

  S<int> b; // S<int> already instantiated and S<int>::foo is instantiated with attribute no_openmp
  b.foo(); // calls function S<int>::foo with no_openmp attribute.
}

What do you think?

[jdoerfert]

clang/include/clang/Sema/Sema.h

10335

Thanks for the extra information. I'll add a test based on this for sure, no matter what we come up with for now.

#pragma omp begin assumes no_openmp
  S<int> s; // S<int> is instantiated here and function S<int>::foo() gets no_openmp attribute.
  s.a = 0;
#pragma omp end assumes

Hm, I guess the question is:
Is S<int>::foo() a function defined in between the begin/end or not. If we say it is, the behavior right now seems consistent but the standard might need some clarifications. If we say it is not, we should not annotate it. Though, if it is not, where is it defined? I could avoid annotating template instantiations I guess. I might start with that. Update is coming.

[ABataev]

clang/include/clang/Sema/Sema.h

10335

Yep, this is what I'm asking about. I would check С++ standard, I believe we shall follow С++ rules here about the instantiation context.

[jdoerfert]

Improve template handling. Do not annotate template instantiations with scoped assumptions for now.

[jdoerfert]

clang/include/clang/Sema/Sema.h

10335

I think the new solution, that is not to use scoped assumptions during template instantiations, is a conservative approximation of what people would assume and also matches the standard wrt. "defintition-declaration-seq". I included your example as a test.

WDYT?

[ABataev]

Why don't yo want to try to implement the scheme similar to the declare target?

clang/include/clang/Sema/Sema.h

10335

const member function

10338

const member function

clang/lib/Parse/ParseOpenMP.cpp

1530

I think you will need just capture this by value here

1549–1554

Why not express it as clauses?

1559

S

1568

Idx

1568–1592

Better to outline it into function/lambda

clang/lib/Sema/SemaOpenMP.cpp

3210

What are you trying to do here? Scan all the decls in the current context and mark all the functions with the assumption attributes?

[jdoerfert]

Addressed comments

[jdoerfert]

In D91980#2416950, @ABataev wrote:

Why don't yo want to try to implement the scheme similar to the declare target?

Because it is not clear that the standard even says that right now. Also, what is the user expectation here.
The scheme now is conservative but consistent. I'd prefer to use something like that first before we clarify edge cases.

clang/lib/Parse/ParseOpenMP.cpp

1530

I've never done that. Could you explain why? Is it just an "optimization"? A simple test doesn't really show a difference https://godbolt.org/z/bfr47f

1549–1554

Various reasons actually:

• It is a lot of boilerplate code, if we would auto-generate them from the tablegen definition the situation might be different.
• There seems to be no immediate benefit.
• We want to be in-sync with the generic assumption handling not develop two systems for the same thing.

1568–1592

Which part? Assuming you mean the 9 line conditional above, I'm not convinced. Even if I do outline it, 4 lines will have to stay here and I don't see any reuse.

clang/lib/Sema/SemaOpenMP.cpp

3210

Pretty much, yeah. Find all function declarations and add the attribute. I'll add a comment.

[ABataev]

In D91980#2417049, @jdoerfert wrote:

In D91980#2416950, @ABataev wrote:

Why don't yo want to try to implement the scheme similar to the declare target?

Because it is not clear that the standard even says that right now. Also, what is the user expectation here.
The scheme now is conservative but consistent. I'd prefer to use something like that first before we clarify edge cases.

I don't think that it will affect the implementation significantly. Still, OpenMP standard must preserve visibility etc. rules of C/C++, so the changes should not be critical (I think).

clang/lib/Parse/ParseOpenMP.cpp

1530

Yes, this is just the optimization.It is not for the current implementation but for possible future changes. There is no difference for this but might be difference if you'll need to capture something else.

1549–1554

Just it is always to have and to maintain just one design rather than 2 or more. But it is up to you.

1568–1592

The part that searches over the table.

clang/lib/Sema/SemaOpenMP.cpp

3210

Does it mean that for this code

void foo();
 #pragma omp assume ...
void bar();
 #pragma omp end assume

foo() must have the same assumptions like bar()? It looks to me, currently we'll end up with the same attributes for both functions. But I'm not sure that this is correct to mark the function outside of the assumes region with the same attributes just like the function inside.

[jdoerfert]

In D91980#2417129, @ABataev wrote:

In D91980#2417049, @jdoerfert wrote:

In D91980#2416950, @ABataev wrote:

Why don't yo want to try to implement the scheme similar to the declare target?

Because it is not clear that the standard even says that right now. Also, what is the user expectation here.
The scheme now is conservative but consistent. I'd prefer to use something like that first before we clarify edge cases.

I don't think that it will affect the implementation significantly. Still, OpenMP standard must preserve visibility etc. rules of C/C++, so the changes should not be critical (I think).

We can add more support later, once it is clear how that will look like. As the person that wrote the OpenMP standard section, I'm trying to tell you it is not clear what the expected behavior is supposed to be. Not applying assumes is *always* correct. Applying them for some cases that is not easy to argue about, is at least confusing. If you feel we should add assumes in more cases, I'm happy to bring up the discussion in the OpenMP standards committee and look at a follow up patch.

clang/lib/Parse/ParseOpenMP.cpp

1568–1592

The part that searches over the table.

I don't follow. Where is a search? The part that builds a string switch? line 1569 - 1579?

clang/lib/Sema/SemaOpenMP.cpp

3210

Your example mixes the global omp assumes with the scoped begin/end omp assumes. The former has the effect to be added to prior declarations, the latter does not have this effect. This is how it is supposed to work:

void foo();  // no assumption
 #pragma omp begin assumes ext_XYZ
void bar();  // XYZ assumption
 #pragma omp end assumes

void foo();  // XYZ assumption
 #pragma omp assumes ext_XYZ
void bar();  // XYZ assumption

[ABataev]

clang/lib/Parse/ParseOpenMP.cpp

1568–1592

Yes.

clang/lib/Sema/SemaOpenMP.cpp

3210

Ah, I see. omp assumes looks weird. :( It may lead to conflicts
PS. What about lambdas? Could add a test that lambdas get annotations too?

3218

I would also add a check that the decl is valid here before checking it and applying the attributes.

3225

if (auto *SubCtx = dyn_cast<DeclContext>(SubDC)) {
  DeclContexts.push_back(cast<DeclContext>(SubCtx));
 continue;
}

?

[jdoerfert]

clang/lib/Sema/SemaOpenMP.cpp

3210

Ah, I see. omp assumes looks weird. :( It may lead to conflicts

The reason we have assumes and assume is to avoid conflicts. Let's hope it works.

PS. What about lambdas? Could add a test that lambdas get annotations too?

Done, assuming lambdas are function declarations. Otherwise we would need to exclude them from the scoped assumptions as we do for template instantiations now.

[jdoerfert]

Addressed comments, lambda handling and more

[ABataev]

clang/include/clang/Basic/DiagnosticParseKinds.td

1296–1297

There is err_omp_no_clause_for_directive defined in Sema. See Sema::ActOnOpenMPTargetDataDirective how to use it, better to reduce the number of diagnostics if possible

clang/lib/Sema/SemaLambda.cpp

999–1002 ↗

(On Diff #308817)

Are there any other function-like constructs that also should have markings? Coroutines maybe?

clang/lib/Sema/SemaOpenMP.cpp

3203

How this comment relates to the code?

3207

Add a message in for the assertion

3217

Maybe, better to use getLexicalParent? getParent returns semantic parent, while getLexicalParent - the lexical parent. Do you need to mark the declarations in the lexical context or in the semantic context?

[jdoerfert]

Addressed comments

clang/lib/Sema/SemaLambda.cpp

999–1002 ↗

(On Diff #308817)

Hm, can we postpone coroutines for now?

clang/lib/Sema/SemaOpenMP.cpp

3203

Leftover.

3207

Good catch!

3217

I go to the top, should not matter. I want all declarations so I start at the top most scope, then traverse everything. I go with lexical now.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

3203–3206

Early exit from the function in case of error?

[jdoerfert]

clang/lib/Sema/SemaOpenMP.cpp

3203–3206

We can't always because we want to create an (empty) scope so we can diagnose the end assumes properly. I'll check in the non-scope case though.

[jdoerfert]

Early exit

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

5926

You can also check if the decl is invalid here

5931–5932

Looks like FD==nullptr is impossible here. Or UTemplDecl->getTemplatedDecl() can return nullptr?

[jdoerfert]

Addressed comments

[ABataev]

LG

[mehdi_amini]

llvm/include/llvm/Frontend/OpenMP/OMPConstants.h

101 ↗

(On Diff #312044)

This is broken on gcc-5:

llvm/include/llvm/Frontend/OpenMP/OMPConstants.h:101:1: error: could not convert '(const char*)"ext_"' from 'const char*' to 'llvm::StringLiteral'
 };
 ^

[jdoerfert]

I guess I revert it for now :(

llvm/include/llvm/Frontend/OpenMP/OMPConstants.h

101 ↗

(On Diff #312044)

You don't happen to have a suggestion on how to fix this, do you? I don't have a gcc-5 lying around to try (I think).

[thakis]

In addition to the build error, it also broke check-clang everywhere as far as I can tell.
mac: http://45.33.8.238/macm1/292/step_6.txt
linux: http://45.33.8.238/linux/35399/step_7.txt
win: http://45.33.8.238/win/29934/step_7.txt

[jdoerfert]

In D91980#2456600, @thakis wrote:

In addition to the build error, it also broke check-clang everywhere as far as I can tell.
mac: http://45.33.8.238/macm1/292/step_6.txt
linux: http://45.33.8.238/linux/35399/step_7.txt
win: http://45.33.8.238/win/29934/step_7.txt

I think I fixed 2 errors but I'm unsure about the linux one. I changed one test and I'll try it again as I find a workaround for gcc-5

[thakis]

In D91980#2456637, @jdoerfert wrote:

In D91980#2456600, @thakis wrote:

In addition to the build error, it also broke check-clang everywhere as far as I can tell.
mac: http://45.33.8.238/macm1/292/step_6.txt
linux: http://45.33.8.238/linux/35399/step_7.txt
win: http://45.33.8.238/win/29934/step_7.txt

I think I fixed 2 errors but I'm unsure about the linux one. I changed one test and I'll try it again as I find a workaround for gcc-5

The same tests fail on linux and mac and win at these links:

Clang :: OpenMP/assumes_codegen.cpp (win, mac/m1)
Clang :: OpenMP/assumes_include_nvptx.cpp (linux, mac/x86, mac/m1)

(mac/x86 link: http://45.33.8.238/mac/25261/step_7.txt)

[jdoerfert]

Fixes

clang/include/clang/Basic/DiagnosticParseKinds.td

def err_omp_map_type_modifier_missing : Error<
   "missing map type modifier">;
 def err_omp_declare_simd_inbranch_notinbranch : Error<
   "unexpected '%0' clause, '%1' is specified already">;
 def err_expected_end_declare_target_or_variant : Error<
   "expected '#pragma omp end declare %select{target|variant}0'">;
 def err_expected_begin_declare_variant
     : Error<"'#pragma omp end declare variant' with no matching '#pragma omp "
             "begin declare variant'">;
+def err_expected_begin_assumes
+    : Error<"'#pragma omp end assumes' with no matching '#pragma omp begin assumes'">;
+def err_expected_assumption_clause
+    : Error<"'%0' directive requires at least one clause">;
+def warn_omp_unknown_assumption_clause_missing_id
+    : Warning<"valid %0 clauses start with %1; %select{token|tokens}2 will be ignored">,
+      InGroup<OpenMPClauses>;
+def warn_omp_unknown_assumption_clause_without_args
+    : Warning<"%0 clause should not be followed by arguments; tokens will be ignored">,
+      InGroup<OpenMPClauses>;
+def note_omp_assumption_clause_continue_here
+    : Note<"the ignored tokens spans until here">;
 def err_omp_declare_target_unexpected_clause: Error<
   "unexpected '%0' clause, only %select{'to' or 'link'|'to', 'link' or 'device_type'}1 clauses expected">;
 def err_omp_expected_clause: Error<
   "expected at least one clause on '#pragma omp %0' directive">;
 def err_omp_mapper_illegal_identifier : Error<
   "illegal OpenMP user-defined mapper identifier">;
 def err_omp_mapper_expected_declarator : Error<
   "expected declarator on 'omp declare mapper' directive">;

clang/include/clang/Parse/Parser.h

private:
   /// if there was an error.
   bool parseOMPDeclareVariantMatchClause(SourceLocation Loc, OMPTraitInfo &TI,
                                          OMPTraitInfo *ParentTI);
 
   /// Parse clauses for '#pragma omp declare variant'.
   void ParseOMPDeclareVariantClauses(DeclGroupPtrTy Ptr, CachedTokens &Toks,
                                      SourceLocation Loc);
 
+  /// Parse 'omp [begin] assume[s]' directive.
+  void ParseOpenMPAssumesDirective(OpenMPDirectiveKind DKind,
+                                   SourceLocation Loc);
+
+  /// Parse 'omp end assumes' directive.
+  void ParseOpenMPEndAssumesDirective(SourceLocation Loc);
+
   /// Parse clauses for '#pragma omp declare target'.
   DeclGroupPtrTy ParseOMPDeclareTargetClauses();
   /// Parse '#pragma omp end declare target'.
   void ParseOMPEndDeclareTargetDirective(OpenMPDirectiveKind DKind,
                                          SourceLocation Loc);
 
   /// Skip tokens until a `annot_pragma_openmp_end` was found. Emit a warning if
   /// it is not the current token.

clang/include/clang/Sema/Sema.h

private:
   OMPTraitInfo *getOMPTraitInfoForSurroundingScope() {
     return OMPDeclareVariantScopes.empty() ? nullptr
                                            : OMPDeclareVariantScopes.back().TI;
   }
 
   /// The current `omp begin/end declare variant` scopes.
   SmallVector<OMPDeclareVariantScope, 4> OMPDeclareVariantScopes;
 
+  /// The current `omp begin/end assumes` scopes.
+  SmallVector<AssumptionAttr *, 4> OMPAssumeScoped;
+
+  /// All `omp assumes` we encountered so far.
+  SmallVector<AssumptionAttr *, 4> OMPAssumeGlobal;
+
+public:
   /// The declarator \p D defines a function in the scope \p S which is nested
   /// in an `omp begin/end declare variant` scope. In this method we create a
   /// declaration for \p D and rename \p D according to the OpenMP context
   /// selector of the surrounding scope. Return all base functions in \p Bases.
   void ActOnStartOfFunctionDefinitionInOpenMPDeclareVariantScope(
       Scope *S, Declarator &D, MultiTemplateParamsArg TemplateParameterLists,
       SmallVectorImpl<FunctionDecl *> &Bases);
 
   /// Register \p D as specialization of all base functions in \p Bases in the
   /// current `omp begin/end declare variant` scope.
   void ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(
       Decl *D, SmallVectorImpl<FunctionDecl *> &Bases);
 
-public:
+  /// Act on \p D, a function definition inside of an `omp [begin/end] assumes`.
+  void ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(Decl *D);
 
-  /// Can we exit a scope at the moment.
-  bool isInOpenMPDeclareVariantScope() {
+  /// Can we exit an OpenMP declare variant scope at the moment.
+  bool isInOpenMPDeclareVariantScope() const {
     return !OMPDeclareVariantScopes.empty();
   }
 
   /// Given the potential call expression \p Call, determine if there is a
   /// specialization via the OpenMP declare variant mechanism available. If
   /// there is, return the specialized call expression, otherwise return the
   /// original \p Call.
   ExprResult ActOnOpenMPCall(ExprResult Call, Scope *Scope,
public:
   /// Builds a new OpenMPThreadPrivateDecl and checks its correctness.
   OMPThreadPrivateDecl *CheckOMPThreadPrivateDecl(SourceLocation Loc,
                                                   ArrayRef<Expr *> VarList);
   /// Called on well-formed '#pragma omp allocate'.
   DeclGroupPtrTy ActOnOpenMPAllocateDirective(SourceLocation Loc,
                                               ArrayRef<Expr *> VarList,
                                               ArrayRef<OMPClause *> Clauses,
                                               DeclContext *Owner = nullptr);
+
+  /// Called on well-formed '#pragma omp [begin] assume[s]'.
+  void ActOnOpenMPAssumesDirective(SourceLocation Loc,
+                                   OpenMPDirectiveKind DKind,
+                                   ArrayRef<StringRef> Assumptions);
+
+  /// Check if there is an active global `omp begin assumes` directive.
+  bool isInOpenMPAssumeScope() const { return !OMPAssumeScoped.empty(); }
+
+  /// Check if there is an active global `omp assumes` directive.
+  bool hasGlobalOpenMPAssumes() const { return !OMPAssumeGlobal.empty(); }
+
+  /// Called on well-formed '#pragma omp end assumes'.
+  void ActOnOpenMPEndAssumesDirective();
+
   /// Called on well-formed '#pragma omp requires'.
   DeclGroupPtrTy ActOnOpenMPRequiresDirective(SourceLocation Loc,
                                               ArrayRef<OMPClause *> ClauseList);
   /// Check restrictions on Requires directive
   OMPRequiresDecl *CheckOMPRequiresDecl(SourceLocation Loc,
                                         ArrayRef<OMPClause *> Clauses);
   /// Check if the specified type is allowed to be used in 'omp declare
   /// reduction' construct.

clang/lib/Parse/ParseOpenMP.cpp

 }
 
 static OpenMPDirectiveKindExWrapper parseOpenMPDirectiveKind(Parser &P) {
   // Array of foldings: F[i][0] F[i][1] ===> F[i][2].
   // E.g.: OMPD_for OMPD_simd ===> OMPD_for_simd
   // TODO: add other combined directives in topological order.
   static const OpenMPDirectiveKindExWrapper F[][3] = {
       {OMPD_begin, OMPD_declare, OMPD_begin_declare},
+      {OMPD_begin, OMPD_assumes, OMPD_begin_assumes},
       {OMPD_end, OMPD_declare, OMPD_end_declare},
+      {OMPD_end, OMPD_assumes, OMPD_end_assumes},
       {OMPD_cancellation, OMPD_point, OMPD_cancellation_point},
       {OMPD_declare, OMPD_reduction, OMPD_declare_reduction},
       {OMPD_declare, OMPD_mapper, OMPD_declare_mapper},
       {OMPD_declare, OMPD_simd, OMPD_declare_simd},
       {OMPD_declare, OMPD_target, OMPD_declare_target},
       {OMPD_declare, OMPD_variant, OMPD_declare_variant},
       {OMPD_begin_declare, OMPD_variant, OMPD_begin_declare_variant},
       {OMPD_end_declare, OMPD_variant, OMPD_end_declare_variant},
for (const OMPTraitSet &ParentSet : ParentTI->Sets) {
     }
     if (!MergedSet)
       TI.Sets.push_back(ParentSet);
   }
 
   return false;
 }
 
+/// `omp assumes` or `omp begin/end assumes` <clause> [[,]<clause>]...
+/// where
+///
+///   clause:
+///     'ext_IMPL_DEFINED'
+///     'absent' '(' directive-name [, directive-name]* ')'
+///     'contains' '(' directive-name [, directive-name]* ')'
+///     'holds' '(' scalar-expression ')'
+///     'no_openmp'
+///     'no_openmp_routines'
+///     'no_parallelism'
+///
+void Parser::ParseOpenMPAssumesDirective(OpenMPDirectiveKind DKind,
+                                         SourceLocation Loc) {
+  SmallVector<StringRef, 4> Assumptions;
+  bool SkippedClauses = false;
+
+  auto SkipBraces = [&](llvm::StringRef Spelling, bool IssueNote) {
+    BalancedDelimiterTracker T(*this, tok::l_paren,
+                               tok::annot_pragma_openmp_end);
+    if (T.expectAndConsume(diag::err_expected_lparen_after, Spelling.data()))
+      return;
+    T.skipToEnd();
+    if (IssueNote && T.getCloseLocation().isValid())
+      Diag(T.getCloseLocation(),
+           diag::note_omp_assumption_clause_continue_here);
+  };
+
+  struct AssumptionClauseMappingInfo {
+    llvm::StringLiteral Identifier;
+    bool StartsWith;
+    bool HasDirectiveList;
+    bool HasExpression;
+  };
+  AssumptionClauseMappingInfo AssumptionClauseMappings[7] = {
+      {"ext_", true, false, false},
+      {"absent", false, true, false},
+      {"contains", false, true, false},
+      {"holds", false, false, true},
+      {"no_openmp", false, false, false},
+      {"no_openmp_routines", false, false, false},
+      {"no_parallelism", false, false, false},
+  };
+
+  // Helper to create a diagnostic listing all valid clause (beginnings).
+  auto GetAllowedClauseStartValues = [&]() {
+    std::string S;
+    for (const AssumptionClauseMappingInfo &ACMI : AssumptionClauseMappings)
+      S += (S.empty() ? "'" : "', '") + ACMI.Identifier.str();
+    return S + "'";
+  };
+
+  while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+    IdentifierInfo *II = nullptr;
+    SourceLocation StartLoc = Tok.getLocation();
+    int Idx = -1;
+    if (Tok.isAnyIdentifier()) {
+      II = Tok.getIdentifierInfo();
+      llvm::StringSwitch<int> SS(II->getName());
+      for (const AssumptionClauseMappingInfo &ACMI : AssumptionClauseMappings) {
+        if (ACMI.StartsWith)
+          SS.StartsWith(ACMI.Identifier, ++Idx);
+        else
+          SS.Case(ACMI.Identifier, ++Idx);
+      }
+      Idx = SS.Default(-1);
+    }
+    ConsumeAnyToken();
+
+    bool NextIsLPar = Tok.is(tok::l_paren);
+    // Handle unknown clauses by skipping them.
+    if (Idx == -1) {
+      Diag(StartLoc, diag::warn_omp_unknown_assumption_clause_missing_id)
+          << llvm::omp::getOpenMPDirectiveName(DKind)
+          << GetAllowedClauseStartValues() << NextIsLPar;
+      if (NextIsLPar)
+        SkipBraces(II ? II->getName() : "", /* IssueNote */ true);
+      SkippedClauses = true;
+      continue;
+    }
+    const AssumptionClauseMappingInfo &ACMI = AssumptionClauseMappings[Idx];
+    if (ACMI.HasDirectiveList || ACMI.HasExpression) {
+      // TODO: We ignore absent, contains, and holds assumptions for now. We
+      //       also do not verify the content in the parenthesis at all.
+      SkippedClauses = true;
+      SkipBraces(II->getName(), /* IssueNote */ false);
+      continue;
+    }
+
+    if (NextIsLPar) {
+      Diag(Tok.getLocation(),
+           diag::warn_omp_unknown_assumption_clause_without_args)
+          << II;
+      SkipBraces(II->getName(), /* IssueNote */ true);
+    }
+
+    assert(II && "Expected an identifier clause!");
+    StringRef Assumption = II->getName();
+    if (ACMI.StartsWith)
+      Assumption = Assumption.substr(ACMI.Identifier.size());
+    Assumptions.push_back(Assumption);
+  }
+
+  if (Assumptions.empty()) {
+    if (!SkippedClauses)
+      Diag(Loc, diag::err_expected_assumption_clause)
+          << llvm::omp::getOpenMPDirectiveName(DKind);
+  }
+
+  Actions.ActOnOpenMPAssumesDirective(Loc, DKind, Assumptions);
+}
+
+void Parser::ParseOpenMPEndAssumesDirective(SourceLocation Loc) {
+  if (Actions.isInOpenMPAssumeScope())
+    Actions.ActOnOpenMPEndAssumesDirective();
+  else
+    Diag(Loc, diag::err_expected_begin_assumes);
+}
+
 /// Parsing of simple OpenMP clauses like 'default' or 'proc_bind'.
 ///
 ///    default-clause:
 ///         'default' '(' 'none' | 'shared'  | 'firstprivate' ')
 ///
 ///    proc_bind-clause:
 ///         'proc_bind' '(' 'master' | 'close' | 'spread' ')
 ///

 ///         annot_pragma_openmp 'declare simd' {<clause> [,]}
 ///         annot_pragma_openmp_end
 ///         <function declaration/definition>
 ///
 ///       requires directive:
 ///         annot_pragma_openmp 'requires' <clause> [[[,] <clause>] ... ]
 ///         annot_pragma_openmp_end
 ///
+///       assumes directive:
+///         annot_pragma_openmp 'assumes' <clause> [[[,] <clause>] ... ]
+///         annot_pragma_openmp_end
+///       or
+///         annot_pragma_openmp 'begin assumes' <clause> [[[,] <clause>] ... ]
+///         annot_pragma_openmp 'end assumes'
+///         annot_pragma_openmp_end
+///
 Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
     AccessSpecifier &AS, ParsedAttributesWithRange &Attrs, bool Delayed,
     DeclSpec::TST TagType, Decl *Tag) {
   assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
   ParsingOpenMPDirectiveRAII DirScope(*this);
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
 
   SourceLocation Loc;
if (Clauses.empty()) {
       Diag(Tok, diag::err_omp_expected_clause)
           << getOpenMPDirectiveName(OMPD_requires);
       ConsumeAnnotationToken();
       return nullptr;
     }
     ConsumeAnnotationToken();
     return Actions.ActOnOpenMPRequiresDirective(StartLoc, Clauses);
   }
+  case OMPD_assumes:
+  case OMPD_begin_assumes:
+    ParseOpenMPAssumesDirective(DKind, ConsumeToken());
+    break;
+  case OMPD_end_assumes:
+    ParseOpenMPEndAssumesDirective(ConsumeToken());
+    break;
   case OMPD_declare_reduction:
     ConsumeToken();
     if (DeclGroupPtrTy Res = ParseOpenMPDeclareReductionDirective(AS)) {
       skipUntilPragmaOpenMPEnd(OMPD_declare_reduction);
       // Skip the last annot_pragma_openmp_end.
       ConsumeAnnotationToken();
       return Res;
     }

clang/lib/Sema/SemaDecl.cpp

if (OtherUnmarkedIter != Previous.end()) {
       Diag((*OtherUnmarkedIter)->getLocation(),
            diag::note_attribute_overloadable_prev_overload)
           << false;
 
       NewFD->addAttr(OverloadableAttr::CreateImplicit(Context));
     }
   }
 
+  if (LangOpts.OpenMP)
+    ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
+
   // Semantic checking for this function declaration (in isolation).
 
   if (getLangOpts().CPlusPlus) {
     // C++-specific checks.
     if (CXXConstructorDecl *Constructor = dyn_cast<CXXConstructorDecl>(NewFD)) {
       CheckConstructor(Constructor);
     } else if (CXXDestructorDecl *Destructor =
                 dyn_cast<CXXDestructorDecl>(NewFD)) {

clang/lib/Sema/SemaOpenMP.cpp

 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
 #include "llvm/ADT/IndexedMap.h"
 #include "llvm/ADT/PointerEmbeddedInt.h"
 #include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/StringExtras.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include <set>
 
 using namespace clang;
 using namespace llvm::omp;
 
 //===----------------------------------------------------------------------===//
 // Stack of data-sharing attributes for variables
if (!CurContext->isFileContext()) {
     if (D) {
       CurContext->addDecl(D);
       DSAStack->addRequiresDecl(D);
     }
   }
   return DeclGroupPtrTy::make(DeclGroupRef(D));
 }
 
+void Sema::ActOnOpenMPAssumesDirective(SourceLocation Loc,
+                                       OpenMPDirectiveKind DKind,
+                                       ArrayRef<StringRef> Assumptions) {
+  auto *AA = AssumptionAttr::Create(Context, llvm::join(Assumptions, ","));
+  // Disable assumes in OpenMP simd mode.
+  if (DKind == llvm::omp::Directive::OMPD_begin_assumes) {
+    OMPAssumeScoped.push_back(AA);
+  } else {
+    assert(DKind == llvm::omp::Directive::OMPD_assumes && "");
+    OMPAssumeGlobal.push_back(AA);
+
+    // The OMPAssumeGlobal scope above will take care of new declarations but
+    // we also want to apply the assumption to existing ones, e.g., to
+    // declarations in included headers. To this end, we traverse all existing
+    // declaration contexts and annotate function declarations here.
+    SmallVector<DeclContext *, 8> DeclContexts;
+    DeclContexts.push_back(CurContext);
+    while (!DeclContexts.empty()) {
+      DeclContext *DC = DeclContexts.pop_back_val();
+      for (auto *SubDC : DC->decls()) {
+        if (auto *CTD = dyn_cast<ClassTemplateDecl>(SubDC)) {
+          DeclContexts.push_back(CTD->getTemplatedDecl());
+          for (auto *S : CTD->specializations())
+            DeclContexts.push_back(S);
+          continue;
+        }
+        if (isa<CXXRecordDecl>(SubDC) || isa<NamespaceDecl>(SubDC)) {
+          DeclContexts.push_back(cast<DeclContext>(SubDC));
+          continue;
+        }
+        if (auto *F = dyn_cast<FunctionDecl>(SubDC)) {
+          F->addAttr(AA);
+          continue;
+        }
+      }
+    }
+  }
+}
+
+void Sema::ActOnOpenMPEndAssumesDirective() {
+  assert(isInOpenMPAssumeScope() && "Not in OpenMP assumes scope!");
+  OMPAssumeScoped.pop_back();
+}
+
 OMPRequiresDecl *Sema::CheckOMPRequiresDecl(SourceLocation Loc,
                                             ArrayRef<OMPClause *> ClauseList) {
   /// For target specific clauses, the requires directive cannot be
   /// specified after the handling of any of the target regions in the
   /// current compilation unit.
   ArrayRef<SourceLocation> TargetLocations =
       DSAStack->getEncounteredTargetLocs();
   SourceLocation AtomicLoc = DSAStack->getAtomicDirectiveLoc();
for (const ParmVarDecl *P : FDWithProto->parameters()) {
     Param->setScopeInfo(0, Params.size());
     Param->setImplicit();
     Params.push_back(Param);
   }
 
   FD->setParams(Params);
 }
 
+void Sema::ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(Decl *D) {
+  FunctionDecl *FD = nullptr;
+  if (auto *UTemplDecl = dyn_cast<FunctionTemplateDecl>(D))
+    FD = UTemplDecl->getTemplatedDecl();
+  else
+    FD = cast<FunctionDecl>(D);
+  if (!FD)
+    return;
+
+  // If we are intantiating templates we do *not* apply scoped assumptions but
+  // only global ones. We apply scoped assumption to the template definition
+  // though.
+  if (!inTemplateInstantiation()) {
+    for (AssumptionAttr *AA : OMPAssumeScoped)
+      FD->addAttr(AA);
+  }
+  for (AssumptionAttr *AA : OMPAssumeGlobal)
+    FD->addAttr(AA);
+}
+
 Sema::OMPDeclareVariantScope::OMPDeclareVariantScope(OMPTraitInfo &TI)
     : TI(&TI), NameSuffix(TI.getMangledName()) {}
 
 void Sema::ActOnStartOfFunctionDefinitionInOpenMPDeclareVariantScope(
     Scope *S, Declarator &D, MultiTemplateParamsArg TemplateParamLists,
     SmallVectorImpl<FunctionDecl *> &Bases) {
   if (!D.getIdentifier())
     return;

clang/test/OpenMP/assumes_codegen.cpp

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s --check-prefix=AST
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -verify -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify=pch %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp -fexceptions -fcxx-exceptions -debug-info-kind=line-tables-only -x c++ -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-enable-irbuilder -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-enable-irbuilder -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -verify -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-enable-irbuilder -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify=pch %s -emit-llvm -o - | FileCheck %s
+
+// pch-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void foo() {
+}
+
+#pragma omp assumes no_openmp_routines warning ext_another_warning(1) ext_after_invalid_clauses // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{'ext_another_warning' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+
+#pragma omp assumes no_openmp
+
+#pragma omp begin assumes ext_range_bar_only
+
+#pragma omp begin assumes ext_range_bar_only_2
+
+class BAR {
+public:
+  BAR() {}
+
+  void bar1() {
+  }
+
+  static void bar2() {
+  }
+};
+
+void bar() { BAR b; }
+
+#pragma omp end assumes
+#pragma omp end assumes
+
+#pragma omp begin assumes ext_not_seen
+#pragma omp end assumes
+
+#pragma omp begin assumes ext_1234
+void baz();
+
+template<typename T>
+class BAZ {
+public:
+  BAZ() {}
+
+  void baz1() {
+  }
+
+  static void baz2() {
+  }
+};
+
+void baz() { BAZ<float> b; }
+#pragma omp end assumes
+
+// AST:      void foo() __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) {
+// AST-NEXT: }
+// AST-NEXT: class BAR {
+// AST-NEXT: public:
+// AST-NEXT:     BAR() __attribute__((assume("range_bar_only"))) __attribute__((assume("range_bar_only_2"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT:     void bar1() __attribute__((assume("range_bar_only"))) __attribute__((assume("range_bar_only_2"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT:     static void bar2() __attribute__((assume("range_bar_only"))) __attribute__((assume("range_bar_only_2"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT: };
+// AST-NEXT: void bar() __attribute__((assume("range_bar_only"))) __attribute__((assume("range_bar_only_2"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) {
+// AST-NEXT:     BAR b;
+// AST-NEXT: }
+// AST-NEXT: void baz() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")));
+// AST-NEXT: template <typename T> class BAZ {
+// AST-NEXT: public:
+// AST-NEXT:     BAZ<T>() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT:     void baz1() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT:     static void baz2() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))     {
+// AST-NEXT:     }
+// AST-NEXT: };
+// AST-NEXT: template<> class BAZ<float> {
+// AST-NEXT: public:
+// AST-NEXT:     BAZ() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")))    {
+// AST-NEXT:     }
+// AST-NEXT:     void baz1() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")));
+// AST-NEXT:     static void baz2() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp")));
+// AST-NEXT: };
+// AST-NEXT: void baz() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) {
+// AST-NEXT:     BAZ<float> b;
+// AST-NEXT: }
+
+#endif
+
+// CHECK: define void @_Z3foov()
+// CHECK-SAME: [[attr0:#[0-9]]]
+// CHECK: define void @_Z3barv()
+// CHECK-SAME: [[attr1:#[0-9]]]
+// CHECK:   call void @_ZN3BARC1Ev(%class.BAR* %b)
+// CHECK-SAME: [[attr7:#[0-9]]]
+// CHECK: define linkonce_odr void @_ZN3BARC1Ev(%class.BAR* %this)
+// CHECK-SAME: [[attr2:#[0-9]]]
+// CHECK:   call void @_ZN3BARC2Ev(%class.BAR* %this1)
+// CHECK-SAME: [[attr7]]
+// CHECK: define linkonce_odr void @_ZN3BARC2Ev(%class.BAR* %this)
+// CHECK-SAME: [[attr3:#[0-9]]]
+// CHECK: define void @_Z3bazv()
+// CHECK-SAME: [[attr4:#[0-9]]]
+// CHECK:   call void @_ZN3BAZIfEC1Ev(%class.BAZ* %b)
+// CHECK-SAME: [[attr8:#[0-9]]]
+// CHECK: define linkonce_odr void @_ZN3BAZIfEC1Ev(%class.BAZ* %this)
+// CHECK-SAME: [[attr5:#[0-9]]]
+// CHECK:   call void @_ZN3BAZIfEC2Ev(%class.BAZ* %this1)
+// CHECK-SAME: [[attr8]]
+// CHECK: define linkonce_odr void @_ZN3BAZIfEC2Ev(%class.BAZ* %this)
+// CHECK-SAME: [[attr6:#[0-9]]]
+
+// CHECK:     attributes [[attr0]]
+// CHECK-SAME:  "llvm.assume"="no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr1]]
+// CHECK-SAME:  "llvm.assume"="range_bar_only,range_bar_only_2,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr2]]
+// CHECK-SAME:  "llvm.assume"="range_bar_only,range_bar_only_2,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr3]]
+// CHECK-SAME:  "llvm.assume"="range_bar_only,range_bar_only_2,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr4]]
+// CHECK-SAME:  "llvm.assume"="1234,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp,1234,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr5]]
+// CHECK-SAME:  "llvm.assume"="1234,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr6]]
+// CHECK-SAME:  "llvm.assume"="1234,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr7]]
+// CHECK-SAME:  "llvm.assume"="range_bar_only,range_bar_only_2,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK:     attributes [[attr8]]
+// CHECK-SAME:  "llvm.assume"="1234,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"

clang/test/OpenMP/assumes_include_nvptx.cpp

+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown  -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+#include <cmath>
+
+// TODO: Think about teaching the OMPIRBuilder about default attributes as well so the __kmpc* declarations are annotated.
+
+// CHECK: define internal void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}_worker() [[attr0:#[0-9]*]]
+// CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0]]
+// CHECK: %call = call float @_Z3sinf(float 0.000000e+00) [[attr5:#[0-9]*]]
+// CHECK: declare i32 @llvm.nvvm.read.ptx.sreg.warpsize() [[attr1:#[0-9]*]]
+// CHECK: declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() [[attr1]]
+// CHECK: declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() [[attr1]]
+// CHECK: declare void @__kmpc_kernel_init(i32, i16)
+// CHECK-NOT: #
+// CHECK: declare void @__kmpc_data_sharing_init_stack()
+// CHECK-NOT: #
+// CHECK: declare float @_Z3sinf(float) [[attr2:#[0-9]*]]
+// CHECK: declare void @__kmpc_kernel_deinit(i16)
+// CHECK-NOT: #
+// CHECK: declare void @__kmpc_barrier_simple_spmd(%struct.ident_t*, i32) [[attr3:#[0-9]*]]
+// CHECK: declare i1 @__kmpc_kernel_parallel(i8**)
+// CHECK-NOT: #
+// CHECK: declare i32 @__kmpc_global_thread_num(%struct.ident_t*) [[attr4:#[0-9]*]]
+// CHECK: declare void @__kmpc_kernel_end_parallel()
+// CHECK-NOT: #
+// CHECK: define internal void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}_worker() [[attr0]]
+// CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]]
+// CHECK: %call = call double @_Z3sind(double 0.000000e+00) [[attr5]]
+// CHECK: declare double @_Z3sind(double) [[attr2]]
+
+// CHECK:       attributes [[attr0]]
+// CHECK-NOT:  "llvm.assume"
+// CHECK:       attributes [[attr1]]
+// CHECK-NOT:  "llvm.assume"
+// CHECK:       attributes [[attr2]]
+// CHECK-SAME:  "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
+// CHECK:       attributes [[attr3]]
+// CHECK-NOT:  "llvm.assume"
+// CHECK:       attributes [[attr4]]
+// CHECK-NOT:  "llvm.assume"
+// CHECK:       attributes [[attr5]]
+// CHECK-SAME:  "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
+
+
+template <typename T>
+void foo() {
+  cos(T(0));
+}
+
+template <typename T>
+void complex_reduction() {
+  foo<T>();
+#pragma omp target
+  sin(T(0));
+}
+
+#pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations
+
+void test() {
+  complex_reduction<float>();
+  complex_reduction<double>();
+}
+#endif

clang/test/OpenMP/assumes_messages.c

+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -x c -std=c99 -fms-extensions -Wno-pragma-pack %s
+
+#pragma omp assumes // expected-error {{'assumes' directive requires at least one clause}}
+#pragma omp begin // expected-error {{expected an OpenMP directive}}
+#pragma omp begin assumes // expected-error {{'begin assumes' directive requires at least one clause}}
+#pragma omp end assumes
+
+#pragma omp assumes foobar // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp begin assumes foobar // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp end assumes
+
+#pragma omp begin assumes foobar(foo 2 no_openmp // expected-error {{expected ')'}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{to match this '('}}
+#pragma omp assumes foobar(foo 2 no_openmp // expected-error {{expected ')'}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{to match this '('}}
+#pragma omp end assumes
+
+#pragma omp begin assumes foobar(foo 2 baz) // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp assumes foobar(foo 2 baz) // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp end assumes
+
+#pragma omp begin assumes foobar foo 2 baz) bar // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp assumes foobar foo 2 baz) bar // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp end assumes
+
+#pragma omp assumes no_openmp(1) // expected-warning {{'no_openmp' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp begin assumes no_openmp(1 2 3) // expected-warning {{'no_openmp' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp end assumes no_openmp(1)
+
+#pragma omp assumes foobar no_openmp bazbaz // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp begin assumes foobar no_openmp bazbaz // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp end assumes
+
+#pragma omp begin assumes foobar(foo 2 baz) no_openmp bazbaz(foo 2 baz) // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{the ignored tokens spans until here}}
+#pragma omp assumes foobar(foo 2 baz) no_openmp bazbaz(foo 2 baz) // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{the ignored tokens spans until here}}
+#pragma omp end assumes
+
+#pragma omp begin assumes foobar(foo (2) baz) no_openmp bazbaz(foo (2)) baz) // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{the ignored tokens spans until here}}
+#pragma omp assumes foobar(foo () baz) no_openmp bazbaz(foo ((2) baz) // expected-error {{expected ')'}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{to match this '('}}
+#pragma omp end assumes
+
+#pragma omp assumes no_openmp foobar no_openmp // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp begin assumes no_openmp foobar no_openmp // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp end assumes
+
+#pragma omp assumes holds(1, 2 3)
+#pragma omp begin assumes holds(1, 2 3)
+#pragma omp end assumes
+
+#pragma omp assumes absent(1, 2 3)
+#pragma omp begin assumes absent(1, 2 3)
+#pragma omp end assumes
+
+#pragma omp assumes contains(1, 2 3)
+#pragma omp begin assumes contains(1, 2 3)
+#pragma omp end assumes
+
+#pragma omp assumes ext // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp begin assumes ext // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+#pragma omp end assumes
+
+#pragma omp assumes ext_123(not allowed) // expected-warning {{'ext_123' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp begin assumes ext_123(not allowed) // expected-warning {{'ext_123' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+#pragma omp end assumes
+
+#pragma omp end assumes // expected-error {{'#pragma omp end assumes' with no matching '#pragma omp begin assumes'}}
+
+// TODO: we should emit a warning at least.
+#pragma omp begin assumes ext_abc

clang/test/OpenMP/assumes_print.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void foo() {
+}
+
+#pragma omp assumes no_openmp_routines
+
+#pragma omp assumes no_openmp
+
+#pragma omp begin assumes ext_range_bar_only
+
+#pragma omp begin assumes ext_range_bar_only_2
+
+void bar() {
+}
+
+#pragma omp end assumes
+#pragma omp end assumes
+
+#pragma omp begin assumes ext_not_seen
+#pragma omp end assumes
+
+#pragma omp begin assumes ext_1234
+void baz() {
+}
+#pragma omp end assumes
+
+// CHECK: void foo() __attribute__((assume("no_openmp_routines"))) __attribute__((assume("no_openmp")))
+// CHECK: void bar() __attribute__((assume("range_bar_only"))) __attribute__((assume("range_bar_only_2"))) __attribute__((assume("no_openmp_routines"))) __attribute__((assume("no_openmp")))
+// CHECK: void baz() __attribute__((assume("1234"))) __attribute__((assume("no_openmp_routines"))) __attribute__((assume("no_openmp")))
+
+#endif

clang/test/OpenMP/assumes_template_print.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+// It is unclear if we want to annotate the template instantiations, e.g., S<int>::foo, or not in the two
+// situations shown below. Since it is always fair to drop assumptions, we do that for now.
+
+#ifndef HEADER
+#define HEADER
+
+template <typename T>
+struct S {
+  int a;
+// CHECK: template <typename T> struct S {
+// CHECK:     void foo() __attribute__((assume("global_assumption")))     {
+  void foo() {
+    #pragma omp parallel
+    {}
+  }
+};
+
+// CHECK: template<> struct S<int> {
+// CHECK:     void foo() __attribute__((assume("global_assumption")))     {
+
+#pragma omp begin assumes no_openmp
+// CHECK: void S_with_assumes_no_call() __attribute__((assume("no_openmp"))) __attribute__((assume("global_assumption"))) {
+void S_with_assumes_no_call() {
+  S<int> s;
+  s.a = 0;
+}
+// CHECK: void S_with_assumes_call() __attribute__((assume("no_openmp"))) __attribute__((assume("global_assumption"))) {
+void S_with_assumes_call() {
+  S<int> s;
+  s.a = 0;
+  // If this is executed we have UB!
+  s.foo();
+}
+#pragma omp end assumes
+
+// CHECK: void S_without_assumes() __attribute__((assume("global_assumption"))) {
+void S_without_assumes() {
+  S<int> s;
+  s.foo();
+}
+
+#pragma omp assumes ext_global_assumption
+
+// Same as the struct S above but the order in which we instantiate P is different, first outside of an assumes.
+template <typename T>
+struct P {
+// CHECK: template <typename T> struct P {
+// CHECK:     void foo() __attribute__((assume("global_assumption")))     {
+  int a;
+  void foo() {
+    #pragma omp parallel
+    {}
+  }
+};
+
+// TODO: Avoid the duplication here:
+
+// CHECK: template<> struct P<int> {
+// CHECK:     void foo() __attribute__((assume("global_assumption"))) __attribute__((assume("global_assumption")))     {
+
+// CHECK: void P_without_assumes() __attribute__((assume("global_assumption"))) {
+void P_without_assumes() {
+  P<int> p;
+  p.foo();
+}
+
+#pragma omp begin assumes no_openmp
+// CHECK: void P_with_assumes_no_call() __attribute__((assume("no_openmp"))) __attribute__((assume("global_assumption"))) {
+void P_with_assumes_no_call() {
+  P<int> p;
+  p.a = 0;
+}
+// CHECK: void P_with_assumes_call() __attribute__((assume("no_openmp"))) __attribute__((assume("global_assumption"))) {
+void P_with_assumes_call() {
+  P<int> p;
+  p.a = 0;
+  // If this is executed we have UB!
+  p.foo();
+}
+#pragma omp end assumes
+
+#endif

llvm/include/llvm/Frontend/OpenMP/OMP.td

def OMP_Depobj : Directive<"depobj"> {
   ];
 }
 def OMP_Scan : Directive<"scan"> {
   let allowedClauses = [
     VersionedClause<OMPC_Inclusive, 50>,
     VersionedClause<OMPC_Exclusive, 50>
   ];
 }
+def OMP_Assumes : Directive<"assumes"> {}
+def OMP_BeginAssumes : Directive<"begin assumes"> {}
+def OMP_EndAssumes : Directive<"end assumes"> {}
 def OMP_BeginDeclareVariant : Directive<"begin declare variant"> {}
 def OMP_EndDeclareVariant : Directive<"end declare variant"> {}
 def OMP_ParallelWorkshare : Directive<"parallel workshare"> {
   let allowedClauses = [
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Copyin>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_FirstPrivate>,