[OPENMP]Allow cancellation constructs in target parallel regions.

[ABataev]

Summary

omp cancellation point parallel and omp cancel parallel directives are
allowed in target paralle regions.

[jdoerfert]

LGTM.

clang/include/clang/AST/StmtOpenMP.h

 /// \code
 /// #pragma omp target parallel if(a)
 /// \endcode
 /// In this example directive '#pragma omp target parallel' has clause 'if' with
 /// condition 'a'.
 ///
 class OMPTargetParallelDirective : public OMPExecutableDirective {
   friend class ASTStmtReader;
+  /// true if the construct has inner cancel directive.
+  bool HasCancel = false;
+
   /// Build directive with the given start and end location.
   ///
   /// \param StartLoc Starting location of the directive kind.
   /// \param EndLoc Ending location of the directive.
   /// \param NumClauses Number of clauses.
   ///
   OMPTargetParallelDirective(SourceLocation StartLoc, SourceLocation EndLoc,
                              unsigned NumClauses)
       : OMPExecutableDirective(this, OMPTargetParallelDirectiveClass,
                                llvm::omp::OMPD_target_parallel, StartLoc,
                                EndLoc, NumClauses, /*NumChildren=*/1) {}
 
   /// Build an empty directive.
   ///
   /// \param NumClauses Number of clauses.
   ///
   explicit OMPTargetParallelDirective(unsigned NumClauses)
       : OMPExecutableDirective(this, OMPTargetParallelDirectiveClass,
                                llvm::omp::OMPD_target_parallel,
                                SourceLocation(), SourceLocation(), NumClauses,
                                /*NumChildren=*/1) {}
 
+  /// Set cancel state.
+  void setHasCancel(bool Has) { HasCancel = Has; }
+
 public:
   /// Creates directive with a list of \a Clauses.
   ///
   /// \param C AST context.
   /// \param StartLoc Starting location of the directive kind.
   /// \param EndLoc Ending Location of the directive.
   /// \param Clauses List of clauses.
   /// \param AssociatedStmt Statement, associated with the directive.
+  /// \param HasCancel true if this directive has inner cancel directive.
   ///
   static OMPTargetParallelDirective *
   Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
-         ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt);
+         ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, bool HasCancel);
 
   /// Creates an empty directive with the place for \a NumClauses
   /// clauses.
   ///
   /// \param C AST context.
   /// \param NumClauses Number of clauses.
   ///
   static OMPTargetParallelDirective *
   CreateEmpty(const ASTContext &C, unsigned NumClauses, EmptyShell);
 
+  /// Return true if current directive has inner cancel directive.
+  bool hasCancel() const { return HasCancel; }
+
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTargetParallelDirectiveClass;
   }
 };
 
 /// This represents '#pragma omp target parallel for' directive.
 ///
 /// \code

clang/lib/AST/StmtOpenMP.cpp

unsigned Size =
       llvm::alignTo(sizeof(OMPTargetDirective), alignof(OMPClause *));
   void *Mem =
       C.Allocate(Size + sizeof(OMPClause *) * NumClauses + sizeof(Stmt *));
   return new (Mem) OMPTargetDirective(NumClauses);
 }
 
 OMPTargetParallelDirective *OMPTargetParallelDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
-    ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt) {
+    ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, bool HasCancel) {
   unsigned Size =
       llvm::alignTo(sizeof(OMPTargetParallelDirective), alignof(OMPClause *));
   void *Mem =
       C.Allocate(Size + sizeof(OMPClause *) * Clauses.size() + sizeof(Stmt *));
   OMPTargetParallelDirective *Dir =
       new (Mem) OMPTargetParallelDirective(StartLoc, EndLoc, Clauses.size());
   Dir->setClauses(Clauses);
   Dir->setAssociatedStmt(AssociatedStmt);
+  Dir->setHasCancel(HasCancel);
   return Dir;
 }
 
 OMPTargetParallelDirective *
 OMPTargetParallelDirective::CreateEmpty(const ASTContext &C,
                                         unsigned NumClauses, EmptyShell) {
   unsigned Size =
       llvm::alignTo(sizeof(OMPTargetParallelDirective), alignof(OMPClause *));

clang/lib/CodeGen/CGOpenMPRuntime.cpp

static llvm::Function *emitParallelOrTeamsOutlinedFunction(
     const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
     const StringRef OutlinedHelperName, const RegionCodeGenTy &CodeGen) {
   assert(ThreadIDVar->getType()->isPointerType() &&
          "thread id variable must be of type kmp_int32 *");
   CodeGenFunction CGF(CGM, true);
   bool HasCancel = false;
   if (const auto *OPD = dyn_cast<OMPParallelDirective>(&D))
     HasCancel = OPD->hasCancel();
+  else if (const auto *OPD = dyn_cast<OMPTargetParallelDirective>(&D))
+    HasCancel = OPD->hasCancel();
   else if (const auto *OPSD = dyn_cast<OMPParallelSectionsDirective>(&D))
     HasCancel = OPSD->hasCancel();
   else if (const auto *OPFD = dyn_cast<OMPParallelForDirective>(&D))
     HasCancel = OPFD->hasCancel();
   else if (const auto *OPFD = dyn_cast<OMPTargetParallelForDirective>(&D))
     HasCancel = OPFD->hasCancel();
   else if (const auto *OPFD = dyn_cast<OMPDistributeParallelForDirective>(&D))
     HasCancel = OPFD->hasCancel();

clang/lib/Sema/SemaOpenMP.cpp

for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_parallel);
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
   setFunctionHasBranchProtectedScope();
 
   return OMPTargetParallelDirective::Create(Context, StartLoc, EndLoc, Clauses,
-                                            AStmt);
+                                            AStmt, DSAStack->isCancelRegion());
 }
 
 StmtResult Sema::ActOnOpenMPTargetParallelForDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 

clang/lib/Serialization/ASTReaderStmt.cpp

void ASTStmtReader::VisitOMPTargetExitDataDirective(
   VisitOMPExecutableDirective(D);
 }
 
 void ASTStmtReader::VisitOMPTargetParallelDirective(
     OMPTargetParallelDirective *D) {
   VisitStmt(D);
   Record.skipInts(1);
   VisitOMPExecutableDirective(D);
+  D->setHasCancel(Record.readBool());
 }
 
 void ASTStmtReader::VisitOMPTargetParallelForDirective(
     OMPTargetParallelForDirective *D) {
   VisitOMPLoopDirective(D);
   D->setHasCancel(Record.readInt());
 }
 

clang/lib/Serialization/ASTWriterStmt.cpp

void ASTStmtWriter::VisitOMPTargetExitDataDirective(
   Code = serialization::STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPTargetParallelDirective(
     OMPTargetParallelDirective *D) {
   VisitStmt(D);
   Record.push_back(D->getNumClauses());
   VisitOMPExecutableDirective(D);
+  Record.writeBool(D->hasCancel());
   Code = serialization::STMT_OMP_TARGET_PARALLEL_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPTargetParallelForDirective(
     OMPTargetParallelForDirective *D) {
   VisitOMPLoopDirective(D);
   Record.push_back(D->hasCancel() ? 1 : 0);
   Code = serialization::STMT_OMP_TARGET_PARALLEL_FOR_DIRECTIVE;

clang/test/OpenMP/target_parallel_ast_print.cpp

 
 #pragma omp target parallel depend(in : argc, argv[i:argc], a[:])
 // CHECK-NEXT: #pragma omp target parallel depend(in : argc,argv[i:argc],a[:])
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target parallel defaultmap(tofrom: scalar) reduction(task, +:argc)
 // CHECK-NEXT: #pragma omp target parallel defaultmap(tofrom: scalar) reduction(task, +: argc)
+  {
   foo();
+#pragma omp cancellation point parallel
+#pragma omp cancel parallel
+  }
+// CHECK-NEXT: {
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp cancellation point parallel
+// CHECK-NEXT: #pragma omp cancel parallel
+// CHECK-NEXT: }
 
   return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
 }
 
 extern template int S<int>::TS;
 extern template char S<char>::TS;
 
 #endif

clang/test/OpenMP/target_parallel_codegen.cpp

int foo(int n) {
   // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
   // CHECK:       [[FAIL]]
   // CHECK:       call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}})
   // CHECK-NEXT:  br label %[[END]]
   // CHECK:       [[END]]
   #pragma omp target parallel if(target: 1)
   {
     aa += 1;
+#pragma omp cancel parallel
   }
 
   // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
   // CHECK:       br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
   // CHECK:       [[IFTHEN]]
   // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT3]], i32 0, i32 0), i32 1, i32 0)
   // CHECK-DAG:   [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]], i32 0, i32 0
   // CHECK-DAG:   [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]], i32 0, i32 0

 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]])* [[OMP_OUTLINED2:@.+]] to void (i32*, i32*, ...)*), i[[SZ]] [[PARAM]])
 //
 //
 // CHECK:       define internal {{.*}}void [[OMP_OUTLINED2]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}})
 // CHECK:       [[AA_ADDR:%.+]] = alloca i[[SZ]], align
 // CHECK:       store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align
 // CHECK:       [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
 // CHECK:       [[AA:%.+]] = load i16, i16* [[AA_CADDR]], align
+// CHECK:       [[IS_CANCEL:%.+]] = call i32 @__kmpc_cancel(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 1)
+// CHECK:       [[CMP:%.+]] = icmp ne i32 [[IS_CANCEL]], 0
+// CHECK:       br i1 [[CMP]], label %[[EXIT:.+]], label %[[CONTINUE:[^,]+]]
+// CHECK:       [[EXIT]]:
+// CHECK:       br label %[[CONTINUE]]
+// CHECK:       [[CONTINUE]]:
 // CHECK:       ret void
 // CHECK-NEXT:  }
 
 // CHECK:       define internal void [[HVT3]]
 // Create stack storage and store argument in there.
 // CHECK:       [[A_ADDR:%.+]] = alloca i[[SZ]], align
 // CHECK:       [[AA_ADDR:%.+]] = alloca i[[SZ]], align
 // CHECK:       [[A_CASTED:%.+]] = alloca i[[SZ]], align

clang/test/OpenMP/target_parallel_messages.cpp

int main(int argc, char **argv) {
 
   [[]] // expected-error {{an attribute list cannot appear here}}
   #pragma omp target parallel
   for (int n = 0; n < 100; ++n) {}
 
   #pragma omp target parallel copyin(pvt) // expected-error {{unexpected OpenMP clause 'copyin' in directive '#pragma omp target parallel'}}
   foo();
 
+  #pragma omp target parallel
+  {
+#pragma omp cancel // expected-error {{one of 'for', 'parallel', 'sections' or 'taskgroup' is expected}}
+#pragma omp cancellation point // expected-error {{one of 'for', 'parallel', 'sections' or 'taskgroup' is expected}}
+#pragma omp cancel for // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancellation point for // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancel sections // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancellation point sections // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancel taskgroup // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancellation point taskgroup // expected-error {{region cannot be closely nested inside 'target parallel' region}}
+#pragma omp cancel parallel
+#pragma omp cancellation point parallel
+  }
+
   return 0;
 }