[OPENMP] Add support for nested 'declare target' directives

[patricklyster]

Summary

Add the capability to nest multiple declare target directives - including header files within a declare target region.

[ABataev]

LG

[RaviNarayanaswamy]

Is there a way to tell if the header files have matching omp declare target/omp end declare target.
The reason is if one of the header files is missing a matching omp end declare target all files which include it will end up having everything marked with declare target

[patricklyster]

In D51378#1216600, @RaviNarayanaswamy wrote:

Is there a way to tell if the header files have matching omp declare target/omp end declare target.
The reason is if one of the header files is missing a matching omp end declare target all files which include it will end up having everything marked with declare target

Are you suggesting to throw a warning/error at the end of the header file should there be a hanging omp declare target? Wouldn't this still be the same issue with each file that includes it still getting the same warning/error?

There is a corner case when the header file is missing the end declare target and the file that includes it has an extra end declare target. In this case no error will be thrown and all code in between the directives will be associated with the declare target. However, I don't know if this is something worth checking for.

[RaviNarayanaswamy]

We should just go with generating an error if the DeclareTargetNestingLevel is not 0 at the end of compilation unit.
Hard to detect if user accidentally forgot to have end declare in header file and had it in the include file or it was intentional.

[Hahnfeld]

In D51378#1218184, @RaviNarayanaswamy wrote:

We should just go with generating an error if the DeclareTargetNestingLevel is not 0 at the end of compilation unit.
Hard to detect if user accidentally forgot to have end declare in header file and had it in the include file or it was intentional.

That will effectively forbid the legacy approach of doing

#pragma omp declare target
#include <...>
#pragma omp end declare target

as in the test because DeclareTargetNestingLevel will be 1 throughout the header file. I think that's still relevant today, so the condition should be "has the same value as when entering this file".

[patricklyster]

In D51378#1218184, @RaviNarayanaswamy wrote:

We should just go with generating an error if the DeclareTargetNestingLevel is not 0 at the end of compilation unit.
Hard to detect if user accidentally forgot to have end declare in header file and had it in the include file or it was intentional.

This is the current behaviour

[RaviNarayanaswamy]

I did not see the code where check is done if Nestingdepth is 0 at end of compilation.

[patricklyster]

In D51378#1219474, @RaviNarayanaswamy wrote:

I did not see the code where check is done if Nestingdepth is 0 at end of compilation.

Sorry I should clarify. You will get an error if you have an unmatched #pragma omp declare target once you reach the end of the compilation unit (not the header file) because execution will exit the while loop in ParseOpenMP.cpp:774 when it encounters a tok::eof and fall into the else statement on ParseOpenMP.cpp:812 which emits the diagnostic. So it is not explicitly checking if DeclareTargetNestingLevel > 0 but you will see an appropriate diagnostic anyways.

[RaviNarayanaswamy]

Ok.
Thanks.

[patricklyster]

Previous implementation only supported immediately nested declare targets such as below:

#pragma omp declare target
  #pragma omp declare target
    int foo();
  #pragma omp end declare target
  int bar();
#pragma omp end declare target

Update allows for nesting in this format:

#pragma omp declare target
  int foo();
  #pragma omp declare target
    int bar();
  #pragma omp end declare target
#pragma omp end declare target

[ABataev]

LG

include/clang/Basic/DiagnosticSemaKinds.td

def err_omp_linear_expected_int_or_ptr : Error<
   "argument of a linear clause should be of integral or pointer "
   "type, not %0">;
 def warn_omp_linear_step_zero : Warning<
   "zero linear step (%0 %select{|and other variables in clause }1should probably be const)">,
   InGroup<OpenMPClauses>;
 def warn_omp_alignment_not_power_of_two : Warning<
   "aligned clause will be ignored because the requested alignment is not a power of 2">,
   InGroup<OpenMPClauses>;
-def err_omp_enclosed_declare_target : Error<
-  "declare target region may not be enclosed within another declare target region">;
 def err_omp_invalid_target_decl : Error<
   "%0 used in declare target directive is not a variable or a function name">;
 def err_omp_declare_target_multiple : Error<
   "%0 appears multiple times in clauses on the same declare target directive">;
 def err_omp_declare_target_to_and_link : Error<
   "%0 must not appear in both clauses 'to' and 'link'">;
 def warn_omp_not_in_target_context : Warning<
   "declaration is not declared in any declare target region">,

include/clang/Sema/Sema.h

public:
   /// \return true if type is disabled.
   bool checkOpenCLDisabledDecl(const NamedDecl &D, const Expr &E);
 
   //===--------------------------------------------------------------------===//
   // OpenMP directives and clauses.
   //
 private:
   void *VarDataSharingAttributesStack;
-  /// Set to true inside '#pragma omp declare target' region.
-  bool IsInOpenMPDeclareTargetContext = false;
+  /// Number of nested '#pragma omp declare target' directives.
+  unsigned DeclareTargetNestingLevel = 0;
   /// Initialization of data-sharing attributes stack.
   void InitDataSharingAttributesStack();
   void DestroyDataSharingAttributesStack();
   ExprResult
   VerifyPositiveIntegerConstantInClause(Expr *Op, OpenMPClauseKind CKind,
                                         bool StrictlyPositive = true);
   /// Returns OpenMP nesting level for current directive.
   unsigned getOpenMPNestingLevel() const;
void ActOnOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec,
                                     OMPDeclareTargetDeclAttr::MapTypeTy MT,
                                     NamedDeclSetType &SameDirectiveDecls);
   /// Check declaration inside target region.
   void
   checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
                                    SourceLocation IdLoc = SourceLocation());
   /// Return true inside OpenMP declare target region.
   bool isInOpenMPDeclareTargetContext() const {
-    return IsInOpenMPDeclareTargetContext;
+    return DeclareTargetNestingLevel > 0;
   }
   /// Return true inside OpenMP target region.
   bool isInOpenMPTargetExecutionDirective() const;
   /// Return true if (un)supported features for the current target should be
   /// diagnosed if OpenMP (offloading) is enabled.
   bool shouldDiagnoseTargetSupportFromOpenMP() const {
     return !getLangOpts().OpenMPIsDevice || isInOpenMPDeclareTargetContext() ||
       isInOpenMPTargetExecutionDirective();

lib/Parse/ParseOpenMP.cpp

case OMPD_declare_target: {
     // Skip the last annot_pragma_openmp_end.
     ConsumeAnyToken();
 
     if (!Actions.ActOnStartOpenMPDeclareTargetDirective(DTLoc))
       return DeclGroupPtrTy();
 
     llvm::SmallVector<Decl *, 4>  Decls;
     DKind = parseOpenMPDirectiveKind(*this);
-    while (DKind != OMPD_end_declare_target && DKind != OMPD_declare_target &&
-           Tok.isNot(tok::eof) && Tok.isNot(tok::r_brace)) {
+    while (DKind != OMPD_end_declare_target && Tok.isNot(tok::eof) &&
+           Tok.isNot(tok::r_brace)) {
       DeclGroupPtrTy Ptr;
       // Here we expect to see some function declaration.
       if (AS == AS_none) {
         assert(TagType == DeclSpec::TST_unspecified);
         MaybeParseCXX11Attributes(Attrs);
         ParsingDeclSpec PDS(*this);
         Ptr = ParseExternalDeclaration(Attrs, &PDS);
       } else {

lib/Sema/SemaOpenMP.cpp

if (!CurLexicalContext->isFileContext() &&
       !CurLexicalContext->isExternCXXContext() &&
       !isa<CXXRecordDecl>(CurLexicalContext) &&
       !isa<ClassTemplateDecl>(CurLexicalContext) &&
       !isa<ClassTemplatePartialSpecializationDecl>(CurLexicalContext) &&
       !isa<ClassTemplateSpecializationDecl>(CurLexicalContext)) {
     Diag(Loc, diag::err_omp_region_not_file_context);
     return false;
   }
-  if (IsInOpenMPDeclareTargetContext) {
-    Diag(Loc, diag::err_omp_enclosed_declare_target);
-    return false;
-  }
-
-  IsInOpenMPDeclareTargetContext = true;
+  ++DeclareTargetNestingLevel;
   return true;
 }
 
 void Sema::ActOnFinishOpenMPDeclareTargetDirective() {
-  assert(IsInOpenMPDeclareTargetContext &&
+  assert(DeclareTargetNestingLevel > 0 &&
          "Unexpected ActOnFinishOpenMPDeclareTargetDirective");
-  IsInOpenMPDeclareTargetContext = false;
+  --DeclareTargetNestingLevel;
 }
 
 void Sema::ActOnOpenMPDeclareTargetName(Scope *CurScope,
                                         CXXScopeSpec &ScopeSpec,
                                         const DeclarationNameInfo &Id,
                                         OMPDeclareTargetDeclAttr::MapTypeTy MT,
                                         NamedDeclSetType &SameDirectiveDecls) {
   LookupResult Lookup(*this, Id, LookupOrdinaryName);

test/OpenMP/Inputs/declare_target_include.h

+#pragma omp declare target
+  void zyx();
+#pragma omp end declare target

test/OpenMP/declare_target_ast_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
+// RUN: %clang_cc1 -verify -fopenmp -I %S/Inputs -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -I %S/Inputs -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 int out_decl_target = 0;
 // CHECK: #pragma omp declare target{{$}}
 // CHECK: int out_decl_target = 0;

 // CHECK: template<> float baz<float>() {
 // CHECK:     return float();
 // CHECK: }
 // CHECK: template<> int baz<int>() {
 // CHECK:     return 1;
 // CHECK: }
 // CHECK: #pragma omp end declare target
 
+#pragma omp declare target
+  #include "declare_target_include.h"
+  void xyz();
+#pragma omp end declare target
+
+// CHECK: #pragma omp declare target
+// CHECK: void zyx();
+// CHECK: #pragma omp end declare target
+// CHECK: #pragma omp declare target
+// CHECK: void xyz();
+// CHECK: #pragma omp end declare target
+
+#pragma omp declare target
+  #pragma omp declare target
+    void abc();
+  #pragma omp end declare target
+  void cba();
+#pragma omp end declare target
+
+// CHECK: #pragma omp declare target
+// CHECK: void abc();
+// CHECK: #pragma omp end declare target
+// CHECK: #pragma omp declare target
+// CHECK: void cba();
+// CHECK: #pragma omp end declare target
+
 int main (int argc, char **argv) {
   foo();
   foo_c();
   foo_cpp();
   test1();
   baz<float>();
   baz<int>();
   return (0);
 }
 
 // CHECK: #pragma omp declare target
 // CHECK-NEXT: int ts = 1;
 // CHECK-NEXT: #pragma omp end declare target
 #endif

test/OpenMP/declare_target_messages.cpp

 void t2() {
 #pragma omp target
   {
     bla2<float>();
     bla2<long>();
   }
 }
 
+#pragma omp declare target
+  void abc();
+#pragma omp end declare target
+void cba();
+#pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}}
 
-#pragma omp declare target // expected-note {{to match this '#pragma omp declare target'}}
+#pragma omp declare target  
+  #pragma omp declare target
+    void def();
+  #pragma omp end declare target
+  void fed();
+
+#pragma omp declare target
 #pragma omp threadprivate(a) // expected-note {{defined as threadprivate or thread local}}
 extern int b;
 int g;
 
 struct T {
   int a;
   virtual int method();
 };
void foo(int p) {
   C object;
   VC object1;
   g = object.method();
   g += object.method1();
   g += object1.method() + p;
   f();
   c();
 }
-#pragma omp declare target // expected-error {{expected '#pragma omp end declare target'}}
+#pragma omp declare target
 void foo1() {}
 #pragma omp end declare target
+
+#pragma omp end declare target
+#pragma omp end declare target
 #pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}}
 
 int C::method() {
   return 0;
 }
 
 struct S {
 #pragma omp declare target