replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[OPENMP] Improve handling of cancel directives in target-based 

constructs, NFC.

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

llvm-svn: 319046
This commit is contained in:
Alexey Bataev 2017-11-27 16:54:08 +00:00
parent 647dd6a602
commit 10a5431239
3 changed files with 13 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -1293,6 +1293,13 @@ static llvm::Value *emitParallelOrTeamsOutlinedFunction(
HasCancel = OPFD->hasCancel(); HasCancel = OPFD->hasCancel();
else if (auto *OPFD = dyn_cast<OMPTargetParallelForDirective>(&D)) else if (auto *OPFD = dyn_cast<OMPTargetParallelForDirective>(&D))
HasCancel = OPFD->hasCancel(); 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, CGOpenMPOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen, InnermostKind,
HasCancel, OutlinedHelperName); HasCancel, OutlinedHelperName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo); CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);

10
clang/lib/CodeGen/CGStmtOpenMP.cpp
View File

@ -2039,8 +2039,7 @@ void CodeGenFunction::EmitOMPDistributeParallelForDirective(
S.getDistInc()); S.getDistInc());
}; };
OMPLexicalScope Scope(*this, S, /*AsInlined=*/true); OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen, CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen);
S.hasCancel());
} }
void CodeGenFunction::EmitOMPDistributeParallelForSimdDirective( void CodeGenFunction::EmitOMPDistributeParallelForSimdDirective(
@ -3201,8 +3200,7 @@ void CodeGenFunction::EmitOMPDistributeDirective(
CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc()); CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc());
}; };
OMPLexicalScope Scope(*this, S, /*AsInlined=*/true); OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen, CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_distribute, CodeGen);
false);
} }
static llvm::Function *emitOutlinedOrderedFunction(CodeGenModule &CGM, static llvm::Function *emitOutlinedOrderedFunction(CodeGenModule &CGM,
@ -3915,8 +3913,8 @@ void CodeGenFunction::EmitOMPTeamsDistributeParallelForDirective(
OMPPrivateScope PrivateScope(CGF); OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope); CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize(); (void)PrivateScope.Privatize();
CGF.CGM.getOpenMPRuntime().emitInlinedDirective( CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
CGF, OMPD_distribute, CodeGenDistribute, S.hasCancel()); CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams); CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
}; };
emitCommonOMPTeamsDirective(*this, S, OMPD_distribute_parallel_for, CodeGen); emitCommonOMPTeamsDirective(*this, S, OMPD_distribute_parallel_for, CodeGen);

2
clang/test/OpenMP/target_parallel_for_codegen.cpp
View File

@ -106,6 +106,7 @@ int foo(int n) {
#pragma omp target parallel for #pragma omp target parallel for
for (int i = 3; i < 32; i += 5) { for (int i = 3; i < 32; i += 5) {
#pragma omp cancel for #pragma omp cancel for
#pragma omp cancellation point for
} }
// CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}, i{{32|64}}{{[*]*}} {{[^)]+}}) // 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: define internal {{.*}}void [[OMP_OUTLINED]](i32* noalias %.global_tid., i32* noalias %.bound_tid.)
// CHECK: call i32 @__kmpc_cancel(%ident_t* @ // CHECK: call i32 @__kmpc_cancel(%ident_t* @
// CHECK: call i32 @__kmpc_cancellationpoint(%ident_t* @
// CHECK: ret void // CHECK: ret void
// CHECK: } // CHECK: }