From: Alexey Bataev Date: Mon, 2 Oct 2017 16:32:39 +0000 (+0000) Subject: [OPENMP] Capture argument of `device` clause for target-based X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=3ac9cc204042907aa7a1665d10ace04995f15456;p=clang [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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@314686 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h index 24d51bc747..7c24e34251 100644 --- a/include/clang/AST/OpenMPClause.h +++ b/include/clang/AST/OpenMPClause.h @@ -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 '('. diff --git a/include/clang/AST/RecursiveASTVisitor.h b/include/clang/AST/RecursiveASTVisitor.h index 7459504703..7fb5221b7d 100644 --- a/include/clang/AST/RecursiveASTVisitor.h +++ b/include/clang/AST/RecursiveASTVisitor.h @@ -3109,6 +3109,7 @@ bool RecursiveASTVisitor::VisitOMPDependClause(OMPDependClause *C) { template bool RecursiveASTVisitor::VisitOMPDeviceClause(OMPDeviceClause *C) { + TRY_TO(VisitOMPClauseWithPreInit(C)); TRY_TO(TraverseStmt(C->getDevice())); return true; } diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp index 497c605cae..a4fa4b07a1 100644 --- a/lib/AST/OpenMPClause.cpp +++ b/lib/AST/OpenMPClause.cpp @@ -60,6 +60,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) { return static_cast(C); case OMPC_thread_limit: return static_cast(C); + case OMPC_device: + return static_cast(C); case OMPC_default: case OMPC_proc_bind: case OMPC_final: @@ -83,7 +85,6 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) { case OMPC_capture: case OMPC_seq_cst: case OMPC_depend: - case OMPC_device: case OMPC_threads: case OMPC_simd: case OMPC_map: diff --git a/lib/CodeGen/CGStmtOpenMP.cpp b/lib/CodeGen/CGStmtOpenMP.cpp index 20f1535209..2f819a14bc 100644 --- a/lib/CodeGen/CGStmtOpenMP.cpp +++ b/lib/CodeGen/CGStmtOpenMP.cpp @@ -2091,6 +2091,7 @@ void CodeGenFunction::EmitOMPTeamsDistributeParallelForDirective( 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::EmitOMPTargetTeamsDistributeDirective( 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::EmitOMPTargetTeamsDistributeParallelForDirective( 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::EmitOMPTargetTeamsDistributeParallelForSimdDirective( 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 &) { diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp index 4d04f5a837..8bded5854a 100644 --- a/lib/Sema/SemaOpenMP.cpp +++ b/lib/Sema/SemaOpenMP.cpp @@ -10708,6 +10708,7 @@ OMPClause *Sema::ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc, 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(Expr *Device, SourceLocation StartLoc, /*StrictlyPositive=*/false)) return nullptr; - return new (Context) OMPDeviceClause(ValExpr, StartLoc, LParenLoc, EndLoc); + OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective(); + if (isOpenMPTargetExecutionDirective(DKind) && + !CurContext->isDependentContext()) { + llvm::MapVector 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, diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp index bff8f9d05f..2b8a10f174 100644 --- a/lib/Serialization/ASTReaderStmt.cpp +++ b/lib/Serialization/ASTReaderStmt.cpp @@ -2347,6 +2347,7 @@ void OMPClauseReader::VisitOMPDependClause(OMPDependClause *C) { } void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) { + VisitOMPClauseWithPreInit(C); C->setDevice(Reader->Record.readSubExpr()); C->setLParenLoc(Reader->ReadSourceLocation()); } diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp index 1bc64b6a17..15d2425713 100644 --- a/lib/Serialization/ASTWriterStmt.cpp +++ b/lib/Serialization/ASTWriterStmt.cpp @@ -2102,6 +2102,7 @@ void OMPClauseWriter::VisitOMPDependClause(OMPDependClause *C) { } void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) { + VisitOMPClauseWithPreInit(C); Record.AddStmt(C->getDevice()); Record.AddSourceLocation(C->getLParenLoc()); } diff --git a/test/OpenMP/target_codegen.cpp b/test/OpenMP/target_codegen.cpp index 5073fd760c..8debf92e8e 100644 --- a/test/OpenMP/target_codegen.cpp +++ b/test/OpenMP/target_codegen.cpp @@ -92,14 +92,17 @@ int foo(int n) { double cn[5][n]; TT 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) { }