[OpenMP] Prevent nesting of target constructs within target code exec…

…ution regions. Summary: This patch enhances Sema to check for the following restriction: OpenMP 4.5 [2.17 Nesting of Regions] If a target, target update, target data, target enter data, or target exit data construct is encountered during execution of a target region, the behavior is unspecified. Reviewers: ABataev Differential Revision: http://reviews.llvm.org/D16758 llvm-svn: 259366
llvm · Feb 1, 2016 · f195862 · f195862
1 parent 04bf91a
commit f195862
Show file tree

Hide file tree

Showing 22 changed files with 423 additions and 74 deletions.
diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h
@@ -156,11 +156,20 @@ bool isOpenMPTaskLoopDirective(OpenMPDirectiveKind DKind);
 /// parallel', otherwise - false.
 bool isOpenMPParallelDirective(OpenMPDirectiveKind DKind);
 
-/// \brief Checks if the specified directive is a target-kind directive.
+/// \brief Checks if the specified directive is a target code offload directive.
 /// \param DKind Specified directive.
-/// \return true - the directive is a target-like directive like 'omp target',
+/// \return true - the directive is a target code offload directive like
+/// 'omp target', 'omp target parallel', 'omp target xxx'
 /// otherwise - false.
-bool isOpenMPTargetDirective(OpenMPDirectiveKind DKind);
+bool isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind);
+
+/// \brief Checks if the specified directive is a target data offload directive.
+/// \param DKind Specified directive.
+/// \return true - the directive is a target data offload directive like
+/// 'omp target data', 'omp target update', 'omp target enter data',
+/// 'omp target exit data'
+/// otherwise - false.
+bool isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind);
 
 /// \brief Checks if the specified directive is a teams-kind directive.
 /// \param DKind Specified directive.

diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
@@ -576,8 +576,15 @@ bool clang::isOpenMPParallelDirective(OpenMPDirectiveKind DKind) {
          // TODO add next directives.
 }
 
-bool clang::isOpenMPTargetDirective(OpenMPDirectiveKind DKind) {
-  return DKind == OMPD_target; // TODO add next directives.
+bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
+  // TODO add next directives.
+  return DKind == OMPD_target || DKind == OMPD_target_parallel;
+}
+
+bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
+  // TODO add target update directive check.
+  return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
+         DKind == OMPD_target_exit_data;
 }
 
 bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
@@ -809,7 +809,7 @@ bool Sema::IsOpenMPCapturedByRef(ValueDecl *D,
   auto DKind = DSAStack->getDirectiveForScope(RSI->TheScope);
   auto Ty = D->getType();
 
-  if (isOpenMPTargetDirective(DKind)) {
+  if (isOpenMPTargetExecutionDirective(DKind)) {
     // This table summarizes how a given variable should be passed to the device
     // given its type and the clauses where it appears. This table is based on
     // the description in OpenMP 4.5 [2.10.4, target Construct] and
@@ -907,7 +907,7 @@ bool Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
         DSAStack->hasDirective(
             [](OpenMPDirectiveKind K, const DeclarationNameInfo &DNI,
                SourceLocation Loc) -> bool {
-              return isOpenMPTargetDirective(K);
+              return isOpenMPTargetExecutionDirective(K);
             },
             false)) {
       return true;
@@ -944,7 +944,8 @@ bool Sema::isOpenMPTargetCapturedDecl(ValueDecl *D, unsigned Level) {
 
   auto *VD = dyn_cast<VarDecl>(D);
   return VD && !VD->hasLocalStorage() &&
-         DSAStack->hasExplicitDirective(isOpenMPTargetDirective, Level);
+         DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
+                                        Level);
 }
 
 void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }
@@ -2313,11 +2314,11 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
   // | target           | flush           | *                                  |
   // | target           | ordered         | *                                  |
   // | target           | atomic          | *                                  |
-  // | target           | target          | *                                  |
-  // | target           | target parallel | *                                  |
-  // | target           | target enter    | *                                  |
+  // | target           | target          |                                    |
+  // | target           | target parallel |                                    |
+  // | target           | target enter    |                                    |
   // |                  | data            |                                    |
-  // | target           | target exit     | *                                  |
+  // | target           | target exit     |                                    |
   // |                  | data            |                                    |
   // | target           | teams           | *                                  |
   // | target           | cancellation    |                                    |
@@ -2347,11 +2348,11 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
   // | target parallel  | flush           | *                                  |
   // | target parallel  | ordered         | *                                  |
   // | target parallel  | atomic          | *                                  |
-  // | target parallel  | target          | *                                  |
-  // | target parallel  | target parallel | *                                  |
-  // | target parallel  | target enter    | *                                  |
+  // | target parallel  | target          |                                    |
+  // | target parallel  | target parallel |                                    |
+  // | target parallel  | target enter    |                                    |
   // |                  | data            |                                    |
-  // | target parallel  | target exit     | *                                  |
+  // | target parallel  | target exit     |                                    |
   // |                  | data            |                                    |
   // | target parallel  | teams           |                                    |
   // | target parallel  | cancellation    |                                    |
@@ -2498,6 +2499,7 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
   // +------------------+-----------------+------------------------------------+
   if (Stack->getCurScope()) {
     auto ParentRegion = Stack->getParentDirective();
+    auto OffendingRegion = ParentRegion;
     bool NestingProhibited = false;
     bool CloseNesting = true;
     enum {
@@ -2658,10 +2660,32 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
       NestingProhibited = !isOpenMPTeamsDirective(ParentRegion);
       Recommend = ShouldBeInTeamsRegion;
     }
+    if (!NestingProhibited &&
+        (isOpenMPTargetExecutionDirective(CurrentRegion) ||
+         isOpenMPTargetDataManagementDirective(CurrentRegion))) {
+      // OpenMP 4.5 [2.17 Nesting of Regions]
+      // If a target, target update, target data, target enter data, or
+      // target exit data construct is encountered during execution of a
+      // target region, the behavior is unspecified.
+      OpenMPDirectiveKind PreviousTargetExecutionDirective;
+      NestingProhibited = Stack->hasDirective(
+          [&PreviousTargetExecutionDirective](OpenMPDirectiveKind K,
+                                              const DeclarationNameInfo &DNI,
+                                              SourceLocation Loc) -> bool {
+            if (isOpenMPTargetExecutionDirective(K)) {
+              PreviousTargetExecutionDirective = K;
+              return true;
+            } else
+              return false;
+          },
+          false /* don't skip top directive */);
+      CloseNesting = false;
+      OffendingRegion = PreviousTargetExecutionDirective;
+    }
     if (NestingProhibited) {
       SemaRef.Diag(StartLoc, diag::err_omp_prohibited_region)
-          << CloseNesting << getOpenMPDirectiveName(ParentRegion) << Recommend
-          << getOpenMPDirectiveName(CurrentRegion);
+          << CloseNesting << getOpenMPDirectiveName(OffendingRegion)
+          << Recommend << getOpenMPDirectiveName(CurrentRegion);
       return true;
     }
   }

diff --git a/clang/test/OpenMP/distribute_private_messages.cpp b/clang/test/OpenMP/distribute_private_messages.cpp
@@ -98,6 +98,7 @@ int main(int argc, char **argv) {
   #pragma omp target
   #pragma omp teams firstprivate(i)
   #pragma omp parallel private(i)
+  {}
   #pragma omp target
   #pragma omp teams reduction(+:i)
   #pragma omp distribute private(i)
@@ -113,20 +114,20 @@ int main(int argc, char **argv) {
   #pragma omp teams
   #pragma omp distribute firstprivate(i)
   for (int k = 0; k < 10; ++k) {
-    #pragma omp target
-    #pragma omp teams firstprivate(i)
-    #pragma omp distribute private(i)
-    for (int x = 0; x < 10; ++x) foo();
   }
   #pragma omp target
+  #pragma omp teams firstprivate(i)
+  #pragma omp distribute private(i)
+  for (int x = 0; x < 10; ++x) foo();
+  #pragma omp target
   #pragma omp teams reduction(+:i)
   #pragma omp distribute
   for (int k = 0; k < 10; ++k) {
-    #pragma omp target
-    #pragma omp teams reduction(+:i)
-    #pragma omp distribute private(i)
-    for (int x = 0; x < 10; ++x) foo();
   }
+  #pragma omp target
+  #pragma omp teams reduction(+:i)
+  #pragma omp distribute private(i)
+  for (int x = 0; x < 10; ++x) foo();
 
   return 0;
 }
diff --git a/clang/test/OpenMP/nesting_of_regions.cpp b/clang/test/OpenMP/nesting_of_regions.cpp
@@ -2666,12 +2666,12 @@ void foo() {
   }
 #pragma omp target
   {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
@@ -2699,11 +2699,11 @@ void foo() {
   }
 #pragma omp target
   {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
   {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 
 // TARGET PARALLEL DIRECTIVE
@@ -2796,12 +2796,12 @@ void foo() {
   }
 #pragma omp target parallel
   {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target parallel' region}}
     ++a;
   }
 #pragma omp target parallel
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target parallel' region}}
     ++a;
   }
 #pragma omp target parallel
@@ -2829,11 +2829,11 @@ void foo() {
   }
 #pragma omp target parallel
   {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 #pragma omp target parallel
   {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 
 // TEAMS DIRECTIVE
@@ -2952,7 +2952,7 @@ void foo() {
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
@@ -3337,28 +3337,28 @@ void foo() {
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
@@ -5807,21 +5807,21 @@ void foo() {
   }
 #pragma omp target
   {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
   {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
   {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
   {
@@ -5937,12 +5937,12 @@ void foo() {
   }
 #pragma omp target parallel
   {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target parallel' region}}
     ++a;
   }
 #pragma omp target parallel
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target parallel' region}}
     ++a;
   }
 #pragma omp target parallel
@@ -5970,11 +5970,11 @@ void foo() {
   }
 #pragma omp target parallel
   {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 #pragma omp target parallel
   {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 
 // TEAMS DIRECTIVE
@@ -6093,7 +6093,7 @@ void foo() {
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
@@ -6477,14 +6477,14 @@ void foo() {
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target parallel
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
@@ -6499,14 +6499,14 @@ void foo() {
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target enter data map(to: a)
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
-#pragma omp target exit data map(from: a)
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 }
diff --git a/clang/test/OpenMP/target_data_device_messages.cpp b/clang/test/OpenMP/target_data_device_messages.cpp
@@ -21,8 +21,8 @@ int main(int argc, char **argv) {
   #pragma omp target data map(to: a) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target data' cannot contain more than one 'device' clause}}
   #pragma omp target data map(to: a) device (S1) // expected-error {{'S1' does not refer to a value}}
   #pragma omp target data map(to: a) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
-  #pragma omp target map(to: a) device (-10u)
-  #pragma omp target map(to: a) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  #pragma omp target data map(to: a) device (-10u)
+  #pragma omp target data map(to: a) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
   foo();
 
   return 0;