[clang] 2d4f80f - [OPENMP50]Full handling of atomic_default_mem_order in requires
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 11 12:54:00 PST 2020
Author: Alexey Bataev
Date: 2020-02-11T15:42:34-05:00
New Revision: 2d4f80f78aa5c25f19c396bf85d022d009706936
URL: https://github.com/llvm/llvm-project/commit/2d4f80f78aa5c25f19c396bf85d022d009706936
DIFF: https://github.com/llvm/llvm-project/commit/2d4f80f78aa5c25f19c396bf85d022d009706936.diff
LOG: [OPENMP50]Full handling of atomic_default_mem_order in requires
directive.
According to OpenMP 5.0, The atomic_default_mem_order clause specifies the default memory ordering behavior for atomic constructs that must be provided by an implementation. If the default memory ordering is specified as seq_cst, all atomic constructs on which memory-order-clause is not specified behave as if the seq_cst clause appears. If the default memory ordering is specified as relaxed, all atomic constructs on which memory-order-clause is not specified behave as if the relaxed clause appears.
If the default memory ordering is specified as acq_rel, atomic constructs on which memory-order-clause is not specified behave as if the release clause appears if the atomic write or atomic update operation is specified, as if the acquire clause appears if the atomic read operation is specified, and as if the acq_rel clause appears if the atomic captured update operation is specified.
Added:
clang/test/OpenMP/requires_acq_rel_codegen.cpp
clang/test/OpenMP/requires_default_atomic_mem_order_messages.cpp
clang/test/OpenMP/requires_relaxed_codegen.cpp
clang/test/OpenMP/requires_seq_cst_codegen.cpp
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/AST/ASTContext.cpp
clang/lib/CodeGen/CGDecl.cpp
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Serialization/ASTReaderDecl.cpp
clang/test/OpenMP/requires_target_messages.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index e5714a744692..37f0acf010d6 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9852,10 +9852,10 @@ def err_omp_requires_clause_redeclaration : Error <
"Only one %0 clause can appear on a requires directive in a single translation unit">;
def note_omp_requires_previous_clause : Note <
"%0 clause previously used here">;
-def err_omp_target_before_requires : Error <
- "target region encountered before requires directive with '%0' clause">;
-def note_omp_requires_encountered_target : Note <
- "target previously encountered here">;
+def err_omp_directive_before_requires : Error <
+ "'%0' region encountered before requires directive with '%1' clause">;
+def note_omp_requires_encountered_directive : Note <
+ "'%0' previously encountered here">;
def err_omp_invalid_scope : Error <
"'#pragma omp %0' directive must appear only in file scope">;
def note_omp_invalid_length_on_this_ptr_mapping : Note <
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index c80d3948d003..50a0c3d76da2 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -10041,6 +10041,8 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
return true;
else if (isa<PragmaDetectMismatchDecl>(D))
return true;
+ else if (isa<OMPRequiresDecl>(D))
+ return true;
else if (isa<OMPThreadPrivateDecl>(D))
return !D->getDeclContext()->isDependentContext();
else if (isa<OMPAllocateDecl>(D))
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 29462592887a..1767e744bac7 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -2537,5 +2537,5 @@ void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
}
void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
- getOpenMPRuntime().checkArchForUnifiedAddressing(D);
+ getOpenMPRuntime().processRequiresDirective(D);
}
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 23d49b23a3b4..c3e2e1e0a5d9 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -21,6 +21,7 @@
#include "clang/AST/StmtOpenMP.h"
#include "clang/AST/StmtVisitor.h"
#include "clang/Basic/BitmaskEnum.h"
+#include "clang/Basic/OpenMPKinds.h"
#include "clang/CodeGen/ConstantInitBuilder.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/SetOperations.h"
@@ -30,6 +31,7 @@
#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/GlobalValue.h"
#include "llvm/IR/Value.h"
+#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/Format.h"
#include "llvm/Support/raw_ostream.h"
#include <cassert>
@@ -9784,16 +9786,33 @@ void CGOpenMPRuntime::adjustTargetSpecificDataForLambdas(
" Expected target-based directive.");
}
-void CGOpenMPRuntime::checkArchForUnifiedAddressing(
- const OMPRequiresDecl *D) {
+void CGOpenMPRuntime::processRequiresDirective(const OMPRequiresDecl *D) {
for (const OMPClause *Clause : D->clauselists()) {
if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
HasRequiresUnifiedSharedMemory = true;
- break;
+ } else if (const auto *AC =
+ dyn_cast<OMPAtomicDefaultMemOrderClause>(Clause)) {
+ switch (AC->getAtomicDefaultMemOrderKind()) {
+ case OMPC_ATOMIC_DEFAULT_MEM_ORDER_acq_rel:
+ RequiresAtomicOrdering = llvm::AtomicOrdering::AcquireRelease;
+ break;
+ case OMPC_ATOMIC_DEFAULT_MEM_ORDER_seq_cst:
+ RequiresAtomicOrdering = llvm::AtomicOrdering::SequentiallyConsistent;
+ break;
+ case OMPC_ATOMIC_DEFAULT_MEM_ORDER_relaxed:
+ RequiresAtomicOrdering = llvm::AtomicOrdering::Monotonic;
+ break;
+ case OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown:
+ break;
+ }
}
}
}
+llvm::AtomicOrdering CGOpenMPRuntime::getDefaultMemoryOrdering() const {
+ return RequiresAtomicOrdering;
+}
+
bool CGOpenMPRuntime::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
LangAS &AS) {
if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index a4fe15eee26a..f559e0d22574 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -26,6 +26,7 @@
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/ValueHandle.h"
+#include "llvm/Support/AtomicOrdering.h"
namespace llvm {
class ArrayType;
@@ -704,6 +705,9 @@ class CGOpenMPRuntime {
/// directive is present.
bool HasRequiresUnifiedSharedMemory = false;
+ /// Atomic ordering from the omp requires directive.
+ llvm::AtomicOrdering RequiresAtomicOrdering = llvm::AtomicOrdering::Monotonic;
+
/// Flag for keeping track of weather a target region has been emitted.
bool HasEmittedTargetRegion = false;
@@ -1700,7 +1704,10 @@ class CGOpenMPRuntime {
/// Perform check on requires decl to ensure that target architecture
/// supports unified addressing
- virtual void checkArchForUnifiedAddressing(const OMPRequiresDecl *D);
+ virtual void processRequiresDirective(const OMPRequiresDecl *D);
+
+ /// Gets default memory ordering as specified in requires directive.
+ llvm::AtomicOrdering getDefaultMemoryOrdering() const;
/// Checks if the variable has associated OMPAllocateDeclAttr attribute with
/// the predefined allocator and translates it into the corresponding address
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
index d00d84b79cfe..867bfb072736 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
@@ -4962,7 +4962,7 @@ static CudaArch getCudaArch(CodeGenModule &CGM) {
/// Check to see if target architecture supports unified addressing which is
/// a restriction for OpenMP requires clause "unified_shared_memory".
-void CGOpenMPRuntimeNVPTX::checkArchForUnifiedAddressing(
+void CGOpenMPRuntimeNVPTX::processRequiresDirective(
const OMPRequiresDecl *D) {
for (const OMPClause *Clause : D->clauselists()) {
if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
@@ -5017,7 +5017,7 @@ void CGOpenMPRuntimeNVPTX::checkArchForUnifiedAddressing(
}
}
}
- CGOpenMPRuntime::checkArchForUnifiedAddressing(D);
+ CGOpenMPRuntime::processRequiresDirective(D);
}
/// Get number of SMs and number of blocks per SM.
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
index 4159af0a622f..834adb3782a0 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
@@ -395,7 +395,7 @@ class CGOpenMPRuntimeNVPTX : public CGOpenMPRuntime {
/// Perform check on requires decl to ensure that target architecture
/// supports unified addressing
- void checkArchForUnifiedAddressing(const OMPRequiresDecl *D) override;
+ void processRequiresDirective(const OMPRequiresDecl *D) override;
/// Returns default address space for the constant firstprivates, __constant__
/// address space by default.
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 16914648a34d..7181374a73fc 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -25,6 +25,7 @@
#include "clang/Basic/PrettyStackTrace.h"
#include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
#include "llvm/IR/Instructions.h"
+#include "llvm/Support/AtomicOrdering.h"
using namespace clang;
using namespace CodeGen;
using namespace llvm::omp;
@@ -4542,16 +4543,23 @@ static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
llvm::AtomicOrdering AO = llvm::AtomicOrdering::Monotonic;
- if (S.getSingleClause<OMPSeqCstClause>())
+ bool MemOrderingSpecified = false;
+ if (S.getSingleClause<OMPSeqCstClause>()) {
AO = llvm::AtomicOrdering::SequentiallyConsistent;
- else if (S.getSingleClause<OMPAcqRelClause>())
+ MemOrderingSpecified = true;
+ } else if (S.getSingleClause<OMPAcqRelClause>()) {
AO = llvm::AtomicOrdering::AcquireRelease;
- else if (S.getSingleClause<OMPAcquireClause>())
+ MemOrderingSpecified = true;
+ } else if (S.getSingleClause<OMPAcquireClause>()) {
AO = llvm::AtomicOrdering::Acquire;
- else if (S.getSingleClause<OMPReleaseClause>())
+ MemOrderingSpecified = true;
+ } else if (S.getSingleClause<OMPReleaseClause>()) {
AO = llvm::AtomicOrdering::Release;
- else if (S.getSingleClause<OMPRelaxedClause>())
+ MemOrderingSpecified = true;
+ } else if (S.getSingleClause<OMPRelaxedClause>()) {
AO = llvm::AtomicOrdering::Monotonic;
+ MemOrderingSpecified = true;
+ }
OpenMPClauseKind Kind = OMPC_unknown;
for (const OMPClause *C : S.clauses()) {
// Find first clause (skip seq_cst|acq_rel|aqcuire|release|relaxed clause,
@@ -4565,6 +4573,23 @@ void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
break;
}
}
+ if (!MemOrderingSpecified) {
+ llvm::AtomicOrdering DefaultOrder =
+ CGM.getOpenMPRuntime().getDefaultMemoryOrdering();
+ if (DefaultOrder == llvm::AtomicOrdering::Monotonic ||
+ DefaultOrder == llvm::AtomicOrdering::SequentiallyConsistent ||
+ (DefaultOrder == llvm::AtomicOrdering::AcquireRelease &&
+ Kind == OMPC_capture)) {
+ AO = DefaultOrder;
+ } else if (DefaultOrder == llvm::AtomicOrdering::AcquireRelease) {
+ if (Kind == OMPC_unknown || Kind == OMPC_update || Kind == OMPC_write) {
+ AO = llvm::AtomicOrdering::Release;
+ } else if (Kind == OMPC_read) {
+ assert(Kind == OMPC_read && "Unexpected atomic kind.");
+ AO = llvm::AtomicOrdering::Acquire;
+ }
+ }
+ }
const Stmt *CS = S.getInnermostCapturedStmt()->IgnoreContainers();
if (const auto *FE = dyn_cast<FullExpr>(CS))
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index c893ffc13f51..fbabe92977c9 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -1473,7 +1473,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
Actions.EndOpenMPClause();
}
// Consume final annot_pragma_openmp_end
- if (Clauses.size() == 0) {
+ if (Clauses.empty()) {
Diag(Tok, diag::err_omp_expected_clause)
<< getOpenMPDirectiveName(OMPD_requires);
ConsumeAnnotationToken();
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 2a13f72f6936..bd1c4777bc0d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -271,6 +271,7 @@ class DSAStackTy {
nullptr};
/// Vector of previously encountered target directives
SmallVector<SourceLocation, 2> TargetLocations;
+ SourceLocation AtomicLocation;
public:
explicit DSAStackTy(Sema &S) : SemaRef(S) {}
@@ -555,7 +556,7 @@ class DSAStackTy {
/// Checks if the defined 'requires' directive has specified type of clause.
template <typename ClauseType>
- bool hasRequiresDeclWithClause() {
+ bool hasRequiresDeclWithClause() const {
return llvm::any_of(RequiresDecls, [](const OMPRequiresDecl *D) {
return llvm::any_of(D->clauselists(), [](const OMPClause *C) {
return isa<ClauseType>(C);
@@ -590,6 +591,18 @@ class DSAStackTy {
TargetLocations.push_back(LocStart);
}
+ /// Add location for the first encountered atomicc directive.
+ void addAtomicDirectiveLoc(SourceLocation Loc) {
+ if (AtomicLocation.isInvalid())
+ AtomicLocation = Loc;
+ }
+
+ /// Returns the location of the first encountered atomic directive in the
+ /// module.
+ SourceLocation getAtomicDirectiveLoc() const {
+ return AtomicLocation;
+ }
+
// Return previously encountered target region locations.
ArrayRef<SourceLocation> getEncounteredTargetLocs() const {
return TargetLocations;
@@ -2830,18 +2843,26 @@ OMPRequiresDecl *Sema::CheckOMPRequiresDecl(SourceLocation Loc,
/// current compilation unit.
ArrayRef<SourceLocation> TargetLocations =
DSAStack->getEncounteredTargetLocs();
- if (!TargetLocations.empty()) {
+ SourceLocation AtomicLoc = DSAStack->getAtomicDirectiveLoc();
+ if (!TargetLocations.empty() || !AtomicLoc.isInvalid()) {
for (const OMPClause *CNew : ClauseList) {
// Check if any of the requires clauses affect target regions.
if (isa<OMPUnifiedSharedMemoryClause>(CNew) ||
isa<OMPUnifiedAddressClause>(CNew) ||
isa<OMPReverseOffloadClause>(CNew) ||
isa<OMPDynamicAllocatorsClause>(CNew)) {
- Diag(Loc, diag::err_omp_target_before_requires)
- << getOpenMPClauseName(CNew->getClauseKind());
+ Diag(Loc, diag::err_omp_directive_before_requires)
+ << "target" << getOpenMPClauseName(CNew->getClauseKind());
for (SourceLocation TargetLoc : TargetLocations) {
- Diag(TargetLoc, diag::note_omp_requires_encountered_target);
+ Diag(TargetLoc, diag::note_omp_requires_encountered_directive)
+ << "target";
}
+ } else if (!AtomicLoc.isInvalid() &&
+ isa<OMPAtomicDefaultMemOrderClause>(CNew)) {
+ Diag(Loc, diag::err_omp_directive_before_requires)
+ << "atomic" << getOpenMPClauseName(CNew->getClauseKind());
+ Diag(AtomicLoc, diag::note_omp_requires_encountered_directive)
+ << "atomic";
}
}
}
@@ -8926,6 +8947,8 @@ StmtResult Sema::ActOnOpenMPAtomicDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt,
SourceLocation StartLoc,
SourceLocation EndLoc) {
+ // Register location of the first atomic directive.
+ DSAStack->addAtomicDirectiveLoc(StartLoc);
if (!AStmt)
return StmtError();
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 093b69ab19d0..3d47274079fa 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -2840,7 +2840,8 @@ static bool isConsumerInterestedIn(ASTContext &Ctx, Decl *D, bool HasBody) {
isa<PragmaDetectMismatchDecl>(D))
return true;
if (isa<OMPThreadPrivateDecl>(D) || isa<OMPDeclareReductionDecl>(D) ||
- isa<OMPDeclareMapperDecl>(D) || isa<OMPAllocateDecl>(D))
+ isa<OMPDeclareMapperDecl>(D) || isa<OMPAllocateDecl>(D) ||
+ isa<OMPRequiresDecl>(D))
return !D->getDeclContext()->isFunctionOrMethod();
if (const auto *Var = dyn_cast<VarDecl>(D))
return Var->isFileVarDecl() &&
diff --git a/clang/test/OpenMP/requires_acq_rel_codegen.cpp b/clang/test/OpenMP/requires_acq_rel_codegen.cpp
new file mode 100644
index 000000000000..b8ba01b0cafe
--- /dev/null
+++ b/clang/test/OpenMP/requires_acq_rel_codegen.cpp
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o -| FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd %s -fopenmp-version=50 -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -emit-llvm -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp requires atomic_default_mem_order(acq_rel)
+
+// CHECK-LABEL: foo
+void foo() {
+ int a = 0, b = 0;
+// CHECK: load atomic i32,{{.*}}acquire
+#pragma omp atomic read
+ a = b;
+// CHECK: store atomic i32{{.*}}release
+#pragma omp atomic write
+ a = b;
+// CHECK: atomicrmw add i32{{.*}}release
+#pragma omp atomic
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}release
+#pragma omp atomic update
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}acq_rel
+#pragma omp atomic capture
+ {
+ b = a;
+ a += 1;
+ }
+}
+
+#endif
diff --git a/clang/test/OpenMP/requires_default_atomic_mem_order_messages.cpp b/clang/test/OpenMP/requires_default_atomic_mem_order_messages.cpp
new file mode 100644
index 000000000000..19f6ede043d8
--- /dev/null
+++ b/clang/test/OpenMP/requires_default_atomic_mem_order_messages.cpp
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo2() {
+ int a;
+ #pragma omp atomic update // expected-note 3 {{'atomic' previously encountered here}}
+ a = a + 1;
+}
+
+#pragma omp requires atomic_default_mem_order(seq_cst) // expected-error {{'atomic' region encountered before requires directive with 'atomic_default_mem_order' clause}} expected-note 2 {{atomic_default_mem_order clause previously used here}}
+#pragma omp requires atomic_default_mem_order(acq_rel) // expected-error {{'atomic' region encountered before requires directive with 'atomic_default_mem_order' clause}} expected-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
+#pragma omp requires atomic_default_mem_order(relaxed) // expected-error {{'atomic' region encountered before requires directive with 'atomic_default_mem_order' clause}} expected-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
+#pragma omp requires atomic_default_mem_order(release) // expected-error {{expected 'seq_cst', 'acq_rel' or 'relaxed' in OpenMP clause 'atomic_default_mem_order'}} expected-error {{expected at least one clause on '#pragma omp requires' directive}}
diff --git a/clang/test/OpenMP/requires_relaxed_codegen.cpp b/clang/test/OpenMP/requires_relaxed_codegen.cpp
new file mode 100644
index 000000000000..e92b55e8f137
--- /dev/null
+++ b/clang/test/OpenMP/requires_relaxed_codegen.cpp
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o -| FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd %s -fopenmp-version=50 -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -emit-llvm -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp requires atomic_default_mem_order(relaxed)
+
+// CHECK-LABEL: foo
+void foo() {
+ int a = 0, b = 0;
+// CHECK: load atomic i32,{{.*}}monotonic
+#pragma omp atomic read
+ a = b;
+// CHECK: store atomic i32{{.*}}monotonic
+#pragma omp atomic write
+ a = b;
+// CHECK: atomicrmw add i32{{.*}}monotonic
+#pragma omp atomic
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}monotonic
+#pragma omp atomic update
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}monotonic
+#pragma omp atomic capture
+ {
+ b = a;
+ a += 1;
+ }
+}
+
+#endif
diff --git a/clang/test/OpenMP/requires_seq_cst_codegen.cpp b/clang/test/OpenMP/requires_seq_cst_codegen.cpp
new file mode 100644
index 000000000000..c2f02665d153
--- /dev/null
+++ b/clang/test/OpenMP/requires_seq_cst_codegen.cpp
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -triple x86_64-apple-darwin10 -x c++ -emit-llvm -o -| FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd %s -fopenmp-version=50 -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s -triple x86_64-apple-darwin10
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -verify %s -emit-llvm -x c++ -emit-llvm -triple x86_64-apple-darwin10 -o -| FileCheck %s --check-prefix SIMD-ONLY0
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp requires atomic_default_mem_order(seq_cst)
+
+// CHECK-LABEL: foo
+void foo() {
+ int a = 0, b = 0;
+// CHECK: load atomic i32,{{.*}}seq_cst
+#pragma omp atomic read
+ a = b;
+// CHECK: store atomic i32{{.*}}seq_cst
+#pragma omp atomic write
+ a = b;
+// CHECK: atomicrmw add i32{{.*}}seq_cst
+#pragma omp atomic
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}seq_cst
+#pragma omp atomic update
+ a += 1;
+// CHECK: atomicrmw add i32{{.*}}seq_cst
+#pragma omp atomic capture
+ {
+ b = a;
+ a += 1;
+ }
+}
+
+#endif
diff --git a/clang/test/OpenMP/requires_target_messages.cpp b/clang/test/OpenMP/requires_target_messages.cpp
index ef65d98fed9c..93f318ea1bb7 100644
--- a/clang/test/OpenMP/requires_target_messages.cpp
+++ b/clang/test/OpenMP/requires_target_messages.cpp
@@ -2,14 +2,14 @@
void foo2() {
int a;
- #pragma omp target // expected-note 4 {{target previously encountered here}}
+ #pragma omp target // expected-note 4 {{'target' previously encountered here}}
{
a = a + 1;
}
}
#pragma omp requires atomic_default_mem_order(seq_cst)
-#pragma omp requires unified_address //expected-error {{target region encountered before requires directive with 'unified_address' clause}}
-#pragma omp requires unified_shared_memory //expected-error {{target region encountered before requires directive with 'unified_shared_memory' clause}}
-#pragma omp requires reverse_offload //expected-error {{target region encountered before requires directive with 'reverse_offload' clause}}
-#pragma omp requires dynamic_allocators //expected-error {{target region encountered before requires directive with 'dynamic_allocators' clause}}
+#pragma omp requires unified_address //expected-error {{'target' region encountered before requires directive with 'unified_address' clause}}
+#pragma omp requires unified_shared_memory //expected-error {{'target' region encountered before requires directive with 'unified_shared_memory' clause}}
+#pragma omp requires reverse_offload //expected-error {{'target' region encountered before requires directive with 'reverse_offload' clause}}
+#pragma omp requires dynamic_allocators //expected-error {{'target' region encountered before requires directive with 'dynamic_allocators' clause}}
More information about the cfe-commits
mailing list