[OpenMP] Add checks for requires and target directives.

[gtbercea]

Summary

The requires directive containing target related clauses must appear before any target region in the compilation unit.

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

9136

Enclose %0 in single quotes for better reading

lib/Sema/SemaOpenMP.cpp

2440

Better to check for TargetLocations.empty() before the loop

4208

Better to check for the dependent context at first, only after that check for the clause

test/OpenMP/requires_target_messages.cpp

12

Add the checks for all the clauses that may cause this error

[gtbercea]

Fix test,

[gtbercea]

• Fix test.

[ABataev]

lib/Sema/SemaOpenMP.cpp

4208

Better to check for the dependent context at first, only after that check for the clause

test/OpenMP/requires_target_messages.cpp

6

Just expected-note 4 {{

[gtbercea]

• Fix.

[gtbercea]

• Fix.

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

9136

The messages must start from lower letter. Other messages must be fixed too in separate patches.

[gtbercea]

• Lower case.

[ABataev]

LG

include/clang/Basic/DiagnosticSemaKinds.td

def err_omp_linear_distribute_var_non_loop_iteration : Error<
   "only loop iteration variables are allowed in 'linear' clause in distribute directives">;
 def warn_omp_non_trivial_type_mapped : Warning<
   "Non-trivial type %0 is mapped, only trivial types are guaranteed to be mapped correctly">,
   InGroup<OpenMPTarget>;
 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_invalid_scope : Error <
   "'#pragma omp %0' directive must appear only in file scope">;
 def note_omp_invalid_length_on_this_ptr_mapping : Note <
   "expected length on mapping of 'this' array section expression to be '1'">;
 def note_omp_invalid_lower_bound_on_this_ptr_mapping : Note <
   "expected lower bound on mapping of 'this' array section expression to be '0' or not specified">;
 def note_omp_invalid_subscript_on_this_ptr_map : Note <
   "expected 'this' subscript expression on map clause to be 'this[0]'">;

lib/Sema/SemaOpenMP.cpp

private:
 
   /// Vector of previously declared requires directives
   SmallVector<const OMPRequiresDecl *, 2> RequiresDecls;
   /// omp_allocator_handle_t type.
   QualType OMPAllocatorHandleT;
   /// Expression for the predefined allocators.
   Expr *OMPPredefinedAllocators[OMPAllocateDeclAttr::OMPUserDefinedMemAlloc] = {
       nullptr};
+  /// Vector of previously encountered target directives
+  SmallVector<SourceLocation, 2> TargetLocations;
 
 public:
   explicit DSAStackTy(Sema &S) : SemaRef(S) {}
 
   /// Sets omp_allocator_handle_t type.
   void setOMPAllocatorHandleT(QualType Ty) { OMPAllocatorHandleT = Ty; }
   /// Gets omp_allocator_handle_t type.
   QualType getOMPAllocatorHandleT() const { return OMPAllocatorHandleT; }
for (OMPClause *CNew : ClauseList) {
             IsDuplicate = true;
           }
         }
       }
     }
     return IsDuplicate;
   }
 
+  /// Add location of previously encountered target to internal vector
+  void addTargetDirLocation(SourceLocation LocStart) {
+    TargetLocations.push_back(LocStart);
+  }
+
+  // Return previously encountered target region locations.
+  ArrayRef<SourceLocation> getEncounteredTargetLocs() const {
+    return TargetLocations;
+  }
+
   /// Set default data sharing attribute to none.
   void setDefaultDSANone(SourceLocation Loc) {
     assert(!isStackEmpty());
     Stack.back().first.back().DefaultAttr = DSA_none;
     Stack.back().first.back().DefaultAttrLoc = Loc;
   }
   /// Set default data sharing attribute to shared.
   void setDefaultDSAShared(SourceLocation Loc) {
if (D) {
       DSAStack->addRequiresDecl(D);
     }
   }
   return DeclGroupPtrTy::make(DeclGroupRef(D));
 }
 
 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();
+  if (!TargetLocations.empty()) {
+    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());
+        for (SourceLocation TargetLoc : TargetLocations) {
+          Diag(TargetLoc, diag::note_omp_requires_encountered_target);
+        }
+      }
+    }
+  }
+
   if (!DSAStack->hasDuplicateRequiresClause(ClauseList))
     return OMPRequiresDecl::Create(Context, getCurLexicalContext(), Loc,
                                    ClauseList);
   return nullptr;
 }
 
 static void reportOriginalDsa(Sema &SemaRef, const DSAStackTy *Stack,
                               const ValueDecl *D,
if (ErrorFound)
     return StmtError();
 
   if (!(Res.getAs<OMPExecutableDirective>()->isStandaloneDirective())) {
     Res.getAs<OMPExecutableDirective>()
         ->getStructuredBlock()
         ->setIsOMPStructuredBlock(true);
   }
 
+  if (!CurContext->isDependentContext() &&
+      isOpenMPTargetExecutionDirective(Kind) &&
+      !(DSAStack->hasRequiresDeclWithClause<OMPUnifiedSharedMemoryClause>() ||
+        DSAStack->hasRequiresDeclWithClause<OMPUnifiedAddressClause>() ||
+        DSAStack->hasRequiresDeclWithClause<OMPReverseOffloadClause>() ||
+        DSAStack->hasRequiresDeclWithClause<OMPDynamicAllocatorsClause>())) {
+    // Register target to DSA Stack.
+    DSAStack->addTargetDirLocation(StartLoc);
+  }
+
   return Res;
 }
 
 Sema::DeclGroupPtrTy Sema::ActOnOpenMPDeclareSimdDirective(
     DeclGroupPtrTy DG, OMPDeclareSimdDeclAttr::BranchStateTy BS, Expr *Simdlen,
     ArrayRef<Expr *> Uniforms, ArrayRef<Expr *> Aligneds,
     ArrayRef<Expr *> Alignments, ArrayRef<Expr *> Linears,
     ArrayRef<unsigned> LinModifiers, ArrayRef<Expr *> Steps, SourceRange SR) {

test/OpenMP/requires_messages.cpp

 // RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100  %s
 
 int a;
 #pragma omp requires unified_address allocate(a) // expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note {{unified_address clause previously used here}} expected-note{{unified_address clause previously used here}} expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp requires'}}
 
 #pragma omp requires unified_shared_memory // expected-note {{unified_shared_memory clause previously used here}} expected-note{{unified_shared_memory clause previously used here}}
 
 #pragma omp requires unified_shared_memory, unified_shared_memory // expected-error {{Only one unified_shared_memory clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'unified_shared_memory' clause}}
 
-#pragma omp requires unified_address // expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}} 
+#pragma omp requires unified_address // expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}}
 
 #pragma omp requires unified_address, unified_address // expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'unified_address' clause}}
 
 #pragma omp requires reverse_offload // expected-note {{reverse_offload clause previously used here}} expected-note {{reverse_offload clause previously used here}}
 
 #pragma omp requires reverse_offload, reverse_offload // expected-error {{Only one reverse_offload clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'reverse_offload' clause}}
 
 #pragma omp requires dynamic_allocators // expected-note {{dynamic_allocators clause previously used here}} expected-note {{dynamic_allocators clause previously used here}}
 
 #pragma omp requires dynamic_allocators, dynamic_allocators // expected-error {{Only one dynamic_allocators clause can appear on a requires directive in a single translation unit}} expected-error {{directive '#pragma omp requires' cannot contain more than one 'dynamic_allocators' clause}}
 
 #pragma omp requires atomic_default_mem_order(seq_cst) // expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}} expected-note {{atomic_default_mem_order clause previously used here}}
 
 #pragma omp requires atomic_default_mem_order(acq_rel) // 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 {{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 // expected-error {{expected '(' after 'atomic_default_mem_order'}} expected-error {{expected at least one clause on '#pragma omp requires' directive}}
 
 #pragma omp requires atomic_default_mem_order( // expected-error {{expected ')'}} expected-note {{to match this '('}} 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}}
 
-#pragma omp requires atomic_default_mem_order(seq_cst // expected-error {{expected ')'}} expected-note {{to match this '('}} 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(seq_cst // expected-error {{expected ')'}} expected-note {{to match this '('}} 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(invalid_modifier) // 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}} 
+#pragma omp requires atomic_default_mem_order(invalid_modifier) // 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}}
 
 #pragma omp requires atomic_default_mem_order(shared) // 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}}
 
-#pragma omp requires atomic_default_mem_order(acq_rel), atomic_default_mem_order(relaxed) // expected-error {{directive '#pragma omp requires' cannot contain more than one 'atomic_default_mem_order' claus}} 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(acq_rel), atomic_default_mem_order(relaxed) // expected-error {{directive '#pragma omp requires' cannot contain more than one 'atomic_default_mem_order' claus}} expected-error {{Only one atomic_default_mem_order clause can appear on a requires directive in a single translation unit}}
 
 #pragma omp requires // expected-error {{expected at least one clause on '#pragma omp requires' directive}}
 
 #pragma omp requires invalid_clause // expected-warning {{extra tokens at the end of '#pragma omp requires' are ignored}} expected-error {{expected at least one clause on '#pragma omp requires' directive}}
 
 #pragma omp requires nowait // expected-error {{unexpected OpenMP clause 'nowait' in directive '#pragma omp requires'}} expected-error {{expected at least one clause on '#pragma omp requires' directive}}
 
 #pragma omp requires unified_address, invalid_clause // expected-warning {{extra tokens at the end of '#pragma omp requires' are ignored}} expected-error {{Only one unified_address clause can appear on a requires directive in a single translation unit}}

test/OpenMP/requires_target_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100  %s
+
+void foo2() {
+  int a;
+  #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}}