From: Alexey Bataev <a.bataev@hotmail.com>
Date: Mon, 27 Nov 2017 16:54:08 +0000 (+0000)
Subject: [OPENMP] Improve handling of cancel directives in target-based
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=61396a3c038285ab7f13984f1fdf80a217faf646;p=clang

[OPENMP] Improve handling of cancel directives in target-based
constructs, NFC.

Improved handling of cancel|cancellation point directives inside
target-based for directives.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@319046 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 0d8b5c2fff..3de5d3d84f 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1293,6 +1293,13 @@ static llvm::Value *emitParallelOrTeamsOutlinedFunction(
     HasCancel = OPFD->hasCancel();
   else if (auto *OPFD = dyn_cast<OMPTargetParallelForDirective>(&D))
     HasCancel = OPFD->hasCancel();
+  else if (auto *OPFD = dyn_cast<OMPDistributeParallelForDirective>(&D))
+    HasCancel = OPFD->hasCancel();
+  else if (auto *OPFD = dyn_cast<OMPTeamsDistributeParallelForDirective>(&D))
+    HasCancel = OPFD->hasCancel();
+  else if (auto *OPFD =
+               dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&D))
+    HasCancel = OPFD->hasCancel();
   CGOpenMPOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen, InnermostKind,
                                     HasCancel, OutlinedHelperName);
   CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
diff --git a/lib/CodeGen/CGStmtOpenMP.cpp b/lib/CodeGen/CGStmtOpenMP.cpp
index af9f7e2eff..22a36eb0b5 100644
--- a/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2039,8 +2039,7 @@ void CodeGenFunction::EmitOMPDistributeParallelForDirective(
                               S.getDistInc());
   };
   OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
-  CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen,
-                                              S.hasCancel());
+  CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen);
 }
 
 void CodeGenFunction::EmitOMPDistributeParallelForSimdDirective(
@@ -3201,8 +3200,7 @@ void CodeGenFunction::EmitOMPDistributeDirective(
     CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc());
   };
   OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
-  CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen,
-                                              false);
+  CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen);
 }
 
 static llvm::Function *emitOutlinedOrderedFunction(CodeGenModule &CGM,
@@ -3915,8 +3913,8 @@ void CodeGenFunction::EmitOMPTeamsDistributeParallelForDirective(
     OMPPrivateScope PrivateScope(CGF);
     CGF.EmitOMPReductionClauseInit(S, PrivateScope);
     (void)PrivateScope.Privatize();
-    CGF.CGM.getOpenMPRuntime().emitInlinedDirective(
-        CGF, OMPD_distribute, CodeGenDistribute, S.hasCancel());
+    CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
+                                                    CodeGenDistribute);
     CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
   };
   emitCommonOMPTeamsDirective(*this, S, OMPD_distribute_parallel_for, CodeGen);
diff --git a/test/OpenMP/target_parallel_for_codegen.cpp b/test/OpenMP/target_parallel_for_codegen.cpp
index 2b725fcbcd..55dd4d2a57 100644
--- a/test/OpenMP/target_parallel_for_codegen.cpp
+++ b/test/OpenMP/target_parallel_for_codegen.cpp
@@ -106,6 +106,7 @@ int foo(int n) {
   #pragma omp target parallel for
   for (int i = 3; i < 32; i += 5) {
 #pragma omp cancel for
+#pragma omp cancellation point for
   }
 
   // CHECK:       call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}, i{{32|64}}{{[*]*}} {{[^)]+}})
@@ -325,6 +326,7 @@ int foo(int n) {
 //
 // CHECK:       define internal {{.*}}void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid.)
 // CHECK:       call i32 @__kmpc_cancel(%ident_t* @
+// CHECK:       call i32 @__kmpc_cancellationpoint(%ident_t* @
 // CHECK:       ret void
 // CHECK:       }