r314686 - [OPENMP] Capture argument of `device` clause for target-based
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 2 09:32:39 PDT 2017
Author: abataev
Date: Mon Oct 2 09:32:39 2017
New Revision: 314686
URL: http://llvm.org/viewvc/llvm-project?rev=314686&view=rev
Log:
[OPENMP] Capture argument of `device` clause for target-based
directives.
The argument of the `device` clause in target-based executable
directives must be captured to support codegen for the `target`
directives with the `depend` clauses.
Modified:
cfe/trunk/include/clang/AST/OpenMPClause.h
cfe/trunk/include/clang/AST/RecursiveASTVisitor.h
cfe/trunk/lib/AST/OpenMPClause.cpp
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/lib/Serialization/ASTReaderStmt.cpp
cfe/trunk/lib/Serialization/ASTWriterStmt.cpp
cfe/trunk/test/OpenMP/target_codegen.cpp
Modified: cfe/trunk/include/clang/AST/OpenMPClause.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/OpenMPClause.h (original)
+++ cfe/trunk/include/clang/AST/OpenMPClause.h Mon Oct 2 09:32:39 2017
@@ -3141,7 +3141,7 @@ public:
/// In this example directive '#pragma omp target' has clause 'device'
/// with single expression 'a'.
///
-class OMPDeviceClause : public OMPClause {
+class OMPDeviceClause : public OMPClause, public OMPClauseWithPreInit {
friend class OMPClauseReader;
/// \brief Location of '('.
SourceLocation LParenLoc;
@@ -3161,16 +3161,19 @@ public:
/// \param LParenLoc Location of '('.
/// \param EndLoc Ending location of the clause.
///
- OMPDeviceClause(Expr *E, SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc)
- : OMPClause(OMPC_device, StartLoc, EndLoc), LParenLoc(LParenLoc),
- Device(E) {}
+ OMPDeviceClause(Expr *E, Stmt *HelperE, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc)
+ : OMPClause(OMPC_device, StartLoc, EndLoc), OMPClauseWithPreInit(this),
+ LParenLoc(LParenLoc), Device(E) {
+ setPreInitStmt(HelperE);
+ }
/// \brief Build an empty clause.
///
OMPDeviceClause()
- : OMPClause(OMPC_device, SourceLocation(), SourceLocation()),
- LParenLoc(SourceLocation()), Device(nullptr) {}
+ : OMPClause(OMPC_device, SourceLocation(), SourceLocation()),
+ OMPClauseWithPreInit(this), LParenLoc(SourceLocation()),
+ Device(nullptr) {}
/// \brief Sets the location of '('.
void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; }
/// \brief Returns the location of '('.
Modified: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/RecursiveASTVisitor.h?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/RecursiveASTVisitor.h (original)
+++ cfe/trunk/include/clang/AST/RecursiveASTVisitor.h Mon Oct 2 09:32:39 2017
@@ -3109,6 +3109,7 @@ bool RecursiveASTVisitor<Derived>::Visit
template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPDeviceClause(OMPDeviceClause *C) {
+ TRY_TO(VisitOMPClauseWithPreInit(C));
TRY_TO(TraverseStmt(C->getDevice()));
return true;
}
Modified: cfe/trunk/lib/AST/OpenMPClause.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/lib/AST/OpenMPClause.cpp (original)
+++ cfe/trunk/lib/AST/OpenMPClause.cpp Mon Oct 2 09:32:39 2017
@@ -60,6 +60,8 @@ const OMPClauseWithPreInit *OMPClauseWit
return static_cast<const OMPNumTeamsClause *>(C);
case OMPC_thread_limit:
return static_cast<const OMPThreadLimitClause *>(C);
+ case OMPC_device:
+ return static_cast<const OMPDeviceClause *>(C);
case OMPC_default:
case OMPC_proc_bind:
case OMPC_final:
@@ -83,7 +85,6 @@ const OMPClauseWithPreInit *OMPClauseWit
case OMPC_capture:
case OMPC_seq_cst:
case OMPC_depend:
- case OMPC_device:
case OMPC_threads:
case OMPC_simd:
case OMPC_map:
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Mon Oct 2 09:32:39 2017
@@ -2091,6 +2091,7 @@ void CodeGenFunction::EmitOMPTeamsDistri
void CodeGenFunction::EmitOMPTargetTeamsDistributeDirective(
const OMPTargetTeamsDistributeDirective &S) {
+ OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(
*this, OMPD_target_teams_distribute,
[&S](CodeGenFunction &CGF, PrePostActionTy &) {
@@ -2101,6 +2102,7 @@ void CodeGenFunction::EmitOMPTargetTeams
void CodeGenFunction::EmitOMPTargetTeamsDistributeParallelForDirective(
const OMPTargetTeamsDistributeParallelForDirective &S) {
+ OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(
*this, OMPD_target_teams_distribute_parallel_for,
[&S](CodeGenFunction &CGF, PrePostActionTy &) {
@@ -2111,6 +2113,7 @@ void CodeGenFunction::EmitOMPTargetTeams
void CodeGenFunction::EmitOMPTargetTeamsDistributeParallelForSimdDirective(
const OMPTargetTeamsDistributeParallelForSimdDirective &S) {
+ OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(
*this, OMPD_target_teams_distribute_parallel_for_simd,
[&S](CodeGenFunction &CGF, PrePostActionTy &) {
@@ -2121,6 +2124,7 @@ void CodeGenFunction::EmitOMPTargetTeams
void CodeGenFunction::EmitOMPTargetTeamsDistributeSimdDirective(
const OMPTargetTeamsDistributeSimdDirective &S) {
+ OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
CGM.getOpenMPRuntime().emitInlinedDirective(
*this, OMPD_target_teams_distribute_simd,
[&S](CodeGenFunction &CGF, PrePostActionTy &) {
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Mon Oct 2 09:32:39 2017
@@ -10708,6 +10708,7 @@ OMPClause *Sema::ActOnOpenMPDeviceClause
SourceLocation LParenLoc,
SourceLocation EndLoc) {
Expr *ValExpr = Device;
+ Stmt *HelperValStmt = nullptr;
// OpenMP [2.9.1, Restrictions]
// The device expression must evaluate to a non-negative integer value.
@@ -10715,7 +10716,16 @@ OMPClause *Sema::ActOnOpenMPDeviceClause
/*StrictlyPositive=*/false))
return nullptr;
- return new (Context) OMPDeviceClause(ValExpr, StartLoc, LParenLoc, EndLoc);
+ OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
+ if (isOpenMPTargetExecutionDirective(DKind) &&
+ !CurContext->isDependentContext()) {
+ llvm::MapVector<Expr *, DeclRefExpr *> Captures;
+ ValExpr = tryBuildCapture(*this, ValExpr, Captures).get();
+ HelperValStmt = buildPreInits(Context, Captures);
+ }
+
+ return new (Context)
+ OMPDeviceClause(ValExpr, HelperValStmt, StartLoc, LParenLoc, EndLoc);
}
static bool CheckTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
Modified: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTReaderStmt.cpp Mon Oct 2 09:32:39 2017
@@ -2347,6 +2347,7 @@ void OMPClauseReader::VisitOMPDependClau
}
void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
+ VisitOMPClauseWithPreInit(C);
C->setDevice(Reader->Record.readSubExpr());
C->setLParenLoc(Reader->ReadSourceLocation());
}
Modified: cfe/trunk/lib/Serialization/ASTWriterStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTWriterStmt.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTWriterStmt.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTWriterStmt.cpp Mon Oct 2 09:32:39 2017
@@ -2102,6 +2102,7 @@ void OMPClauseWriter::VisitOMPDependClau
}
void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) {
+ VisitOMPClauseWithPreInit(C);
Record.AddStmt(C->getDevice());
Record.AddSourceLocation(C->getLParenLoc());
}
Modified: cfe/trunk/test/OpenMP/target_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_codegen.cpp?rev=314686&r1=314685&r2=314686&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_codegen.cpp Mon Oct 2 09:32:39 2017
@@ -92,14 +92,17 @@ int foo(int n) {
double cn[5][n];
TT<long long, char> d;
- // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i32* null)
+ // CHECK: [[ADD:%.+]] = add nsw i32
+ // CHECK: store i32 [[ADD]], i32* [[CAPTURE:%.+]],
+ // CHECK: [[LD:%.+]] = load i32, i32* [[CAPTURE]],
+ // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i32 [[LD]], i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i32* null)
// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: call void [[HVT0:@.+]]()
// CHECK-NEXT: br label %[[END]]
// CHECK: [[END]]
- #pragma omp target
+ #pragma omp target device(global + a)
{
}
More information about the cfe-commits
mailing list