diff options
author | Samuel Antao <sfantao@us.ibm.com> | 2016-05-26 18:30:22 +0000 |
---|---|---|
committer | Samuel Antao <sfantao@us.ibm.com> | 2016-05-26 18:30:22 +0000 |
commit | 8d2d730f2a50bad6812f16369edb57b3333d7db9 (patch) | |
tree | 003964c69c0745d4fe3c3220305a276694160f1a /clang/lib/CodeGen/CGStmtOpenMP.cpp | |
parent | 143f684a79a28a02a85eb450ad5cee7a86062b1a (diff) | |
download | bcm5719-llvm-8d2d730f2a50bad6812f16369edb57b3333d7db9.tar.gz bcm5719-llvm-8d2d730f2a50bad6812f16369edb57b3333d7db9.zip |
[OpenMP] Codegen for target update directive.
Summary: This patch implements the code generation for the `target update` directive. The implemntation relies on the logic already in place for target data standalone directives, i.e. target enter/exit data.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D20650
llvm-svn: 270886
Diffstat (limited to 'clang/lib/CodeGen/CGStmtOpenMP.cpp')
-rw-r--r-- | clang/lib/CodeGen/CGStmtOpenMP.cpp | 23 |
1 files changed, 18 insertions, 5 deletions
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index e20ac0aebc0..2f4a2c7f7b6 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -3366,8 +3366,7 @@ void CodeGenFunction::EmitOMPTargetEnterDataDirective( if (auto *C = S.getSingleClause<OMPDeviceClause>()) Device = C->getDevice(); - CGM.getOpenMPRuntime().emitTargetEnterOrExitDataCall(*this, S, IfCond, - Device); + CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device); } void CodeGenFunction::EmitOMPTargetExitDataDirective( @@ -3387,8 +3386,7 @@ void CodeGenFunction::EmitOMPTargetExitDataDirective( if (auto *C = S.getSingleClause<OMPDeviceClause>()) Device = C->getDevice(); - CGM.getOpenMPRuntime().emitTargetEnterOrExitDataCall(*this, S, IfCond, - Device); + CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device); } void CodeGenFunction::EmitOMPTargetParallelDirective( @@ -3550,5 +3548,20 @@ void CodeGenFunction::EmitOMPTaskLoopSimdDirective( // Generate the instructions for '#pragma omp target update' directive. void CodeGenFunction::EmitOMPTargetUpdateDirective( const OMPTargetUpdateDirective &S) { - // TODO: codegen for target update + // If we don't have target devices, don't bother emitting the data mapping + // code. + if (CGM.getLangOpts().OMPTargetTriples.empty()) + return; + + // Check if we have any if clause associated with the directive. + const Expr *IfCond = nullptr; + if (auto *C = S.getSingleClause<OMPIfClause>()) + IfCond = C->getCondition(); + + // Check if we have any device clause associated with the directive. + const Expr *Device = nullptr; + if (auto *C = S.getSingleClause<OMPDeviceClause>()) + Device = C->getDevice(); + + CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device); } |