[llvm-branch-commits] [clang] a5a14cb - [OpenMP] Add initial support for `omp [begin/end] assumes`
Johannes Doerfert via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Dec 15 14:56:40 PST 2020
Author: Johannes Doerfert
Date: 2020-12-15T16:51:34-06:00
New Revision: a5a14cbe7f87e01882ecaa14df5d596cbf38823a
URL: https://github.com/llvm/llvm-project/commit/a5a14cbe7f87e01882ecaa14df5d596cbf38823a
DIFF: https://github.com/llvm/llvm-project/commit/a5a14cbe7f87e01882ecaa14df5d596cbf38823a.diff
LOG: [OpenMP] Add initial support for `omp [begin/end] assumes`
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`.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D91980
Added:
clang/test/OpenMP/assumes_codegen.cpp
clang/test/OpenMP/assumes_include_nvptx.cpp
clang/test/OpenMP/assumes_messages.c
clang/test/OpenMP/assumes_print.cpp
clang/test/OpenMP/assumes_template_print.cpp
Modified:
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/Sema.h
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaLambda.cpp
clang/lib/Sema/SemaOpenMP.cpp
llvm/include/llvm/Frontend/OpenMP/OMP.td
llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index c0a5a2b4c144..8f78bbfc4e70 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1291,6 +1291,16 @@ def err_expected_end_declare_target_or_variant : Error<
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 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<
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 73943ed5f0a6..02aab3b43be0 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3104,6 +3104,13 @@ class Parser : public CodeCompletionHandler {
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'.
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index eb6429912bf6..ff0257634d9d 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10128,6 +10128,13 @@ class Sema final {
/// 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
@@ -10141,10 +10148,11 @@ class Sema final {
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();
}
@@ -10260,6 +10268,22 @@ class Sema final {
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,
+ bool SkippedClauses);
+
+ /// 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);
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 6f5ec060e5ff..c4aa361b8262 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -21,6 +21,7 @@
#include "clang/Parse/RAIIObjectsForParser.h"
#include "clang/Sema/Scope.h"
#include "llvm/ADT/PointerIntPair.h"
+#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/UniqueVector.h"
#include "llvm/Frontend/OpenMP/OMPContext.h"
@@ -115,7 +116,9 @@ static OpenMPDirectiveKindExWrapper parseOpenMPDirectiveKind(Parser &P) {
// 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},
@@ -1508,6 +1511,103 @@ bool Parser::parseOMPDeclareVariantMatchClause(SourceLocation Loc,
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);
+ };
+
+ /// Helper to determine which AssumptionClauseMapping (ACM) in the
+ /// AssumptionClauseMappings table matches \p RawString. The return value is
+ /// the index of the matching ACM into the table or -1 if there was no match.
+ auto MatchACMClause = [&](StringRef RawString) {
+ llvm::StringSwitch<int> SS(RawString);
+ unsigned ACMIdx = 0;
+ for (const AssumptionClauseMappingInfo &ACMI : AssumptionClauseMappings) {
+ if (ACMI.StartsWith)
+ SS.StartsWith(ACMI.Identifier, ACMIdx++);
+ else
+ SS.Case(ACMI.Identifier, ACMIdx++);
+ }
+ return SS.Default(-1);
+ };
+
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ IdentifierInfo *II = nullptr;
+ SourceLocation StartLoc = Tok.getLocation();
+ int Idx = -1;
+ if (Tok.isAnyIdentifier()) {
+ II = Tok.getIdentifierInfo();
+ Idx = MatchACMClause(II->getName());
+ }
+ 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)
+ << llvm::omp::getAllAssumeClauseOptions() << 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);
+ }
+
+ Actions.ActOnOpenMPAssumesDirective(Loc, DKind, Assumptions, SkippedClauses);
+}
+
+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:
@@ -1716,6 +1816,14 @@ void Parser::ParseOMPEndDeclareTargetDirective(OpenMPDirectiveKind DKind,
/// 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) {
@@ -1853,6 +1961,13 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
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)) {
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 0031f874c05a..d6be0d257329 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -10842,6 +10842,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
}
}
+ if (LangOpts.OpenMP)
+ ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
+
// Semantic checking for this function declaration (in isolation).
if (getLangOpts().CPlusPlus) {
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index ec1db50cfaaa..af61c82c2002 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -998,6 +998,10 @@ void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
if (getLangOpts().CUDA)
CUDASetLambdaAttrs(Method);
+ // OpenMP lambdas might get assumumption attributes.
+ if (LangOpts.OpenMP)
+ ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(Method);
+
// Number the lambda for linkage purposes if necessary.
handleLambdaNumbering(Class, Method);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index b01184c402db..78707484f588 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -35,6 +35,7 @@
#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>
@@ -3195,6 +3196,64 @@ Sema::ActOnOpenMPRequiresDirective(SourceLocation Loc,
return DeclGroupPtrTy::make(DeclGroupRef(D));
}
+void Sema::ActOnOpenMPAssumesDirective(SourceLocation Loc,
+ OpenMPDirectiveKind DKind,
+ ArrayRef<StringRef> Assumptions,
+ bool SkippedClauses) {
+ if (!SkippedClauses && Assumptions.empty())
+ Diag(Loc, diag::err_omp_no_clause_for_directive)
+ << llvm::omp::getAllAssumeClauseOptions()
+ << llvm::omp::getOpenMPDirectiveName(DKind);
+
+ auto *AA = AssumptionAttr::Create(Context, llvm::join(Assumptions, ","), Loc);
+ if (DKind == llvm::omp::Directive::OMPD_begin_assumes) {
+ OMPAssumeScoped.push_back(AA);
+ return;
+ }
+
+ // Global assumes without assumption clauses are ignored.
+ if (Assumptions.empty())
+ return;
+
+ assert(DKind == llvm::omp::Directive::OMPD_assumes &&
+ "Unexpected omp assumption directive!");
+ 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;
+ auto *Ctx = CurContext;
+ while (Ctx->getLexicalParent())
+ Ctx = Ctx->getLexicalParent();
+ DeclContexts.push_back(Ctx);
+ while (!DeclContexts.empty()) {
+ DeclContext *DC = DeclContexts.pop_back_val();
+ for (auto *SubDC : DC->decls()) {
+ if (SubDC->isInvalidDecl())
+ continue;
+ if (auto *CTD = dyn_cast<ClassTemplateDecl>(SubDC)) {
+ DeclContexts.push_back(CTD->getTemplatedDecl());
+ for (auto *S : CTD->specializations())
+ DeclContexts.push_back(S);
+ continue;
+ }
+ if (auto *DC = dyn_cast<DeclContext>(SubDC))
+ DeclContexts.push_back(DC);
+ 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
@@ -5936,6 +5995,27 @@ static void setPrototype(Sema &S, FunctionDecl *FD, FunctionDecl *FDWithProto,
FD->setParams(Params);
}
+void Sema::ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(Decl *D) {
+ if (D->isInvalidDecl())
+ return;
+ FunctionDecl *FD = nullptr;
+ if (auto *UTemplDecl = dyn_cast<FunctionTemplateDecl>(D))
+ FD = UTemplDecl->getTemplatedDecl();
+ else
+ FD = cast<FunctionDecl>(D);
+ assert(FD && "Expected a function declaration!");
+
+ // 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()) {}
diff --git a/clang/test/OpenMP/assumes_codegen.cpp b/clang/test/OpenMP/assumes_codegen.cpp
new file mode 100644
index 000000000000..45935b5a2332
--- /dev/null
+++ b/clang/test/OpenMP/assumes_codegen.cpp
@@ -0,0 +1,165 @@
+// 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
+
+#pragma omp begin assumes ext_lambda_assumption
+int lambda_outer() {
+ auto lambda_inner = []() { return 42; };
+ return lambda_inner();
+}
+#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: }
+// AST-NEXT: int lambda_outer() __attribute__((assume("lambda_assumption"))) __attribute__((assume("no_openmp_routines,another_warning,after_invalid_clauses"))) __attribute__((assume("no_openmp"))) {
+// AST-NEXT: auto lambda_inner = []() {
+// AST-NEXT: return 42;
+// AST-NEXT: };
+// AST-NEXT: return lambda_inner();
+// 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: [[attr9:#[0-9]]]
+// CHECK: define{{.*}} void @_ZN3BARC1Ev(%class.BAR*{{.*}} %this)
+// CHECK-SAME: [[attr2:#[0-9]]]
+// CHECK: call void @_ZN3BARC2Ev(%class.BAR*{{.*}} %this1)
+// CHECK-SAME: [[attr9]]
+// CHECK: define{{.*}} 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: [[attr10:#[0-9]]]
+// CHECK: define{{.*}} void @_ZN3BAZIfEC1Ev(%class.BAZ*{{.*}} %this)
+// CHECK-SAME: [[attr5:#[0-9]]]
+// CHECK: call void @_ZN3BAZIfEC2Ev(%class.BAZ*{{.*}} %this1)
+// CHECK-SAME: [[attr10]]
+// CHECK: define{{.*}} void @_ZN3BAZIfEC2Ev(%class.BAZ*{{.*}} %this)
+// CHECK-SAME: [[attr6:#[0-9]]]
+// CHECK: define{{.*}} i32 @_Z12lambda_outerv()
+// CHECK-SAME: [[attr7:#[0-9]]]
+// CHECK: call i32 @"_ZZ12lambda_outervENK3$_0clEv"
+// CHECK-SAME: [[attr11:#[0-9]]]
+// CHECK: define{{.*}} i32 @"_ZZ12lambda_outervENK3$_0clEv"(%class.anon*{{.*}} %this)
+// CHECK-SAME: [[attr8:#[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"="lambda_assumption,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK: attributes [[attr8]]
+// CHECK-SAME: "llvm.assume"="lambda_assumption,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK: attributes [[attr9]]
+// CHECK-SAME: "llvm.assume"="range_bar_only,range_bar_only_2,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
+// CHECK: attributes [[attr10]]
+// 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 [[attr11]]
+// CHECK-SAME: "llvm.assume"="lambda_assumption,no_openmp_routines,another_warning,after_invalid_clauses,no_openmp"
diff --git a/clang/test/OpenMP/assumes_include_nvptx.cpp b/clang/test/OpenMP/assumes_include_nvptx.cpp
new file mode 100644
index 000000000000..188726601c00
--- /dev/null
+++ b/clang/test/OpenMP/assumes_include_nvptx.cpp
@@ -0,0 +1,70 @@
+// 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
diff --git a/clang/test/OpenMP/assumes_messages.c b/clang/test/OpenMP/assumes_messages.c
new file mode 100644
index 000000000000..99bbe93dc1fc
--- /dev/null
+++ b/clang/test/OpenMP/assumes_messages.c
@@ -0,0 +1,69 @@
+// 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 {{expected at least one 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism' clause for '#pragma omp assumes'}}
+#pragma omp begin // expected-error {{expected an OpenMP directive}}
+#pragma omp begin assumes // expected-error {{expected at least one 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism' clause for '#pragma omp begin assumes'}}
+#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
diff --git a/clang/test/OpenMP/assumes_print.cpp b/clang/test/OpenMP/assumes_print.cpp
new file mode 100644
index 000000000000..1112d99bcb94
--- /dev/null
+++ b/clang/test/OpenMP/assumes_print.cpp
@@ -0,0 +1,44 @@
+// 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
+
+namespace inner {
+ #pragma omp assumes no_openmp
+} // namespace inner
+
+#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
diff --git a/clang/test/OpenMP/assumes_template_print.cpp b/clang/test/OpenMP/assumes_template_print.cpp
new file mode 100644
index 000000000000..7b4090a043d2
--- /dev/null
+++ b/clang/test/OpenMP/assumes_template_print.cpp
@@ -0,0 +1,91 @@
+// 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
diff erent, 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
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index 59198f8a64fb..58aa1bf23b68 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -1593,6 +1593,9 @@ def OMP_Scan : Directive<"scan"> {
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"> {
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
index 3ad13ddc5136..36ce3fc0f66f 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
@@ -16,6 +16,7 @@
#include "llvm/ADT/BitmaskEnum.h"
+#include "llvm/ADT/StringRef.h"
#include "llvm/Frontend/OpenMP/OMP.h.inc"
namespace llvm {
@@ -79,6 +80,33 @@ enum class IdentFlag {
#define OMP_IDENT_FLAG(Enum, ...) constexpr auto Enum = omp::IdentFlag::Enum;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
+/// Helper to describe assume clauses.
+struct AssumptionClauseMappingInfo {
+ /// The identifier describing the (beginning of the) clause.
+ llvm::StringLiteral Identifier;
+ /// Flag to determine if the identifier is a full name or the start of a name.
+ bool StartsWith;
+ /// Flag to determine if a directive lists follows.
+ bool HasDirectiveList;
+ /// Flag to determine if an expression follows.
+ bool HasExpression;
+};
+
+/// All known assume clauses.
+static constexpr AssumptionClauseMappingInfo AssumptionClauseMappings[] = {
+#define OMP_ASSUME_CLAUSE(Identifier, StartsWith, HasDirectiveList, \
+ HasExpression) \
+ {Identifier, StartsWith, HasDirectiveList, HasExpression},
+#include "llvm/Frontend/OpenMP/OMPKinds.def"
+};
+
+inline std::string getAllAssumeClauseOptions() {
+ std::string S;
+ for (const AssumptionClauseMappingInfo &ACMI : AssumptionClauseMappings)
+ S += (S.empty() ? "'" : "', '") + ACMI.Identifier.str();
+ return S + "'";
+}
+
} // end namespace omp
} // end namespace llvm
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 947b0c8d5ef2..ad4f8bd5b079 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -1224,3 +1224,27 @@ OMP_LAST_TRAIT_PROPERTY(
#undef __OMP_REQUIRES_TRAIT
#undef OMP_REQUIRES_TRAIT
///}
+
+
+/// Assumption clauses
+///
+///{
+
+#ifdef OMP_ASSUME_CLAUSE
+#define __OMP_ASSUME_CLAUSE(Identifier, StartsWith, HasDirectiveList, HasExpression) \
+OMP_ASSUME_CLAUSE(Identifier, StartsWith, HasDirectiveList, HasExpression)
+#else
+#define __OMP_ASSUME_CLAUSE(...)
+#endif
+
+__OMP_ASSUME_CLAUSE("ext_", true, false, false)
+__OMP_ASSUME_CLAUSE("absent", false, true, false)
+__OMP_ASSUME_CLAUSE("contains", false, true, false)
+__OMP_ASSUME_CLAUSE("holds", false, false, true)
+__OMP_ASSUME_CLAUSE("no_openmp", false, false, false)
+__OMP_ASSUME_CLAUSE("no_openmp_routines", false, false, false)
+__OMP_ASSUME_CLAUSE("no_parallelism", false, false, false)
+
+#undef __OMP_ASSUME_CLAUSE
+#undef OMP_ASSUME_CLAUSE
+///}
More information about the llvm-branch-commits
mailing list